diff options
| author | aamir <[email protected]> | 2018-08-14 17:46:21 -0700 |
|---|---|---|
| committer | aamir <[email protected]> | 2018-08-14 17:46:21 -0700 |
| commit | cdd42ac0384cb31bcee05ac72d026f1bc8133d26 (patch) | |
| tree | fd63bd4ee1998b32e9ff30098587bc46cfb0bb14 /cuda-kernels/v16p_genericMatrixMultiply.cu | |
| parent | 44f0114ad2c208f69c0c1baa980a5b3bda37e16b (diff) | |
parsing changes for timing model
Diffstat (limited to 'cuda-kernels/v16p_genericMatrixMultiply.cu')
| -rw-r--r-- | cuda-kernels/v16p_genericMatrixMultiply.cu | 8 |
1 files changed, 4 insertions, 4 deletions
diff --git a/cuda-kernels/v16p_genericMatrixMultiply.cu b/cuda-kernels/v16p_genericMatrixMultiply.cu index fd5a0f8..c36d257 100644 --- a/cuda-kernels/v16p_genericMatrixMultiply.cu +++ b/cuda-kernels/v16p_genericMatrixMultiply.cu @@ -20,9 +20,9 @@ void curandErrCheck_(curandStatus_t stat, const char *file, int line) { using namespace nvcuda; // Must be multiples of 16 for wmma code to work -#define MATRIX_M (256) -#define MATRIX_N (256) -#define MATRIX_K (256) +#define MATRIX_M (32) +#define MATRIX_N (32) +#define MATRIX_K (32) // The only dimensions currently supported by WMMA @@ -92,7 +92,7 @@ __global__ void vp_example(int *a, int *b, int *c, int M, int N, int K ) { //vp::mma_sync(acc_frag, a_frag, b_frag, acc_frag); asm("/*"); asm("CPTX_BEGIN"); - asm("vp.mma.sync.row.row.m16n16k16.s32 {%0, %1, %2, %3, %4, %5, %6, %7}, {%8, %9, %10, %11, %12, %13, %14, %15}, {%16, %17, %18, %19}, { %20, %21, %22, %23, %24, %25, %26,%27};" : + asm("vp.mma16.sync.row.row.m16n16k16.s32 {%0, %1, %2, %3, %4, %5, %6, %7}, {%8, %9, %10, %11, %12, %13, %14, %15}, {%16, %17, %18, %19}, { %20, %21, %22, %23, %24, %25, %26,%27};" : "=r"(acc_frag[0]), "=r"(acc_frag[1]),"=r"(acc_frag[2]),"=r"(acc_frag[3]), "=r"(acc_frag[4]),"=r"(acc_frag[5]),"=r"(acc_frag[6]),"=r"(acc_frag[7]): "r"(a_frag[0]),"r"(a_frag[1]),"r"(a_frag[2]),"r"(a_frag[3]), |