From e5f532a3b65e17f49991ed08a275f87ac2d68d0a Mon Sep 17 00:00:00 2001 From: aamir Date: Fri, 1 Jun 2018 09:52:12 -0700 Subject: wmma load working --- src/cuda-sim/cuda-math.h | 1 + 1 file changed, 1 insertion(+) (limited to 'src/cuda-sim/cuda-math.h') diff --git a/src/cuda-sim/cuda-math.h b/src/cuda-sim/cuda-math.h index 4721e8a..e38e499 100644 --- a/src/cuda-sim/cuda-math.h +++ b/src/cuda-sim/cuda-math.h @@ -64,6 +64,7 @@ * the above Disclaimer and U.S. Government End Users Notice. */ + #ifndef CUDA_MATH #define CUDA_MATH -- cgit v1.3 From a089ac9ed1c798108bd0866e422392d6873f6681 Mon Sep 17 00:00:00 2001 From: aamir Date: Wed, 6 Jun 2018 16:31:03 -0700 Subject: added profilling xls and changes for regression --- cuda-kernels/TensorCoreMatrixCArrangement.xlsx | Bin 0 -> 76737 bytes cuda-kernels/tensorcore_type16_32.cu | 226 +++++++++++++++++++++++++ src/cuda-sim/cuda-math.h | 4 +- src/cuda-sim/instructions.cc | 32 ++-- src/cuda-sim/ptx_sim.h | 11 ++ 5 files changed, 255 insertions(+), 18 deletions(-) create mode 100644 cuda-kernels/TensorCoreMatrixCArrangement.xlsx create mode 100644 cuda-kernels/tensorcore_type16_32.cu (limited to 'src/cuda-sim/cuda-math.h') diff --git a/cuda-kernels/TensorCoreMatrixCArrangement.xlsx b/cuda-kernels/TensorCoreMatrixCArrangement.xlsx new file mode 100644 index 0000000..785ecc3 Binary files /dev/null and b/cuda-kernels/TensorCoreMatrixCArrangement.xlsx differ diff --git a/cuda-kernels/tensorcore_type16_32.cu b/cuda-kernels/tensorcore_type16_32.cu new file mode 100644 index 0000000..7084bf9 --- /dev/null +++ b/cuda-kernels/tensorcore_type16_32.cu @@ -0,0 +1,226 @@ +#include +#include + +// Define some error checking macros. +#define cudaErrCheck(stat) { cudaErrCheck_((stat), __FILE__, __LINE__); } +void cudaErrCheck_(cudaError_t stat, const char *file, int line) { + if (stat != cudaSuccess) { + fprintf(stderr, "CUDA Error: %s %s %d\n", cudaGetErrorString(stat), file, line); + } +} + +#define curandErrCheck(stat) { curandErrCheck_((stat), __FILE__, __LINE__); } +void curandErrCheck_(curandStatus_t stat, const char *file, int line) { + if (stat != CURAND_STATUS_SUCCESS) { + fprintf(stderr, "cuRand Error: %d %s %d\n", stat, file, line); + } +} + +#include +using namespace nvcuda; + +// Must be multiples of 16 for wmma code to work +#define MATRIX_M (16) +#define MATRIX_N (16) +#define MATRIX_K (16) + + +// The only dimensions currently supported by WMMA +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,half *d_fp16, int M, int N, int K) { + //unsigned int start_time=0,end_time=0; + //start_time=clock(); + + // Declare the fragments + wmma::fragment a_frag; + wmma::fragment b_frag; + wmma::fragment c_frag; + wmma::fragment d_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); + +// for(int i=0; i < c_frag.num_elements; i++) { +//// c_frag.x[i]=c_frag.x[i]+c_frag.x[i]; +// float temp=c_frag.x[i]; +// printf("THREAD%d:%d: %f \n",threadIdx.x,i,temp ); +// } + wmma::mma_sync(c_frag, a_frag, b_frag, c_frag); + + for(int i=0; i < c_frag.num_elements; i++) { + d_frag.x[i]=c_frag.x[i]; + } + wmma::store_matrix_sync(d_fp16, d_frag, N, wmma::mem_col_major); + //printf("clock=%d",end_time-start_time); +} + +__global__ void convertFp32ToFp16 (half *out, float *in, int n) { + int idx = blockDim.x * blockIdx.x + threadIdx.x; + if (idx < n) { + out[idx] = in[idx]; + } +} +__global__ void convertFp16ToFp32 (float *out, half *in, int n) { + int idx = blockDim.x * blockIdx.x + threadIdx.x; + if (idx < n) { + out[idx] = in[idx]; + } +} + +int main(int argc, char* argv[]) { + float *a_fp32; + float *b_fp32; + float *c_fp32; + float *d_fp32; + + half *a_fp16; + half *b_fp16; + half *c_fp16; + half *d_fp16; + + float *a_host_wmma; + float *b_host_wmma; + float *c_host_wmma; + float *d_host_wmma; + float *d_cal_host_wmma; + + cudaEvent_t startWMMA; + cudaEvent_t stopWMMA; + + + cudaErrCheck(cudaEventCreate(&startWMMA)); + cudaErrCheck(cudaEventCreate(&stopWMMA)); + + // Use tensor cores + cudaErrCheck(cudaMalloc((void**)&a_fp32, MATRIX_M * MATRIX_K * sizeof(float))); + cudaErrCheck(cudaMalloc((void**)&b_fp32, MATRIX_K * MATRIX_N * sizeof(float))); + cudaErrCheck(cudaMalloc((void**)&c_fp32, MATRIX_K * MATRIX_N * sizeof(float))); + cudaErrCheck(cudaMalloc((void**)&d_fp32, MATRIX_K * MATRIX_N * sizeof(float))); + cudaErrCheck(cudaMalloc((void**)&a_fp16, MATRIX_M * MATRIX_K * sizeof(half))); + cudaErrCheck(cudaMalloc((void**)&b_fp16, MATRIX_K * MATRIX_N * sizeof(half))); + cudaErrCheck(cudaMalloc((void**)&c_fp16, MATRIX_K * MATRIX_N * sizeof(half))); + cudaErrCheck(cudaMalloc((void**)&d_fp16, MATRIX_K * MATRIX_N * sizeof(half))); + + + a_host_wmma = (float*)malloc(MATRIX_M * MATRIX_K * sizeof(float)); + b_host_wmma = (float*)malloc(MATRIX_K * MATRIX_N * sizeof(float)); + c_host_wmma = (float*)malloc(MATRIX_M * MATRIX_N * sizeof(float)); + d_host_wmma = (float*)malloc(MATRIX_M * MATRIX_N * sizeof(float)); + d_cal_host_wmma = (float*)malloc(MATRIX_M * MATRIX_N * sizeof(float)); + + //printf("a_fp32\n"); + for(int m=0;m>> (a_fp16, a_fp32, MATRIX_M * MATRIX_K); + convertFp32ToFp16 <<< (MATRIX_K * MATRIX_N + 255) / 256, 256 >>> (b_fp16, b_fp32, MATRIX_K * MATRIX_N); + + printf("\nM = %d, N = %d, K = %d. \n", MATRIX_M, MATRIX_N, MATRIX_K); + + printf("Running with wmma...\n"); + cudaErrCheck(cudaEventRecord(startWMMA)); + wmma_example <<< 1, 32>>> (a_fp16, b_fp16, c_fp32, d_fp16 , MATRIX_M, MATRIX_N, MATRIX_K); + cudaErrCheck(cudaEventRecord(stopWMMA)); + cudaErrCheck(cudaEventSynchronize(stopWMMA)); + + convertFp16ToFp32 <<< (MATRIX_M * MATRIX_N + 255) / 256, 256 >>> (d_fp32, d_fp16, MATRIX_K * MATRIX_N); + // Error checking + printf("\nChecking results...\n"); + cudaErrCheck(cudaMemcpy(d_host_wmma, d_fp32, MATRIX_M * MATRIX_N * sizeof(float), cudaMemcpyDeviceToHost)); + + printf("Results verified: cublas and WMMA agree.\n\n"); + float wmmaTime; + cudaErrCheck(cudaEventElapsedTime(&wmmaTime, startWMMA, stopWMMA)); + printf("wmma took %fms\n", wmmaTime); + + cudaErrCheck(cudaEventDestroy(startWMMA)); + cudaErrCheck(cudaEventDestroy(stopWMMA)); + int t=600000; + while(t-->0); + + printf("D_CALCULATED\n"); + for(int m=0;m1) + { + printf("ERROR:\n"); + suc=0; + } + } + } + if(suc==1) + printf("COMPLETED_SUCCESSFULLY\n"); + + cudaErrCheck(cudaFree(a_fp32)); + cudaErrCheck(cudaFree(b_fp32)); + cudaErrCheck(cudaFree(c_fp32)); + cudaErrCheck(cudaFree(d_fp32)); + cudaErrCheck(cudaFree(a_fp16)); + cudaErrCheck(cudaFree(b_fp16)); + cudaErrCheck(cudaFree(c_fp16)); + cudaErrCheck(cudaFree(d_fp16)); + + free(a_host_wmma); + free(b_host_wmma); + free(c_host_wmma); + free(d_host_wmma); + cudaErrCheck(cudaDeviceReset()); + return 0; +} + + diff --git a/src/cuda-sim/cuda-math.h b/src/cuda-sim/cuda-math.h index e38e499..314f813 100644 --- a/src/cuda-sim/cuda-math.h +++ b/src/cuda-sim/cuda-math.h @@ -322,7 +322,7 @@ float __internal_accurate_fdividef(float a, float b) float __saturatef(float a) { float b; - if (isnan(a)) b = 0.0f; + if (std::isnan(a)) b = 0.0f; else if (a >= 1.0f) b = 1.0f; else if (a <= 0.0f) b = 0.0f; else b = a; @@ -358,7 +358,7 @@ int __signbitd(double d) #ifdef __APPLE__ int isnanf(float a) { - return (isnan(a)); + return (std::isnan(a)); } #endif diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index f314e62..59415f4 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -1644,9 +1644,9 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) for(k=0;k<8;k++){ if(k%2==0) - hex_val=((v[k/2].s64&0xffff0000)>>16); - else hex_val=(v[k/2].s64&0xffff); + else + hex_val=((v[k/2].s64&0xffff0000)>>16); nw_v[k].f16 =*((half *)&hex_val); matrix_c[row+k][offset]=nw_v[k]; } @@ -1767,10 +1767,10 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) } else if(type==F16_TYPE){ ptx_reg_t nw_data1, nw_data2, nw_data3, nw_data4; - nw_data1.s64=((matrix_d[row][offset].s64 & 0xffff)<<16)|((matrix_d[row+1][offset].s64&0xffff)); - nw_data2.s64=((matrix_d[row+2][offset].s64 & 0xffff)<<16)|((matrix_d[row+3][offset].s64&0xffff)); - nw_data3.s64=((matrix_d[row+4][offset].s64 & 0xffff)<<16)|((matrix_d[row+5][offset].s64&0xffff)); - nw_data4.s64=((matrix_d[row+6][offset].s64 & 0xffff)<<16)|((matrix_d[row+7][offset].s64&0xffff)); + nw_data1.s64=((matrix_d[row][offset].s64 & 0xffff))|((matrix_d[row+1][offset].s64&0xffff)<<16); + nw_data2.s64=((matrix_d[row+2][offset].s64 & 0xffff))|((matrix_d[row+3][offset].s64&0xffff)<<16); + nw_data3.s64=((matrix_d[row+4][offset].s64 & 0xffff))|((matrix_d[row+5][offset].s64&0xffff)<<16); + nw_data4.s64=((matrix_d[row+6][offset].s64 & 0xffff))|((matrix_d[row+7][offset].s64&0xffff)<<16); thread->set_vector_operand_values(dst,nw_data1,nw_data2,nw_data3,nw_data4); printf("thread%d=%x,%x,%x,%x",thrd,nw_data1.s64,nw_data2.s64,nw_data3.s64,nw_data4.s64); @@ -2295,7 +2295,7 @@ ptx_reg_t d2d( ptx_reg_t x, unsigned from_width, unsigned to_width, int to_sign, y.f64 = x.f64; break; } - if (isnan(y.f64)) { + if (std::isnan(y.f64)) { y.u64 = 0xfff8000000000000ull; } else if (saturation_mode) { y.f64 = cuda_math::__saturatef(y.f64); @@ -2426,7 +2426,7 @@ void ptx_round(ptx_reg_t& data, int rounding_mode, int type) } } if ((type == F64_TYPE)||(type == FF64_TYPE)) { - if (isnan(data.f64)) { + if (std::isnan(data.f64)) { data.u64 = 0xfff8000000000000ull; } } @@ -2828,9 +2828,9 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) ptx_reg_t nw_v[8]; for(k=0;k<8;k++){ if(k%2==0) - nw_v[k].s64=((v[k/2].s64&0xffff0000)>>16); - else nw_v[k].s64=(v[k/2].s64&0xffff); + else + nw_v[k].s64=((v[k/2].s64&0xffff0000)>>16); } for(k=0;k<8;k++){ @@ -2905,10 +2905,10 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) } else{ ptx_reg_t nw_data1, nw_data2, nw_data3, nw_data4; - nw_data1.s64=((data1.s64 & 0xffff)<<16)|((data2.s64&0xffff)); - nw_data2.s64=((data3.s64 & 0xffff)<<16)|((data4.s64&0xffff)); - nw_data3.s64=((data5.s64 & 0xffff)<<16)|((data6.s64&0xffff)); - nw_data4.s64=((data7.s64 & 0xffff)<<16)|((data8.s64&0xffff)); + nw_data1.s64=((data1.s64 & 0xffff))|((data2.s64&0xffff)<<16); + nw_data2.s64=((data3.s64 & 0xffff))|((data4.s64&0xffff)<<16); + nw_data3.s64=((data5.s64 & 0xffff))|((data6.s64&0xffff)<<16); + nw_data4.s64=((data7.s64 & 0xffff))|((data8.s64&0xffff)<<16); printf("wmma_load:data1.s64=%x,data2.s64=%x,new_data1.s64=%x\n",data1.s64,data2.s64,nw_data1.s64); printf("wmma_load:data3.s64=%x,data4.s64=%x,new_data2.s64=%x\n",data3.s64,data4.s64,nw_data2.s64); printf("wmma_load:data5.s64=%x,data6.s64=%x,new_data3.s64=%x\n",data5.s64,data6.s64,nw_data3.s64); @@ -3139,12 +3139,12 @@ void mad_def( const ptx_instruction *pI, ptx_thread_info *thread, bool use_carry bool isNaN(float x) { - return isnan(x); + return std::isnan(x); } bool isNaN(double x) { - return isnan(x); + return std::isnan(x); } void max_impl( const ptx_instruction *pI, ptx_thread_info *thread ) diff --git a/src/cuda-sim/ptx_sim.h b/src/cuda-sim/ptx_sim.h index fbd7881..ea87e59 100644 --- a/src/cuda-sim/ptx_sim.h +++ b/src/cuda-sim/ptx_sim.h @@ -42,6 +42,12 @@ #include "memory.h" +#define GCC_VERSION (__GNUC__ * 10000 \ + + __GNUC_MINOR__ * 100 \ + + __GNUC_PATCHLEVEL__) + + + struct param_t { const void *pdata; int type; @@ -128,7 +134,12 @@ union ptx_reg_t { unsigned short u16; unsigned int u32; unsigned long long u64; + //gcc 4.7.0 + #if GCC_VERSION >= 40700 half f16; + #else + float f16; + #endif float f32; double f64; struct { -- cgit v1.3