From 46aad91327a265c2fea2cfe629cc38eadb629200 Mon Sep 17 00:00:00 2001 From: speverel Date: Thu, 2 Jun 2016 11:28:15 -0700 Subject: Added handling of .cc option for arithmetic instructions. NOTE: Only made changes to parse instructions. Carry functionality NOT fully implemented; .cc instructions function like their unmodified ueqivelents. Also modified GTX750Ti config to model L1 data cache as simply not being used for global loads (instead of not existing at all). Changed ptxinfo parsing to avoid crashing when info includes texture information. --- src/cuda-sim/instructions.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'src/cuda-sim/instructions.cc') diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index cf7f04a..4dd5ed8 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -2421,7 +2421,7 @@ void mad_impl( const ptx_instruction *pI, ptx_thread_info *thread ) mad_def(pI, thread, false); } -void madp_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +void madc_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { mad_def(pI, thread, true); } -- cgit v1.3 From 8ed1522aede92bcafc60f69d7e03e6d48c44a86c Mon Sep 17 00:00:00 2001 From: speverel Date: Mon, 6 Jun 2016 14:48:13 -0700 Subject: Added support for BFE (Bit field extract) instruction. --- src/cuda-sim/instructions.cc | 77 +++++++++++++++++++++++++++++++++++++++++++- 1 file changed, 76 insertions(+), 1 deletion(-) (limited to 'src/cuda-sim/instructions.cc') diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 4dd5ed8..02ce01c 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -1336,7 +1336,82 @@ void bar_impl( const ptx_instruction *pIin, ptx_thread_info *thread ) thread->m_last_dram_callback.instruction = pIin; } -void bfe_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void bfe_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +{ + unsigned i_type = pI->get_type(); + unsigned msb = (i_type == U32_TYPE || i_type == S32_TYPE) ? 31 : 63; + const operand_info &dst = pI->dst(); + const operand_info &src1 = pI->src1(); + const operand_info &src2 = pI->src2(); + const operand_info &src3 = pI->src3(); + ptx_reg_t a = thread->get_operand_value(src1, dst, i_type, thread, 1); + ptx_reg_t b = thread->get_operand_value(src2, dst, i_type, thread, 1); + ptx_reg_t c = thread->get_operand_value(src3, dst, i_type, thread, 1); + unsigned pos = b.u32 & 0xFF; + unsigned len = c.u32 & 0xFF; + unsigned d = 0; + switch (i_type) + { + case U32_TYPE: + { + unsigned mask; + d = a.u32 >> pos; + mask = 0xFFFFFFFF >> (32 - len); + d &= mask; + break; + } + case U64_TYPE: + { + unsigned long mask; + d = a.u64 >> pos; + mask = 0xFFFFFFFFFFFFFFFF >> (64 - len); + d &= mask; + break; + } + case S32_TYPE: + { + unsigned mask; + unsigned min = MY_MIN_I(pos + len - 1, msb); + unsigned sbit = len == 0 ? 0 : (a.s32 >> min) & 0x1; + d = a.s32 >> pos; + if (sbit > 0) + { + mask = 0xFFFFFFFF << len; + d |= mask; + } + else + { + mask = 0xFFFFFFFF >> (32 - len); + d &= mask; + } + break; + } + case S64_TYPE: + { + unsigned long mask; + unsigned min = MY_MIN_I(pos + len - 1, msb); + unsigned sbit = len == 0 ? 0 : (a.s64 >> min) & 0x1; + d = a.s64 >> pos; + if (sbit > 0) + { + mask = 0xFFFFFFFFFFFFFFFF << len; + d |= mask; + } + else + { + mask = 0xFFFFFFFFFFFFFFFF >> (64 - len); + d &= mask; + } + break; + } + default: + printf("Operand type not supported for BFE instruction.\n"); + abort(); + return; + } + thread->set_operand_value(dst,d, i_type, thread, pI); +} + void bfi_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } void bfind_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } -- cgit v1.3 From 281798191f9bc37a75592d34a5e38cc5d6c41b6d Mon Sep 17 00:00:00 2001 From: speverel Date: Thu, 16 Jun 2016 15:51:17 -0700 Subject: Added the ability to inject arbitrary PTX instructions. This will be used to add custom instructions in the future; the imaginary instructions 'spr' and 'ama' have been added as samples. --- src/cuda-sim/instructions.cc | 10 ++++++++++ src/cuda-sim/opcodes.def | 2 ++ src/cuda-sim/ptx.l | 6 +++++- 3 files changed, 17 insertions(+), 1 deletion(-) (limited to 'src/cuda-sim/instructions.cc') diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 02ce01c..922e14a 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -816,6 +816,11 @@ void add_impl( const ptx_instruction *pI, ptx_thread_info *thread ) void addc_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } +void ama_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +{ + printf("AMA instruction found.\n"); +} + void and_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { ptx_reg_t src1_data, src2_data, data; @@ -3698,6 +3703,11 @@ void slct_impl( const ptx_instruction *pI, ptx_thread_info *thread ) thread->set_operand_value(dst,d, i_type, thread, pI); } +void spr_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +{ + printf("SPR instruction found.\n"); +} + void sqrt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { ptx_reg_t a, d; diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def index 874acc7..33ee0ca 100644 --- a/src/cuda-sim/opcodes.def +++ b/src/cuda-sim/opcodes.def @@ -41,6 +41,7 @@ OP_DEF(ABS_OP,abs_impl,"abs",1,1) OP_DEF(ADD_OP,add_impl,"add",1,1) OP_DEF(ADDP_OP,addp_impl,"addp",1,1) OP_DEF(ADDC_OP,addc_impl,"addc",1,1) +OP_DEF(AMA_OP,ama_impl,"ama",1,2) OP_DEF(AND_OP,and_impl,"and",1,1) OP_DEF(ANDN_OP,andn_impl,"andn",1,1) OP_DEF(ATOM_OP,atom_impl,"atom",1,3) @@ -101,6 +102,7 @@ OP_DEF(SHL_OP,shl_impl,"shl",1,1) OP_DEF(SHR_OP,shr_impl,"shr",1,1) OP_DEF(SIN_OP,sin_impl,"sin",1,4) OP_DEF(SLCT_OP,slct_impl,"slct",1,1) +OP_DEF(SPR_OP,spr_impl,"spr",1,1) OP_DEF(SQRT_OP,sqrt_impl,"sqrt",1,4) OP_DEF(SSY_OP,ssy_impl,"ssy",0,3) OP_DEF(ST_OP,st_impl,"st",0,5) diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index a44177b..026270a 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -145,6 +145,8 @@ xor TC; ptx_lval.int_value = XOR_OP; return OPCODE; nop TC; ptx_lval.int_value = NOP_OP; return OPCODE; break TC; ptx_lval.int_value = BREAK_OP; return OPCODE; breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; +spr TC; ptx_lval.int_value = SPR_OP; return OPCODE; +ama TC; ptx_lval.int_value = AMA_OP; return OPCODE; { @@ -390,7 +392,9 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; } { "*/" BEGIN(INITIAL); -[^*\n]+ // eat comment in chunks +"CPTX_BEGIN" printf("BEGINNING CUSTOM PTX.\n"); BEGIN(INITIAL); +[^C*\n]+ // eat comment in chunks +"C" "*" // eat the lone star \n TC; } -- cgit v1.3 From f7c57e76c086ce417626f37ffc91097c839c687d Mon Sep 17 00:00:00 2001 From: sspenst Date: Mon, 4 Jul 2016 09:43:36 -0700 Subject: Reverted part of the previous commit so that our new changes related to DNNs can be done in a different branch --- src/cuda-sim/instructions.cc | 10 ---------- src/cuda-sim/opcodes.def | 2 -- src/cuda-sim/ptx.l | 2 -- 3 files changed, 14 deletions(-) (limited to 'src/cuda-sim/instructions.cc') diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 922e14a..02ce01c 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -816,11 +816,6 @@ void add_impl( const ptx_instruction *pI, ptx_thread_info *thread ) void addc_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } -void ama_impl( const ptx_instruction *pI, ptx_thread_info *thread ) -{ - printf("AMA instruction found.\n"); -} - void and_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { ptx_reg_t src1_data, src2_data, data; @@ -3703,11 +3698,6 @@ void slct_impl( const ptx_instruction *pI, ptx_thread_info *thread ) thread->set_operand_value(dst,d, i_type, thread, pI); } -void spr_impl( const ptx_instruction *pI, ptx_thread_info *thread ) -{ - printf("SPR instruction found.\n"); -} - void sqrt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { ptx_reg_t a, d; diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def index 33ee0ca..874acc7 100644 --- a/src/cuda-sim/opcodes.def +++ b/src/cuda-sim/opcodes.def @@ -41,7 +41,6 @@ OP_DEF(ABS_OP,abs_impl,"abs",1,1) OP_DEF(ADD_OP,add_impl,"add",1,1) OP_DEF(ADDP_OP,addp_impl,"addp",1,1) OP_DEF(ADDC_OP,addc_impl,"addc",1,1) -OP_DEF(AMA_OP,ama_impl,"ama",1,2) OP_DEF(AND_OP,and_impl,"and",1,1) OP_DEF(ANDN_OP,andn_impl,"andn",1,1) OP_DEF(ATOM_OP,atom_impl,"atom",1,3) @@ -102,7 +101,6 @@ OP_DEF(SHL_OP,shl_impl,"shl",1,1) OP_DEF(SHR_OP,shr_impl,"shr",1,1) OP_DEF(SIN_OP,sin_impl,"sin",1,4) OP_DEF(SLCT_OP,slct_impl,"slct",1,1) -OP_DEF(SPR_OP,spr_impl,"spr",1,1) OP_DEF(SQRT_OP,sqrt_impl,"sqrt",1,4) OP_DEF(SSY_OP,ssy_impl,"ssy",0,3) OP_DEF(ST_OP,st_impl,"st",0,5) diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index 026270a..b8ce497 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -145,8 +145,6 @@ xor TC; ptx_lval.int_value = XOR_OP; return OPCODE; nop TC; ptx_lval.int_value = NOP_OP; return OPCODE; break TC; ptx_lval.int_value = BREAK_OP; return OPCODE; breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; -spr TC; ptx_lval.int_value = SPR_OP; return OPCODE; -ama TC; ptx_lval.int_value = AMA_OP; return OPCODE; { -- cgit v1.3 From 2683b8bd7ba9950e0aa174915ef9ff64e0a20421 Mon Sep 17 00:00:00 2001 From: speverel Date: Mon, 4 Jul 2016 15:07:50 -0700 Subject: Restored madp instruction. --- src/cuda-sim/cuda-sim.cc | 2 +- src/cuda-sim/instructions.cc | 5 +++++ src/cuda-sim/opcodes.def | 1 + src/cuda-sim/ptx.l | 1 + 4 files changed, 8 insertions(+), 1 deletion(-) (limited to 'src/cuda-sim/instructions.cc') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index fba3a59..09e9a81 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -719,7 +719,7 @@ void ptx_instruction::set_opcode_and_latency() break; } break; - case MAD_OP: case MADC_OP: + case MAD_OP: case MADC_OP: case MADP_OP: //MAD latency switch(get_type()){ case F32_TYPE: diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 02ce01c..7b0f4fa 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -2496,6 +2496,11 @@ void mad_impl( const ptx_instruction *pI, ptx_thread_info *thread ) mad_def(pI, thread, false); } +void madp_impl( const ptx_instruction *pI, ptx_thread_info *thread ) +{ + mad_def(pI, thread, true); +} + void madc_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { mad_def(pI, thread, true); diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def index 874acc7..2ee6976 100644 --- a/src/cuda-sim/opcodes.def +++ b/src/cuda-sim/opcodes.def @@ -70,6 +70,7 @@ OP_DEF(LG2_OP,lg2_impl,"lg2",1,4) OP_DEF(MAD24_OP,mad24_impl,"mad24",1,2) OP_DEF(MAD_OP,mad_impl,"mad",1,2) OP_DEF(MADC_OP,madc_impl,"madc",1,2) +OP_DEF(MADP_OP,madp_impl,"madp",1,2) OP_DEF(MAX_OP,max_impl,"max",1,1) OP_DEF(MEMBAR_OP,membar_impl,"membar",1,3) OP_DEF(MIN_OP,min_impl,"min",1,1) diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index b8ce497..88ccf6a 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -87,6 +87,7 @@ lg2 TC; ptx_lval.int_value = LG2_OP; return OPCODE; mad24 TC; ptx_lval.int_value = MAD24_OP; return OPCODE; mad TC; ptx_lval.int_value = MAD_OP; return OPCODE; madc TC; ptx_lval.int_value = MADC_OP; return OPCODE; +madp TC; ptx_lval.int_value = MADP_OP; return OPCODE; max TC; ptx_lval.int_value = MAX_OP; return OPCODE; membar TC; ptx_lval.int_value = MEMBAR_OP; return OPCODE; min TC; ptx_lval.int_value = MIN_OP; return OPCODE; -- cgit v1.3 From 5cc6392a98bf736a05ffa7c7c557c4556753c8c2 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Wed, 1 Oct 2014 23:23:51 -0400 Subject: ADD: handle child kernel name in mov instruction. ADD: detect call cudaGetParameterBufferV2 and call cudaLaunchDeviceV2 --- libcuda/cuda_runtime_api.cc | 1 - src/cuda-sim/instructions.cc | 13 ++++++++++++- 2 files changed, 12 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim/instructions.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 5dc9cc1..5310a52 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -266,7 +266,6 @@ struct CUctx_st { m_kernel_lookup[hostFun] = f; } else { - m_code[fat_cubin_handle]->dump(); printf("Warning: cannot find deviceFun %s\n", deviceFun); m_kernel_lookup[hostFun] = NULL; } diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index cf7f04a..254427b 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -148,6 +148,9 @@ ptx_reg_t ptx_thread_info::get_operand_value( const operand_info &op, operand_in result.u64 = op.get_symbol()->get_address(); } else if ( op.is_local() ) { result.u64 = op.get_symbol()->get_address(); + } else if ( op.is_function_address() ) { + result.u64 = op.get_symbol()->get_pc()->get_start_PC(); + printf("Get pc for kernel function %u\n", op.get_symbol()->get_pc()->get_start_PC()); } else { const char *name = op.name().c_str(); printf("GPGPU-Sim PTX: ERROR ** get_operand_value : unknown operand type for %s\n", name ); @@ -1407,7 +1410,15 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) if( fname == "vprintf" ) { gpgpusim_cuda_vprintf(pI, thread, target_func); return; - } + } + else if(fname == "cudaGetParameterBufferV2") { + printf("calling cudaGetParameterBufferV2\n"); + return; + } + else if(fname == "cudaLaunchDeviceV2") { + printf("calling cudaLaunchDeviceV2\n"); + return; + } // read source arguements into register specified in declaration of function arg_buffer_list_t arg_values; -- cgit v1.3 From 70e02ee5283cb96f0edcb46a15edf0ab6e1d0697 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Mon, 13 Oct 2014 18:58:37 -0400 Subject: ADD: add cudaGetParameterBufferV2 and add cudaLaunchDeviceV2 implementation. Kernel launch to stream not yet implemented --- src/cuda-sim/Makefile | 3 +- src/cuda-sim/cuda-sim.cc | 10 +++ src/cuda-sim/cuda_device_runtime.cc | 175 ++++++++++++++++++++++++++++++++++++ src/cuda-sim/cuda_device_runtime.h | 7 ++ src/cuda-sim/instructions.cc | 12 ++- src/cuda-sim/ptx_ir.h | 4 +- src/cuda-sim/ptx_parser.cc | 14 +++ 7 files changed, 219 insertions(+), 6 deletions(-) create mode 100644 src/cuda-sim/cuda_device_runtime.cc create mode 100644 src/cuda-sim/cuda_device_runtime.h (limited to 'src/cuda-sim/instructions.cc') diff --git a/src/cuda-sim/Makefile b/src/cuda-sim/Makefile index 166e256..f479294 100644 --- a/src/cuda-sim/Makefile +++ b/src/cuda-sim/Makefile @@ -62,7 +62,7 @@ ifeq ($(GNUC_CPP0X),1) endif endif -OBJS := $(OUTPUT_DIR)/ptx_parser.o $(OUTPUT_DIR)/ptx_loader.o $(OUTPUT_DIR)/cuda_device_printf.o $(OUTPUT_DIR)/instructions.o $(OUTPUT_DIR)/cuda-sim.o $(OUTPUT_DIR)/ptx_ir.o $(OUTPUT_DIR)/ptx_sim.o $(OUTPUT_DIR)/memory.o $(OUTPUT_DIR)/ptx-stats.o $(OUTPUT_DIR)/decuda_pred_table/decuda_pred_table.o $(OUTPUT_DIR)/ptx.tab.o $(OUTPUT_DIR)/lex.ptx_.o $(OUTPUT_DIR)/ptxinfo.tab.o $(OUTPUT_DIR)/lex.ptxinfo_.o +OBJS := $(OUTPUT_DIR)/ptx_parser.o $(OUTPUT_DIR)/ptx_loader.o $(OUTPUT_DIR)/cuda_device_printf.o $(OUTPUT_DIR)/instructions.o $(OUTPUT_DIR)/cuda-sim.o $(OUTPUT_DIR)/ptx_ir.o $(OUTPUT_DIR)/ptx_sim.o $(OUTPUT_DIR)/memory.o $(OUTPUT_DIR)/ptx-stats.o $(OUTPUT_DIR)/decuda_pred_table/decuda_pred_table.o $(OUTPUT_DIR)/ptx.tab.o $(OUTPUT_DIR)/lex.ptx_.o $(OUTPUT_DIR)/ptxinfo.tab.o $(OUTPUT_DIR)/lex.ptxinfo_.o $(OUTPUT_DIR)/cuda_device_runtime.o OPT += -DCUDART_VERSION=$(CUDART_VERSION) @@ -145,5 +145,6 @@ $(OUTPUT_DIR)/ptx_sim.o: $(OUTPUT_DIR)/ptx.tab.c $(OUTPUT_DIR)/cuda-sim.o: $(OUTPUT_DIR)/ptx.tab.c $(OUTPUT_DIR)/lex.ptxinfo_.o: $(OUTPUT_DIR)/ptx.tab.c $(OUTPUT_DIR)/lex.ptx_.o: $(OUTPUT_DIR)/ptx.tab.c +$(OUTPUT_DIR)/cuda_device_runtime.o: $(OUTPUT_DIR)/ptx.tab.c include $(OUTPUT_DIR)/Makefile.makedepend diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 715be98..4933029 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1065,6 +1065,16 @@ void function_info::add_param_data( unsigned argn, struct gpgpu_ptx_sim_arg *arg } } +unsigned function_info::get_args_aligned_size() { + unsigned int align_size = 4; // a word + unsigned int total_size = 0; + for(unsigned int i = 0; i < num_args(); i++) { + total_size += ((m_args[i]->get_size_in_bytes() + align_size - 1) / align_size) * align_size; + } + return total_size; +} + + void function_info::finalize( memory_space *param_mem ) { unsigned param_address = 0; diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc new file mode 100644 index 0000000..937eec8 --- /dev/null +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -0,0 +1,175 @@ +//Jin: cuda_device_runtime.cc +//Defines CUDA device runtime APIs for CDP support + +#include +#include + +#define __CUDA_RUNTIME_API_H__ + +#include +#include +#include "../gpgpu-sim/gpu-sim.h" +#include "cuda-sim.h" +#include "ptx_ir.h" +#include "../stream_manager.h" +#include "cuda_device_runtime.h" + +#define DEV_RUNTIME_REPORT(a) \ + if( g_debug_execution ) { \ + std::cout << __FILE__ << ", " << __LINE__ << ": " << a << "\n"; \ + std::cout.flush(); \ + } + +std::map g_cuda_device_launch_map; +struct CUstream_st * g_device_default_stream = NULL; +extern stream_manager *g_stream_manager; + +//Handling device runtime api: +//void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize) +void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func) +{ + DEV_RUNTIME_REPORT("Calling cudaGetParameterBufferV2"); + + unsigned n_return = target_func->has_return(); + assert(n_return); + unsigned n_args = target_func->num_args(); + assert( n_args == 4 ); + + function_info * child_kernel_entry; + struct dim3 gridDim, blockDim; + unsigned int sharedMem; + + for( unsigned arg=0; arg < n_args; arg ++ ) { + const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg); //param# + const symbol *formal_param = target_func->get_arg(arg); //cudaGetParameterBufferV2_param_# + unsigned size=formal_param->get_size_in_bytes(); + assert( formal_param->is_param_local() ); + assert( actual_param_op.is_param_local() ); + addr_t from_addr = actual_param_op.get_symbol()->get_address(); + + if(arg == 0) {//function_info* for the child kernel + unsigned long long buf; + assert(size == sizeof(function_info *)); + thread->m_local_mem->read(from_addr, size, &buf); + child_kernel_entry = (function_info *)buf; + assert(child_kernel_entry); + DEV_RUNTIME_REPORT("child kernel name " << child_kernel_entry->get_name()); + } + else if(arg == 1) { //dim3 gridDim for the child kernel + assert(size == sizeof(struct dim3)); + thread->m_local_mem->read(from_addr, size, & gridDim); + DEV_RUNTIME_REPORT("grid (" << gridDim.x << ", " << gridDim.y << ", " << gridDim.z << ")"); + } + else if(arg == 2) { //dim3 blockDim for the child kernel + assert(size == sizeof(struct dim3)); + thread->m_local_mem->read(from_addr, size, & blockDim); + DEV_RUNTIME_REPORT("block (" << blockDim.x << ", " << blockDim.y << ", " << blockDim.z << ")"); + } + else if(arg == 3) { //unsigned int sharedMem + assert(size == sizeof(unsigned int)); + thread->m_local_mem->read(from_addr, size, & sharedMem); + DEV_RUNTIME_REPORT("shared memory " << sharedMem); + } + } + + //get total child kernel argument size and malloc buffer in global memory + unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size(); + void * param_buffer = thread->get_gpu()->gpu_malloc(child_kernel_arg_size); + DEV_RUNTIME_REPORT("child kernel arg size total " << child_kernel_arg_size << ", parameter buffer allocated at " << param_buffer); + + //create child kernel_info_t and index it with parameter_buffer address + kernel_info_t * child_grid = new kernel_info_t(gridDim, blockDim, child_kernel_entry); + assert(g_cuda_device_launch_map.find(param_buffer) == g_cuda_device_launch_map.end()); + g_cuda_device_launch_map[param_buffer] = child_grid; + + //copy the buffer address to retval0 + const operand_info &actual_return_op = pI->operand_lookup(0); //retval0 + const symbol *formal_return = target_func->get_return_var(); //void * + unsigned int return_size = formal_return->get_size_in_bytes(); + DEV_RUNTIME_REPORT("cudaGetParameterBufferV2 return value has size of " << return_size); + assert(actual_return_op.is_param_local()); + assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size && return_size == sizeof(void *)); + addr_t ret_param_addr = actual_return_op.get_symbol()->get_address(); + thread->m_local_mem->write(ret_param_addr, return_size, ¶m_buffer, NULL, NULL); + +} + +//Handling device runtime api: +//cudaError_t cudaLaunchDeviceV2(void *parameterBuffer, cudaStream_t stream) +void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func) { + DEV_RUNTIME_REPORT("Calling cudaLaunchDeviceV2"); + + unsigned n_return = target_func->has_return(); + assert(n_return); + unsigned n_args = target_func->num_args(); + assert( n_args == 2 ); + + kernel_info_t * child_grid = NULL; + function_info * child_kernel_entry = NULL; + void * parameter_buffer; + struct CUstream_st * child_stream; + for( unsigned arg=0; arg < n_args; arg ++ ) { + const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg); //param# + const symbol *formal_param = target_func->get_arg(arg); //cudaLaunchDeviceV2_param_# + unsigned size=formal_param->get_size_in_bytes(); + assert( formal_param->is_param_local() ); + assert( actual_param_op.is_param_local() ); + addr_t from_addr = actual_param_op.get_symbol()->get_address(); + + if(arg == 0) {//paramter buffer for child kernel (in global memory) + //get parameter_buffer from the cudaDeviceLaunchV2_param0 + assert(size == sizeof(void *)); + thread->m_local_mem->read(from_addr, size, ¶meter_buffer); + assert((size_t)parameter_buffer >= GLOBAL_HEAP_START); + DEV_RUNTIME_REPORT("Parameter buffer locating at global memory " << parameter_buffer); + + //get child grid info through parameter_buffer address + assert(g_cuda_device_launch_map.find(parameter_buffer) != g_cuda_device_launch_map.end()); + child_grid = g_cuda_device_launch_map[parameter_buffer]; + child_kernel_entry = child_grid->entry(); + DEV_RUNTIME_REPORT("find child kernel " << child_kernel_entry->get_name()); + + //copy data in parameter_buffer to child kernel param memory + unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size(); + DEV_RUNTIME_REPORT("child_kernel_arg_size " << child_kernel_arg_size); + memory_space *child_kernel_param_mem = child_grid->get_param_memory(); + size_t param_start_address = 0; + for(unsigned n = 0; n < child_kernel_arg_size; n++) { + unsigned char one_byte; + thread->get_gpu()->get_global_memory()->read((size_t)parameter_buffer + n, 1, &one_byte); + child_kernel_param_mem->write(param_start_address + n, 1, &one_byte, NULL, NULL); + } + } + else if(arg == 1) { //cudaStream for the child kernel + assert(size == sizeof(cudaStream_t)); + thread->m_local_mem->read(from_addr, size, &child_stream); + if(child_stream == 0) { //default stream on device + if(!g_device_default_stream) { + //g_device_default_stream = new struct CUstream_st(); + //g_stream_manager->add_stream(g_device_default_stream); + } + child_stream = g_device_default_stream; + } +// DEV_RUNTIME_REPORT("launching child kernel to stream " << child_stream->get_uid()); + } + + } + + //launch child kernel + stream_operation op(child_grid, g_ptx_sim_mode, child_stream); +// g_stream_manager->push(op); +// g_cuda_device_launch_map.erase(parameter_buffer); + + //set retval0 + const operand_info &actual_return_op = pI->operand_lookup(0); //retval0 + const symbol *formal_return = target_func->get_return_var(); //cudaError_t + unsigned int return_size = formal_return->get_size_in_bytes(); + DEV_RUNTIME_REPORT("cudaLaunchDeviceV2 return value has size of " << return_size); + assert(actual_return_op.is_param_local()); + assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size + && return_size == sizeof(cudaError_t)); + cudaError_t error = cudaSuccess; + addr_t ret_param_addr = actual_return_op.get_symbol()->get_address(); + thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL); + +} diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h new file mode 100644 index 0000000..1b10407 --- /dev/null +++ b/src/cuda-sim/cuda_device_runtime.h @@ -0,0 +1,7 @@ +//Jin: cuda_device_runtime.h +//Defines CUDA device runtime APIs for CDP support + +#pragma once + +void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); +void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 254427b..1c47ad3 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -41,6 +41,9 @@ #include "../gpgpu-sim/gpu-sim.h" #include "../gpgpu-sim/shader.h" +//Jin: include device runtime for CDP +#include "cuda_device_runtime.h" + #include unsigned ptx_instruction::g_num_ptx_inst_uid=0; @@ -149,8 +152,7 @@ ptx_reg_t ptx_thread_info::get_operand_value( const operand_info &op, operand_in } else if ( op.is_local() ) { result.u64 = op.get_symbol()->get_address(); } else if ( op.is_function_address() ) { - result.u64 = op.get_symbol()->get_pc()->get_start_PC(); - printf("Get pc for kernel function %u\n", op.get_symbol()->get_pc()->get_start_PC()); + result.u64 = (size_t)op.get_symbol()->get_pc(); } else { const char *name = op.name().c_str(); printf("GPGPU-Sim PTX: ERROR ** get_operand_value : unknown operand type for %s\n", name ); @@ -1411,12 +1413,14 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) gpgpusim_cuda_vprintf(pI, thread, target_func); return; } + + //Jin: handle device runtime apis for CDP else if(fname == "cudaGetParameterBufferV2") { - printf("calling cudaGetParameterBufferV2\n"); + gpgpusim_cuda_getParameterBufferV2(pI, thread, target_func); return; } else if(fname == "cudaLaunchDeviceV2") { - printf("calling cudaLaunchDeviceV2\n"); + gpgpusim_cuda_launchDeviceV2(pI, thread, target_func); return; } diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 7325e5f..a7ca27e 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -699,7 +699,7 @@ public: } bool is_immediate_address() const { - return m_immediate_address; + return m_immediate_address; } bool is_literal() const { return m_type == int_t || @@ -1209,6 +1209,8 @@ public: { return m_args.size(); } + unsigned get_args_aligned_size(); + const symbol* get_arg( unsigned n ) const { assert( n < m_args.size() ); diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 39257da..baa3bcd 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -120,6 +120,20 @@ symbol_table *init_parser( const char *ptx_filename ) #define DEF(X,Y) g_ptx_token_decode[X] = Y; #include "ptx_parser_decode.def" #undef DEF + g_ptx_token_decode[undefined_space] = "undefined_space"; + g_ptx_token_decode[undefined_space] = "undefined_space=0"; + g_ptx_token_decode[reg_space] = "reg_space"; + g_ptx_token_decode[local_space] = "local_space"; + g_ptx_token_decode[shared_space] = "shared_space"; + g_ptx_token_decode[param_space_unclassified] = "param_space_unclassified"; + g_ptx_token_decode[param_space_kernel] = "param_space_kernel"; + g_ptx_token_decode[param_space_local] = "param_space_local"; + g_ptx_token_decode[const_space] = "const_space"; + g_ptx_token_decode[tex_space] = "tex_space"; + g_ptx_token_decode[surf_space] = "surf_space"; + g_ptx_token_decode[global_space] = "global_space"; + g_ptx_token_decode[generic_space] = "generic_space"; + g_ptx_token_decode[instruction_space] = "instruction_space"; return g_global_symbol_table; } -- cgit v1.3 From 3bb32d87175d873e7089ad50a0069acc195edb34 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Tue, 21 Oct 2014 01:24:26 -0400 Subject: ADD: add support for cudaStreamCreateWithFlags --- src/cuda-sim/cuda_device_runtime.cc | 61 ++++++++++++++++++++++++++++++++++++- src/cuda-sim/cuda_device_runtime.h | 1 + src/cuda-sim/instructions.cc | 4 +++ src/cuda-sim/ptx_ir.cc | 3 +- 4 files changed, 67 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim/instructions.cc') diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 148471b..be87f6a 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -152,11 +152,12 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * kernel_info_t & parent_kernel = thread->get_kernel(); if(child_stream == 0) { //default stream on device for current CTA child_stream = parent_kernel.get_default_stream_cta(thread->get_ctaid()); + DEV_RUNTIME_REPORT("launching child kernel to default stream of the cta " << child_stream->get_uid() << ": " << child_stream); } else { assert(parent_kernel.cta_has_stream(thread->get_ctaid(), child_stream)); + DEV_RUNTIME_REPORT("launching child kernel to stream " << child_stream->get_uid() << ": " << child_stream); } - DEV_RUNTIME_REPORT("launching child kernel to stream " << child_stream->get_uid() << " " << child_stream); } } @@ -179,3 +180,61 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL); } + + +//Handling device runtime api: +//cudaError_t cudaStreamCreateWithFlags ( cudaStream_t* pStream, unsigned int flags) +//flags can only be cudaStreamNonBlocking +void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func) { + DEV_RUNTIME_REPORT("Calling cudaStreamCreateWithFlags"); + + unsigned n_return = target_func->has_return(); + assert(n_return); + unsigned n_args = target_func->num_args(); + assert( n_args == 2 ); + + size_t generic_pStream_addr; + addr_t pStream_addr; + unsigned int flags; + for( unsigned arg=0; arg < n_args; arg ++ ) { + const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg); //param# + const symbol *formal_param = target_func->get_arg(arg); //cudaStreamCreateWithFlags_param_# + unsigned size=formal_param->get_size_in_bytes(); + assert( formal_param->is_param_local() ); + assert( actual_param_op.is_param_local() ); + addr_t from_addr = actual_param_op.get_symbol()->get_address(); + + if(arg == 0) {//cudaStream_t * pStream, address of cudaStream_t + assert(size == sizeof(cudaStream_t *)); + thread->m_local_mem->read(from_addr, size, &generic_pStream_addr); + + //pStream should be non-zero address in local memory + pStream_addr = generic_to_local(thread->get_hw_sid(), thread->get_hw_tid(), generic_pStream_addr); + + DEV_RUNTIME_REPORT("pStream locating at local memory " << pStream_addr); + } + else if(arg == 1) { //unsigned int flags, should be cudaStreamNonBlocking + assert(size == sizeof(unsigned int)); + thread->m_local_mem->read(from_addr, size, &flags); + assert(flags == cudaStreamNonBlocking); + } + } + + //create stream and write back to param0 + CUstream_st * stream = thread->get_kernel().create_stream_cta(thread->get_ctaid()); + DEV_RUNTIME_REPORT("Create stream " << stream->get_uid() << ": " << stream); + thread->m_local_mem->write(pStream_addr, sizeof(cudaStream_t), &stream, NULL, NULL); + + //set retval0 + const operand_info &actual_return_op = pI->operand_lookup(0); //retval0 + const symbol *formal_return = target_func->get_return_var(); //cudaError_t + unsigned int return_size = formal_return->get_size_in_bytes(); + DEV_RUNTIME_REPORT("cudaStreamCreateWithFlags return value has size of " << return_size); + assert(actual_return_op.is_param_local()); + assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size + && return_size == sizeof(cudaError_t)); + cudaError_t error = cudaSuccess; + addr_t ret_param_addr = actual_return_op.get_symbol()->get_address(); + thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL); + +} diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h index 1b10407..7447d5a 100644 --- a/src/cuda-sim/cuda_device_runtime.h +++ b/src/cuda-sim/cuda_device_runtime.h @@ -5,3 +5,4 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); +void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 1c47ad3..5d909d3 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -1423,6 +1423,10 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) gpgpusim_cuda_launchDeviceV2(pI, thread, target_func); return; } + else if(fname == "cudaStreamCreateWithFlags") { + gpgpusim_cuda_streamCreateWithFlags(pI, thread, target_func); + return; + } // read source arguements into register specified in declaration of function arg_buffer_list_t arg_values; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 861f0dc..bdc8381 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1242,7 +1242,8 @@ ptx_instruction::ptx_instruction( int opcode, m_is_printf = true; } if (fname == "cudaGetParameterBufferV2" - || fname == "cudaLaunchDeviceV2") + || fname == "cudaLaunchDeviceV2" + || fname == "cudaStreamCreateWithFlags") m_is_cdp = true; } -- cgit v1.3 From 68336f112117bcef5b943650819a6764e9ebf4ce Mon Sep 17 00:00:00 2001 From: sspenst Date: Wed, 24 Aug 2016 15:24:19 -0700 Subject: Added shfl instruction --- src/abstract_hardware_model.h | 1 + src/cuda-sim/cuda-sim.cc | 20 +++++++++++- src/cuda-sim/instructions.cc | 75 +++++++++++++++++++++++++++++++++++++++++++ src/cuda-sim/opcodes.def | 1 + src/cuda-sim/opcodes.h | 4 ++- src/cuda-sim/ptx.l | 6 ++++ src/cuda-sim/ptx.y | 8 +++++ src/cuda-sim/ptx_ir.cc | 6 ++++ src/cuda-sim/ptx_ir.h | 2 ++ src/cuda-sim/ptx_sim.cc | 21 ++++++++++++ src/cuda-sim/ptx_sim.h | 12 +++++++ 11 files changed, 154 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim/instructions.cc') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index b29f918..d0c807d 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -1051,6 +1051,7 @@ class core_t { warp_inst_t getExecuteWarp(unsigned warpId); void get_pdom_stack_top_info( unsigned warpId, unsigned *pc, unsigned *rpc ) const; kernel_info_t * get_kernel_info(){ return m_kernel;} + class ptx_thread_info ** get_thread_info() { return m_thread; } unsigned get_warp_size() const { return m_warp_size; } void and_reduction(unsigned ctaid, unsigned barid, bool value) { reduction_storage[ctaid][barid] &= value; } void or_reduction(unsigned ctaid, unsigned barid, bool value) { reduction_storage[ctaid][barid] |= value; } diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 09e9a81..8bf4ec8 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -769,6 +769,10 @@ void ptx_instruction::set_opcode_and_latency() initiation_interval = dp_init[2]; op = SFU_OP; break; + case SHFL_OP: + latency = 32; + initiation_interval = 15; + break; default: break; } @@ -845,8 +849,10 @@ void ptx_instruction::pre_decode() switch ( get_opcode() ) { #define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: has_dst = (DST!=0); break; +#define OP_W_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: has_dst = (DST!=0); break; #include "opcodes.def" #undef OP_DEF +#undef OP_W_DEF default: printf( "Execution error: Invalid opcode (0x%x)\n", get_opcode() ); break; @@ -1240,8 +1246,10 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) } switch ( pI->get_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" #undef OP_DEF +#undef OP_W_DEF default: printf( "Execution error: Invalid opcode (0x%x)\n", pI->get_opcode() ); break; } delete pJ; @@ -1408,6 +1416,7 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, static std::map shared_memory_lookup; static std::map ptx_cta_lookup; + static std::map ptx_warp_lookup; static std::map > local_memory_lookup; if ( *thread_info != NULL ) { @@ -1486,7 +1495,16 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, kernel.increment_thread_id(); new_tid += tid; ptx_thread_info *thd = new ptx_thread_info(kernel); - + + ptx_warp_info *warp_info = NULL; + if ( ptx_warp_lookup.find(hw_warp_id) == ptx_warp_lookup.end() ) { + warp_info = new ptx_warp_info(); // num_threads should be threads in the warp + ptx_warp_lookup[hw_warp_id] = warp_info; + } else { + warp_info = ptx_warp_lookup[hw_warp_id]; + } + thd->m_warp_info = warp_info; + memory_space *local_mem = NULL; std::map::iterator l = local_mem_lookup.find(new_tid); if ( l != local_mem_lookup.end() ) { diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 7b0f4fa..05ba78f 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -47,8 +47,10 @@ unsigned ptx_instruction::g_num_ptx_inst_uid=0; const char *g_opcode_string[NUM_OPCODES] = { #define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) STR, +#define OP_W_DEF(OP,FUNC,STR,DST,CLASSIFICATION) STR, #include "opcodes.def" #undef OP_DEF +#undef OP_W_DEF }; void inst_not_implemented( const ptx_instruction * pI ) ; @@ -3516,6 +3518,79 @@ void set_impl( const ptx_instruction *pI, ptx_thread_info *thread ) } +void shfl_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) +{ + unsigned i_type = pI->get_type(); + int tid = inst.warp_id() * core->get_warp_size(); + ptx_thread_info *thread = core->get_thread_info()[tid]; + ptx_warp_info *warp_info = thread->m_warp_info; + int lane = warp_info->get_done_threads(); + thread = core->get_thread_info()[tid + lane]; + + const operand_info &dst = pI->dst(); + const operand_info &src1 = pI->src1(); + const operand_info &src2 = pI->src2(); + const operand_info &src3 = pI->src3(); + int bval = (thread->get_operand_value(src2, dst, i_type, thread, 1)).u32; + int cval = (thread->get_operand_value(src3, dst, i_type, thread, 1)).u32; + int mask = cval >> 8; + cval &= 0x1F; + + int maxLane = (lane & mask) | (cval & ~mask); + int minLane = lane & mask; + + int src_idx; + int p; + switch(pI->shfl_op()) { + case UP_OPTION: + src_idx = lane - bval; + p = (src_idx >= maxLane); + break; + case DOWN_OPTION: + src_idx = lane + bval; + p = (src_idx <= maxLane); + break; + case BFLY_OPTION: + src_idx = lane ^ bval; + p = (src_idx <= maxLane); + break; + case IDX_OPTION: + src_idx = minLane | (bval & ~mask); + p = (src_idx <= maxLane); + break; + default: + printf("GPGPU-Sim PTX: ERROR: Unrecognized shfl option\n"); + assert(0); + break; + } + // copy from own lane + if (!p) src_idx = lane; + // copy input from lane src_idx + ptx_reg_t data; + /*if (inst.active(src_idx) && i_type == PRED_TYPE) { + ptx_thread_info *source = core->get_thread_info()[tid + src_idx]; + data = source->get_operand_value(src1, dst, i_type, source, 1); + data.pred = p; + } else { + printf("GPGPU-Sim PTX: WARNING: shfl input value unpredictable for inactive/predicated-off threads in a warp\n"); + data.u32 = 0; + }*/ + if (inst.active(src_idx)) { + ptx_thread_info *source = core->get_thread_info()[tid + src_idx]; + data = source->get_operand_value(src1, dst, i_type, source, 1); + } + if (i_type == PRED_TYPE) { + data.pred = p; + } + thread->set_operand_value(dst, data, i_type, thread, pI); + + // keep track of the number of threads that have executed in the warp + warp_info->inc_done_threads(); + if (warp_info->get_done_threads() == inst.active_count()) { + warp_info->reset_done_threads(); + } +} + void shl_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { ptx_reg_t a, b, d; diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def index 2ee6976..e1b1422 100644 --- a/src/cuda-sim/opcodes.def +++ b/src/cuda-sim/opcodes.def @@ -98,6 +98,7 @@ OP_DEF(SAD_OP,sad_impl,"sad",1,1) OP_DEF(SELP_OP,selp_impl,"selp",1,1) OP_DEF(SETP_OP,setp_impl,"setp",1,1) OP_DEF(SET_OP,set_impl,"set",1,1) +OP_W_DEF(SHFL_OP,shfl_impl,"shfl",1,10) OP_DEF(SHL_OP,shl_impl,"shl",1,1) OP_DEF(SHR_OP,shr_impl,"shr",1,1) OP_DEF(SIN_OP,sin_impl,"sin",1,4) diff --git a/src/cuda-sim/opcodes.h b/src/cuda-sim/opcodes.h index 871091c..aa133da 100644 --- a/src/cuda-sim/opcodes.h +++ b/src/cuda-sim/opcodes.h @@ -30,9 +30,11 @@ enum opcode_t { #define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) OP, +#define OP_W_DEF(OP,FUNC,STR,DST,CLASSIFICATION) OP, #include "opcodes.def" - NUM_OPCODES + NUM_OPCODES #undef OP_DEF +#undef OP_W_DEF }; enum special_regs { diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index 88ccf6a..8fac4ac 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -115,6 +115,7 @@ sad TC; ptx_lval.int_value = SAD_OP; return OPCODE; selp TC; ptx_lval.int_value = SELP_OP; return OPCODE; setp TC; ptx_lval.int_value = SETP_OP; return OPCODE; set TC; ptx_lval.int_value = SET_OP; return OPCODE; +shfl TC; ptx_lval.int_value = SHFL_OP; return OPCODE; shl TC; ptx_lval.int_value = SHL_OP; return OPCODE; shr TC; ptx_lval.int_value = SHR_OP; return OPCODE; sin TC; ptx_lval.int_value = SIN_OP; return OPCODE; @@ -329,6 +330,11 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; \.nc TC; return NC_OPTION; +\.up TC; return UP_OPTION; +\.down TC; return DOWN_OPTION; +\.bfly TC; return BFLY_OPTION; +\.idx TC; return IDX_OPTION; + \.popc TC; return ATOMIC_POPC; \.and TC; return ATOMIC_AND; \.or TC; return ATOMIC_OR; diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y index 4de39d1..166b15d 100644 --- a/src/cuda-sim/ptx.y +++ b/src/cuda-sim/ptx.y @@ -194,6 +194,10 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. %token WB_OPTION; %token WT_OPTION; %token NC_OPTION; +%token UP_OPTION; +%token DOWN_OPTION; +%token BFLY_OPTION; +%token IDX_OPTION; %type function_decl_header %type function_decl @@ -451,6 +455,10 @@ option: type_spec | WB_OPTION { add_option(WB_OPTION); } | WT_OPTION { add_option(WT_OPTION); } | NC_OPTION { add_option(NC_OPTION); } + | UP_OPTION { add_option(UP_OPTION); } + | DOWN_OPTION { add_option(DOWN_OPTION); } + | BFLY_OPTION { add_option(BFLY_OPTION); } + | IDX_OPTION { add_option(IDX_OPTION); } ; atomic_operation_spec: ATOMIC_AND { add_option(ATOMIC_AND); } diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 2eccabc..4cfe1b9 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1171,6 +1171,12 @@ ptx_instruction::ptx_instruction( int opcode, 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; default: assert(0); break; diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 601a13d..0abbc83 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -993,6 +993,7 @@ public: unsigned saturation_mode() const { return m_saturation_mode;} unsigned dimension() const { return m_geom_spec;} unsigned barrier_op() const {return m_barrier_op;} + unsigned shfl_op() const {return m_shfl_op;} enum vote_mode_t { vote_any, vote_all, vote_uni, vote_ballot }; enum vote_mode_t vote_mode() const { return m_vote_mode; } @@ -1058,6 +1059,7 @@ private: unsigned m_compare_op; unsigned m_saturation_mode; unsigned m_barrier_op; + unsigned m_shfl_op; std::list m_scalar_type; memory_space_t m_space_spec; diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc index 09844ae..a3e43aa 100644 --- a/src/cuda-sim/ptx_sim.cc +++ b/src/cuda-sim/ptx_sim.cc @@ -128,6 +128,26 @@ unsigned ptx_cta_info::get_sm_idx() const return m_sm_idx; } +ptx_warp_info::ptx_warp_info() +{ + reset_done_threads(); +} + +unsigned ptx_warp_info::get_done_threads() const +{ + return m_done_threads; +} + +void ptx_warp_info::inc_done_threads() +{ + m_done_threads++; +} + +void ptx_warp_info::reset_done_threads() +{ + m_done_threads = 0; +} + unsigned g_ptx_thread_info_uid_next=1; unsigned g_ptx_thread_info_delete_count=0; @@ -153,6 +173,7 @@ ptx_thread_info::ptx_thread_info( kernel_info_t &kernel ) m_last_memory_space = undefined_space; m_branch_taken = 0; m_shared_mem = NULL; + m_warp_info = NULL; m_cta_info = NULL; m_local_mem = NULL; m_symbol_table = NULL; diff --git a/src/cuda-sim/ptx_sim.h b/src/cuda-sim/ptx_sim.h index f926e6d..449511f 100644 --- a/src/cuda-sim/ptx_sim.h +++ b/src/cuda-sim/ptx_sim.h @@ -167,6 +167,17 @@ private: std::set m_dangling_pointers; }; +class ptx_warp_info { +public: + ptx_warp_info(); + unsigned get_done_threads() const; + void inc_done_threads(); + void reset_done_threads(); + +private: + unsigned m_done_threads; +}; + class symbol; struct stack_entry { @@ -425,6 +436,7 @@ public: dram_callback_t m_last_dram_callback; memory_space *m_shared_mem; memory_space *m_local_mem; + ptx_warp_info *m_warp_info; ptx_cta_info *m_cta_info; ptx_reg_t m_last_set_operand_value; -- cgit v1.3 From bffac032402b35acb81200540f6f808bd6f851d8 Mon Sep 17 00:00:00 2001 From: sspenst Date: Wed, 24 Aug 2016 15:30:16 -0700 Subject: Cleanup --- src/cuda-sim/cuda-sim.cc | 2 +- src/cuda-sim/instructions.cc | 8 -------- 2 files changed, 1 insertion(+), 9 deletions(-) (limited to 'src/cuda-sim/instructions.cc') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 8bf4ec8..d3c928d 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1498,7 +1498,7 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, ptx_warp_info *warp_info = NULL; if ( ptx_warp_lookup.find(hw_warp_id) == ptx_warp_lookup.end() ) { - warp_info = new ptx_warp_info(); // num_threads should be threads in the warp + warp_info = new ptx_warp_info(); ptx_warp_lookup[hw_warp_id] = warp_info; } else { warp_info = ptx_warp_lookup[hw_warp_id]; diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 05ba78f..4c9392b 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -3567,14 +3567,6 @@ void shfl_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) if (!p) src_idx = lane; // copy input from lane src_idx ptx_reg_t data; - /*if (inst.active(src_idx) && i_type == PRED_TYPE) { - ptx_thread_info *source = core->get_thread_info()[tid + src_idx]; - data = source->get_operand_value(src1, dst, i_type, source, 1); - data.pred = p; - } else { - printf("GPGPU-Sim PTX: WARNING: shfl input value unpredictable for inactive/predicated-off threads in a warp\n"); - data.u32 = 0; - }*/ if (inst.active(src_idx)) { ptx_thread_info *source = core->get_thread_info()[tid + src_idx]; data = source->get_operand_value(src1, dst, i_type, source, 1); -- cgit v1.3 From 6a88a8fb947f6256dcf5dd9f315834f12322ad38 Mon Sep 17 00:00:00 2001 From: sspenst Date: Thu, 25 Aug 2016 15:54:38 -0700 Subject: Fixed minor shfl bugs --- src/cuda-sim/instructions.cc | 22 ++++++++++++++++------ 1 file changed, 16 insertions(+), 6 deletions(-) (limited to 'src/cuda-sim/instructions.cc') diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 4c9392b..78d6450 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -3534,13 +3534,14 @@ void shfl_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) int bval = (thread->get_operand_value(src2, dst, i_type, thread, 1)).u32; int cval = (thread->get_operand_value(src3, dst, i_type, thread, 1)).u32; int mask = cval >> 8; + bval &= 0x1F; cval &= 0x1F; int maxLane = (lane & mask) | (cval & ~mask); int minLane = lane & mask; int src_idx; - int p; + unsigned p; switch(pI->shfl_op()) { case UP_OPTION: src_idx = lane - bval; @@ -3559,23 +3560,32 @@ void shfl_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) p = (src_idx <= maxLane); break; default: - printf("GPGPU-Sim PTX: ERROR: Unrecognized shfl option\n"); + printf("GPGPU-Sim PTX: ERROR: Invalid shfl option\n"); assert(0); break; } // copy from own lane if (!p) src_idx = lane; + // copy input from lane src_idx + ptx_thread_info *source = core->get_thread_info()[tid + src_idx]; ptx_reg_t data; if (inst.active(src_idx)) { - ptx_thread_info *source = core->get_thread_info()[tid + src_idx]; data = source->get_operand_value(src1, dst, i_type, source, 1); - } - if (i_type == PRED_TYPE) { - data.pred = p; + } else { + printf("GPGPU-Sim PTX: WARNING: shfl input value unpredictable for inactive threads in a warp\n"); + data.u32 = 0; } thread->set_operand_value(dst, data, i_type, thread, pI); + /* + TODO: deal with predicates appropriately using the following pseudocode: + if (!isGuardPredicateTrue(src_idx)) { + printf("GPGPU-Sim PTX: WARNING: shfl input value unpredictable for predicated-off threads in a warp\n"); + } + if (dest predicate selected) data.pred = p; + */ + // keep track of the number of threads that have executed in the warp warp_info->inc_done_threads(); if (warp_info->get_done_threads() == inst.active_count()) { -- cgit v1.3 From 0afe825c212560c3b473366d005a50c2568cad1a Mon Sep 17 00:00:00 2001 From: sspenst Date: Thu, 25 Aug 2016 15:56:02 -0700 Subject: OCD --- src/cuda-sim/instructions.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'src/cuda-sim/instructions.cc') diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 78d6450..26d6285 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -3568,9 +3568,9 @@ void shfl_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) if (!p) src_idx = lane; // copy input from lane src_idx - ptx_thread_info *source = core->get_thread_info()[tid + src_idx]; ptx_reg_t data; if (inst.active(src_idx)) { + ptx_thread_info *source = core->get_thread_info()[tid + src_idx]; data = source->get_operand_value(src1, dst, i_type, source, 1); } else { printf("GPGPU-Sim PTX: WARNING: shfl input value unpredictable for inactive threads in a warp\n"); -- cgit v1.3 From 623a88e5d5c6c3edb94404ef6e5ea100caec9deb Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Mon, 29 Aug 2016 18:10:00 -0700 Subject: MOD: Add macros to turn off cuda_device_runtime for CUDA < 5.0 --- src/cuda-sim/cuda-sim.cc | 2 ++ src/cuda-sim/cuda_device_runtime.cc | 10 +++++++--- src/cuda-sim/cuda_device_runtime.h | 3 ++- src/cuda-sim/instructions.cc | 2 ++ src/gpgpu-sim/gpu-sim.cc | 2 ++ src/gpgpu-sim/shader.cc | 2 +- 6 files changed, 16 insertions(+), 5 deletions(-) (limited to 'src/cuda-sim/instructions.cc') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 3f5af7e..ec51779 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1792,7 +1792,9 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) ); cta.execute(); +#if (CUDART_VERSION >= 5000) launch_all_device_kernels(); +#endif } //registering this kernel as done diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 12c83d2..4a8ffe5 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -1,9 +1,15 @@ //Jin: cuda_device_runtime.cc //Defines CUDA device runtime APIs for CDP support + #include #include +unsigned long long g_total_param_size = 0; +unsigned long long g_max_total_param_size = 0; + + +#if (CUDART_VERSION >= 5000) #define __CUDA_RUNTIME_API_H__ #include @@ -59,8 +65,6 @@ public: std::map g_cuda_device_launch_param_map; std::list g_cuda_device_launch_op; extern stream_manager *g_stream_manager; -unsigned long long g_total_param_size = 0; -unsigned long long g_max_total_param_size = 0; //Handling device runtime api: //void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize) @@ -313,4 +317,4 @@ void launch_all_device_kernels() { launch_one_device_kernel(); } } - +#endif diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h index 385d605..6dbcd71 100644 --- a/src/cuda-sim/cuda_device_runtime.h +++ b/src/cuda-sim/cuda_device_runtime.h @@ -1,6 +1,6 @@ //Jin: cuda_device_runtime.h //Defines CUDA device runtime APIs for CDP support - +#if (CUDART_VERSION >= 5000) #pragma once void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); @@ -8,3 +8,4 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); void launch_all_device_kernels(); void launch_one_device_kernel(); +#endif diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 5d909d3..e68f9fd 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -1414,6 +1414,7 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) return; } +#if (CUDART_VERSION >= 5000) //Jin: handle device runtime apis for CDP else if(fname == "cudaGetParameterBufferV2") { gpgpusim_cuda_getParameterBufferV2(pI, thread, target_func); @@ -1427,6 +1428,7 @@ void call_impl( const ptx_instruction *pI, ptx_thread_info *thread ) gpgpusim_cuda_streamCreateWithFlags(pI, thread, target_func); return; } +#endif // read source arguements into register specified in declaration of function arg_buffer_list_t arg_values; diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 0b4b2f6..363fe5a 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1513,8 +1513,10 @@ void gpgpu_sim::cycle() try_snap_shot(gpu_sim_cycle); spill_log_to_file (stdout, 0, gpu_sim_cycle); +#if (CUDART_VERSION >= 5000) //launch device kernel launch_one_device_kernel(); +#endif } } diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index b9caf18..59a2d8b 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -846,7 +846,7 @@ void scheduler_unit::cycle() while( !warp(warp_id).waiting() && !warp(warp_id).ibuffer_empty() && (checked < max_issue) && (checked <= issued) && (issued < max_issue) ) { const warp_inst_t *pI = warp(warp_id).ibuffer_next_inst(); //Jin: handle cdp latency; - if(pI->m_is_cdp && warp(warp_id).m_cdp_latency > 0) { + if(pI && pI->m_is_cdp && warp(warp_id).m_cdp_latency > 0) { assert(warp(warp_id).m_cdp_dummy); warp(warp_id).m_cdp_latency--; break; -- cgit v1.3