debug prints

5 files changed, 108 insertions, 23 deletions

diff --git a/cuda-kernels/Makefile b/cuda-kernels/Makefile
index 8effd11..b33519f 100755
--- a/cuda-kernels/Makefile
+++ b/cuda-kernels/Makefile
@@ -7,4 +7,5 @@ clean:
 	rm _cuob*
 	rm gpgpusim_power* 
 	rm gpgpu_inst_stats.txt
+	rm gpgpusim_visualizer*
 #	nvcc -arch=sm_70 --gpu-architecture=compute_50 --gpu-code=compute_50 -lcudart -g -o tensor_core tensor_core.cu
diff --git a/cuda-kernels/genericMatrixMultiply.cu b/cuda-kernels/genericMatrixMultiply.cu
index 6a5d33f..8b96483 100644
--- a/cuda-kernels/genericMatrixMultiply.cu
+++ b/cuda-kernels/genericMatrixMultiply.cu
@@ -40,9 +40,9 @@ void cudaErrCheck_(cudaError_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 (64)
+#define MATRIX_N (64)
+#define MATRIX_K (64)
 
 
 
diff --git a/cuda-kernels/gpgpusim.config b/cuda-kernels/gpgpusim.config
index 3daa539..5f1be25 100755
--- a/cuda-kernels/gpgpusim.config
+++ b/cuda-kernels/gpgpusim.config
@@ -138,7 +138,7 @@
 -gpgpu_memlatency_stat 14 
 -gpgpu_runtime_stat 500
 -enable_ptx_file_line_stats 1
--visualizer_enabled 0
+-visualizer_enabled 1
 
 # power model configs
 -power_simulation_enabled 1
diff --git a/cuda-kernels/scripts b/cuda-kernels/scripts
index b9a3b50..a24fb72 100755
--- a/cuda-kernels/scripts
+++ b/cuda-kernels/scripts
@@ -10,11 +10,16 @@ nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o vp4_256
 ./vp4_128>vp4_log128
 ./vp4_256>vp4_log256
 
-grep "kernel_name.*vp_example" -A 589 vp4_log16   >vp4_16_summary
-grep "kernel_name.*vp_example" -A 589 vp4_log32   >vp4_32_summary
-grep "kernel_name.*vp_example" -A 589 vp4_log64   >vp4_64_summary
-grep "kernel_name.*vp_example" -A 589 vp4_log128  >vp4_128_summary
-grep "kernel_name.*vp_example" -A 589 vp4_log256  >vp4_256_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp4_log16  	>vp4_16_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp4_log32  	>vp4_32_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp4_log64  	>vp4_64_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp4_log128 	>vp4_128_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp4_log256 	>vp4_256_summary
+###grep "kernel_name.*vp_example" -A 589 vp4_log16      	>>vp4_16_summary
+###grep "kernel_name.*vp_example" -A 589 vp4_log32      	>>vp4_32_summary
+###grep "kernel_name.*vp_example" -A 589 vp4_log64      	>>vp4_64_summary
+#grep "kernel_name.*vp_example" -A 589 vp4_log128     	>>vp4_128_summary
+#grep "kernel_name.*vp_example" -A 589 vp4_log256     	>>vp4_256_summary
 
 
 nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o vp8_16   v8p_genericMatrixMultiply.cu   -DSIZE=16
@@ -28,11 +33,16 @@ nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o vp8_256
 ./vp8_128>vp8_log128
 ./vp8_256>vp8_log256
 
-grep "kernel_name.*vp_example" -A 589 vp8_log16   >vp8_16_summary
-grep "kernel_name.*vp_example" -A 589 vp8_log32   >vp8_32_summary
-grep "kernel_name.*vp_example" -A 589 vp8_log64   >vp8_64_summary
-grep "kernel_name.*vp_example" -A 589 vp8_log128  >vp8_128_summary
-grep "kernel_name.*vp_example" -A 589 vp8_log256  >vp8_256_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp8_log16   >vp8_16_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp8_log32   >vp8_32_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp8_log64   >vp8_64_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp8_log128  >vp8_128_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp8_log256  >vp8_256_summary
+##grep "kernel_name.*vp_example" -A 589 vp8_log16       >>vp8_16_summary
+##grep "kernel_name.*vp_example" -A 589 vp8_log32       >>vp8_32_summary
+##grep "kernel_name.*vp_example" -A 589 vp8_log64       >>vp8_64_summary
+#grep "kernel_name.*vp_example" -A 589 vp8_log128      >>vp8_128_summary
+#grep "kernel_name.*vp_example" -A 589 vp8_log256      >>vp8_256_summary
 
 nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o vp16_16  v16p_genericMatrixMultiply.cu   -DSIZE=16
 nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o vp16_32  v16p_genericMatrixMultiply.cu   -DSIZE=32
@@ -45,11 +55,85 @@ nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o vp16_256
 ./vp16_128>vp16_log128
 ./vp16_256>vp16_log256
 
-grep "kernel_name.*vp_example" -A 589 vp16_log16   >vp16_16_summary
-grep "kernel_name.*vp_example" -A 589 vp16_log32   >vp16_32_summary
-grep "kernel_name.*vp_example" -A 589 vp16_log64   >vp16_64_summary
-grep "kernel_name.*vp_example" -A 589 vp16_log128  >vp16_128_summary
-grep "kernel_name.*vp_example" -A 589 vp16_log256  >vp16_256_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp16_log16   >vp16_16_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp16_log32   >vp16_32_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp16_log64   >vp16_64_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp16_log128  >vp16_128_summary
+#grep "kernel_name.*convertInt32To" -A 589 vp16_log256  >vp16_256_summary
+##grep "kernel_name.*vp_example" -A 589 vp16_log16       >>vp16_16_summary
+##grep "kernel_name.*vp_example" -A 589 vp16_log32       >>vp16_32_summary
+##grep "kernel_name.*vp_example" -A 589 vp16_log64       >>vp16_64_summary
+#grep "kernel_name.*vp_example" -A 589 vp16_log128      >>vp16_128_summary
+#grep "kernel_name.*vp_example" -A 589 vp16_log256      >>vp16_256_summary
 
-mv vp* VPlog/
-make clean
+###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp4_log16  > kernel_log_vp4_log16
+###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp4_log32  > kernel_log_vp4_log32
+###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp4_log64  > kernel_log_vp4_log64
+####sed -n '/bind.*to.*kernel.*vp_example/,$p' vp4_log128 > kernel_log_vp4_log128
+####sed -n '/bind.*to.*kernel.*vp_example/,$p' vp4_log256 > kernel_log_vp4_log256
+###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp8_log16  > kernel_log_vp8_log16
+###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp8_log32  > kernel_log_vp8_log32
+###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp8_log64  > kernel_log_vp8_log64
+####sed -n '/bind.*to.*kernel.*vp_example/,$p' vp8_log128 > kernel_log_vp8_log128
+####sed -n '/bind.*to.*kernel.*vp_example/,$p' vp8_log256 > kernel_log_vp8_log256
+###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp16_log16  > kernel_log_vp16_log16
+###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp16_log32  > kernel_log_vp16_log32
+###sed -n '/bind.*to.*kernel.*vp_example/,$p' vp16_log64  > kernel_log_vp16_log64
+####sed -n '/bind.*to.*kernel.*vp_example/,$p' vp16_log128 > kernel_log_vp16_log128
+####sed -n '/bind.*to.*kernel.*vp_example/,$p' vp16_log256 > kernel_log_vp16_log256
+###
+###grep memtocore:Icnt kernel_log_vp4_log16>memtocore_vp4_log16
+###grep coretomem:Icnt kernel_log_vp4_log16>coretomem_vp4_log16
+###grep MEM_TXN        kernel_log_vp4_log16>  gen_mem_vp4_log16
+###grep memtocore:Icnt kernel_log_vp4_log32>memtocore_vp4_log32
+###grep coretomem:Icnt kernel_log_vp4_log32>coretomem_vp4_log32
+###grep MEM_TXN        kernel_log_vp4_log32>  gen_mem_vp4_log32
+###grep memtocore:Icnt kernel_log_vp4_log64>memtocore_vp4_log64
+###grep coretomem:Icnt kernel_log_vp4_log64>coretomem_vp4_log64
+###grep MEM_TXN        kernel_log_vp4_log64>  gen_mem_vp4_log64
+####grep memtocore:Icnt kernel_log_vp4_log128>memtocore_vp4_log128
+####grep coretomem:Icnt kernel_log_vp4_log128>coretomem_vp4_log128
+####grep MEM_TXN        kernel_log_vp4_log128>  gen_mem_vp4_log128
+####grep memtocore:Icnt kernel_log_vp4_log256>memtocore_vp4_log256
+####grep coretomem:Icnt kernel_log_vp4_log256>coretomem_vp4_log256
+####grep MEM_TXN        kernel_log_vp4_log256>  gen_mem_vp4_log256
+###
+###grep memtocore:Icnt kernel_log_vp8_log16>memtocore_vp8_log16
+###grep coretomem:Icnt kernel_log_vp8_log16>coretomem_vp8_log16
+###grep MEM_TXN        kernel_log_vp8_log16>  gen_mem_vp8_log16
+###grep memtocore:Icnt kernel_log_vp8_log32>memtocore_vp8_log32
+###grep coretomem:Icnt kernel_log_vp8_log32>coretomem_vp8_log32
+###grep MEM_TXN        kernel_log_vp8_log32>  gen_mem_vp8_log32
+###grep memtocore:Icnt kernel_log_vp8_log64>memtocore_vp8_log64
+###grep coretomem:Icnt kernel_log_vp8_log64>coretomem_vp8_log64
+###grep MEM_TXN        kernel_log_vp8_log64>  gen_mem_vp8_log64
+####grep memtocore:Icnt kernel_log_vp8_log128>memtocore_vp8_log128
+####grep coretomem:Icnt kernel_log_vp8_log128>coretomem_vp8_log128
+####grep MEM_TXN        kernel_log_vp8_log128>  gen_mem_vp8_log128
+####grep memtocore:Icnt kernel_log_vp8_log256>memtocore_vp8_log256
+####grep coretomem:Icnt kernel_log_vp8_log256>coretomem_vp8_log256
+####grep MEM_TXN        kernel_log_vp8_log256>  gen_mem_vp8_log256
+###
+###grep memtocore:Icnt kernel_log_vp16_log16>memtocore_vp16_log16
+###grep coretomem:Icnt kernel_log_vp16_log16>coretomem_vp16_log16
+###grep MEM_TXN        kernel_log_vp16_log16>  gen_mem_vp16_log16
+###grep memtocore:Icnt kernel_log_vp16_log32>memtocore_vp16_log32
+###grep coretomem:Icnt kernel_log_vp16_log32>coretomem_vp16_log32
+###grep MEM_TXN        kernel_log_vp16_log32>  gen_mem_vp16_log32
+###grep memtocore:Icnt kernel_log_vp16_log64>memtocore_vp16_log64
+###grep coretomem:Icnt kernel_log_vp16_log64>coretomem_vp16_log64
+###grep MEM_TXN        kernel_log_vp16_log64>  gen_mem_vp16_log64
+####grep memtocore:Icnt kernel_log_vp16_log128>memtocore_vp16_log128
+####grep coretomem:Icnt kernel_log_vp16_log128>coretomem_vp16_log128
+####grep MEM_TXN        kernel_log_vp16_log128>  gen_mem_vp16_log128
+####grep memtocore:Icnt kernel_log_vp16_log256>memtocore_vp16_log256
+####grep coretomem:Icnt kernel_log_vp16_log256>coretomem_vp16_log256
+####grep MEM_TXN        kernel_log_vp16_log256>  gen_mem_vp16_log256
+###
+###mv kernel_* MemTraffic/
+###mv memtocore* MemTraffic/
+###mv coretomem* MemTraffic/
+###mv gen_mem* MemTraffic/
+###
+###mv vp* VPlog/
+###make clean
diff --git a/cuda-kernels/v16p_kernel.cu b/cuda-kernels/v16p_kernel.cu
index 864597d..d84328f 100644
--- a/cuda-kernels/v16p_kernel.cu
+++ b/cuda-kernels/v16p_kernel.cu
@@ -31,7 +31,7 @@ const int WMMA_M = 16;
 const int WMMA_N = 16;
 const int WMMA_K = 16;
 
-__global__ void v4p_example(int *a_int32, int *b_int4, int *c,int *d_int32, int M, int N, int K) {
+__global__ void v16p_example(int *a_int32, int *b_int4, int *c,int *d_int32, int M, int N, int K) {
 	
 	int registers_a[8];
 	int registers_b[8];
@@ -294,7 +294,7 @@ int main(int argc, char* argv[]) {
 //AAMIR   
 //AAMIR   printf("Running with wmma...\n");
    cudaErrCheck(cudaEventRecord(startWMMA));
-   v4p_example <<< 1, 32>>> (a_int32, b_int16, c_int32, d_int32, MATRIX_M, MATRIX_N, MATRIX_K);
+   v16p_example <<< 1, 32>>> (a_int32, b_int16, c_int32, d_int32, MATRIX_M, MATRIX_N, MATRIX_K);
    cudaErrCheck(cudaEventRecord(stopWMMA));
    cudaErrCheck(cudaEventSynchronize(stopWMMA));