From 556191bcbe6cfb2a8a1008166aa51b7a8567be82 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Mon, 10 Jun 2019 11:29:38 -0400 Subject: adding new wmma instruction cogif --- src/cuda-sim/opcodes.h | 6 +++++- src/cuda-sim/ptx.l | 2 ++ src/cuda-sim/ptx_ir.cc | 2 ++ 3 files changed, 9 insertions(+), 1 deletion(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/opcodes.h b/src/cuda-sim/opcodes.h index b91d92f..86d3b99 100644 --- a/src/cuda-sim/opcodes.h +++ b/src/cuda-sim/opcodes.h @@ -68,6 +68,10 @@ enum wmma_type{ MMA, ROW, COL, - M16N16K16 + M16N16K16, + M32N8K16, + M8N32K16 + + }; #endif diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index 3232361..36a8a4d 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -169,6 +169,8 @@ 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; +\.m32n8k16 TC; ptx_lval.int_value = M32N8K16; return CONFIGURATION; +\.m8n32k16 TC; ptx_lval.int_value = M8N32K16; return CONFIGURATION; \.f4e TC; return PRMT_F4E_MODE; \.b4e TC; return PRMT_B4E_MODE; \.rc8 TC; return PRMT_RC8_MODE; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 73db24d..4ad9ddf 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1144,6 +1144,8 @@ ptx_instruction::ptx_instruction( int opcode, m_wmma_layout[rr++]=last_ptx_inst_option; break; case M16N16K16: + case M32N8K16: + case M8N32K16: break; default: assert(0); -- cgit v1.3 From 8b276989af2f82ffc51306016409e01156678273 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Tue, 18 Jun 2019 19:23:06 -0400 Subject: fixing shuffle inst bug --- src/cuda-sim/cuda-sim.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index f7bb9cc..4c0fc58 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -873,7 +873,7 @@ void ptx_instruction::set_opcode_and_latency() op=TENSOR_CORE_OP; break; case SHFL_OP: - latency = 32; + latency = 4; initiation_interval = 4; break; default: -- cgit v1.3 From f23a5de2c7eec680fc8f5c6ba45c64fcd9544e65 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Sun, 23 Jun 2019 15:22:54 -0400 Subject: fixing the buffer limit for function names --- libcuda/cuda_runtime_api.cc | 4 ++++ src/cuda-sim/ptx_ir.cc | 4 ++-- 2 files changed, 6 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 17a5c96..97396ac 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -987,6 +987,10 @@ cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int * context->get_device()->get_gpgpu()->get_config().num_shader()); dim3 blockDim(blockSize); kernel_info_t result(gridDim, blockDim, entry); + //if(entry == NULL){ + // *numBlocks = 1; + // return g_last_cudaError = cudaErrorUnknown; + //} *numBlocks = context->get_device()->get_gpgpu()->get_max_cta(result); printf("Maximum size is %d with gridDim %d and blockDim %d\n", *numBlocks, gridDim.x, blockDim.x); return g_last_cudaError = cudaSuccess; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 4ad9ddf..c4d5a6c 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -177,8 +177,8 @@ void symbol_table::add_function( function_info *func, const char *filename, unsi //Jin: handle instruction group for cdp symbol_table* symbol_table::start_inst_group() { - char inst_group_name[1024]; - snprintf(inst_group_name, 1024, "%s_inst_group_%u", m_scope_name.c_str(), m_inst_group_id); + char inst_group_name[2048]; + snprintf(inst_group_name, 2048, "%s_inst_group_%u", m_scope_name.c_str(), m_inst_group_id); //previous added assert(m_inst_group_symtab.find(std::string(inst_group_name)) == m_inst_group_symtab.end()); -- cgit v1.3 From c9fb34c320a5adbd9649a410b417be70937345ad Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Thu, 27 Jun 2019 18:40:57 -0400 Subject: increase the function buffer again --- src/cuda-sim/ptx_ir.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim') diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index c4d5a6c..e4e0c09 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -177,8 +177,8 @@ void symbol_table::add_function( function_info *func, const char *filename, unsi //Jin: handle instruction group for cdp symbol_table* symbol_table::start_inst_group() { - char inst_group_name[2048]; - snprintf(inst_group_name, 2048, "%s_inst_group_%u", m_scope_name.c_str(), m_inst_group_id); + char inst_group_name[4096]; + snprintf(inst_group_name, 4096, "%s_inst_group_%u", m_scope_name.c_str(), m_inst_group_id); //previous added assert(m_inst_group_symtab.find(std::string(inst_group_name)) == m_inst_group_symtab.end()); -- cgit v1.3