From 7dfa2ae2e6f8ccaaf133318265a7ab00de546e82 Mon Sep 17 00:00:00 2001 From: aamir Date: Sun, 27 May 2018 14:18:53 -0700 Subject: added wmma parsing but execution getting aborted --- src/cuda-sim/ptx_ir.cc | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim/ptx_ir.cc') diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 8ebdcf8..9a4d8d3 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -995,7 +995,7 @@ static std::list check_operands( int opcode, const std::list &operands ) { static int g_warn_literal_operands_two_type_inst; - if( (opcode == CVT_OP) || (opcode == SET_OP) || (opcode == SLCT_OP) || (opcode == TEX_OP) ) { + if( (opcode == CVT_OP) || (opcode == SET_OP) || (opcode == SLCT_OP) || (opcode == TEX_OP) || (opcode==MMA_OP)) { // just make sure these do not have have const operands... if( !g_warn_literal_operands_two_type_inst ) { std::list::const_iterator o; @@ -1043,6 +1043,7 @@ ptx_instruction::ptx_instruction( int opcode, const std::list &operands, const operand_info &return_var, const std::list &options, + const std::list &wmma_options, const std::list &scalar_type, memory_space_t space_spec, const char *file, @@ -1061,6 +1062,7 @@ ptx_instruction::ptx_instruction( int opcode, m_operands.insert(m_operands.begin(), checked_operands.begin(), checked_operands.end() ); m_return_var = return_var; m_options = options; + m_wmma_options = wmma_options; m_wide = false; m_hi = false; m_lo = false; @@ -1078,7 +1080,7 @@ ptx_instruction::ptx_instruction( int opcode, m_atomic_spec = 0; m_membar_level = 0; m_inst_size = 8; // bytes - + int rr=0; std::list::const_iterator i; unsigned n=1; for ( i=options.begin(); i!= options.end(); i++, n++ ) { -- cgit v1.3 From 4161ccba0d4a99157afed3cdccef0e9c2a6d89e6 Mon Sep 17 00:00:00 2001 From: aamir Date: Tue, 5 Jun 2018 12:50:57 -0700 Subject: added support for wmma:load_c:f16_type --- cuda-kernels/tensorcore_type16_16.cu | 217 ++++++++++++++++++++++++++ cuda-kernels/tensorcore_type32_16.cu | 218 ++++++++++++++++++++++++++ cuda-kernels/tensorcore_type32_32.cu | 15 ++ src/cuda-sim/instructions.cc | 287 +++++++++++++++++++++++++---------- src/cuda-sim/ptx_ir.cc | 21 +++ src/cuda-sim/ptx_ir.h | 9 +- 6 files changed, 686 insertions(+), 81 deletions(-) create mode 100644 cuda-kernels/tensorcore_type16_16.cu create mode 100644 cuda-kernels/tensorcore_type32_16.cu (limited to 'src/cuda-sim/ptx_ir.cc') diff --git a/cuda-kernels/tensorcore_type16_16.cu b/cuda-kernels/tensorcore_type16_16.cu new file mode 100644 index 0000000..2b93bf5 --- /dev/null +++ b/cuda-kernels/tensorcore_type16_16.cu @@ -0,0 +1,217 @@ +#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, half *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; + + // 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); + wmma::mma_sync(c_frag, a_frag, b_frag, c_frag); + + wmma::store_matrix_sync(d_fp16, c_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); + convertFp32ToFp16 <<< (MATRIX_M * MATRIX_N + 255) / 256, 256 >>> (c_fp16, c_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_fp16, 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=200000; + 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/cuda-kernels/tensorcore_type32_16.cu b/cuda-kernels/tensorcore_type32_16.cu new file mode 100644 index 0000000..c66d8f8 --- /dev/null +++ b/cuda-kernels/tensorcore_type32_16.cu @@ -0,0 +1,218 @@ +#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, half *c,float *d_fp32, 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); + wmma::mma_sync(d_frag, a_frag, b_frag, c_frag); + + wmma::store_matrix_sync(d_fp32, 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); + convertFp32ToFp16 <<< (MATRIX_M * MATRIX_N + 255) / 256, 256 >>> (c_fp16, c_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_fp16, d_fp32 , 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/cuda-kernels/tensorcore_type32_32.cu b/cuda-kernels/tensorcore_type32_32.cu index 0d26163..73386f9 100644 --- a/cuda-kernels/tensorcore_type32_32.cu +++ b/cuda-kernels/tensorcore_type32_32.cu @@ -167,7 +167,10 @@ int main(int argc, char* argv[]) { cudaErrCheck(cudaEventDestroy(startWMMA)); cudaErrCheck(cudaEventDestroy(stopWMMA)); + int t=200000; + 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)); diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 70aee35..f314e62 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -57,6 +57,42 @@ const char *g_opcode_string[NUM_OPCODES] = { #undef OP_W_DEF }; +unsigned thread_group_offset(int thread){ + unsigned thread_group=thread/4; + unsigned in_tg_index=thread%4; + unsigned offset; + switch(thread_group){ + case 0: + offset=0; + break; + case 1: + offset=8; + break; + + case 2: + offset=128; + break; + case 3: + offset=136; + break; + case 4: + offset=4; + break; + case 5: + offset=12; + break; + case 6: + offset=132; + break; + case 7: + offset=140; + break; + default: + abort(); + + } + return offset+in_tg_index; +} void inst_not_implemented( const ptx_instruction * pI ) ; ptx_reg_t srcOperandModifiers(ptx_reg_t opData, operand_info opInfo, operand_info dstInfo, unsigned type, ptx_thread_info *thread); @@ -655,7 +691,7 @@ void ptx_thread_info::set_wmma_vector_operand_values( const operand_info &dst, const ptx_reg_t &data8 ) { unsigned num_elements = dst.get_vect_nelem(); - if (num_elements > 7) { + if (num_elements == 8) { set_reg(dst.vec_symbol(0), data1); set_reg(dst.vec_symbol(1), data2); set_reg(dst.vec_symbol(2), data3); @@ -1534,6 +1570,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) ptx_reg_t src_data; ptx_thread_info *thread; + unsigned wmma_type = pI->get_wmma_type(); unsigned type = pI->get_type(); unsigned type2 = pI->get_type2(); int tid = inst.warp_id_func() * core->get_warp_size(); @@ -1543,9 +1580,8 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) //NOT WOR const operand_info &src_a= pI->operand_lookup(1); //NOT WOR src_data= (thread->get_operand_value(src_a, dst, type, thread, 1)); //NOT WOR thread->set_operand_value(dst, src_data, type, thread, pI); + unsigned thread_group_index; for (thrd=0; thrd < core->get_warp_size(); thrd++){ - row=thrd/2; - offset=8*(thrd%2); thread = core->get_thread_info()[tid+thrd]; printf("thread=%d:",thrd); for(i=1;i<=3;i++){ @@ -1554,40 +1590,84 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) unsigned nelem = src_a.get_vect_nelem(); ptx_reg_t v[8]; thread->get_vector_operand_values( src_a, v, nelem ); - if(i!=3||((i==3)&&(type==F16_TYPE))){ - printf("%x ",v[0].f16); - printf("%x ",v[1].f16); - printf("%x ",v[2].f16); - printf("%x ",v[3].f16); - printf("%x ",v[4].f16); - printf("%x ",v[5].f16); - printf("%x ",v[6].f16); - printf("%x ",v[7].f16); + if(i!=3){ + printf("%x ",v[0].f16); + printf("%x ",v[1].f16); + printf("%x ",v[2].f16); + printf("%x ",v[3].f16); + printf("%x ",v[4].f16); + printf("%x ",v[5].f16); + printf("%x ",v[6].f16); + printf("%x ",v[7].f16); + } else{ - printf("%f ",v[0].f32); - printf("%f ",v[1].f32); - printf("%f ",v[2].f32); - printf("%f ",v[3].f32); - printf("%f ",v[4].f32); - printf("%f ",v[5].f32); - printf("%f ",v[6].f32); - printf("%f ",v[7].f32); + if(type2==F32_TYPE){ + printf("%f ",v[0].f32); + printf("%f ",v[1].f32); + printf("%f ",v[2].f32); + printf("%f ",v[3].f32); + printf("%f ",v[4].f32); + printf("%f ",v[5].f32); + printf("%f ",v[6].f32); + printf("%f ",v[7].f32); + } + else{ + printf("%x ",v[0].s64); + printf("%x ",v[1].s64); + printf("%x ",v[2].s64); + printf("%x ",v[3].s64); + } } - + thread_group_index=thread_group_offset(thrd); + row=(thread_group_index/16); + offset=thread_group_index%16; switch(i) { case 1 ://operand 1 for(k=0;k<8;k++) - matrix_a[row][offset+k]=v[k]; + matrix_a[row+k][offset]=v[k]; break; case 2 ://operand 2 for(k=0;k<8;k++) - matrix_b[row][offset+k]=v[k]; + matrix_b[row+k][offset]=v[k]; break; case 3 ://operand 3 - for(k=0;k<8;k++) - matrix_c[row][offset+k]=v[k]; - break; + if(type2!=F16_TYPE){ + for(k=0;k<8;k++) + matrix_c[row+k][offset]=v[k]; + } + else { + ptx_reg_t nw_v[8]; + unsigned int n = 0x41933333; + float f = *((float*)&n); + int hex_val; + + 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); + nw_v[k].f16 =*((half *)&hex_val); + matrix_c[row+k][offset]=nw_v[k]; + } + printf("%x ",nw_v[0].f16); + printf("%x ",nw_v[1].f16); + printf("%x ",nw_v[2].f16); + printf("%x ",nw_v[3].f16); + printf("%x ",nw_v[4].f16); + printf("%x ",nw_v[5].f16); + printf("%x ",nw_v[6].f16); + printf("%x ",nw_v[7].f16); + //float t; + //int m; + //printf("\n"); + //for(m=0;m<8;m++){ + // t=nw_v[m].f16; + // printf(" %f ",t); + //} + //printf("\n"); + } + break; default : printf("Invalid Operand Index\n" ); } @@ -1612,10 +1692,10 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) printf("MATRIX_C\n"); for (i=0;i<16;i++){ for(j=0;j<16;j++){ - if(type==F16_TYPE) + if(type2==F16_TYPE) printf("%x ",matrix_c[i][j].f16); else - printf("%f ",matrix_c[i][j].f32); + printf("%f ",matrix_c[i][j].f32); } printf("\n"); } @@ -1628,18 +1708,33 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) printf("MATRIX_D\n"); for (i=0;i<16;i++){ for(j=0;j<16;j++){ - printf("%x ",matrix_d[i][j].f16); + if(type==F16_TYPE) + printf("%x ",matrix_d[i][j].f16); + else + printf("%.2f ",matrix_d[i][j].f32); + } printf("\n"); } float temp; + half temp2; for (i=0;i<16;i++){ for(j=0;j<16;j++){ for(k=0;k<16;k++){ matrix_d[i][j].f16=matrix_d[i][j].f16+matrix_a[i][k].f16*matrix_b[k][j].f16; } - if(type==F16_TYPE) + if((type==F16_TYPE)&&(type2==F16_TYPE)) matrix_d[i][j].f16+=matrix_c[i][j].f16; + else if((type==F32_TYPE)&&(type2==F16_TYPE)){ + temp2=matrix_d[i][j].f16+matrix_c[i][j].f16; + temp=temp2; + matrix_d[i][j].f32=temp; + } + else if((type==F16_TYPE)&&(type2==F32_TYPE)){ + temp=matrix_d[i][j].f16; + temp+=matrix_c[i][j].f32; + matrix_d[i][j].f16=half(temp); + } else{ temp=matrix_d[i][j].f16; temp+=matrix_c[i][j].f32; @@ -1658,16 +1753,33 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) printf("\n"); } for (thrd=0; thrd < core->get_warp_size(); thrd++){ + thread_group_index=thread_group_offset(thrd); + row=(thread_group_index/16); + offset=thread_group_index%16; thread = core->get_thread_info()[tid+thrd]; - row=thrd/2; - offset=8*(thrd%2); - //r2=dst.get_symbol(); - //printf("thrd=%d,i=%d,register%s, data=%f\n",thrd,i,(r2->name()).c_str(),matrix_d[row][offset+i].f32); - //thread->set_operand_value(dst, matrix_d[row][offset+i], type, thread, pI); - thread->set_wmma_vector_operand_values(dst,matrix_d[row][offset],matrix_d[row][offset+1],matrix_d[row][offset+2],matrix_d[row][offset+3],matrix_d[row][offset+4],matrix_d[row][offset+5],matrix_d[row][offset+6],matrix_d[row][offset+7]); - printf("thread%d=%x,%x,%x,%x",thrd,matrix_d[row][offset].f16,matrix_d[row][offset+1].f16,matrix_d[row][offset+2].f16,matrix_d[row][offset+3].f16); - printf(",%x,%x,%x,%x\n",matrix_d[row][offset+4].f16,matrix_d[row][offset+5].f16,matrix_d[row][offset+6].f16,matrix_d[row][offset+7].f16); - } + //r2=dst.get_symbol(); + //printf("thrd=%d,i=%d,register%s, data=%f\n",thrd,i,(r2->name()).c_str(),matrix_d[row][offset+i].f32); + //thread->set_operand_value(dst, matrix_d[row][offset+i], type, thread, pI); + if(type==F32_TYPE){ + thread->set_wmma_vector_operand_values(dst,matrix_d[row][offset],matrix_d[row+1][offset],matrix_d[row+2][offset],matrix_d[row+3][offset],matrix_d[row+4][offset],matrix_d[row+5][offset],matrix_d[row+6][offset],matrix_d[row+7][offset]); + printf("thread%d=%x,%x,%x,%x",thrd,matrix_d[row][offset].f16,matrix_d[row+1][offset].f16,matrix_d[row+2][offset].f16,matrix_d[row+3][offset].f16); + printf(",%x,%x,%x,%x\n",matrix_d[row+4][offset].f16,matrix_d[row+5][offset].f16,matrix_d[row+6][offset].f16,matrix_d[row+7][offset].f16); + } + 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)); + 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); + + } + else{ + printf("wmma:mma:wrong type\n"); + abort(); + } + } } void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) @@ -1910,7 +2022,7 @@ ptx_reg_t f2x( ptx_reg_t x, unsigned from_width, unsigned to_width, int to_sign, { half mytemp; float myfloat; - assert( from_width == 32); + //assert( from_width == 32); enum cuda_math::cudaRoundMode mode = cuda_math::cudaRoundZero; switch (rounding_mode) { @@ -1959,7 +2071,10 @@ ptx_reg_t f2x( ptx_reg_t x, unsigned from_width, unsigned to_width, int to_sign, //y.f16 = half(x.f32); printf("f2x: %f\n",myfloat); break; - case 32: assert(0); break; // handled by f2f + case 32: + y.f32=float(x.f16); + + break; // handled by f2f case 64: y.f64 = x.f32; break; @@ -2673,6 +2788,7 @@ void ldu_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { ld_exec(pI,thread); } + void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) { size_t size; @@ -2685,6 +2801,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) const operand_info &src2 = pI->operand_lookup(2); int tid = inst.warp_id_func()*core->get_warp_size(); unsigned type = pI->get_type(); + unsigned wmma_type = pI->get_wmma_type(); for (thrd=0; thrd < core->get_warp_size(); thrd++) { thread = core->get_thread_info()[tid+thrd]; @@ -2706,23 +2823,28 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) type_info_key::type_decode(type,size,t); printf("mma_st: thrd=%d,addr=%d, fp(size=%d), stride=%d\n",thrd,addr_reg.u32,size,src2_data.u32); - if(type==F16_TYPE){ - for(k=0;k<8;k++){ - mem->write(addr+inx*2*src2_data.u32+odd*16+k*size/8,size/8,&v[k].s64,thread,pI); + addr_t new_addr = addr+thread_group_offset(thrd)*size/8; + + 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); + } + + for(k=0;k<8;k++){ + if(type==F32_TYPE){ + mem->write(new_addr+k*2*size,size/8,&v[k].s64,thread,pI); + printf("wmma:store:thread%d=%x,%x,%x,%x,%x,%x,%x,%x\n",thrd,v[0].s64,v[1].s64,v[2].s64,v[3].s64,v[4].s64,v[5].s64,v[6].s64,v[7].s64); } - } - else if(type==F32_TYPE){ - for(k=0;k<8;k++){ - mem->write(addr+inx*4*src2_data.u32+odd*32+k*size/8,size/8,&v[k].s64,thread,pI); + else if(type==F16_TYPE){ + mem->write(new_addr+k*2*size,size/8,&nw_v[k].s64,thread,pI); + printf("wmma:store:thread%d=%x,%x,%x,%x,%x,%x,%x,%x\n",thrd,nw_v[0].s64,nw_v[1].s64,nw_v[2].s64,nw_v[3].s64,nw_v[4].s64,nw_v[5].s64,nw_v[6].s64,nw_v[7].s64); } - } - else{ - printf("wmma:wrong error type\n"); - } - printf("wmma:store:thread%d=%x,%x,%x,%x,%x,%x,%x,%x\n",thrd,v[0].s64,v[1].s64,v[2].s64,v[3].s64,v[4].s64,v[5].s64,v[6].s64,v[7].s64); - - delete [] v; + + delete [] v; thread->m_last_effective_address = addr; thread->m_last_memory_space = space; } @@ -2736,6 +2858,7 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) const operand_info &src2 = pI->src2(); unsigned type = pI->get_type(); + unsigned wmma_type = pI->get_wmma_type(); int tid = inst.warp_id_func()*core->get_warp_size(); int thrd,odd,inx; @@ -2758,32 +2881,40 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) ptx_reg_t data1, data2, data3, data4; ptx_reg_t data5, data6, data7, data8; printf("mma_ld: thrd=%d,addr=%d, fp16(size=%d), stride=%d\n",thrd,src1_data.u32,size,src2_data.u32); - if(type==F16_TYPE){ - mem->read(addr+inx*2*src2_data.u32+odd*16,size/8,&data1.s64); - mem->read(addr+inx*2*src2_data.u32+odd*16+size/8,size/8,&data2.s64); - mem->read(addr+inx*2*src2_data.u32+odd*16+2*size/8,size/8,&data3.s64); - mem->read(addr+inx*2*src2_data.u32+odd*16+3*size/8,size/8,&data4.s64); - mem->read(addr+inx*2*src2_data.u32+odd*16+4*size/8,size/8,&data5.s64); - mem->read(addr+inx*2*src2_data.u32+odd*16+5*size/8,size/8,&data6.s64); - mem->read(addr+inx*2*src2_data.u32+odd*16+6*size/8,size/8,&data7.s64); - mem->read(addr+inx*2*src2_data.u32+odd*16+7*size/8,size/8,&data8.s64); - printf("thread%d=%x,%x,%x,%x,%x,%x,%x,%x\n",0,data1.s64,data2.s64,data3.s64,data4.s64,data5.s64,data6.s64,data7.s64,data8.s64); - } - else if(type==F32_TYPE){ - mem->read(addr+inx*4*src2_data.u32+odd*32,size/8,&data1.s64); - mem->read(addr+inx*4*src2_data.u32+odd*32+size/8,size/8,&data2.s64); - mem->read(addr+inx*4*src2_data.u32+odd*32+2*size/8,size/8,&data3.s64); - mem->read(addr+inx*4*src2_data.u32+odd*32+3*size/8,size/8,&data4.s64); - mem->read(addr+inx*4*src2_data.u32+odd*32+4*size/8,size/8,&data5.s64); - mem->read(addr+inx*4*src2_data.u32+odd*32+5*size/8,size/8,&data6.s64); - mem->read(addr+inx*4*src2_data.u32+odd*32+6*size/8,size/8,&data7.s64); - mem->read(addr+inx*4*src2_data.u32+odd*32+7*size/8,size/8,&data8.s64); - printf("thread%d=%f,%f,%f,%f,%f,%f,%f,%f\n",thrd,data1.f32,data2.f32,data3.f32,data4.f32,data5.f32,data6.f32,data7.f32,data8.f32); + + addr_t new_addr = addr+thread_group_offset(thrd)*size/8; + mem->read(new_addr,size/8,&data1.s64); + mem->read(new_addr+2*size,size/8,&data2.s64); + mem->read(new_addr+4*size,size/8,&data3.s64); + mem->read(new_addr+6*size,size/8,&data4.s64); + mem->read(new_addr+8*size,size/8,&data5.s64); + mem->read(new_addr+10*size,size/8,&data6.s64); + mem->read(new_addr+12*size,size/8,&data7.s64); + mem->read(new_addr+14*size,size/8,&data8.s64); + + if(type==F16_TYPE) + printf("thread%d=%x,%x,%x,%x,%x,%x,%x,%x\n",thrd,data1.s64,data2.s64,data3.s64,data4.s64,data5.s64,data6.s64,data7.s64,data8.s64); + + else if(type==F32_TYPE) + printf("thread%d=%f,%f,%f,%f,%f,%f,%f,%f\n",thrd,data1.f32,data2.f32,data3.f32,data4.f32,data5.f32,data6.f32,data7.f32,data8.f32); + else + printf("wmma_ld:wrong type\n"); + + if(!((wmma_type==LOAD_C)&&(type==F16_TYPE))){ + thread->set_wmma_vector_operand_values(dst,data1,data2,data3,data4,data5,data6,data7,data8); } else{ - printf("wmma_ld:wrong type\n"); + 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)); + 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); + printf("wmma_load:data7.s64=%x,data8.s64=%x,new_data4.s64=%x\n",data7.s64,data8.s64,nw_data4.s64); + thread->set_vector_operand_values(dst,nw_data1,nw_data2,nw_data3,nw_data4); } - thread->set_wmma_vector_operand_values(dst,data1,data2,data3,data4,data5,data6,data7,data8); thread->m_last_effective_address = addr; thread->m_last_memory_space = space; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 9a4d8d3..fb9adca 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1083,6 +1083,27 @@ ptx_instruction::ptx_instruction( int opcode, int rr=0; std::list::const_iterator i; unsigned n=1; + for ( i=wmma_options.begin(); i!= wmma_options.end(); i++, n++ ) { + int last_ptx_inst_option = *i; + switch ( last_ptx_inst_option ) { + case SYNC_OPTION: + case LOAD_A: + case LOAD_B: + case LOAD_C: + case STORE_D: + case MMA: + m_wmma_type=last_ptx_inst_option; + break; + case ROW: + case COL: + case M16N16K16: + break; + default: + assert(0); + break; + } + } + n=1; for ( i=options.begin(); i!= options.end(); i++, n++ ) { int last_ptx_inst_option = *i; switch ( last_ptx_inst_option ) { diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 6bba717..7bc7522 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1025,6 +1025,9 @@ public: unsigned get_vector() const { return m_vector_spec;} unsigned get_atomic() const { return m_atomic_spec;} + int get_wmma_type() const { + return m_wmma_type; + } int get_type() const { assert( !m_scalar_type.empty() ); @@ -1134,9 +1137,9 @@ private: bool m_uni; //if branch instruction, this evaluates to true for uniform branches (ie jumps) bool m_to_option; unsigned m_cache_option; - unsigned m_wmma_type; - unsigned m_wmma_layout[2]; - unsigned m_wmma_configuration; + int m_wmma_type; + int m_wmma_layout[2]; + int m_wmma_configuration; unsigned m_rounding_mode; unsigned m_compare_op; unsigned m_saturation_mode; -- cgit v1.3 From a86a03769bbdd5ea99d194704086f9fbe82104c3 Mon Sep 17 00:00:00 2001 From: aamir Date: Mon, 11 Jun 2018 10:39:07 -0700 Subject: added all the configuration --- src/cuda-sim/instructions.cc | 543 ++++++++++++++++++++++++++++++------------- src/cuda-sim/ptx_ir.cc | 3 + src/cuda-sim/ptx_ir.h | 3 + 3 files changed, 382 insertions(+), 167 deletions(-) (limited to 'src/cuda-sim/ptx_ir.cc') diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 92c8529..e03cbce 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -56,43 +56,102 @@ const char *g_opcode_string[NUM_OPCODES] = { #undef OP_DEF #undef OP_W_DEF }; +//Using profiled information::check the TensorCoreMatrixArrangement.xls for details +unsigned thread_group_offset(int thread,unsigned wmma_type,unsigned wmma_layout,unsigned type){ -unsigned thread_group_offset(int thread){ + unsigned offset; + unsigned load_a_row[8]={0,128,0,128,64,192,64,192}; + unsigned load_a_col[8]={0,8,0,8,4,12,4,12}; + unsigned load_b_row[8]={0,8,0,8,4,12,4,12}; + unsigned load_b_col[8]={0,128,0,128,64,192,64,192}; + unsigned load_c_float_row[8]={0,128,8,136,64,192,72,200}; + unsigned load_c_float_col[8]={0,8,128,136,4,12,132,140}; + unsigned load_c_half_row[8]={0,128,8,136,64,192,72,200}; + unsigned load_c_half_col[8]={0,8,128,136,4,12,132,140}; unsigned thread_group=thread/4; unsigned in_tg_index=thread%4; - unsigned offset; - switch(thread_group){ - case 0: - offset=0; - break; - case 1: - offset=8; - break; - - case 2: - offset=128; - break; - case 3: - offset=136; - break; - case 4: - offset=4; - break; - case 5: - offset=12; - break; - case 6: - offset=132; + + switch(wmma_type){ + case LOAD_A: + if(wmma_layout==ROW) + offset=load_a_row[thread_group]+16*in_tg_index; + else + offset=load_a_col[thread_group]+16*in_tg_index; + break; + + + case LOAD_B: + if(wmma_layout==ROW) + offset=load_b_row[thread_group]+16*in_tg_index; + else + offset=load_b_col[thread_group]+16*in_tg_index; break; - case 7: - offset=140; + + case LOAD_C: + case STORE_D: + if(type==F16_TYPE){ + if(wmma_layout==ROW) + offset=load_c_half_row[thread_group]+16*in_tg_index; + else + offset=load_c_half_col[thread_group]+in_tg_index; + } + else{ + if(wmma_layout==ROW) + offset=load_c_float_row[thread_group]; + else + offset=load_c_float_col[thread_group]; + + switch(in_tg_index){ + case 0: + break; + case 1: + if(wmma_layout==ROW) + offset+=16; + else + offset+=1; + break; + case 2: + if(wmma_layout==ROW) + offset+=2; + else + offset+=32; + break; + case 3: + if(wmma_layout==ROW) + offset+=18; + else + offset+=33; + break; + default: + abort(); + } + } break; + default: abort(); } - return offset+in_tg_index; + + return offset; +} + +int acc_float_offset(int index,int wmma_layout){ + + int c_row_offset[]={0,1,32,33,4,5,36,37}; + int c_col_offset[]={0,16,2,18,64,80,66,82}; + + if(wmma_layout==ROW) + return c_row_offset[index]; + else if(wmma_layout==COL) + return c_col_offset[index]; + else{ + printf("wrong layout"); + abort(); + } + } + void inst_not_implemented( const ptx_instruction * pI ) ; ptx_reg_t srcOperandModifiers(ptx_reg_t opData, operand_info opInfo, operand_info dstInfo, unsigned type, ptx_thread_info *thread); @@ -1546,11 +1605,62 @@ unsigned trunc(unsigned num, unsigned precision) { } return num; } +void mapping(int thread,int wmma_type,int wmma_layout,int type,int index,int &row,int &col,int &assg_offset){ + int offset; + int c_row_offset[]={0,8,0,8,4,12,4,12}; + int c_col_offset[]={0,0,8,8,0,0,8,8}; + int c_tg_inside_row_offset[]={0,1,0,1}; + int c_tg_inside_col_offset[]={0,0,2,2}; + int c_inside_row_offset[]={0,0,2,2,0,0,2,2}; + int c_inside_col_offset[]={0,1,0,1,4,5,4,5}; + + offset=thread_group_offset(thread,wmma_type,wmma_layout,type); + + if(wmma_type==LOAD_A){ + if(wmma_layout==ROW){ + offset+=index+8*((thread%16)/8); + } + else{ + offset+=64*(index/4)+index%4+128*((thread%16)/8); + } + assg_offset=index+8*((thread%16)/8); + } + else if(wmma_type==LOAD_B){ + if(wmma_layout==ROW){ + offset+=64*(index/4)+index%4+128*((thread%16)/8); + } + else{ + offset+=index+8*((thread%16)/8); + } + assg_offset=index+8*((thread%16)/8); + } + else if( wmma_type==LOAD_C){ + if(type==F16_TYPE){ + row=c_row_offset[thread/4]+thread%4; + col=c_col_offset[thread/4]+index; + } + else{ + row=c_row_offset[thread/4]+c_tg_inside_row_offset[thread%4]+c_inside_row_offset[index]; + col=c_col_offset[thread/4]+c_tg_inside_col_offset[thread%4]+c_inside_col_offset[index]; + } + assg_offset=index; + } + if(wmma_type==LOAD_A||wmma_type==LOAD_B){ + if(wmma_layout==ROW){ + row=offset/16; + col=offset%16; + } + else{ + col=offset/16; + row=offset%16; + } + } +} void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) { int i,j,k,thrd; - int row,offset; + int row,col,offset; printf("mmaWorld\n"); ptx_reg_t matrix_a[16][16]; ptx_reg_t matrix_b[16][16]; @@ -1560,101 +1670,82 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) ptx_thread_info *thread; unsigned wmma_type = pI->get_wmma_type(); + unsigned a_layout = pI->get_wmma_layout(0); + unsigned b_layout = pI->get_wmma_layout(1); unsigned type = pI->get_type(); unsigned type2 = pI->get_type2(); int tid = inst.warp_id_func() * core->get_warp_size(); const operand_info &dst = pI->operand_lookup(0); -//NOT WOR thread = core->get_thread_info()[tid]; -//NOT WOR const operand_info &src_a= pI->operand_lookup(1); -//NOT WOR src_data= (thread->get_operand_value(src_a, dst, type, thread, 1)); -//NOT WOR thread->set_operand_value(dst, src_data, type, thread, pI); unsigned thread_group_index; + float temp; + half temp2; + for (thrd=0; thrd < core->get_warp_size(); thrd++){ thread = core->get_thread_info()[tid+thrd]; - printf("thread=%d:",thrd); + printf("THREAD=%d\n:",thrd); for(i=1;i<=3;i++){ - int k; const operand_info &src_a= pI->operand_lookup(i); unsigned nelem = src_a.get_vect_nelem(); ptx_reg_t v[8]; thread->get_vector_operand_values( src_a, v, nelem ); - if(i!=3){ - printf("%x ",v[0].f16); - printf("%x ",v[1].f16); - printf("%x ",v[2].f16); - printf("%x ",v[3].f16); - printf("%x ",v[4].f16); - printf("%x ",v[5].f16); - printf("%x ",v[6].f16); - printf("%x ",v[7].f16); - + + printf("Thread%d_Iteration=%d\n:",thrd,i); + for(k=0;k>16); + nw_v[k].f16 =*((half *)&hex_val); } - else{ - printf("%x ",v[0].s64); - printf("%x ",v[1].s64); - printf("%x ",v[2].s64); - printf("%x ",v[3].s64); + } + if(!((i==3)&&(type2==F32_TYPE))){ + for(k=0;k<2*nelem;k++){ + temp=nw_v[k].f16; + printf("%f ",temp); } + printf("\n"); + } + else{ + for(k=0;k<8;k++){ + printf("%f ",v[k].f32); + } + printf("\n"); } - thread_group_index=thread_group_offset(thrd); - row=(thread_group_index/16); - offset=thread_group_index%16; switch(i) { case 1 ://operand 1 - for(k=0;k<8;k++) - matrix_a[row+k][offset]=v[k]; - break; + for(k=0;k<8;k++){ + mapping(thrd,LOAD_A,a_layout,F16_TYPE,k,row,col,offset); + printf("A:thread=%d,row=%d,col=%d,offset=%d\n",thrd,row,col,offset); + matrix_a[row][col]=nw_v[offset]; + } + break; case 2 ://operand 2 - for(k=0;k<8;k++) - matrix_b[row+k][offset]=v[k]; - break; + for(k=0;k<8;k++){ + mapping(thrd,LOAD_B,b_layout,F16_TYPE,k,row,col,offset); + printf("B:thread=%d,row=%d,col=%d,offset=%d\n",thrd,row,col,offset); + matrix_b[row][col]=nw_v[offset]; + } + break; case 3 ://operand 3 - if(type2!=F16_TYPE){ - for(k=0;k<8;k++) - matrix_c[row+k][offset]=v[k]; - } - else { - ptx_reg_t nw_v[8]; - unsigned int n = 0x41933333; - float f = *((float*)&n); - int hex_val; - - for(k=0;k<8;k++){ - if(k%2==0) - 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]; + for(k=0;k<8;k++){ + mapping(thrd,LOAD_C,ROW,type2,k,row,col,offset); + printf("C:thread=%d,row=%d,col=%d,offset=%d\n",thrd,row,col,offset); + if(type2!=F16_TYPE){ + matrix_c[row][col]=v[offset]; + } + else { + matrix_c[row][col]=nw_v[offset]; } - printf("%x ",nw_v[0].f16); - printf("%x ",nw_v[1].f16); - printf("%x ",nw_v[2].f16); - printf("%x ",nw_v[3].f16); - printf("%x ",nw_v[4].f16); - printf("%x ",nw_v[5].f16); - printf("%x ",nw_v[6].f16); - printf("%x ",nw_v[7].f16); - //float t; - //int m; - //printf("\n"); - //for(m=0;m<8;m++){ - // t=nw_v[m].f16; - // printf(" %f ",t); - //} - //printf("\n"); } break; default : @@ -1667,22 +1758,26 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) printf("MATRIX_A\n"); for (i=0;i<16;i++){ for(j=0;j<16;j++){ - printf("%x ",matrix_a[i][j].f16); + temp=matrix_a[i][j].f16; + printf("%f ",temp); } printf("\n"); } printf("MATRIX_B\n"); for (i=0;i<16;i++){ for(j=0;j<16;j++){ - printf("%x ",matrix_b[i][j].f16); + temp=matrix_b[i][j].f16; + printf("%f ",temp); } printf("\n"); } printf("MATRIX_C\n"); for (i=0;i<16;i++){ for(j=0;j<16;j++){ - if(type2==F16_TYPE) - printf("%x ",matrix_c[i][j].f16); + if(type2==F16_TYPE){ + temp=matrix_c[i][j].f16; + printf("%f ",temp); + } else printf("%f ",matrix_c[i][j].f32); } @@ -1697,16 +1792,16 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) printf("MATRIX_D\n"); for (i=0;i<16;i++){ for(j=0;j<16;j++){ - if(type==F16_TYPE) - printf("%x ",matrix_d[i][j].f16); + if(type==F16_TYPE){ + temp=matrix_d[i][j].f16; + printf("%f ",temp); + } else printf("%.2f ",matrix_d[i][j].f32); } printf("\n"); } - float temp; - half temp2; for (i=0;i<16;i++){ for(j=0;j<16;j++){ for(k=0;k<16;k++){ @@ -1734,32 +1829,54 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) printf("MATRIX_D\n"); for (i=0;i<16;i++){ for(j=0;j<16;j++){ - if(type==F16_TYPE) - printf("%x ",matrix_d[i][j].f16); + if(type==F16_TYPE){ + temp=matrix_d[i][j].f16; + printf("%f ",temp); + } else printf("%.2f ",matrix_d[i][j].f32); } printf("\n"); } for (thrd=0; thrd < core->get_warp_size(); thrd++){ - thread_group_index=thread_group_offset(thrd); - row=(thread_group_index/16); - offset=thread_group_index%16; + int row_t[8]; + int col_t[8]; + for(k=0;k<8;k++){ + mapping(thrd,LOAD_C,ROW,type,k,row_t[k],col_t[k],offset); + printf("mma:store:row:%d,col%d\n",row_t[k],col_t[k]); + } thread = core->get_thread_info()[tid+thrd]; - //r2=dst.get_symbol(); - //printf("thrd=%d,i=%d,register%s, data=%f\n",thrd,i,(r2->name()).c_str(),matrix_d[row][offset+i].f32); - //thread->set_operand_value(dst, matrix_d[row][offset+i], type, thread, pI); + + if(type==F32_TYPE){ - thread->set_wmma_vector_operand_values(dst,matrix_d[row][offset],matrix_d[row+1][offset],matrix_d[row+2][offset],matrix_d[row+3][offset],matrix_d[row+4][offset],matrix_d[row+5][offset],matrix_d[row+6][offset],matrix_d[row+7][offset]); - printf("thread%d=%x,%x,%x,%x",thrd,matrix_d[row][offset].f16,matrix_d[row+1][offset].f16,matrix_d[row+2][offset].f16,matrix_d[row+3][offset].f16); - printf(",%x,%x,%x,%x\n",matrix_d[row+4][offset].f16,matrix_d[row+5][offset].f16,matrix_d[row+6][offset].f16,matrix_d[row+7][offset].f16); - } + thread->set_wmma_vector_operand_values(dst,matrix_d[row_t[0]][col_t[0]],matrix_d[row_t[1]][col_t[1]],matrix_d[row_t[2]][col_t[2]],matrix_d[row_t[3]][col_t[3]],matrix_d[row_t[4]][col_t[4]],matrix_d[row_t[5]][col_t[5]],matrix_d[row_t[6]][col_t[6]],matrix_d[row_t[7]][col_t[7]]); + + printf("thread%d:",thrd); + for(k=0;k<8;k++){ + printf("%f ",matrix_d[row_t[k]][col_t[k]].f32); + } + printf("\n"); + } else if(type==F16_TYPE){ + + printf("thread%d:",thrd); + for(k=0;k<8;k++){ + temp=matrix_d[row_t[k]][col_t[k]].f16; + printf("%f ",temp); + } + printf("\n"); + + printf("thread%d:",thrd); + for(k=0;k<8;k++){ + printf("%x ",matrix_d[row_t[k]][col_t[k]].f16); + } + printf("\n"); + ptx_reg_t nw_data1, nw_data2, nw_data3, nw_data4; - 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); + nw_data1.s64=((matrix_d[row_t[0]][col_t[0]].s64 & 0xffff))|((matrix_d[row_t[1]][col_t[1]].s64&0xffff)<<16); + nw_data2.s64=((matrix_d[row_t[2]][col_t[2]].s64 & 0xffff))|((matrix_d[row_t[3]][col_t[3]].s64&0xffff)<<16); + nw_data3.s64=((matrix_d[row_t[4]][col_t[4]].s64 & 0xffff))|((matrix_d[row_t[5]][col_t[5]].s64&0xffff)<<16); + nw_data4.s64=((matrix_d[row_t[6]][col_t[6]].s64 & 0xffff))|((matrix_d[row_t[7]][col_t[7]].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); @@ -2791,9 +2908,10 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) int tid = inst.warp_id_func()*core->get_warp_size(); unsigned type = pI->get_type(); unsigned wmma_type = pI->get_wmma_type(); + unsigned wmma_layout = pI->get_wmma_layout(0); for (thrd=0; thrd < core->get_warp_size(); thrd++) { - thread = core->get_thread_info()[tid+thrd]; + thread = core->get_thread_info()[tid+thrd]; odd=thrd%2; inx=thrd/2; ptx_reg_t addr_reg = thread->get_operand_value(src1, src, type, thread, 1); @@ -2812,7 +2930,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) type_info_key::type_decode(type,size,t); printf("mma_st: thrd=%d,addr=%d, fp(size=%d), stride=%d\n",thrd,addr_reg.u32,size,src2_data.u32); - addr_t new_addr = addr+thread_group_offset(thrd)*size/8; + addr_t new_addr = addr+thread_group_offset(thrd,wmma_type,wmma_layout,type)*size/8; ptx_reg_t nw_v[8]; for(k=0;k<8;k++){ @@ -2824,11 +2942,24 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) for(k=0;k<8;k++){ if(type==F32_TYPE){ - mem->write(new_addr+k*2*size,size/8,&v[k].s64,thread,pI); + mem->write(new_addr+4*acc_float_offset(k,wmma_layout),size/8,&v[k].s64,thread,pI); + printf("wmma:store:thread%d=%x,%x,%x,%x,%x,%x,%x,%x\n",thrd,v[0].s64,v[1].s64,v[2].s64,v[3].s64,v[4].s64,v[5].s64,v[6].s64,v[7].s64); + float temp; + int l; + printf("thread=%d:",thrd); + for(l=0;l<8;l++){ + temp=v[0].f32; + printf("%f",temp); + } + printf("\n"); + } else if(type==F16_TYPE){ - mem->write(new_addr+k*2*size,size/8,&nw_v[k].s64,thread,pI); + if(wmma_layout==ROW) + mem->write(new_addr+k*2,size/8,&nw_v[k].s64,thread,pI); + else if(wmma_layout==COL) + mem->write(new_addr+k*32,size/8,&nw_v[k].s64,thread,pI); printf("wmma:store:thread%d=%x,%x,%x,%x,%x,%x,%x,%x\n",thrd,nw_v[0].s64,nw_v[1].s64,nw_v[2].s64,nw_v[3].s64,nw_v[4].s64,nw_v[5].s64,nw_v[6].s64,nw_v[7].s64); } } @@ -2838,24 +2969,24 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) thread->m_last_memory_space = space; } } + void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) { size_t size; - int t; + int t,i; const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); const operand_info &src2 = pI->src2(); unsigned type = pI->get_type(); unsigned wmma_type = pI->get_wmma_type(); - + unsigned wmma_layout = pI->get_wmma_layout(0); int tid = inst.warp_id_func()*core->get_warp_size(); - int thrd,odd,inx; + int thrd; ptx_thread_info *thread; + for (thrd=0; thrd < core->get_warp_size(); thrd++){ thread = core->get_thread_info()[tid+thrd]; - odd=thrd%2; - inx=thrd/2; ptx_reg_t src1_data = thread->get_operand_value(src1, dst, U32_TYPE, thread, 1); ptx_reg_t src2_data = thread->get_operand_value(src2, dst, U32_TYPE, thread, 1); @@ -2863,46 +2994,124 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) memory_space *mem = NULL; addr_t addr = src1_data.u32; - decode_space(space,thread,src1,mem,addr); - type_info_key::type_decode(type,size,t); - ptx_reg_t data1, data2, data3, data4; - ptx_reg_t data5, data6, data7, data8; - printf("mma_ld: thrd=%d,addr=%d, fp16(size=%d), stride=%d\n",thrd,src1_data.u32,size,src2_data.u32); - addr_t new_addr = addr+thread_group_offset(thrd)*size/8; - mem->read(new_addr,size/8,&data1.s64); - mem->read(new_addr+2*size,size/8,&data2.s64); - mem->read(new_addr+4*size,size/8,&data3.s64); - mem->read(new_addr+6*size,size/8,&data4.s64); - mem->read(new_addr+8*size,size/8,&data5.s64); - mem->read(new_addr+10*size,size/8,&data6.s64); - mem->read(new_addr+12*size,size/8,&data7.s64); - mem->read(new_addr+14*size,size/8,&data8.s64); - - if(type==F16_TYPE) - printf("thread%d=%x,%x,%x,%x,%x,%x,%x,%x\n",thrd,data1.s64,data2.s64,data3.s64,data4.s64,data5.s64,data6.s64,data7.s64,data8.s64); + ptx_reg_t data[16]; + printf("mma_ld: thrd=%d,addr=%d, fpsize=%d, stride=%d\n",thrd,src1_data.u32,size,src2_data.u32); + addr_t new_addr = addr+thread_group_offset(thrd,wmma_type,wmma_layout,type)*size/8; + + if(wmma_type==LOAD_A){ + for(i=0;i<16;i++){ + if(wmma_layout==ROW) + mem->read(new_addr+2*i,size/8,&data[i].s64); + else if(wmma_layout==COL){ + mem->read(new_addr+2*(i%4)+128*(i/4),size/8,&data[i].s64); + } + else{ + printf("mma_ld:wrong_layout_type\n"); + abort(); + } + } + } + else if(wmma_type==LOAD_B){ + for(i=0;i<16;i++){ + if(wmma_layout==COL) + mem->read(new_addr+2*i,size/8,&data[i].s64); + else if(wmma_layout==ROW){ + mem->read(new_addr+2*(i%4)+128*(i/4),size/8,&data[i].s64); + } + else{ + printf("mma_ld:wrong_layout_type\n"); + abort(); + } + } + } + else if(wmma_type==LOAD_C){ + for(i=0;i<8;i++){ + if(type==F16_TYPE){ + if(wmma_layout==ROW) + mem->read(new_addr+2*i,size/8,&data[i].s64); + else if(wmma_layout==COL) + mem->read(new_addr+32*i,size/8,&data[i].s64); + else{ + printf("mma_ld:wrong_type\n"); + abort(); + } + } + else if(type==F32_TYPE){ + mem->read(new_addr+4*acc_float_offset(i,wmma_layout),size/8,&data[i].s64); + } + else{ + printf("wrong type"); + abort(); + } + } + } + else{ + printf("wrong wmma type\n");; + abort(); + } - else if(type==F32_TYPE) - printf("thread%d=%f,%f,%f,%f,%f,%f,%f,%f\n",thrd,data1.f32,data2.f32,data3.f32,data4.f32,data5.f32,data6.f32,data7.f32,data8.f32); - else - printf("wmma_ld:wrong type\n"); + if(type==F16_TYPE){ + printf("\nthread%d= ",thrd); + for(i=0;i<16;i++){ + printf("%x ",data[i].u64); + } + printf("\n"); + + printf("\nthread%d= ",thrd); + float temp; + for(i=0;i<16;i++){ + temp=data[i].f16; + printf("%f ",temp); + } + printf("\n"); + } + else{ + printf("\nthread%d= ",thrd); + for(i=0;i<8;i++){ + printf("%f ",data[i].f32); + } + printf("\n"); + printf("\nthread%d= ",thrd); + for(i=0;i<8;i++){ + printf("%x ",data[i].u64); + } + printf("\n"); + } - if(!((wmma_type==LOAD_C)&&(type==F16_TYPE))){ - thread->set_wmma_vector_operand_values(dst,data1,data2,data3,data4,data5,data6,data7,data8); + if((wmma_type==LOAD_C)&&(type==F32_TYPE)){ + thread->set_wmma_vector_operand_values(dst,data[0],data[1],data[2],data[3],data[4],data[5],data[6],data[7]); } else{ - ptx_reg_t nw_data1, nw_data2, nw_data3, nw_data4; - 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); - printf("wmma_load:data7.s64=%x,data8.s64=%x,new_data4.s64=%x\n",data7.s64,data8.s64,nw_data4.s64); - thread->set_vector_operand_values(dst,nw_data1,nw_data2,nw_data3,nw_data4); + ptx_reg_t nw_data[8]; + int num_reg; + + if(wmma_type==LOAD_C) + num_reg=4; + else + num_reg=8; + + for(i=0;iset_vector_operand_values(dst,nw_data[0],nw_data[1],nw_data[2],nw_data[3]); + else + thread->set_wmma_vector_operand_values(dst,nw_data[0],nw_data[1],nw_data[2],nw_data[3],nw_data[4],nw_data[5],nw_data[6],nw_data[7]); + + printf("wmma_load:data[0].s64=%x,data[1].s64=%x,new_data[0].s64=%x\n",data[0].u64,data[1].u64,nw_data[0].u64); + printf("wmma_load:data[2].s64=%x,data[3].s64=%x,new_data[1].s64=%x\n",data[2].u64,data[3].u64,nw_data[1].u64); + printf("wmma_load:data[4].s64=%x,data[5].s64=%x,new_data[2].s64=%x\n",data[4].u64,data[5].u64,nw_data[2].u64); + printf("wmma_load:data[6].s64=%x,data[7].s64=%x,new_data[3].s64=%x\n",data[6].u64,data[7].u64,nw_data[3].u64); + if(wmma_type!=LOAD_C){ + printf("wmma_load:data[8].s64=%x,data[9].s64=%x,new_data[4].s64=%x\n",data[8].u64,data[9].u64,nw_data[4].s64); + printf("wmma_load:data[10].s64=%x,data[11].s64=%x,new_data[5].s64=%x\n",data[10].u64,data[11].u64,nw_data[5].u64); + printf("wmma_load:data[12].s64=%x,data[13].s64=%x,new_data[6].s64=%x\n",data[12].u64,data[13].u64,nw_data[6].u64); + printf("wmma_load:data[14].s64=%x,data[15].s64=%x,new_data[7].s64=%x\n",data[14].u64,data[15].u64,nw_data[3].u64); + } } thread->m_last_effective_address = addr; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index fb9adca..9c2ac69 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1096,6 +1096,8 @@ ptx_instruction::ptx_instruction( int opcode, break; case ROW: case COL: + m_wmma_layout[rr++]=last_ptx_inst_option; + break; case M16N16K16: break; default: @@ -1103,6 +1105,7 @@ ptx_instruction::ptx_instruction( int opcode, break; } } + rr=0; n=1; for ( i=options.begin(); i!= options.end(); i++, n++ ) { int last_ptx_inst_option = *i; diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 7bc7522..62d7c7c 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1028,6 +1028,9 @@ public: int get_wmma_type() const { return m_wmma_type; } + int get_wmma_layout(int index) const { + return m_wmma_layout[index];//0->Matrix D,1->Matrix C + } int get_type() const { assert( !m_scalar_type.empty() ); -- cgit v1.3 From 1c74dcc29176cb3f6464d9088511216ba0e12c8d Mon Sep 17 00:00:00 2001 From: aamir Date: Tue, 7 Aug 2018 18:53:26 -0700 Subject: implemented prmt and started working on variable precision mul inst --- cuda-kernels/v4p_kernel.cu | 232 +++++++++++++++++++++++++++++++++++++++++++ src/cuda-sim/instructions.cc | 91 ++++++++++++++++- src/cuda-sim/ptx.l | 7 +- src/cuda-sim/ptx.y | 16 ++- src/cuda-sim/ptx_ir.cc | 25 +++-- src/cuda-sim/ptx_ir.h | 2 + src/cuda-sim/ptx_parser.cc | 2 +- 7 files changed, 362 insertions(+), 13 deletions(-) create mode 100644 cuda-kernels/v4p_kernel.cu (limited to 'src/cuda-sim/ptx_ir.cc') diff --git a/cuda-kernels/v4p_kernel.cu b/cuda-kernels/v4p_kernel.cu new file mode 100644 index 0000000..bb9064b --- /dev/null +++ b/cuda-kernels/v4p_kernel.cu @@ -0,0 +1,232 @@ +#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,float *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; + + // 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); + wmma::mma_sync(c_frag, a_frag, b_frag, c_frag); + + wmma::store_matrix_sync(d_fp16, c_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]; + } +} + +__global__ void convertInt32ToInt8 (int *out, int *in, int n) { + int idx = blockDim.x * blockIdx.x + threadIdx.x; + if (idx < n/4) { + out[idx] =(in[4*idx]&0xff)|(in[4*idx+1]&0xff)<<8|(in[4*idx+2]&0xff)<<16|(in[4*idx+3]&0xff)<<24; + } +} + +__global__ void convertInt8ToInt32 (int *out, int *in, int n) { + int idx = blockDim.x * blockIdx.x + threadIdx.x; + int shft_amt=8*(idx%4); + int shft_mask=0xff<>shft_amt; + } +} + +int main(int argc, char* argv[]) { + int *a_int32; + int *b_int32; + int *c_int32; + int *d_int32; + + int *a_int8; + int *b_int8; + + int *a_host_wmma; + int *b_host_wmma; + int *c_host_wmma; + int *d_host_wmma; + int *d_cal_host_wmma; + + cudaEvent_t startWMMA; + cudaEvent_t stopWMMA; + + + cudaErrCheck(cudaEventCreate(&startWMMA)); + cudaErrCheck(cudaEventCreate(&stopWMMA)); + + // Use tensor cores + cudaErrCheck(cudaMalloc((void**)&a_int32, MATRIX_M * MATRIX_K * sizeof(int))); + cudaErrCheck(cudaMalloc((void**)&b_int32, MATRIX_K * MATRIX_N * sizeof(int))); + cudaErrCheck(cudaMalloc((void**)&c_int32, MATRIX_K * MATRIX_N * sizeof(int))); + cudaErrCheck(cudaMalloc((void**)&d_int32, MATRIX_K * MATRIX_N * sizeof(int))); + cudaErrCheck(cudaMalloc((void**)&a_int8, MATRIX_M * MATRIX_K * sizeof(int)/4)); + cudaErrCheck(cudaMalloc((void**)&b_int8, MATRIX_K * MATRIX_N * sizeof(int)/4)); + + + a_host_wmma = (int *)malloc(MATRIX_M * MATRIX_K * sizeof(int)); + b_host_wmma = (int *)malloc(MATRIX_K * MATRIX_N * sizeof(int)); + c_host_wmma = (int *)malloc(MATRIX_M * MATRIX_N * sizeof(int)); + d_host_wmma = (int *)malloc(MATRIX_M * MATRIX_N * sizeof(int)); + d_cal_host_wmma = (int *)malloc(MATRIX_M * MATRIX_N * sizeof(int)); + + printf("a_int32\n"); + for(int m=0;m>> (a_int8, a_int32, MATRIX_M * MATRIX_K); + convertInt8ToInt32 <<< (MATRIX_M * MATRIX_K + 255) / 256, 256 >>> (d_int32, a_int8, MATRIX_M * MATRIX_K); + //convertFp32ToFp16 <<< (MATRIX_K * MATRIX_N + 255) / 256, 256 >>> (b_fp16, b_fp32, MATRIX_K * MATRIX_N); + //convertFp32ToFp16 <<< (MATRIX_M * MATRIX_N + 255) / 256, 256 >>> (c_fp16, c_fp32, MATRIX_K * MATRIX_N); + cudaErrCheck(cudaMemcpy(d_host_wmma, d_int32, MATRIX_M * MATRIX_N * sizeof(float), cudaMemcpyDeviceToHost)); + + +//AAMIR printf("\nM = %d, N = %d, K = %d. \n", MATRIX_M, MATRIX_N, MATRIX_K); +//AAMIR +//AAMIR printf("Running with wmma...\n"); +//AAMIR cudaErrCheck(cudaEventRecord(startWMMA)); +//AAMIR wmma_example <<< 1, 32>>> (a_fp16, b_fp16, c_fp32, d_fp32 , MATRIX_M, MATRIX_N, MATRIX_K); +//AAMIR cudaErrCheck(cudaEventRecord(stopWMMA)); +//AAMIR cudaErrCheck(cudaEventSynchronize(stopWMMA)); +//AAMIR +//AAMIR // Error checking +//AAMIR printf("\nChecking results...\n"); +//AAMIR cudaErrCheck(cudaMemcpy(d_host_wmma, d_fp32, MATRIX_M * MATRIX_N * sizeof(float), cudaMemcpyDeviceToHost)); +//AAMIR +//AAMIR printf("Results verified: cublas and WMMA agree.\n\n"); +//AAMIR float wmmaTime; +//AAMIR cudaErrCheck(cudaEventElapsedTime(&wmmaTime, startWMMA, stopWMMA)); +//AAMIR printf("wmma took %fms\n", wmmaTime); +//AAMIR +//AAMIR cudaErrCheck(cudaEventDestroy(startWMMA)); +//AAMIR cudaErrCheck(cudaEventDestroy(stopWMMA)); +//AAMIR +//AAMIR int t=200000; +//AAMIR while(t-->0); +//AAMIR printf("D_CALCULATED\n"); +//AAMIR +//AAMIR for(int m=0;m1) +//AAMIR { +//AAMIR printf("ERROR:\n"); +//AAMIR suc=0; +//AAMIR } +//AAMIR } +//AAMIR } +//AAMIR if(suc==1) +//AAMIR printf("COMPLETED_SUCCESSFULLY\n"); +//AAMIR + + cudaErrCheck(cudaFree(a_int32)); + cudaErrCheck(cudaFree(b_int32)); + cudaErrCheck(cudaFree(c_int32)); + cudaErrCheck(cudaFree(d_int32)); + cudaErrCheck(cudaFree(a_int8)); + cudaErrCheck(cudaFree(b_int8)); + + 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/instructions.cc b/src/cuda-sim/instructions.cc index 7af157f..aaee2a2 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -48,7 +48,7 @@ using half_float::half; unsigned ptx_instruction::g_num_ptx_inst_uid=0; -bool g_debug_instruction = 1; +bool g_debug_instruction = 0; const char *g_opcode_string[NUM_OPCODES] = { @@ -3911,7 +3911,94 @@ void popc_impl( const ptx_instruction *pI, ptx_thread_info *thread ) } void prefetch_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } void prefetchu_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } -void prmt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } + +int prmt_mode_present(int mode) +{ + int returnval=0; + switch(mode){ + case PRMT_F4E_MODE: + case PRMT_B4E_MODE: + case PRMT_RC8_MODE: + case PRMT_RC16_MODE: + case PRMT_ECL_MODE: + case PRMT_ECR_MODE: + returnval=1; + break; + default: + break; + } + return returnval; +} +int read_byte(int mode,int control,int d_sel_index,signed long long value){ + + int returnval; + int prmt_f4e_mode[4][4]={{0,1,2,3},{1,2,3,4},{2,3,4,5},{3,4,5,6}}; + int prmt_b4e_mode[4][4]={{0,7,6,5},{1,0,7,6},{2,1,0,7},{3,2,1,0}}; + int prmt_rc8_mode[4][4]={{0,0,0,0},{1,1,1,1},{2,2,2,2},{3,3,3,3}}; + int prmt_ecl_mode[4][4]={{0,1,2,3},{1,1,2,3},{2,2,2,3},{3,3,3,3}}; + int prmt_ecr_mode[4][4]={{0,0,0,0},{0,1,1,1},{0,1,2,2},{0,1,2,3}}; + int prmt_rc16_mode[4][4]={{0,1,0,1},{2,3,2,3},{0,1,0,1},{2,3,2,3}}; + + if(!prmt_mode_present(mode)){ + if(control&0x8){ + returnval=0xff; + } + else{ + returnval= (value>>(8*control)) & 0xff; + } + } + else{ + switch(mode){ + case PRMT_F4E_MODE: returnval=prmt_f4e_mode[control][d_sel_index];break; + case PRMT_B4E_MODE: returnval=prmt_b4e_mode[control][d_sel_index];break; + case PRMT_RC8_MODE: returnval=prmt_rc8_mode[control][d_sel_index];break; + case PRMT_ECL_MODE: returnval=prmt_ecl_mode[control][d_sel_index];break; + case PRMT_ECR_MODE: returnval=prmt_ecr_mode[control][d_sel_index];break; + case PRMT_RC16_MODE: returnval=prmt_rc16_mode[control][d_sel_index];break; + default: printf("ERROR\n");break; + } + } + return (returnval<<8*d_sel_index); +} + +void prmt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { + + ptx_reg_t src1_data, src2_data, src3_data,tmpdata,data; + const operand_info &dst = pI->dst(); + const operand_info &src1 = pI->src1(); + const operand_info &src2 = pI->src2(); + const operand_info &src3 = pI->src3(); + + unsigned mode = pI->prmt_op(); + unsigned i_type = pI->get_type(); + + src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1); + src2_data = thread->get_operand_value(src2, dst, i_type, thread, 1); + src3_data = thread->get_operand_value(src3, dst, i_type, thread, 1); + + tmpdata.s64=src1_data.s32|(src2_data.s64<<32); + int ctl[4]; + + if(!prmt_mode_present(mode)){ + ctl[0]=(src3_data.s32>>0)&0xf; + ctl[1]=(src3_data.s32>>4)&0xf; + ctl[2]=(src3_data.s32>>8)&0xf; + ctl[3]=(src3_data.s32>>12)&0xf; + } + else{ + ctl[0]=ctl[1]=ctl[2]=ctl[3]=(src3_data.s32>>0)&0x3; + } + + data.s32=0; + data.s32=data.s32|read_byte(mode,ctl[0],0,tmpdata.s64); //First byte-0 + data.s32=data.s32|read_byte(mode,ctl[1],1,tmpdata.s64); //Second byte-1 + data.s32=data.s32|read_byte(mode,ctl[2],2,tmpdata.s64); //Third byte-2 + data.s32=data.s32|read_byte(mode,ctl[3],3,tmpdata.s64); //Fourth byte-3 + + thread->set_operand_value(dst,data, i_type, thread, pI); + + +} void rcp_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index 03d6838..a6b6fcc 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -164,7 +164,12 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; \.row TC; ptx_lval.int_value = ROW; return LAYOUT; \.col TC; ptx_lval.int_value = COL; return LAYOUT; \.m16n16k16 TC; ptx_lval.int_value = M16N16K16; return CONFIGURATION; - +\.f4e TC; return PRMT_F4E_MODE; +\.b4e TC; return PRMT_B4E_MODE; +\.rc8 TC; return PRMT_RC8_MODE; +\.ecl TC; return PRMT_ECL_MODE; +\.ecr TC; return PRMT_ECR_MODE; +\.rc16 TC; return PRMT_RC16_MODE; \.align TC; return ALIGN_DIRECTIVE; \.branchtargets TC; return BRANCHTARGETS_DIRECTIVE; diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y index 8744663..cd455dd 100644 --- a/src/cuda-sim/ptx.y +++ b/src/cuda-sim/ptx.y @@ -50,7 +50,6 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. %token PTR_DIRECTIVE %token ENTRY_DIRECTIVE %token EXTERN_DIRECTIVE -%token WEAK_DIRECTIVE %token FILE_DIRECTIVE %token FUNC_DIRECTIVE %token GLOBAL_DIRECTIVE @@ -203,6 +202,12 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. %token DOWN_OPTION; %token BFLY_OPTION; %token IDX_OPTION; +%token PRMT_F4E_MODE; +%token PRMT_B4E_MODE; +%token PRMT_RC8_MODE; +%token PRMT_RC16_MODE; +%token PRMT_ECL_MODE; +%token PRMT_ECR_MODE; %type function_decl_header %type function_decl @@ -432,6 +437,7 @@ option: type_spec | addressable_spec | rounding_mode | wmma_spec + | prmt_spec | SYNC_OPTION { add_option(SYNC_OPTION); } | ARRIVE_OPTION { add_option(ARRIVE_OPTION); } | RED_OPTION { add_option(RED_OPTION); } @@ -520,6 +526,14 @@ compare_spec:EQ_OPTION { add_option(EQ_OPTION); } | NAN_OPTION { add_option(NAN_OPTION); } ; +prmt_spec: PRMT_F4E_MODE { add_option( PRMT_F4E_MODE); } + | PRMT_B4E_MODE { add_option( PRMT_B4E_MODE); } + | PRMT_RC8_MODE { add_option( PRMT_RC8_MODE); } + | PRMT_RC16_MODE{ add_option( PRMT_RC16_MODE);} + | PRMT_ECL_MODE { add_option( PRMT_ECL_MODE); } + | PRMT_ECR_MODE { add_option( PRMT_ECR_MODE); } + ; + wmma_spec: WMMA_DIRECTIVE LAYOUT CONFIGURATION{add_space_spec(global_space,0);add_ptr_spec(global_space); add_wmma_option($1);add_wmma_option($2);add_wmma_option($3);} | WMMA_DIRECTIVE LAYOUT LAYOUT CONFIGURATION{add_wmma_option($1);add_wmma_option($2);add_wmma_option($3);add_wmma_option($4);} ; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 9c2ac69..d12c741 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1233,16 +1233,25 @@ ptx_instruction::ptx_instruction( int opcode, case HALF_OPTION: m_inst_size = 4; // bytes break; - case EXTP_OPTION: - break; - case NC_OPTION: - break; - case UP_OPTION: - case DOWN_OPTION: - case BFLY_OPTION: - case IDX_OPTION: + case EXTP_OPTION: + break; + case NC_OPTION: + break; + case UP_OPTION: + case DOWN_OPTION: + case BFLY_OPTION: + case IDX_OPTION: m_shfl_op = last_ptx_inst_option; break; + + case PRMT_F4E_MODE: + case PRMT_B4E_MODE: + case PRMT_RC8_MODE: + case PRMT_ECL_MODE: + case PRMT_ECR_MODE: + case PRMT_RC16_MODE: + m_prmt_op = last_ptx_inst_option; + break; default: assert(0); break; diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index cff312b..cb4556e 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1086,6 +1086,7 @@ public: unsigned dimension() const { return m_geom_spec;} unsigned barrier_op() const {return m_barrier_op;} unsigned shfl_op() const {return m_shfl_op;} + unsigned prmt_op() const {return m_prmt_op;} enum vote_mode_t { vote_any, vote_all, vote_uni, vote_ballot }; enum vote_mode_t vote_mode() const { return m_vote_mode; } @@ -1156,6 +1157,7 @@ private: unsigned m_saturation_mode; unsigned m_barrier_op; unsigned m_shfl_op; + unsigned m_prmt_op; std::list m_scalar_type; memory_space_t m_space_spec; diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index eb81961..6757091 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -39,7 +39,7 @@ void set_ptx_warp_size(const struct core_config * warp_size) g_shader_core_config=warp_size; } -static bool g_debug_ir_generation=true; +static bool g_debug_ir_generation=false; const char *g_filename; unsigned g_max_regs_per_thread = 0; -- cgit v1.3 From 12a8816c32a134693011d8b9e587f109e4d7e7f9 Mon Sep 17 00:00:00 2001 From: aamir Date: Thu, 9 Aug 2018 20:07:20 -0700 Subject: added load --- cuda-kernels/v4p_kernel.cu | 109 +++++++++++++++++++++++------- src/cuda-sim/cuda-sim.cc | 7 +- src/cuda-sim/instructions.cc | 155 +++++++++++++++++++++++++++++++++++++++++-- src/cuda-sim/opcodes.def | 3 + src/cuda-sim/opcodes.h | 4 ++ src/cuda-sim/ptx.l | 8 +++ src/cuda-sim/ptx.y | 6 ++ src/cuda-sim/ptx_ir.cc | 4 ++ src/cuda-sim/ptx_parser.cc | 2 +- 9 files changed, 266 insertions(+), 32 deletions(-) (limited to 'src/cuda-sim/ptx_ir.cc') diff --git a/cuda-kernels/v4p_kernel.cu b/cuda-kernels/v4p_kernel.cu index bb9064b..2e84eda 100644 --- a/cuda-kernels/v4p_kernel.cu +++ b/cuda-kernels/v4p_kernel.cu @@ -30,23 +30,30 @@ 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,float *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; - - // 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); - wmma::mma_sync(c_frag, a_frag, b_frag, c_frag); - - wmma::store_matrix_sync(d_fp16, c_frag, N, wmma::mem_col_major); - //printf("clock=%d",end_time-start_time); +__global__ void v4p_example(int *a_int32, int *b_int8, int *c,int *d_int32, int M, int N, int K) { + + int registers_a[8]; + int register_b; //contains 8 4bit b elements + int idx = blockDim.x * blockIdx.x + threadIdx.x; + + asm("/*"); + asm("CPTX_BEGIN"); + asm("vp.load.b4.sync.row.m16n16k16.s32 {%0,%1,%2,%3,%4,%5,%6,%7},[%8],%9;" : + "=r"(registers_a[0]), "=r"(registers_a[1]),"=r"(registers_a[2]),"=r"(registers_a[3]), + "=r"(registers_a[4]),"=r"(registers_a[5]),"=r"(registers_a[6]),"=r"(registers_a[7]): + "l"(b_int8),"r"(M) + ); + asm("CPTX_END"); + asm("*/"); + + b_int8[0]=registers_a[7]; + b_int8[1]=registers_a[7]; + b_int8[2]=registers_a[7]; + b_int8[3]=registers_a[7]; + b_int8[4]=registers_a[7]; + b_int8[5]=registers_a[7]; + b_int8[6]=registers_a[7]; + b_int8[7]=registers_a[7]; } __global__ void convertFp32ToFp16 (half *out, float *in, int n) { @@ -62,13 +69,35 @@ __global__ void convertFp16ToFp32 (float *out, half *in, int n) { } } +__global__ void convertInt32ToInt4 (int *out, int *in, int n) { + int idx = blockDim.x * blockIdx.x + threadIdx.x; + if (idx < n/8) { + out[idx] =(in[8*idx]&0xf)|(in[8*idx+1]&0xf)<<4|(in[8*idx+2]&0xf)<<8|(in[8*idx+3]&0xf)<<12| + (in[8*idx+4]&0xf)<<16|(in[8*idx+5]&0xf)<<20|(in[8*idx+6]&0xf)<<24|(in[8*idx+7]&0xf)<<28; +// printf("thread%d:%x\n",idx,out[idx]); + } +} __global__ void convertInt32ToInt8 (int *out, int *in, int n) { int idx = blockDim.x * blockIdx.x + threadIdx.x; if (idx < n/4) { out[idx] =(in[4*idx]&0xff)|(in[4*idx+1]&0xff)<<8|(in[4*idx+2]&0xff)<<16|(in[4*idx+3]&0xff)<<24; } } +__global__ void convertInt32ToInt16 (int *out, int *in, int n) { + int idx = blockDim.x * blockIdx.x + threadIdx.x; + if (idx < n/2) { + out[idx] =(in[2*idx]&0xffff)|(in[2*idx+1]&0xffff)<<16; + } +} +__global__ void convertInt4ToInt32 (int *out, int *in, int n) { + int idx = blockDim.x * blockIdx.x + threadIdx.x; + int shft_amt=4*(idx%8); + int shft_mask=0xf<>shft_amt; + } +} __global__ void convertInt8ToInt32 (int *out, int *in, int n) { int idx = blockDim.x * blockIdx.x + threadIdx.x; int shft_amt=8*(idx%4); @@ -77,6 +106,14 @@ __global__ void convertInt8ToInt32 (int *out, int *in, int n) { out[idx]= (in[idx/4]&shft_mask)>>shft_amt; } } +__global__ void convertInt16ToInt32 (int *out, int *in, int n) { + int idx = blockDim.x * blockIdx.x + threadIdx.x; + int shft_amt=16*(idx%2); + int shft_mask=0xffff<>shft_amt; + } +} int main(int argc, char* argv[]) { int *a_int32; @@ -84,8 +121,12 @@ int main(int argc, char* argv[]) { int *c_int32; int *d_int32; + int *a_int4; + int *b_int4; int *a_int8; int *b_int8; + int *a_int16; + int *b_int16; int *a_host_wmma; int *b_host_wmma; @@ -105,15 +146,19 @@ int main(int argc, char* argv[]) { cudaErrCheck(cudaMalloc((void**)&b_int32, MATRIX_K * MATRIX_N * sizeof(int))); cudaErrCheck(cudaMalloc((void**)&c_int32, MATRIX_K * MATRIX_N * sizeof(int))); cudaErrCheck(cudaMalloc((void**)&d_int32, MATRIX_K * MATRIX_N * sizeof(int))); + cudaErrCheck(cudaMalloc((void**)&a_int4, MATRIX_M * MATRIX_K * sizeof(int)/8)); + cudaErrCheck(cudaMalloc((void**)&b_int4, MATRIX_K * MATRIX_N * sizeof(int)/8)); cudaErrCheck(cudaMalloc((void**)&a_int8, MATRIX_M * MATRIX_K * sizeof(int)/4)); cudaErrCheck(cudaMalloc((void**)&b_int8, MATRIX_K * MATRIX_N * sizeof(int)/4)); + cudaErrCheck(cudaMalloc((void**)&a_int16, MATRIX_M * MATRIX_K * sizeof(int)/2)); + cudaErrCheck(cudaMalloc((void**)&b_int16, MATRIX_K * MATRIX_N * sizeof(int)/2)); a_host_wmma = (int *)malloc(MATRIX_M * MATRIX_K * sizeof(int)); b_host_wmma = (int *)malloc(MATRIX_K * MATRIX_N * sizeof(int)); c_host_wmma = (int *)malloc(MATRIX_M * MATRIX_N * sizeof(int)); d_host_wmma = (int *)malloc(MATRIX_M * MATRIX_N * sizeof(int)); - d_cal_host_wmma = (int *)malloc(MATRIX_M * MATRIX_N * sizeof(int)); + d_cal_host_wmma = (int *)malloc(MATRIX_M * MATRIX_N * sizeof(int)); printf("a_int32\n"); for(int m=0;m>> (a_int8, a_int32, MATRIX_M * MATRIX_K); - convertInt8ToInt32 <<< (MATRIX_M * MATRIX_K + 255) / 256, 256 >>> (d_int32, a_int8, MATRIX_M * MATRIX_K); + #define TEST8 + #ifdef TEST16 + convertInt32ToInt16 <<< (MATRIX_M * MATRIX_K + 255) / 256, 256 >>> (a_int16, a_int32, MATRIX_M * MATRIX_K); + convertInt16ToInt32 <<< (MATRIX_M * MATRIX_K + 255) / 256, 256 >>> (d_int32, a_int16, MATRIX_M * MATRIX_K); + cudaErrCheck(cudaMemcpy(d_host_wmma, d_int32, MATRIX_M * MATRIX_N * sizeof(int), cudaMemcpyDeviceToHost)); + #endif + #ifdef TEST8 + convertInt32ToInt8 <<< (MATRIX_M * MATRIX_K + 255) / 256, 256 >>> (a_int8, a_int32, MATRIX_M * MATRIX_K); + convertInt8ToInt32 <<< (MATRIX_M * MATRIX_K + 255) / 256, 256 >>> (d_int32, a_int8, MATRIX_M * MATRIX_K); + cudaErrCheck(cudaMemcpy(d_host_wmma, d_int32, MATRIX_M * MATRIX_N * sizeof(int), cudaMemcpyDeviceToHost)); + #endif + #ifdef TEST4 + convertInt32ToInt4 <<< (MATRIX_M * MATRIX_K + 255) / 256, 256 >>> (a_int4, a_int32, MATRIX_M * MATRIX_K); + convertInt4ToInt32 <<< (MATRIX_M * MATRIX_K + 255) / 256, 256 >>> (d_int32, a_int4, MATRIX_M * MATRIX_K); + cudaErrCheck(cudaMemcpy(d_host_wmma, d_int32, MATRIX_M * MATRIX_N * sizeof(int), cudaMemcpyDeviceToHost)); + #endif //convertFp32ToFp16 <<< (MATRIX_K * MATRIX_N + 255) / 256, 256 >>> (b_fp16, b_fp32, MATRIX_K * MATRIX_N); //convertFp32ToFp16 <<< (MATRIX_M * MATRIX_N + 255) / 256, 256 >>> (c_fp16, c_fp32, MATRIX_K * MATRIX_N); - cudaErrCheck(cudaMemcpy(d_host_wmma, d_int32, MATRIX_M * MATRIX_N * sizeof(float), cudaMemcpyDeviceToHost)); //AAMIR printf("\nM = %d, N = %d, K = %d. \n", MATRIX_M, MATRIX_N, MATRIX_K); //AAMIR //AAMIR printf("Running with wmma...\n"); -//AAMIR cudaErrCheck(cudaEventRecord(startWMMA)); -//AAMIR wmma_example <<< 1, 32>>> (a_fp16, b_fp16, c_fp32, d_fp32 , MATRIX_M, MATRIX_N, MATRIX_K); -//AAMIR cudaErrCheck(cudaEventRecord(stopWMMA)); -//AAMIR cudaErrCheck(cudaEventSynchronize(stopWMMA)); + //cudaErrCheck(cudaEventRecord(startWMMA)); + //v4p_example <<< 1, 32>>> (a_int32, a_int8, a_int32, d_int32, MATRIX_M, MATRIX_N, MATRIX_K); + //cudaErrCheck(cudaEventRecord(stopWMMA)); + //cudaErrCheck(cudaEventSynchronize(stopWMMA)); + //AAMIR //AAMIR // Error checking //AAMIR printf("\nChecking results...\n"); diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 8284ad5..2fe5667 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1323,8 +1323,11 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) *((warp_inst_t*)pJ) = inst; // copy active mask information pI = pJ; } - if(((pI->get_opcode()!=MMA_OP)&&(pI->get_opcode()!=MMA_LD_OP)&&(pI->get_opcode()!=MMA_ST_OP))||((pI->get_opcode()==MMA_OP||pI->get_opcode()==MMA_LD_OP||pI->get_opcode()==MMA_ST_OP)&&(lane_id==0))){ - switch ( pI->get_opcode() ) { + + int inst_opcode=pI->get_opcode(); + + if(((inst_opcode!=MMA_OP)&&(inst_opcode!=MMA_LD_OP)&&(inst_opcode!=MMA_ST_OP)&&(inst_opcode!=VP_LD_OP)&&(inst_opcode!=VP_ST_OP)&&(inst_opcode!=VP_MMA_OP))||((inst_opcode==MMA_OP||inst_opcode==MMA_LD_OP||inst_opcode==MMA_ST_OP||inst_opcode==VP_MMA_OP||inst_opcode==VP_LD_OP||inst_opcode==VP_ST_OP)&&(lane_id==0))){ + switch ( inst_opcode ) { #define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,this); op_classification = CLASSIFICATION; break; #define OP_W_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,get_core(),inst); op_classification = CLASSIFICATION; break; #include "opcodes.def" diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index aaee2a2..d0396a5 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -48,7 +48,7 @@ using half_float::half; unsigned ptx_instruction::g_num_ptx_inst_uid=0; -bool g_debug_instruction = 0; +bool g_debug_instruction = 1; const char *g_opcode_string[NUM_OPCODES] = { @@ -129,6 +129,12 @@ unsigned thread_group_offset(int thread,unsigned wmma_type,unsigned wmma_layout } } break; + case VP_MMA: + if(wmma_layout==ROW) + offset=load_c_float_row[thread_group]+16*in_tg_index; + else + offset=load_c_float_col[thread_group]+16*in_tg_index; + break; default: abort(); @@ -1709,6 +1715,14 @@ void mapping(int thread,int wmma_type,int wmma_layout,int type,int index,int str } } +void vp_mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) +{ + unsigned wmma_type = pI->get_wmma_type(); + unsigned a_layout = pI->get_wmma_layout(0); + unsigned b_layout = pI->get_wmma_layout(1); + + +} void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) { int i,j,k,thrd; @@ -1752,7 +1766,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) int hex_val; if(!((i==3)&&(type2==F32_TYPE))){ - for(k=0;k<2*nelem;k++){ + for(k=0;k<2*nelem;k++){ if(k%2==1) hex_val=(v[k/2].s64&0xffff); else @@ -2958,6 +2972,12 @@ void ldu_impl( const ptx_instruction *pI, ptx_thread_info *thread ) ld_exec(pI,thread); } +void vp_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) +{ + const operand_info &src = pI->operand_lookup(1); + const operand_info &src1 = pI->operand_lookup(0); + const operand_info &src2 = pI->operand_lookup(2); +} void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) { size_t size; @@ -3000,7 +3020,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) type_info_key::type_decode(type,size,t); if(g_debug_instruction) - printf("mma_st: thrd=%d,addr=%d, fp(size=%d), stride=%d\n",thrd,addr_reg.u32,size,src2_data.u32); + printf("mma_st: thrd=%d,addr=%x, fp(size=%d), stride=%d\n",thrd,addr_reg.u32,size,src2_data.u32); addr_t new_addr = addr+thread_group_offset(thrd,wmma_type,wmma_layout,type,stride)*size/8; ptx_reg_t nw_v[8]; @@ -3041,6 +3061,133 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) thread->m_last_effective_address = addr; thread->m_last_memory_space = space; } +} +void vp_ld_impl(const ptx_instruction *pI, core_t *core, warp_inst_t inst) +{ + size_t size; + int t,i; + unsigned smid; + const operand_info &dst = pI->dst(); + const operand_info &src1 = pI->src1(); + const operand_info &src2 = pI->src2(); + unsigned type = pI->get_type(); + unsigned wmma_type = pI->get_wmma_type(); + unsigned wmma_layout = pI->get_wmma_layout(0); + int tid = inst.warp_id_func()*core->get_warp_size(); + int thrd,stride; + ptx_thread_info *thread; + for (thrd=0; thrd < core->get_warp_size(); thrd++){ + thread = core->get_thread_info()[tid+thrd]; + ptx_reg_t src1_data = thread->get_operand_value(src1, dst, U32_TYPE, thread, 1); + ptx_reg_t src2_data = thread->get_operand_value(src2, dst, U32_TYPE, thread, 1); + stride=src2_data.u32; + memory_space_t space = pI->get_space(); + + memory_space *mem = NULL; + addr_t addr = src1_data.u32; + + smid = thread->get_hw_sid(); + if( whichspace(addr) == shared_space ) { + addr= generic_to_shared(smid,addr); + space = shared_space; + } + + decode_space(space,thread,src1,mem,addr); + type_info_key::type_decode(type,size,t); + + ptx_reg_t data[8]; + addr_t new_addr; + //note we are using distribution of VP_MMA for every type of load! + if(wmma_type==LOAD_A||wmma_type==LOAD_C){ + new_addr = addr+thread_group_offset(thrd,VP_MMA,wmma_layout,type,stride)*size/8; + } + else if(wmma_type==LOAD_B4){ + new_addr = addr+thread_group_offset(thrd,VP_MMA,wmma_layout,type,stride)/8; + } + else if (wmma_type==LOAD_B8){ + new_addr = addr+thread_group_offset(thrd,VP_MMA,wmma_layout,type,stride)/4; + } + else if (wmma_type==LOAD_B16){ + new_addr = addr+thread_group_offset(thrd,VP_MMA,wmma_layout,type,stride)/2; + } + + if(g_debug_instruction) + printf("vp_ld: thrx=%d,addr=%x, base_addr=%x, size=%d, stride=%d\n",thrd,new_addr,addr,size,src2_data.u32); + + if(wmma_type==LOAD_A||wmma_type==LOAD_C){ + for(i=0;i<8;i++){ + mem->read(new_addr+4*i,size/8,&data[i].s64); + } + } + else if(wmma_type==LOAD_B4){ + mem->read(new_addr,size/8,&data[0].s64); + } + else if(wmma_type==LOAD_B8){ + mem->read(new_addr,size/8,&data[0].s64); + mem->read(new_addr+4,size/8,&data[1].s64); + } + else if(wmma_type==LOAD_B16){ + mem->read(new_addr,size/8,&data[0].s64); + mem->read(new_addr+4,size/8,&data[1].s64); + mem->read(new_addr+8,size/8,&data[2].s64); + mem->read(new_addr+12,size/8,&data[3].s64); + } + else{ + printf("wrong vp_load type\n");; + abort(); + } + + int num_reg; + if(g_debug_instruction){ + printf("\nvp_ld:thread%d= ",thrd); + if((wmma_type==LOAD_A)||(wmma_type==LOAD_C)){ + num_reg=8; + } + else if(wmma_type==LOAD_B4){ + num_reg=1; + } + else if(wmma_type==LOAD_B8){ + num_reg=2; + } + else if(wmma_type==LOAD_B16){ + num_reg=4; + } + + for(i=0;iset_wmma_vector_operand_values(dst,data[0],data[1],data[2],data[3],data[4],data[5],data[6],data[7]); + } + else if(wmma_type==LOAD_B4){ + thread->set_operand_value(dst,data[0], type, thread, pI); + } + else if(wmma_type==LOAD_B8){ + thread->set_vector_operand_values(dst,data[0],data[1],data[1],data[1]); + } + else if (wmma_type==LOAD_B16){ + thread->set_vector_operand_values(dst,data[0],data[1],data[2],data[3]); + } + else + abort(); + + + + if(g_debug_instruction){ + for(int i=0;im_last_effective_address = addr; + thread->m_last_memory_space = space; + } + + + } void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) @@ -3080,7 +3227,7 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) ptx_reg_t data[16]; if(g_debug_instruction) - printf("mma_ld: thrd=%d,addr=%d, fpsize=%d, stride=%d\n",thrd,src1_data.u32,size,src2_data.u32); + printf("mma_ld: thrd=%d,addr=%x, fpsize=%d, stride=%d\n",thrd,src1_data.u32,size,src2_data.u32); addr_t new_addr = addr+thread_group_offset(thrd,wmma_type,wmma_layout,type,stride)*size/8; diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def index e6f957a..73f6161 100644 --- a/src/cuda-sim/opcodes.def +++ b/src/cuda-sim/opcodes.def @@ -55,6 +55,9 @@ OP_DEF(BRKPT_OP,brkpt_impl,"brkpt",1,9) OP_W_DEF(MMA_OP,mma_impl,"mma",1,1) OP_W_DEF(MMA_LD_OP,mma_ld_impl,"mma_load",1,5) OP_W_DEF(MMA_ST_OP,mma_st_impl,"mma_store",0,5) +OP_W_DEF(VP_MMA_OP,vp_mma_impl,"vp_mma",1,1) +OP_W_DEF(VP_LD_OP,vp_ld_impl,"vp_load",1,5) +OP_W_DEF(VP_ST_OP,vp_st_impl,"vp_store",0,5) OP_DEF(CALL_OP,call_impl,"call",1,3) OP_DEF(CALLP_OP,callp_impl,"callp",1,3) OP_DEF(CLZ_OP,clz_impl,"clz",1,1) diff --git a/src/cuda-sim/opcodes.h b/src/cuda-sim/opcodes.h index b91d92f..31b71d0 100644 --- a/src/cuda-sim/opcodes.h +++ b/src/cuda-sim/opcodes.h @@ -63,9 +63,13 @@ enum special_regs { enum wmma_type{ LOAD_A, LOAD_B, + LOAD_B4,//vp + LOAD_B8,//vp + LOAD_B16,//vp LOAD_C, STORE_D, MMA, + VP_MMA, ROW, COL, M16N16K16 diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index a6b6fcc..fea2420 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -68,9 +68,13 @@ bra TC; ptx_lval.int_value = BRA_OP; return OPCODE; brx TC; ptx_lval.int_value = BRX_OP; return OPCODE; brev TC; ptx_lval.int_value = BREV_OP; return OPCODE; brkpt TC; ptx_lval.int_value = BRKPT_OP; return OPCODE; + wmma TC; ptx_lval.int_value = MMA_OP; return OPCODE; wmma\.load TC; ptx_lval.int_value = MMA_LD_OP; return OPCODE; wmma\.store TC; ptx_lval.int_value = MMA_ST_OP; return OPCODE; +vp\.load TC; ptx_lval.int_value=VP_LD_OP; return OPCODE; +vp\.store TC; ptx_lval.int_value=VP_ST_OP; return OPCODE; +vp TC; ptx_lval.int_value=VP_MMA_OP; return OPCODE; call TC; BEGIN(NOT_OPCODE); ptx_lval.int_value = CALL_OP; return OPCODE; // blocking opcode token in case the callee has the same name as an opcode callp TC; BEGIN(NOT_OPCODE); ptx_lval.int_value = CALLP_OP; return OPCODE; @@ -161,6 +165,10 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; \.c\.sync TC; ptx_lval.int_value = LOAD_C; return WMMA_DIRECTIVE; \.d\.sync TC; ptx_lval.int_value = STORE_D; return WMMA_DIRECTIVE; \.mma\.sync TC;ptx_lval.int_value=MMA; return WMMA_DIRECTIVE; +\.b4\.sync TC; ptx_lval.int_value=LOAD_B4; return WMMA_DIRECTIVE; +\.b8\.sync TC; ptx_lval.int_value=LOAD_B8; return WMMA_DIRECTIVE; +\.b16\.sync TC; ptx_lval.int_value=LOAD_B16; return WMMA_DIRECTIVE; + \.row TC; ptx_lval.int_value = ROW; return LAYOUT; \.col TC; ptx_lval.int_value = COL; return LAYOUT; \.m16n16k16 TC; ptx_lval.int_value = M16N16K16; return CONFIGURATION; diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y index cd455dd..54ac2bc 100644 --- a/src/cuda-sim/ptx.y +++ b/src/cuda-sim/ptx.y @@ -538,6 +538,12 @@ wmma_spec: WMMA_DIRECTIVE LAYOUT CONFIGURATION{add_space_spec(global_space,0);ad | WMMA_DIRECTIVE LAYOUT LAYOUT CONFIGURATION{add_wmma_option($1);add_wmma_option($2);add_wmma_option($3);add_wmma_option($4);} ; +vp_spec: WMMA_DIRECTIVE LAYOUT CONFIGURATION{add_space_spec(global_space,0);add_ptr_spec(global_space);add_wmma_option($1);add_wmma_option($2);add_wmma_option($3);} + | WMMA_DIRECTIVE LAYOUT LAYOUT CONFIGURATION{add_wmma_option($1);add_wmma_option($2);add_wmma_option($3);add_wmma_option($4);} + ; + + + operand_list: operand | operand COMMA operand_list; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index d12c741..55b8e11 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1089,9 +1089,13 @@ ptx_instruction::ptx_instruction( int opcode, case SYNC_OPTION: case LOAD_A: case LOAD_B: + case LOAD_B4: + case LOAD_B8: + case LOAD_B16: case LOAD_C: case STORE_D: case MMA: + case VP_MMA: m_wmma_type=last_ptx_inst_option; break; case ROW: diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 6757091..eb81961 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -39,7 +39,7 @@ void set_ptx_warp_size(const struct core_config * warp_size) g_shader_core_config=warp_size; } -static bool g_debug_ir_generation=false; +static bool g_debug_ir_generation=true; const char *g_filename; unsigned g_max_regs_per_thread = 0; -- cgit v1.3 From cdd42ac0384cb31bcee05ac72d026f1bc8133d26 Mon Sep 17 00:00:00 2001 From: aamir Date: Tue, 14 Aug 2018 17:46:21 -0700 Subject: parsing changes for timing model --- cuda-kernels/v16p_genericMatrixMultiply.cu | 8 ++++---- cuda-kernels/v16p_kernel.cu | 2 +- cuda-kernels/v4p_genericMatrixMultiply.cu | 8 ++++---- cuda-kernels/v4p_kernel.cu | 2 +- cuda-kernels/v8p_genericMatrixMultiply.cu | 2 +- cuda-kernels/v8p_kernel.cu | 2 +- src/cuda-sim/instructions.cc | 13 ++++++++----- src/cuda-sim/opcodes.h | 3 +++ src/cuda-sim/ptx.l | 3 +++ src/cuda-sim/ptx_ir.cc | 4 +++- 10 files changed, 29 insertions(+), 18 deletions(-) (limited to 'src/cuda-sim/ptx_ir.cc') diff --git a/cuda-kernels/v16p_genericMatrixMultiply.cu b/cuda-kernels/v16p_genericMatrixMultiply.cu index fd5a0f8..c36d257 100644 --- a/cuda-kernels/v16p_genericMatrixMultiply.cu +++ b/cuda-kernels/v16p_genericMatrixMultiply.cu @@ -20,9 +20,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 (256) -#define MATRIX_N (256) -#define MATRIX_K (256) +#define MATRIX_M (32) +#define MATRIX_N (32) +#define MATRIX_K (32) // The only dimensions currently supported by WMMA @@ -92,7 +92,7 @@ __global__ void vp_example(int *a, int *b, int *c, int M, int N, int K ) { //vp::mma_sync(acc_frag, a_frag, b_frag, acc_frag); asm("/*"); asm("CPTX_BEGIN"); - asm("vp.mma.sync.row.row.m16n16k16.s32 {%0, %1, %2, %3, %4, %5, %6, %7}, {%8, %9, %10, %11, %12, %13, %14, %15}, {%16, %17, %18, %19}, { %20, %21, %22, %23, %24, %25, %26,%27};" : + asm("vp.mma16.sync.row.row.m16n16k16.s32 {%0, %1, %2, %3, %4, %5, %6, %7}, {%8, %9, %10, %11, %12, %13, %14, %15}, {%16, %17, %18, %19}, { %20, %21, %22, %23, %24, %25, %26,%27};" : "=r"(acc_frag[0]), "=r"(acc_frag[1]),"=r"(acc_frag[2]),"=r"(acc_frag[3]), "=r"(acc_frag[4]),"=r"(acc_frag[5]),"=r"(acc_frag[6]),"=r"(acc_frag[7]): "r"(a_frag[0]),"r"(a_frag[1]),"r"(a_frag[2]),"r"(a_frag[3]), diff --git a/cuda-kernels/v16p_kernel.cu b/cuda-kernels/v16p_kernel.cu index 011fdfd..31a1460 100644 --- a/cuda-kernels/v16p_kernel.cu +++ b/cuda-kernels/v16p_kernel.cu @@ -100,7 +100,7 @@ __global__ void v4p_example(int *a_int32, int *b_int4, int *c,int *d_int32, int //B16 asm("/*"); asm("CPTX_BEGIN"); - asm("vp.mma.sync.row.row.m16n16k16.s32 {%0, %1, %2, %3, %4, %5, %6, %7}, {%8, %9, %10, %11, %12, %13, %14, %15}, {%16, %17, %18, %19}, { %20, %21, %22, %23, %24, %25, %26, %27};" : + asm("vp.mma16.sync.row.row.m16n16k16.s32 {%0, %1, %2, %3, %4, %5, %6, %7}, {%8, %9, %10, %11, %12, %13, %14, %15}, {%16, %17, %18, %19}, { %20, %21, %22, %23, %24, %25, %26, %27};" : "=r"(registers_d[0]), "=r"(registers_d[1]),"=r"(registers_d[2]),"=r"(registers_d[3]), "=r"(registers_d[4]),"=r"(registers_d[5]),"=r"(registers_d[6]),"=r"(registers_d[7]): "r"(registers_a[0]),"r"(registers_a[1]),"r"(registers_a[2]),"r"(registers_a[3]), diff --git a/cuda-kernels/v4p_genericMatrixMultiply.cu b/cuda-kernels/v4p_genericMatrixMultiply.cu index 1b56eb2..abcab8e 100644 --- a/cuda-kernels/v4p_genericMatrixMultiply.cu +++ b/cuda-kernels/v4p_genericMatrixMultiply.cu @@ -20,9 +20,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 (1024) -#define MATRIX_N (1024) -#define MATRIX_K (1024) +#define MATRIX_M (32) +#define MATRIX_N (32) +#define MATRIX_K (32) // The only dimensions currently supported by WMMA @@ -92,7 +92,7 @@ __global__ void vp_example(int *a, int *b, int *c, int M, int N, int K ) { //vp::mma_sync(acc_frag, a_frag, b_frag, acc_frag); asm("/*"); asm("CPTX_BEGIN"); - asm("vp.mma.sync.row.row.m16n16k16.s32 {%0, %1, %2, %3, %4, %5, %6, %7}, {%8, %9, %10, %11, %12, %13, %14, %15}, {%16}, {%17, %18, %19, %20, %21, %22, %23, %24};" : + asm("vp.mma4.sync.row.row.m16n16k16.s32 {%0, %1, %2, %3, %4, %5, %6, %7}, {%8, %9, %10, %11, %12, %13, %14, %15}, {%16}, {%17, %18, %19, %20, %21, %22, %23, %24};" : "=r"(acc_frag[0]), "=r"(acc_frag[1]),"=r"(acc_frag[2]),"=r"(acc_frag[3]), "=r"(acc_frag[4]),"=r"(acc_frag[5]),"=r"(acc_frag[6]),"=r"(acc_frag[7]): "r"(a_frag[0]),"r"(a_frag[1]),"r"(a_frag[2]),"r"(a_frag[3]), diff --git a/cuda-kernels/v4p_kernel.cu b/cuda-kernels/v4p_kernel.cu index 8c9bbd4..f1a9dbe 100644 --- a/cuda-kernels/v4p_kernel.cu +++ b/cuda-kernels/v4p_kernel.cu @@ -70,7 +70,7 @@ __global__ void v4p_example(int *a_int32, int *b_int4, int *c,int *d_int32, int //B4 asm("/*"); asm("CPTX_BEGIN"); - asm("vp.mma.sync.row.row.m16n16k16.s32 {%0, %1, %2, %3, %4, %5, %6, %7}, {%8, %9, %10, %11, %12, %13, %14, %15}, {%16}, {%17, %18, %19, %20, %21, %22, %23, %24};" : + asm("vp.mma4.sync.row.row.m16n16k16.s32 {%0, %1, %2, %3, %4, %5, %6, %7}, {%8, %9, %10, %11, %12, %13, %14, %15}, {%16}, {%17, %18, %19, %20, %21, %22, %23, %24};" : "=r"(registers_d[0]), "=r"(registers_d[1]),"=r"(registers_d[2]),"=r"(registers_d[3]), "=r"(registers_d[4]),"=r"(registers_d[5]),"=r"(registers_d[6]),"=r"(registers_d[7]): "r"(registers_a[0]),"r"(registers_a[1]),"r"(registers_a[2]),"r"(registers_a[3]), diff --git a/cuda-kernels/v8p_genericMatrixMultiply.cu b/cuda-kernels/v8p_genericMatrixMultiply.cu index 2e487e8..208c369 100644 --- a/cuda-kernels/v8p_genericMatrixMultiply.cu +++ b/cuda-kernels/v8p_genericMatrixMultiply.cu @@ -92,7 +92,7 @@ __global__ void vp_example(int *a, int *b, int *c, int M, int N, int K ) { //vp::mma_sync(acc_frag, a_frag, b_frag, acc_frag); asm("/*"); asm("CPTX_BEGIN"); - asm("vp.mma.sync.row.row.m16n16k16.s32 {%0, %1, %2, %3, %4, %5, %6, %7}, {%8, %9, %10, %11, %12, %13, %14, %15}, {%16, %17}, {%18, %19, %20, %21, %22, %23, %24, %25};" : + asm("vp.mma8.sync.row.row.m16n16k16.s32 {%0, %1, %2, %3, %4, %5, %6, %7}, {%8, %9, %10, %11, %12, %13, %14, %15}, {%16, %17}, {%18, %19, %20, %21, %22, %23, %24, %25};" : "=r"(acc_frag[0]), "=r"(acc_frag[1]),"=r"(acc_frag[2]),"=r"(acc_frag[3]), "=r"(acc_frag[4]),"=r"(acc_frag[5]),"=r"(acc_frag[6]),"=r"(acc_frag[7]): "r"(a_frag[0]),"r"(a_frag[1]),"r"(a_frag[2]),"r"(a_frag[3]), diff --git a/cuda-kernels/v8p_kernel.cu b/cuda-kernels/v8p_kernel.cu index b1b0eba..f824eac 100644 --- a/cuda-kernels/v8p_kernel.cu +++ b/cuda-kernels/v8p_kernel.cu @@ -85,7 +85,7 @@ __global__ void v4p_example(int *a_int32, int *b_int4, int *c,int *d_int32, int //B8 asm("/*"); asm("CPTX_BEGIN"); - asm("vp.mma.sync.row.row.m16n16k16.s32 {%0, %1, %2, %3, %4, %5, %6, %7}, {%8, %9, %10, %11, %12, %13, %14, %15}, {%16, %17}, {%18, %19, %20, %21, %22, %23, %24, %25};" : + asm("vp.mma8.sync.row.row.m16n16k16.s32 {%0, %1, %2, %3, %4, %5, %6, %7}, {%8, %9, %10, %11, %12, %13, %14, %15}, {%16, %17}, {%18, %19, %20, %21, %22, %23, %24, %25};" : "=r"(registers_d[0]), "=r"(registers_d[1]),"=r"(registers_d[2]),"=r"(registers_d[3]), "=r"(registers_d[4]),"=r"(registers_d[5]),"=r"(registers_d[6]),"=r"(registers_d[7]): "r"(registers_a[0]),"r"(registers_a[1]),"r"(registers_a[2]),"r"(registers_a[3]), diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index b26adfd..39b8ba5 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -130,6 +130,9 @@ unsigned thread_group_offset(int thread,unsigned wmma_type,unsigned wmma_layout } break; case VP_MMA: + case VP_MMA4: + case VP_MMA8: + case VP_MMA16: if(wmma_layout==ROW) offset=load_c_float_row[thread_group]+16*in_tg_index; else @@ -1703,7 +1706,7 @@ void mapping(int thread,int wmma_type,int wmma_layout,int type,int index,int str } assg_offset=index; } - else if( wmma_type==VP_MMA){ + else if(wmma_type==VP_MMA || wmma_type==VP_MMA4 || wmma_type==VP_MMA8 || wmma_type==VP_MMA16){ row=c_row_offset[thread/4]+thread%4; col=c_col_offset[thread/4]+index; assg_offset=index; @@ -1761,7 +1764,7 @@ void vp_mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) switch(i) { case 1 ://operand 1 for(k=0;k<8;k++){ - mapping(thrd,VP_MMA,a_layout,S32_TYPE,k,16,row,col,offset); + mapping(thrd,wmma_type,a_layout,S32_TYPE,k,16,row,col,offset); if(g_debug_instruction) printf("A:thread=%d,row=%d,col=%d,offset=%d\n",thrd,row,col,offset); matrix_a[row][col]=v[offset]; @@ -1769,7 +1772,7 @@ void vp_mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) break; case 2 ://operand 2 for(k=0;k<8;k++){ - mapping(thrd,VP_MMA,b_layout,S32_TYPE,k,16,row,col,offset); + mapping(thrd,wmma_type,b_layout,S32_TYPE,k,16,row,col,offset); if(g_debug_instruction) printf("B:thread=%d,row=%d,col=%d,offset=%d\n",thrd,row,col,offset); if(nelem==1){ @@ -1785,7 +1788,7 @@ void vp_mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) break; case 3 ://operand 3 for(k=0;k<8;k++){ - mapping(thrd,VP_MMA,ROW,S32_TYPE,k,16,row,col,offset); + mapping(thrd,wmma_type,ROW,S32_TYPE,k,16,row,col,offset); if(g_debug_instruction) printf("C:thread=%d,row=%d,col=%d,offset=%d\n",thrd,row,col,offset); matrix_c[row][col]=v[offset]; @@ -1848,7 +1851,7 @@ void vp_mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) int row_t[8]; int col_t[8]; for(k=0;k<8;k++){ - mapping(thrd,VP_MMA,ROW,type,k,16,row_t[k],col_t[k],offset); + mapping(thrd,wmma_type,ROW,type,k,16,row_t[k],col_t[k],offset); if(g_debug_instruction) printf("vp_mma:row:%d,col%d\n",row_t[k],col_t[k]); } diff --git a/src/cuda-sim/opcodes.h b/src/cuda-sim/opcodes.h index 31b71d0..ad8d8f1 100644 --- a/src/cuda-sim/opcodes.h +++ b/src/cuda-sim/opcodes.h @@ -70,6 +70,9 @@ enum wmma_type{ STORE_D, MMA, VP_MMA, + VP_MMA4, + VP_MMA8, + VP_MMA16, ROW, COL, M16N16K16 diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index fea2420..9c78e9f 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -165,6 +165,9 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; \.c\.sync TC; ptx_lval.int_value = LOAD_C; return WMMA_DIRECTIVE; \.d\.sync TC; ptx_lval.int_value = STORE_D; return WMMA_DIRECTIVE; \.mma\.sync TC;ptx_lval.int_value=MMA; return WMMA_DIRECTIVE; +\.mma4\.sync TC;ptx_lval.int_value=VP_MMA4; return WMMA_DIRECTIVE; +\.mma8\.sync TC;ptx_lval.int_value=VP_MMA8; return WMMA_DIRECTIVE; +\.mma16\.sync TC;ptx_lval.int_value=VP_MMA16; return WMMA_DIRECTIVE; \.b4\.sync TC; ptx_lval.int_value=LOAD_B4; return WMMA_DIRECTIVE; \.b8\.sync TC; ptx_lval.int_value=LOAD_B8; return WMMA_DIRECTIVE; \.b16\.sync TC; ptx_lval.int_value=LOAD_B16; return WMMA_DIRECTIVE; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 55b8e11..9800e1e 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1095,7 +1095,9 @@ ptx_instruction::ptx_instruction( int opcode, case LOAD_C: case STORE_D: case MMA: - case VP_MMA: + case VP_MMA4: + case VP_MMA8: + case VP_MMA16: m_wmma_type=last_ptx_inst_option; break; case ROW: -- cgit v1.3