added load

9 files changed, 265 insertions, 31 deletions

diff --git a/cuda-kernels/v4p_kernel.cu b/cuda-kernels/v4p_kernel.cu
index bb9064b..2e84eda 100644
--- a/cuda-kernels/v4p_kernel.cu
+++ b/cuda-kernels/v4p_kernel.cu
@@ -30,23 +30,30 @@ const int WMMA_M = 16;
 const int WMMA_N = 16;
 const int WMMA_K = 16;
 
-__global__ void wmma_example(half *a, half *b, float *c,float *d_fp16, int M, int N, int K) {
-   //unsigned int start_time=0,end_time=0;
-   //start_time=clock();
+__global__ void v4p_example(int *a_int32, int *b_int8, int *c,int *d_int32, int M, int N, int K) {
+	
+	int registers_a[8];
+	int register_b;		//contains 8 4bit b elements
+   	int idx = blockDim.x * blockIdx.x + threadIdx.x;
+	
+	asm("/*");
+	asm("CPTX_BEGIN");
+	asm("vp.load.b4.sync.row.m16n16k16.s32 {%0,%1,%2,%3,%4,%5,%6,%7},[%8],%9;" : 
+	"=r"(registers_a[0]), "=r"(registers_a[1]),"=r"(registers_a[2]),"=r"(registers_a[3]),
+	"=r"(registers_a[4]),"=r"(registers_a[5]),"=r"(registers_a[6]),"=r"(registers_a[7]):
+	"l"(b_int8),"r"(M)
+	);
+	asm("CPTX_END");
+	asm("*/");
 
-   // Declare the fragments
-   wmma::fragment<wmma::matrix_a, WMMA_M, WMMA_N, WMMA_K, half, wmma::col_major> a_frag;
-   wmma::fragment<wmma::matrix_b, WMMA_M, WMMA_N, WMMA_K, half, wmma::col_major> b_frag;
-   wmma::fragment<wmma::accumulator, WMMA_M, WMMA_N, WMMA_K, float> c_frag;
-
-   // Bounds checking
-   wmma::load_matrix_sync(a_frag, a, K);
-   wmma::load_matrix_sync(b_frag, b, K);
-   wmma::load_matrix_sync(c_frag, c, N,wmma::mem_col_major);
-   wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
-
-   wmma::store_matrix_sync(d_fp16, c_frag, N, wmma::mem_col_major);
-   //printf("clock=%d",end_time-start_time);
+	b_int8[0]=registers_a[7];
+	b_int8[1]=registers_a[7];
+	b_int8[2]=registers_a[7];
+	b_int8[3]=registers_a[7];
+	b_int8[4]=registers_a[7];
+	b_int8[5]=registers_a[7];
+	b_int8[6]=registers_a[7];
+	b_int8[7]=registers_a[7];
 }
 
 __global__ void convertFp32ToFp16 (half *out, float *in, int n) {
@@ -62,13 +69,35 @@ __global__ void convertFp16ToFp32 (float *out, half *in, int n) {
    }
 }
 
+__global__ void convertInt32ToInt4 (int *out, int *in, int n) {
+   int idx = blockDim.x * blockIdx.x + threadIdx.x;
+   if (idx < n/8) {
+      		out[idx] =(in[8*idx]&0xf)|(in[8*idx+1]&0xf)<<4|(in[8*idx+2]&0xf)<<8|(in[8*idx+3]&0xf)<<12|
+      			  (in[8*idx+4]&0xf)<<16|(in[8*idx+5]&0xf)<<20|(in[8*idx+6]&0xf)<<24|(in[8*idx+7]&0xf)<<28;
+//		printf("thread%d:%x\n",idx,out[idx]);
+   }
+}
 __global__ void convertInt32ToInt8 (int *out, int *in, int n) {
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    if (idx < n/4) {
       		out[idx] =(in[4*idx]&0xff)|(in[4*idx+1]&0xff)<<8|(in[4*idx+2]&0xff)<<16|(in[4*idx+3]&0xff)<<24;
    }
 }
+__global__ void convertInt32ToInt16 (int *out, int *in, int n) {
+   int idx = blockDim.x * blockIdx.x + threadIdx.x;
+   if (idx < n/2) {
+      		out[idx] =(in[2*idx]&0xffff)|(in[2*idx+1]&0xffff)<<16;
+   }
+}
 
+__global__ void convertInt4ToInt32 (int *out, int *in, int n) {
+   int idx = blockDim.x * blockIdx.x + threadIdx.x;
+   int shft_amt=4*(idx%8);
+   int shft_mask=0xf<<shft_amt;
+   if (idx < n) {
+      	out[idx]= (in[idx/8]&shft_mask)>>shft_amt;
+   }
+}
 __global__ void convertInt8ToInt32 (int *out, int *in, int n) {
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    int shft_amt=8*(idx%4);
@@ -77,6 +106,14 @@ __global__ void convertInt8ToInt32 (int *out, int *in, int n) {
       	out[idx]= (in[idx/4]&shft_mask)>>shft_amt;
    }
 }
+__global__ void convertInt16ToInt32 (int *out, int *in, int n) {
+   int idx = blockDim.x * blockIdx.x + threadIdx.x;
+   int shft_amt=16*(idx%2);
+   int shft_mask=0xffff<<shft_amt;
+   if (idx < n) {
+      	out[idx]= (in[idx/2]&shft_mask)>>shft_amt;
+   }
+}
 
 int main(int argc, char* argv[]) {
    int *a_int32;
@@ -84,8 +121,12 @@ int main(int argc, char* argv[]) {
    int *c_int32;
    int *d_int32;
 
+   int *a_int4;
+   int *b_int4;
    int *a_int8;
    int *b_int8;
+   int *a_int16;
+   int *b_int16;
    
    int  *a_host_wmma;
    int  *b_host_wmma;
@@ -105,15 +146,19 @@ int main(int argc, char* argv[]) {
    cudaErrCheck(cudaMalloc((void**)&b_int32, MATRIX_K * MATRIX_N * sizeof(int)));
    cudaErrCheck(cudaMalloc((void**)&c_int32, MATRIX_K * MATRIX_N * sizeof(int)));
    cudaErrCheck(cudaMalloc((void**)&d_int32, MATRIX_K * MATRIX_N * sizeof(int)));
+   cudaErrCheck(cudaMalloc((void**)&a_int4, MATRIX_M * MATRIX_K * sizeof(int)/8));
+   cudaErrCheck(cudaMalloc((void**)&b_int4, MATRIX_K * MATRIX_N * sizeof(int)/8));
    cudaErrCheck(cudaMalloc((void**)&a_int8, MATRIX_M * MATRIX_K * sizeof(int)/4));
    cudaErrCheck(cudaMalloc((void**)&b_int8, MATRIX_K * MATRIX_N * sizeof(int)/4));
+   cudaErrCheck(cudaMalloc((void**)&a_int16, MATRIX_M * MATRIX_K * sizeof(int)/2));
+   cudaErrCheck(cudaMalloc((void**)&b_int16, MATRIX_K * MATRIX_N * sizeof(int)/2));
 
 
    a_host_wmma      = (int *)malloc(MATRIX_M * MATRIX_K * sizeof(int));
    b_host_wmma      = (int *)malloc(MATRIX_K * MATRIX_N * sizeof(int));
    c_host_wmma      = (int *)malloc(MATRIX_M * MATRIX_N * sizeof(int));
    d_host_wmma      = (int *)malloc(MATRIX_M * MATRIX_N * sizeof(int));
-   d_cal_host_wmma      = (int *)malloc(MATRIX_M * MATRIX_N * sizeof(int));
+   d_cal_host_wmma  = (int *)malloc(MATRIX_M * MATRIX_N * sizeof(int));
 
    printf("a_int32\n");
    for(int m=0;m<MATRIX_M;m++){
@@ -156,20 +201,34 @@ int main(int argc, char* argv[]) {
    cudaErrCheck(cudaMemcpy(b_int32,b_host_wmma,  MATRIX_K * MATRIX_N * sizeof(int), cudaMemcpyHostToDevice));
    cudaErrCheck(cudaMemcpy(c_int32,c_host_wmma,  MATRIX_M * MATRIX_N * sizeof(int), cudaMemcpyHostToDevice));
 
-   convertInt32ToInt8 <<< (MATRIX_M * MATRIX_K + 255) / 256, 256 >>> (a_int8, a_int32, MATRIX_M * MATRIX_K);
-   convertInt8ToInt32 <<< (MATRIX_M * MATRIX_K + 255) / 256, 256 >>> (d_int32, a_int8, MATRIX_M * MATRIX_K);
+   #define TEST8
+   #ifdef TEST16
+   	convertInt32ToInt16 <<< (MATRIX_M * MATRIX_K + 255) / 256, 256 >>> (a_int16, a_int32, MATRIX_M * MATRIX_K);
+   	convertInt16ToInt32 <<< (MATRIX_M * MATRIX_K + 255) / 256, 256 >>> (d_int32, a_int16, MATRIX_M * MATRIX_K);
+   	cudaErrCheck(cudaMemcpy(d_host_wmma, d_int32, MATRIX_M * MATRIX_N * sizeof(int), cudaMemcpyDeviceToHost));
+   #endif
+   #ifdef TEST8
+   	convertInt32ToInt8 <<< (MATRIX_M * MATRIX_K + 255) / 256, 256 >>> (a_int8, a_int32, MATRIX_M * MATRIX_K);
+  	convertInt8ToInt32 <<< (MATRIX_M * MATRIX_K + 255) / 256, 256 >>> (d_int32, a_int8, MATRIX_M * MATRIX_K);
+ 	cudaErrCheck(cudaMemcpy(d_host_wmma, d_int32, MATRIX_M * MATRIX_N * sizeof(int), cudaMemcpyDeviceToHost));
+   #endif
+   #ifdef TEST4
+   	convertInt32ToInt4 <<< (MATRIX_M * MATRIX_K + 255) / 256, 256 >>> (a_int4, a_int32, MATRIX_M * MATRIX_K);
+  	convertInt4ToInt32 <<< (MATRIX_M * MATRIX_K + 255) / 256, 256 >>> (d_int32, a_int4, MATRIX_M * MATRIX_K);
+ 	cudaErrCheck(cudaMemcpy(d_host_wmma, d_int32, MATRIX_M * MATRIX_N * sizeof(int), cudaMemcpyDeviceToHost));
+   #endif
    //convertFp32ToFp16 <<< (MATRIX_K * MATRIX_N + 255) / 256, 256 >>> (b_fp16, b_fp32, MATRIX_K * MATRIX_N);
    //convertFp32ToFp16 <<< (MATRIX_M * MATRIX_N + 255) / 256, 256 >>> (c_fp16, c_fp32, MATRIX_K * MATRIX_N);
-   cudaErrCheck(cudaMemcpy(d_host_wmma, d_int32, MATRIX_M * MATRIX_N * sizeof(float), cudaMemcpyDeviceToHost));
 
 
 //AAMIR   printf("\nM = %d, N = %d, K = %d. \n", MATRIX_M, MATRIX_N, MATRIX_K);
 //AAMIR   
 //AAMIR   printf("Running with wmma...\n");
-//AAMIR   cudaErrCheck(cudaEventRecord(startWMMA));
-//AAMIR   wmma_example <<< 1, 32>>> (a_fp16, b_fp16, c_fp32, d_fp32 , MATRIX_M, MATRIX_N, MATRIX_K);
-//AAMIR   cudaErrCheck(cudaEventRecord(stopWMMA));
-//AAMIR   cudaErrCheck(cudaEventSynchronize(stopWMMA));
+   //cudaErrCheck(cudaEventRecord(startWMMA));
+   //v4p_example <<< 1, 32>>> (a_int32, a_int8, a_int32, d_int32, MATRIX_M, MATRIX_N, MATRIX_K);
+   //cudaErrCheck(cudaEventRecord(stopWMMA));
+   //cudaErrCheck(cudaEventSynchronize(stopWMMA));
+
 //AAMIR
 //AAMIR   // Error checking
 //AAMIR   printf("\nChecking results...\n");
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 8284ad5..2fe5667 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -1323,8 +1323,11 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
          *((warp_inst_t*)pJ) = inst; // copy active mask information
          pI = pJ;
       }
-      if(((pI->get_opcode()!=MMA_OP)&&(pI->get_opcode()!=MMA_LD_OP)&&(pI->get_opcode()!=MMA_ST_OP))||((pI->get_opcode()==MMA_OP||pI->get_opcode()==MMA_LD_OP||pI->get_opcode()==MMA_ST_OP)&&(lane_id==0))){
-      switch ( pI->get_opcode() ) {
+     
+      int inst_opcode=pI->get_opcode();
+
+      if(((inst_opcode!=MMA_OP)&&(inst_opcode!=MMA_LD_OP)&&(inst_opcode!=MMA_ST_OP)&&(inst_opcode!=VP_LD_OP)&&(inst_opcode!=VP_ST_OP)&&(inst_opcode!=VP_MMA_OP))||((inst_opcode==MMA_OP||inst_opcode==MMA_LD_OP||inst_opcode==MMA_ST_OP||inst_opcode==VP_MMA_OP||inst_opcode==VP_LD_OP||inst_opcode==VP_ST_OP)&&(lane_id==0))){
+      switch ( inst_opcode ) {
 #define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,this); op_classification = CLASSIFICATION; break;
 #define OP_W_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,get_core(),inst); op_classification = CLASSIFICATION; break;
 #include "opcodes.def"
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index aaee2a2..d0396a5 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -48,7 +48,7 @@
 using half_float::half;
 
 unsigned ptx_instruction::g_num_ptx_inst_uid=0;
-bool g_debug_instruction = 0;
+bool g_debug_instruction = 1;
 
 
 const char *g_opcode_string[NUM_OPCODES] = {
@@ -129,6 +129,12 @@ unsigned thread_group_offset(int thread,unsigned  wmma_type,unsigned wmma_layout
 				}
 			}
 			break;	
+		case VP_MMA:
+				if(wmma_layout==ROW)
+					offset=load_c_float_row[thread_group]+16*in_tg_index;
+				else	
+					offset=load_c_float_col[thread_group]+16*in_tg_index;
+			break;	
 
          	default:
          	   abort();
@@ -1709,6 +1715,14 @@ void mapping(int thread,int wmma_type,int wmma_layout,int type,int index,int str
 	}
 }
 
+void vp_mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
+{
+	unsigned wmma_type = pI->get_wmma_type();
+	unsigned a_layout = pI->get_wmma_layout(0);
+	unsigned b_layout = pI->get_wmma_layout(1);
+	
+
+}
 void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
 {
 	int i,j,k,thrd;
@@ -1752,7 +1766,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
 			int hex_val;
 
 			if(!((i==3)&&(type2==F32_TYPE))){
-				for(k=0;k<2*nelem;k++){
+					for(k=0;k<2*nelem;k++){
 					if(k%2==1)
 						hex_val=(v[k/2].s64&0xffff);
 					else
@@ -2958,6 +2972,12 @@ void ldu_impl( const ptx_instruction *pI, ptx_thread_info *thread )
    ld_exec(pI,thread);
 }
 
+void vp_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
+{
+   const operand_info &src  = pI->operand_lookup(1);
+   const operand_info &src1 = pI->operand_lookup(0);
+   const operand_info &src2 = pI->operand_lookup(2);
+}
 void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
 {
    size_t size;
@@ -3000,7 +3020,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
 
    	type_info_key::type_decode(type,size,t);
    	if(g_debug_instruction)
-		printf("mma_st: thrd=%d,addr=%d, fp(size=%d), stride=%d\n",thrd,addr_reg.u32,size,src2_data.u32);
+		printf("mma_st: thrd=%d,addr=%x, fp(size=%d), stride=%d\n",thrd,addr_reg.u32,size,src2_data.u32);
 	addr_t new_addr = addr+thread_group_offset(thrd,wmma_type,wmma_layout,type,stride)*size/8;  
 
 	ptx_reg_t nw_v[8];
@@ -3042,6 +3062,133 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
    	thread->m_last_memory_space = space;
    } 
 }
+void vp_ld_impl(const ptx_instruction *pI, core_t *core, warp_inst_t inst)
+{
+   size_t size;
+   int t,i;
+   unsigned smid;
+   const operand_info &dst = pI->dst();
+   const operand_info &src1 = pI->src1();
+   const operand_info &src2 = pI->src2();
+   unsigned type = pI->get_type();
+   unsigned wmma_type = pI->get_wmma_type();
+   unsigned wmma_layout = pI->get_wmma_layout(0);
+   int tid = inst.warp_id_func()*core->get_warp_size();
+   int thrd,stride;
+   ptx_thread_info *thread;
+   for (thrd=0; thrd < core->get_warp_size(); thrd++){
+   	thread = core->get_thread_info()[tid+thrd];
+   	ptx_reg_t src1_data = thread->get_operand_value(src1, dst, U32_TYPE, thread, 1);
+   	ptx_reg_t src2_data = thread->get_operand_value(src2, dst, U32_TYPE, thread, 1);
+	stride=src2_data.u32;
+   	memory_space_t space = pI->get_space();
+
+   	memory_space *mem = NULL;
+   	addr_t addr = src1_data.u32;
+   	
+        smid = thread->get_hw_sid();
+        if( whichspace(addr) == shared_space ) {
+          addr= generic_to_shared(smid,addr);
+          space = shared_space;
+	}
+
+	decode_space(space,thread,src1,mem,addr);
+   	type_info_key::type_decode(type,size,t);
+	
+	ptx_reg_t data[8];
+	addr_t new_addr;	
+	//note we are using distribution of VP_MMA for every type of load!
+	if(wmma_type==LOAD_A||wmma_type==LOAD_C){
+		 new_addr = addr+thread_group_offset(thrd,VP_MMA,wmma_layout,type,stride)*size/8;  
+	}
+	else if(wmma_type==LOAD_B4){
+		new_addr = addr+thread_group_offset(thrd,VP_MMA,wmma_layout,type,stride)/8;  
+	}
+	else if (wmma_type==LOAD_B8){
+		new_addr = addr+thread_group_offset(thrd,VP_MMA,wmma_layout,type,stride)/4;  
+	}
+	else if (wmma_type==LOAD_B16){
+		new_addr = addr+thread_group_offset(thrd,VP_MMA,wmma_layout,type,stride)/2;  
+	}
+
+   	if(g_debug_instruction)	
+		printf("vp_ld: thrx=%d,addr=%x, base_addr=%x, size=%d, stride=%d\n",thrd,new_addr,addr,size,src2_data.u32);
+
+	if(wmma_type==LOAD_A||wmma_type==LOAD_C){
+		for(i=0;i<8;i++){
+			mem->read(new_addr+4*i,size/8,&data[i].s64);
+		}
+	}
+	else if(wmma_type==LOAD_B4){
+			mem->read(new_addr,size/8,&data[0].s64);
+	}
+	else if(wmma_type==LOAD_B8){
+			mem->read(new_addr,size/8,&data[0].s64);
+			mem->read(new_addr+4,size/8,&data[1].s64);
+	}
+	else if(wmma_type==LOAD_B16){
+			mem->read(new_addr,size/8,&data[0].s64);
+			mem->read(new_addr+4,size/8,&data[1].s64);
+			mem->read(new_addr+8,size/8,&data[2].s64);
+			mem->read(new_addr+12,size/8,&data[3].s64);
+	}
+	else{
+		printf("wrong vp_load type\n");;
+		abort();
+	}
+
+	int num_reg;
+	if(g_debug_instruction){
+		printf("\nvp_ld:thread%d= ",thrd);
+		if((wmma_type==LOAD_A)||(wmma_type==LOAD_C)){	
+			num_reg=8;	
+		}	
+		else if(wmma_type==LOAD_B4){
+			num_reg=1;	
+		}
+		else if(wmma_type==LOAD_B8){
+			num_reg=2;	
+		}
+		else if(wmma_type==LOAD_B16){
+			num_reg=4;	
+		}
+	
+		for(i=0;i<num_reg;i++){
+			printf("%x ",data[i].s32);
+		}
+		printf("\n");
+	}
+
+	if((wmma_type==LOAD_A)||(wmma_type==LOAD_C)){
+   		thread->set_wmma_vector_operand_values(dst,data[0],data[1],data[2],data[3],data[4],data[5],data[6],data[7]);
+	}
+	else if(wmma_type==LOAD_B4){
+      		thread->set_operand_value(dst,data[0], type, thread, pI);
+	}
+	else if(wmma_type==LOAD_B8){
+   		thread->set_vector_operand_values(dst,data[0],data[1],data[1],data[1]);
+	}
+	else if (wmma_type==LOAD_B16){
+   		thread->set_vector_operand_values(dst,data[0],data[1],data[2],data[3]);
+	}
+	else
+		abort();	
+
+	
+
+	if(g_debug_instruction){
+		for(int i=0;i<num_reg;i++)
+			printf("vp_ld:data[%i].s64=%x\n",i,data[i].u64);	
+	}
+	
+
+   	thread->m_last_effective_address = addr;
+   	thread->m_last_memory_space = space;
+    }
+
+
+
+}
 
 void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
 {
@@ -3080,7 +3227,7 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
 	
 	ptx_reg_t data[16];
    	if(g_debug_instruction)	
-		printf("mma_ld: thrd=%d,addr=%d, fpsize=%d, stride=%d\n",thrd,src1_data.u32,size,src2_data.u32);
+		printf("mma_ld: thrd=%d,addr=%x, fpsize=%d, stride=%d\n",thrd,src1_data.u32,size,src2_data.u32);
 	
 	addr_t new_addr = addr+thread_group_offset(thrd,wmma_type,wmma_layout,type,stride)*size/8;  
 
diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def
index e6f957a..73f6161 100644
--- a/src/cuda-sim/opcodes.def
+++ b/src/cuda-sim/opcodes.def
@@ -55,6 +55,9 @@ OP_DEF(BRKPT_OP,brkpt_impl,"brkpt",1,9)
 OP_W_DEF(MMA_OP,mma_impl,"mma",1,1)
 OP_W_DEF(MMA_LD_OP,mma_ld_impl,"mma_load",1,5)
 OP_W_DEF(MMA_ST_OP,mma_st_impl,"mma_store",0,5)
+OP_W_DEF(VP_MMA_OP,vp_mma_impl,"vp_mma",1,1)
+OP_W_DEF(VP_LD_OP,vp_ld_impl,"vp_load",1,5)
+OP_W_DEF(VP_ST_OP,vp_st_impl,"vp_store",0,5)
 OP_DEF(CALL_OP,call_impl,"call",1,3)
 OP_DEF(CALLP_OP,callp_impl,"callp",1,3)
 OP_DEF(CLZ_OP,clz_impl,"clz",1,1)
diff --git a/src/cuda-sim/opcodes.h b/src/cuda-sim/opcodes.h
index b91d92f..31b71d0 100644
--- a/src/cuda-sim/opcodes.h
+++ b/src/cuda-sim/opcodes.h
@@ -63,9 +63,13 @@ enum special_regs {
 enum wmma_type{
 	LOAD_A,
 	LOAD_B,
+	LOAD_B4,//vp
+	LOAD_B8,//vp
+	LOAD_B16,//vp
 	LOAD_C,
 	STORE_D,
 	MMA,
+	VP_MMA,
 	ROW,
 	COL,
 	M16N16K16
diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index a6b6fcc..fea2420 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -68,9 +68,13 @@ bra     TC; ptx_lval.int_value = BRA_OP; return OPCODE;
 brx     TC; ptx_lval.int_value = BRX_OP; return OPCODE;
 brev    TC; ptx_lval.int_value = BREV_OP; return OPCODE;
 brkpt   TC; ptx_lval.int_value = BRKPT_OP; return OPCODE;
+
 wmma	TC; ptx_lval.int_value = MMA_OP; return OPCODE;
 wmma\.load	TC; ptx_lval.int_value = MMA_LD_OP; return OPCODE;
 wmma\.store 	TC; ptx_lval.int_value = MMA_ST_OP; return OPCODE;
+vp\.load TC; ptx_lval.int_value=VP_LD_OP; return OPCODE;
+vp\.store TC; ptx_lval.int_value=VP_ST_OP; return OPCODE;
+vp 	TC; ptx_lval.int_value=VP_MMA_OP; return OPCODE;
 
 call	TC; BEGIN(NOT_OPCODE); ptx_lval.int_value = CALL_OP; return OPCODE; // blocking opcode token in case the callee has the same name as an opcode
 callp   TC; BEGIN(NOT_OPCODE); ptx_lval.int_value = CALLP_OP; return OPCODE;
@@ -161,6 +165,10 @@ 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;
+\.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;
+
 \.row TC; ptx_lval.int_value = ROW; return LAYOUT;
 \.col TC; ptx_lval.int_value = COL; return LAYOUT;
 \.m16n16k16 TC; ptx_lval.int_value = M16N16K16; return CONFIGURATION;
diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y
index cd455dd..54ac2bc 100644
--- a/src/cuda-sim/ptx.y
+++ b/src/cuda-sim/ptx.y
@@ -538,6 +538,12 @@ wmma_spec: WMMA_DIRECTIVE LAYOUT CONFIGURATION{add_space_spec(global_space,0);ad
 	| WMMA_DIRECTIVE LAYOUT LAYOUT CONFIGURATION{add_wmma_option($1);add_wmma_option($2);add_wmma_option($3);add_wmma_option($4);}
 	;
 
+vp_spec: WMMA_DIRECTIVE LAYOUT CONFIGURATION{add_space_spec(global_space,0);add_ptr_spec(global_space);add_wmma_option($1);add_wmma_option($2);add_wmma_option($3);}
+	| WMMA_DIRECTIVE LAYOUT LAYOUT CONFIGURATION{add_wmma_option($1);add_wmma_option($2);add_wmma_option($3);add_wmma_option($4);}
+	;
+
+
+
 operand_list: operand
 	| operand COMMA operand_list;
 
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index d12c741..55b8e11 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -1089,9 +1089,13 @@ ptx_instruction::ptx_instruction( int opcode,
       		case SYNC_OPTION:
       		case LOAD_A:
       		case LOAD_B:
+      		case LOAD_B4:
+      		case LOAD_B8:
+      		case LOAD_B16:
       		case LOAD_C:
       		case STORE_D:
       		case MMA:
+      		case VP_MMA:
       		  m_wmma_type=last_ptx_inst_option;
       		  break;
       		case ROW:
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 6757091..eb81961 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -39,7 +39,7 @@ void set_ptx_warp_size(const struct core_config * warp_size)
    g_shader_core_config=warp_size;
 }
 
-static bool g_debug_ir_generation=false;
+static bool g_debug_ir_generation=true;
 const char *g_filename;
 unsigned g_max_regs_per_thread = 0;