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/ptx_ir.cc | 2 ++ 1 file changed, 2 insertions(+) (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 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 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/ptx_ir.cc') 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/ptx_ir.cc') 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