corrected the bug in vpcore functional model

8 files changed, 63 insertions, 45 deletions

diff --git a/cuda-kernels/v16p_genericMatrixMultiply.cu b/cuda-kernels/v16p_genericMatrixMultiply.cu
index c36d257..284bd7e 100644
--- a/cuda-kernels/v16p_genericMatrixMultiply.cu
+++ b/cuda-kernels/v16p_genericMatrixMultiply.cu
@@ -1,5 +1,6 @@
 #include <stdio.h>
 #include <curand.h>
+#include <stdlib.h>
 
 // Define some error checking macros.
 #define cudaErrCheck(stat) { cudaErrCheck_((stat), __FILE__, __LINE__); }
@@ -20,9 +21,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 (32)
-#define MATRIX_N (32)
-#define MATRIX_K (32)
+#define MATRIX_M (SIZE)
+#define MATRIX_N (SIZE)
+#define MATRIX_K (SIZE)
 
 
 // The only dimensions currently supported by WMMA
@@ -83,7 +84,7 @@ __global__ void vp_example(int *a, int *b, int *c, int M, int N, int K ) {
 	asm("CPTX_BEGIN");
 	asm("vp.load.b16.sync.row.m16n16k16.s32 {%0,%1,%2,%3},[%4],%5;" : 
 	"=r"(b_frag[0]),"=r"(b_frag[1]),"=r"(b_frag[2]),"=r"(b_frag[3]):
-	"l"(b+bRow*ldb/2+bCol),"r"(ldb/2)
+	"l"(b+bRow*ldb/2+bCol/2),"r"(ldb)//NOTE pointer arithematic is happenning over here so be extracareful
 	);
 	asm("CPTX_END");
 	asm("*/");
@@ -118,7 +119,7 @@ __global__ void vp_example(int *a, int *b, int *c, int M, int N, int K ) {
 	asm("vp.load.c.sync.row.m16n16k16.s32 {%0,%1,%2,%3,%4,%5,%6,%7},[%8],%9;" : 
 	"=r"(c_frag[0]), "=r"(c_frag[1]),"=r"(c_frag[2]),"=r"(c_frag[3]),
 	"=r"(c_frag[4]),"=r"(c_frag[5]),"=r"(c_frag[6]),"=r"(c_frag[7]):
-	"l"(c+cRow*ldc),"r"(ldc)
+	"l"(c+cRow*ldc+cCol),"r"(ldc)
 	);
 	asm("CPTX_END");
 	asm("*/");
@@ -223,7 +224,7 @@ int main(int argc, char* argv[]) {
    cudaEvent_t startWMMA;
    cudaEvent_t stopWMMA;
    
-   
+   srand(time(NULL)); 
    cudaErrCheck(cudaEventCreate(&startWMMA));
    cudaErrCheck(cudaEventCreate(&stopWMMA));
    
@@ -249,7 +250,8 @@ int main(int argc, char* argv[]) {
    printf("a_int32\n");
    for(int m=0;m<MATRIX_M;m++){
 	for(int n=0;n<MATRIX_K;n++){
-		a_host_wmma[m*MATRIX_K+n]=(m*MATRIX_K+n)%10;
+		a_host_wmma[m*MATRIX_K+n]= random()%2;
+		//a_host_wmma[m*MATRIX_K+n]= m*MATRIX_K+n; 
 		printf("%d ",a_host_wmma[m*MATRIX_K+n]);
 	}
 	printf(";\n");
@@ -258,7 +260,8 @@ int main(int argc, char* argv[]) {
    printf("b_int32\n");
    for(int m=0;m<MATRIX_K;m++){
 	for(int n=0;n<MATRIX_N;n++){
-		b_host_wmma[m*MATRIX_N+n]=(m*MATRIX_N+n)%4;
+		b_host_wmma[m*MATRIX_N+n]= random()%4;
+		//b_host_wmma[m*MATRIX_N+n]= m*MATRIX_N+n; 
 		printf("%d ",b_host_wmma[m*MATRIX_N+n]);
 	}
 		printf(";\n");
@@ -267,7 +270,8 @@ int main(int argc, char* argv[]) {
    printf("c_int32\n");
    for(int m=0;m<MATRIX_M;m++){
 	for(int n=0;n<MATRIX_N;n++){
-		c_host_wmma[m*MATRIX_N+n]=(m*MATRIX_N+n)%4;
+		c_host_wmma[m*MATRIX_N+n]= random()%32;
+		//c_host_wmma[m*MATRIX_N+n]= m*MATRIX_N+n; 
 		d_cal_host_wmma[m*MATRIX_N+n]=0;
 		printf("%d ",c_host_wmma[m*MATRIX_N+n]);
 	}
@@ -335,7 +339,7 @@ int main(int argc, char* argv[]) {
    cudaErrCheck(cudaEventDestroy(startWMMA));
    cudaErrCheck(cudaEventDestroy(stopWMMA));
 
-   int t=200000;
+   int t=1000000;
    while(t-->0);
    printf("D_CALCULATED\n");
 
diff --git a/cuda-kernels/v16p_kernel.cu b/cuda-kernels/v16p_kernel.cu
index 31a1460..864597d 100644
--- a/cuda-kernels/v16p_kernel.cu
+++ b/cuda-kernels/v16p_kernel.cu
@@ -1,5 +1,6 @@
 #include <stdio.h>
 #include <curand.h>
+#include <stdlib.h>
 
 // Define some error checking macros.
 #define cudaErrCheck(stat) { cudaErrCheck_((stat), __FILE__, __LINE__); }
@@ -205,7 +206,7 @@ int main(int argc, char* argv[]) {
    cudaEvent_t startWMMA;
    cudaEvent_t stopWMMA;
    
-   
+   srand(time(NULL)); 
    cudaErrCheck(cudaEventCreate(&startWMMA));
    cudaErrCheck(cudaEventCreate(&stopWMMA));
    
@@ -231,7 +232,7 @@ int main(int argc, char* argv[]) {
    printf("a_int32\n");
    for(int m=0;m<MATRIX_M;m++){
 	for(int n=0;n<MATRIX_K;n++){
-		a_host_wmma[m*MATRIX_K+n]=(m*MATRIX_K+n)%4;
+		a_host_wmma[m*MATRIX_K+n]= random()%5;
 		printf("%d ",a_host_wmma[m*MATRIX_K+n]);
 	}
 	printf(";\n");
@@ -240,7 +241,7 @@ int main(int argc, char* argv[]) {
    printf("b_int32\n");
    for(int m=0;m<MATRIX_K;m++){
 	for(int n=0;n<MATRIX_N;n++){
-		b_host_wmma[m*MATRIX_N+n]=(m*MATRIX_N+n)%256;
+		b_host_wmma[m*MATRIX_N+n]=random()%32;
 		printf("%d ",b_host_wmma[m*MATRIX_N+n]);
 	}
 		printf(";\n");
@@ -249,7 +250,7 @@ int main(int argc, char* argv[]) {
    printf("c_int32\n");
    for(int m=0;m<MATRIX_M;m++){
 	for(int n=0;n<MATRIX_N;n++){
-		c_host_wmma[m*MATRIX_N+n]=(m*MATRIX_N+n)%8;
+		c_host_wmma[m*MATRIX_N+n]=random()%32;
 		d_cal_host_wmma[m*MATRIX_N+n]=0;
 		printf("%d ",c_host_wmma[m*MATRIX_N+n]);
 	}
diff --git a/cuda-kernels/v4p_genericMatrixMultiply.cu b/cuda-kernels/v4p_genericMatrixMultiply.cu
index abcab8e..42283c4 100644
--- a/cuda-kernels/v4p_genericMatrixMultiply.cu
+++ b/cuda-kernels/v4p_genericMatrixMultiply.cu
@@ -1,5 +1,6 @@
 #include <stdio.h>
 #include <curand.h>
+#include<stdlib.h>
 
 // Define some error checking macros.
 #define cudaErrCheck(stat) { cudaErrCheck_((stat), __FILE__, __LINE__); }
@@ -20,9 +21,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 (32)
-#define MATRIX_N (32)
-#define MATRIX_K (32)
+#define MATRIX_M (SIZE)
+#define MATRIX_N (SIZE)
+#define MATRIX_K (SIZE)
 
 
 // The only dimensions currently supported by WMMA
@@ -83,7 +84,7 @@ __global__ void vp_example(int *a, int *b, int *c, int M, int N, int K ) {
 	asm("CPTX_BEGIN");
 	asm("vp.load.b4.sync.row.m16n16k16.s32 {%0},[%1],%2;" : 
 	"=r"(b_frag[0]):
-	"l"(b+bRow*ldb/8+bCol),"r"(ldb/8)
+	"l"(b+bRow*ldb/8+bCol/8),"r"(ldb)
 	);
 	asm("CPTX_END");
 	asm("*/");
@@ -118,7 +119,7 @@ __global__ void vp_example(int *a, int *b, int *c, int M, int N, int K ) {
 	asm("vp.load.c.sync.row.m16n16k16.s32 {%0,%1,%2,%3,%4,%5,%6,%7},[%8],%9;" : 
 	"=r"(c_frag[0]), "=r"(c_frag[1]),"=r"(c_frag[2]),"=r"(c_frag[3]),
 	"=r"(c_frag[4]),"=r"(c_frag[5]),"=r"(c_frag[6]),"=r"(c_frag[7]):
-	"l"(c+cRow*ldc),"r"(ldc)
+	"l"(c+cRow*ldc+cCol),"r"(ldc)
 	);
 	asm("CPTX_END");
 	asm("*/");
@@ -223,7 +224,8 @@ int main(int argc, char* argv[]) {
    cudaEvent_t startWMMA;
    cudaEvent_t stopWMMA;
    
-   
+  
+   srand (time(NULL));
    cudaErrCheck(cudaEventCreate(&startWMMA));
    cudaErrCheck(cudaEventCreate(&stopWMMA));
    
@@ -249,7 +251,8 @@ int main(int argc, char* argv[]) {
    printf("a_int32\n");
    for(int m=0;m<MATRIX_M;m++){
 	for(int n=0;n<MATRIX_K;n++){
-		a_host_wmma[m*MATRIX_K+n]=(m*MATRIX_K+n)%1024;
+		a_host_wmma[m*MATRIX_K+n]= random()%4;
+		//a_host_wmma[m*MATRIX_K+n]=m*MATRIX_K+n; 
 		printf("%d ",a_host_wmma[m*MATRIX_K+n]);
 	}
 	printf(";\n");
@@ -258,7 +261,7 @@ int main(int argc, char* argv[]) {
    printf("b_int32\n");
    for(int m=0;m<MATRIX_K;m++){
 	for(int n=0;n<MATRIX_N;n++){
-		b_host_wmma[m*MATRIX_N+n]=(m*MATRIX_N+n)%4;
+		b_host_wmma[m*MATRIX_N+n]=random()%2;
 		printf("%d ",b_host_wmma[m*MATRIX_N+n]);
 	}
 		printf(";\n");
@@ -267,7 +270,7 @@ int main(int argc, char* argv[]) {
    printf("c_int32\n");
    for(int m=0;m<MATRIX_M;m++){
 	for(int n=0;n<MATRIX_N;n++){
-		c_host_wmma[m*MATRIX_N+n]=(m*MATRIX_N+n)%4;
+		c_host_wmma[m*MATRIX_N+n]= random()%64;
 		d_cal_host_wmma[m*MATRIX_N+n]=0;
 		printf("%d ",c_host_wmma[m*MATRIX_N+n]);
 	}
@@ -335,7 +338,7 @@ int main(int argc, char* argv[]) {
    cudaErrCheck(cudaEventDestroy(startWMMA));
    cudaErrCheck(cudaEventDestroy(stopWMMA));
 
-   int t=200000;
+   int t=1000000;
    while(t-->0);
    printf("D_CALCULATED\n");
 
diff --git a/cuda-kernels/v4p_kernel.cu b/cuda-kernels/v4p_kernel.cu
index f1a9dbe..053c07b 100644
--- a/cuda-kernels/v4p_kernel.cu
+++ b/cuda-kernels/v4p_kernel.cu
@@ -1,4 +1,5 @@
 #include <stdio.h>
+#include <stdlib.h>
 #include <curand.h>
 
 // Define some error checking macros.
@@ -205,7 +206,7 @@ int main(int argc, char* argv[]) {
    cudaEvent_t startWMMA;
    cudaEvent_t stopWMMA;
    
-   
+   srand (time(NULL));  
    cudaErrCheck(cudaEventCreate(&startWMMA));
    cudaErrCheck(cudaEventCreate(&stopWMMA));
    
@@ -231,7 +232,7 @@ int main(int argc, char* argv[]) {
    printf("a_int32\n");
    for(int m=0;m<MATRIX_M;m++){
 	for(int n=0;n<MATRIX_K;n++){
-		a_host_wmma[m*MATRIX_K+n]=(m*MATRIX_K+n)%4;
+		a_host_wmma[m*MATRIX_K+n]=random()%64;
 		printf("%d ",a_host_wmma[m*MATRIX_K+n]);
 	}
 	printf(";\n");
@@ -240,7 +241,7 @@ int main(int argc, char* argv[]) {
    printf("b_int32\n");
    for(int m=0;m<MATRIX_K;m++){
 	for(int n=0;n<MATRIX_N;n++){
-		b_host_wmma[m*MATRIX_N+n]=(m*MATRIX_N+n)%4;
+		b_host_wmma[m*MATRIX_N+n]=random()%16;
 		printf("%d ",b_host_wmma[m*MATRIX_N+n]);
 	}
 		printf(";\n");
@@ -249,7 +250,7 @@ int main(int argc, char* argv[]) {
    printf("c_int32\n");
    for(int m=0;m<MATRIX_M;m++){
 	for(int n=0;n<MATRIX_N;n++){
-		c_host_wmma[m*MATRIX_N+n]=(m*MATRIX_N+n)%4;
+		c_host_wmma[m*MATRIX_N+n]= random()%64;
 		d_cal_host_wmma[m*MATRIX_N+n]=0;
 		printf("%d ",c_host_wmma[m*MATRIX_N+n]);
 	}
diff --git a/cuda-kernels/v8p_genericMatrixMultiply.cu b/cuda-kernels/v8p_genericMatrixMultiply.cu
index 208c369..d96b07c 100644
--- a/cuda-kernels/v8p_genericMatrixMultiply.cu
+++ b/cuda-kernels/v8p_genericMatrixMultiply.cu
@@ -1,5 +1,6 @@
 #include <stdio.h>
 #include <curand.h>
+#include <stdlib.h>
 
 // Define some error checking macros.
 #define cudaErrCheck(stat) { cudaErrCheck_((stat), __FILE__, __LINE__); }
@@ -20,9 +21,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 (32)
-#define MATRIX_N (32)
-#define MATRIX_K (32)
+#define MATRIX_M (SIZE)
+#define MATRIX_N (SIZE)
+#define MATRIX_K (SIZE)
 
 
 // The only dimensions currently supported by WMMA
@@ -83,7 +84,7 @@ __global__ void vp_example(int *a, int *b, int *c, int M, int N, int K ) {
 	asm("CPTX_BEGIN");
 	asm("vp.load.b8.sync.row.m16n16k16.s32 {%0,%1},[%2],%3;" : 
 	"=r"(b_frag[0]),"=r"(b_frag[1]):
-	"l"(b+bRow*ldb/4+bCol),"r"(ldb/4)
+	"l"(b+bRow*ldb/4+bCol/4),"r"(ldb)
 	);
 	asm("CPTX_END");
 	asm("*/");
@@ -118,7 +119,7 @@ __global__ void vp_example(int *a, int *b, int *c, int M, int N, int K ) {
 	asm("vp.load.c.sync.row.m16n16k16.s32 {%0,%1,%2,%3,%4,%5,%6,%7},[%8],%9;" : 
 	"=r"(c_frag[0]), "=r"(c_frag[1]),"=r"(c_frag[2]),"=r"(c_frag[3]),
 	"=r"(c_frag[4]),"=r"(c_frag[5]),"=r"(c_frag[6]),"=r"(c_frag[7]):
-	"l"(c+cRow*ldc),"r"(ldc)
+	"l"(c+cRow*ldc+cCol),"r"(ldc)
 	);
 	asm("CPTX_END");
 	asm("*/");
@@ -223,7 +224,7 @@ int main(int argc, char* argv[]) {
    cudaEvent_t startWMMA;
    cudaEvent_t stopWMMA;
    
-   
+   srand(time(NULL)); 
    cudaErrCheck(cudaEventCreate(&startWMMA));
    cudaErrCheck(cudaEventCreate(&stopWMMA));
    
@@ -249,7 +250,7 @@ int main(int argc, char* argv[]) {
    printf("a_int32\n");
    for(int m=0;m<MATRIX_M;m++){
 	for(int n=0;n<MATRIX_K;n++){
-		a_host_wmma[m*MATRIX_K+n]=(m*MATRIX_K+n)%1024;
+		a_host_wmma[m*MATRIX_K+n]=random()%5;
 		printf("%d ",a_host_wmma[m*MATRIX_K+n]);
 	}
 	printf(";\n");
@@ -258,7 +259,7 @@ int main(int argc, char* argv[]) {
    printf("b_int32\n");
    for(int m=0;m<MATRIX_K;m++){
 	for(int n=0;n<MATRIX_N;n++){
-		b_host_wmma[m*MATRIX_N+n]=(m*MATRIX_N+n)%4;
+		b_host_wmma[m*MATRIX_N+n]=random()%32;
 		printf("%d ",b_host_wmma[m*MATRIX_N+n]);
 	}
 		printf(";\n");
@@ -267,7 +268,7 @@ int main(int argc, char* argv[]) {
    printf("c_int32\n");
    for(int m=0;m<MATRIX_M;m++){
 	for(int n=0;n<MATRIX_N;n++){
-		c_host_wmma[m*MATRIX_N+n]=(m*MATRIX_N+n)%4;
+		c_host_wmma[m*MATRIX_N+n]=random()%32;
 		d_cal_host_wmma[m*MATRIX_N+n]=0;
 		printf("%d ",c_host_wmma[m*MATRIX_N+n]);
 	}
@@ -335,7 +336,7 @@ int main(int argc, char* argv[]) {
    cudaErrCheck(cudaEventDestroy(startWMMA));
    cudaErrCheck(cudaEventDestroy(stopWMMA));
 
-   int t=200000;
+   int t=1000000;
    while(t-->0);
    printf("D_CALCULATED\n");
 
diff --git a/cuda-kernels/v8p_kernel.cu b/cuda-kernels/v8p_kernel.cu
index f824eac..9bfcea7 100644
--- a/cuda-kernels/v8p_kernel.cu
+++ b/cuda-kernels/v8p_kernel.cu
@@ -1,5 +1,6 @@
 #include <stdio.h>
 #include <curand.h>
+#include<stdlib.h>
 
 // Define some error checking macros.
 #define cudaErrCheck(stat) { cudaErrCheck_((stat), __FILE__, __LINE__); }
@@ -204,7 +205,7 @@ int main(int argc, char* argv[]) {
 
    cudaEvent_t startWMMA;
    cudaEvent_t stopWMMA;
-   
+   srand(time(NULL)); 
    
    cudaErrCheck(cudaEventCreate(&startWMMA));
    cudaErrCheck(cudaEventCreate(&stopWMMA));
@@ -231,7 +232,7 @@ int main(int argc, char* argv[]) {
    printf("a_int32\n");
    for(int m=0;m<MATRIX_M;m++){
 	for(int n=0;n<MATRIX_K;n++){
-		a_host_wmma[m*MATRIX_K+n]=(m*MATRIX_K+n)%8;
+		a_host_wmma[m*MATRIX_K+n]=random()%10;
 		printf("%d ",a_host_wmma[m*MATRIX_K+n]);
 	}
 	printf(";\n");
@@ -240,7 +241,7 @@ int main(int argc, char* argv[]) {
    printf("b_int32\n");
    for(int m=0;m<MATRIX_K;m++){
 	for(int n=0;n<MATRIX_N;n++){
-		b_host_wmma[m*MATRIX_N+n]=(m*MATRIX_N+n)%4;
+		b_host_wmma[m*MATRIX_N+n]=random()%5;
 		printf("%d ",b_host_wmma[m*MATRIX_N+n]);
 	}
 		printf(";\n");
@@ -249,7 +250,7 @@ int main(int argc, char* argv[]) {
    printf("c_int32\n");
    for(int m=0;m<MATRIX_M;m++){
 	for(int n=0;n<MATRIX_N;n++){
-		c_host_wmma[m*MATRIX_N+n]=(m*MATRIX_N+n)%8;
+		c_host_wmma[m*MATRIX_N+n]=random()%64;
 		d_cal_host_wmma[m*MATRIX_N+n]=0;
 		printf("%d ",c_host_wmma[m*MATRIX_N+n]);
 	}
diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc
index b24a77e..83e76fe 100644
--- a/src/abstract_hardware_model.cc
+++ b/src/abstract_hardware_model.cc
@@ -389,7 +389,8 @@ void warp_inst_t::memory_coalescing_arch_13( bool is_write, mem_access_type acce
 
             assert(num_accesses <= MAX_ACCESSES_PER_INSN_PER_THREAD);
 
-            for(unsigned access=0; access<num_accesses; access++) {
+//            for(unsigned access=0; access<num_accesses; access++) {
+            for(unsigned access=0; (access<MAX_ACCESSES_PER_INSN_PER_THREAD)&&(m_per_scalar_thread[thread].memreqaddr[access]!=0); access++) {
                 new_addr_type addr = m_per_scalar_thread[thread].memreqaddr[access];
                 unsigned block_address = line_size_based_tag_func(addr,segment_size);
                 unsigned chunk = (addr&127)/32; // which 32-byte chunk within in a 128-byte chunk does this thread access?
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index 6f1b5d3..42c63ca 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -3363,7 +3363,7 @@ void vp_ld_impl(const ptx_instruction *pI, core_t *core, warp_inst_t &inst)
 			else if(wmma_layout==COL){
 				//mem->read(new_addr+4*stride*i,size/8,&data[i].s64);
 				mem->read(new_addr+4*stride*i,size/8,&data[i].s64);
-				mem_txn_addr[num_mem_txn++]=new_addr+4*i;			
+				mem_txn_addr[num_mem_txn++]=new_addr+4*stride*i;			
 			}
 		}
 	}
@@ -3410,7 +3410,13 @@ void vp_ld_impl(const ptx_instruction *pI, core_t *core, warp_inst_t &inst)
    	inst.set_addr(thrd, (new_addr_type *)mem_txn_addr , num_mem_txn);
    	inst.data_size = 4; // 4 byte transaction 
    	assert( inst.memory_op == insn_memory_op );
-
+	printf("\nVPLOAD_THREAD%d:",thrd);
+	for(int lll=0;lll<num_mem_txn;lll++)
+	{
+		printf("%x ",mem_txn_addr[lll]);
+	}
+	printf("\n");
+	
 	int num_reg;
 	if(g_debug_instruction){
 		printf("\nvp_ld:thread%d= ",thrd);