From cc14e087c164a25cd9e2374777fdc4e57d9cec39 Mon Sep 17 00:00:00 2001
From: aamir <araihan@ece.ubc.ca>
Date: Fri, 14 Sep 2018 01:42:35 -0700
Subject: corrected the bug in vpcore functional model

---
 cuda-kernels/v16p_genericMatrixMultiply.cu | 24 ++++++++++++++----------
 cuda-kernels/v16p_kernel.cu                |  9 +++++----
 cuda-kernels/v4p_genericMatrixMultiply.cu  | 23 +++++++++++++----------
 cuda-kernels/v4p_kernel.cu                 |  9 +++++----
 cuda-kernels/v8p_genericMatrixMultiply.cu  | 21 +++++++++++----------
 cuda-kernels/v8p_kernel.cu                 |  9 +++++----
 6 files changed, 53 insertions(+), 42 deletions(-)

(limited to 'cuda-kernels')

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]);
 	}
-- 
cgit v1.3