6 files changed, 12 insertions, 12 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]),
diff --git a/cuda-kernels/v16p_kernel.cu b/cuda-kernels/v16p_kernel.cu
index 011fdfd..31a1460 100644
--- a/cuda-kernels/v16p_kernel.cu
+++ b/cuda-kernels/v16p_kernel.cu
@@ -100,7 +100,7 @@ __global__ void v4p_example(int *a_int32, int *b_int4, int *c,int *d_int32, int
 	//B16
 	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"(registers_d[0]), "=r"(registers_d[1]),"=r"(registers_d[2]),"=r"(registers_d[3]),
 	"=r"(registers_d[4]),"=r"(registers_d[5]),"=r"(registers_d[6]),"=r"(registers_d[7]):
 	"r"(registers_a[0]),"r"(registers_a[1]),"r"(registers_a[2]),"r"(registers_a[3]),
diff --git a/cuda-kernels/v4p_genericMatrixMultiply.cu b/cuda-kernels/v4p_genericMatrixMultiply.cu
index 1b56eb2..abcab8e 100644
--- a/cuda-kernels/v4p_genericMatrixMultiply.cu
+++ b/cuda-kernels/v4p_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 (1024)
-#define MATRIX_N (1024)
-#define MATRIX_K (1024)
+#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};" : 
+	asm("vp.mma4.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};" : 
 	"=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]),
diff --git a/cuda-kernels/v4p_kernel.cu b/cuda-kernels/v4p_kernel.cu
index 8c9bbd4..f1a9dbe 100644
--- a/cuda-kernels/v4p_kernel.cu
+++ b/cuda-kernels/v4p_kernel.cu
@@ -70,7 +70,7 @@ __global__ void v4p_example(int *a_int32, int *b_int4, int *c,int *d_int32, int
 	//B4
 	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};" : 
+	asm("vp.mma4.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};" : 
 	"=r"(registers_d[0]), "=r"(registers_d[1]),"=r"(registers_d[2]),"=r"(registers_d[3]),
 	"=r"(registers_d[4]),"=r"(registers_d[5]),"=r"(registers_d[6]),"=r"(registers_d[7]):
 	"r"(registers_a[0]),"r"(registers_a[1]),"r"(registers_a[2]),"r"(registers_a[3]),
diff --git a/cuda-kernels/v8p_genericMatrixMultiply.cu b/cuda-kernels/v8p_genericMatrixMultiply.cu
index 2e487e8..208c369 100644
--- a/cuda-kernels/v8p_genericMatrixMultiply.cu
+++ b/cuda-kernels/v8p_genericMatrixMultiply.cu
@@ -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};" : 
+	asm("vp.mma8.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};" : 
 	"=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]),
diff --git a/cuda-kernels/v8p_kernel.cu b/cuda-kernels/v8p_kernel.cu
index b1b0eba..f824eac 100644
--- a/cuda-kernels/v8p_kernel.cu
+++ b/cuda-kernels/v8p_kernel.cu
@@ -85,7 +85,7 @@ __global__ void v4p_example(int *a_int32, int *b_int4, int *c,int *d_int32, int
 	//B8
 	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};" : 
+	asm("vp.mma8.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};" : 
 	"=r"(registers_d[0]), "=r"(registers_d[1]),"=r"(registers_d[2]),"=r"(registers_d[3]),
 	"=r"(registers_d[4]),"=r"(registers_d[5]),"=r"(registers_d[6]),"=r"(registers_d[7]):
 	"r"(registers_a[0]),"r"(registers_a[1]),"r"(registers_a[2]),"r"(registers_a[3]),