parsing changes for timing model

10 files changed, 29 insertions, 18 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]),
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index b26adfd..39b8ba5 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -130,6 +130,9 @@ unsigned thread_group_offset(int thread,unsigned  wmma_type,unsigned wmma_layout
 			}
 			break;	
 		case VP_MMA:
+		case VP_MMA4:
+		case VP_MMA8:
+		case VP_MMA16:
 				if(wmma_layout==ROW)
 					offset=load_c_float_row[thread_group]+16*in_tg_index;
 				else	
@@ -1703,7 +1706,7 @@ void mapping(int thread,int wmma_type,int wmma_layout,int type,int index,int str
 		}
 		assg_offset=index;
 	}
-	else if( wmma_type==VP_MMA){
+	else if(wmma_type==VP_MMA || wmma_type==VP_MMA4 || wmma_type==VP_MMA8 || wmma_type==VP_MMA16){
 		row=c_row_offset[thread/4]+thread%4;	
 		col=c_col_offset[thread/4]+index;
 		assg_offset=index;
@@ -1761,7 +1764,7 @@ void vp_mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
 			switch(i) {
 			      case 1 ://operand 1
 				for(k=0;k<8;k++){
-					mapping(thrd,VP_MMA,a_layout,S32_TYPE,k,16,row,col,offset);
+					mapping(thrd,wmma_type,a_layout,S32_TYPE,k,16,row,col,offset);
 					if(g_debug_instruction)
 					printf("A:thread=%d,row=%d,col=%d,offset=%d\n",thrd,row,col,offset);
 				 	matrix_a[row][col]=v[offset];
@@ -1769,7 +1772,7 @@ void vp_mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
 			      break;
 			      case 2 ://operand 2
 				for(k=0;k<8;k++){
-					mapping(thrd,VP_MMA,b_layout,S32_TYPE,k,16,row,col,offset);
+					mapping(thrd,wmma_type,b_layout,S32_TYPE,k,16,row,col,offset);
 					if(g_debug_instruction)
 					printf("B:thread=%d,row=%d,col=%d,offset=%d\n",thrd,row,col,offset);
 					if(nelem==1){
@@ -1785,7 +1788,7 @@ void vp_mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
 			      break;
 			      case 3 ://operand 3
 				for(k=0;k<8;k++){
-					mapping(thrd,VP_MMA,ROW,S32_TYPE,k,16,row,col,offset);
+					mapping(thrd,wmma_type,ROW,S32_TYPE,k,16,row,col,offset);
 					if(g_debug_instruction)
 					printf("C:thread=%d,row=%d,col=%d,offset=%d\n",thrd,row,col,offset);
 				 	matrix_c[row][col]=v[offset];
@@ -1848,7 +1851,7 @@ void vp_mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
 		int row_t[8];
 		int col_t[8];	
 		for(k=0;k<8;k++){
-			mapping(thrd,VP_MMA,ROW,type,k,16,row_t[k],col_t[k],offset);
+			mapping(thrd,wmma_type,ROW,type,k,16,row_t[k],col_t[k],offset);
 			if(g_debug_instruction)
 			printf("vp_mma:row:%d,col%d\n",row_t[k],col_t[k]);
 		}
diff --git a/src/cuda-sim/opcodes.h b/src/cuda-sim/opcodes.h
index 31b71d0..ad8d8f1 100644
--- a/src/cuda-sim/opcodes.h
+++ b/src/cuda-sim/opcodes.h
@@ -70,6 +70,9 @@ enum wmma_type{
 	STORE_D,
 	MMA,
 	VP_MMA,
+	VP_MMA4,
+	VP_MMA8,
+	VP_MMA16,
 	ROW,
 	COL,
 	M16N16K16
diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index fea2420..9c78e9f 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -165,6 +165,9 @@ breakaddr  TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE;
 \.c\.sync TC; ptx_lval.int_value = LOAD_C; return WMMA_DIRECTIVE;
 \.d\.sync TC; ptx_lval.int_value = STORE_D; return WMMA_DIRECTIVE;
 \.mma\.sync TC;ptx_lval.int_value=MMA; return WMMA_DIRECTIVE;
+\.mma4\.sync TC;ptx_lval.int_value=VP_MMA4; return WMMA_DIRECTIVE;
+\.mma8\.sync TC;ptx_lval.int_value=VP_MMA8; return WMMA_DIRECTIVE;
+\.mma16\.sync TC;ptx_lval.int_value=VP_MMA16; return WMMA_DIRECTIVE;
 \.b4\.sync TC; ptx_lval.int_value=LOAD_B4;	return WMMA_DIRECTIVE;
 \.b8\.sync TC; ptx_lval.int_value=LOAD_B8;	return WMMA_DIRECTIVE;
 \.b16\.sync TC; ptx_lval.int_value=LOAD_B16;	return WMMA_DIRECTIVE;
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index 55b8e11..9800e1e 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -1095,7 +1095,9 @@ ptx_instruction::ptx_instruction( int opcode,
       		case LOAD_C:
       		case STORE_D:
       		case MMA:
-      		case VP_MMA:
+      		case VP_MMA4:
+      		case VP_MMA8:
+      		case VP_MMA16:
       		  m_wmma_type=last_ptx_inst_option;
       		  break;
       		case ROW: