From 09e6092ace5213a5d5a49bf80b052802c06a4268 Mon Sep 17 00:00:00 2001 From: aamir Date: Sat, 22 Sep 2018 15:28:40 -0700 Subject: debug prints --- cuda-kernels/Makefile | 1 + cuda-kernels/genericMatrixMultiply.cu | 6 +- cuda-kernels/gpgpusim.config | 2 +- cuda-kernels/scripts | 118 +++++++++++++++++++++++++++++----- cuda-kernels/v16p_kernel.cu | 4 +- 5 files changed, 108 insertions(+), 23 deletions(-) (limited to 'cuda-kernels') 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)); -- cgit v1.3