From 742c4dc4c2c85329754043d38c60b2a37fefdaa1 Mon Sep 17 00:00:00 2001 From: Amruth Date: Fri, 23 Mar 2018 19:13:00 -0700 Subject: dynamic pdom analysis at runtime --- src/cuda-sim/ptx_ir.h | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) (limited to 'src/cuda-sim/ptx_ir.h') diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 9ad1571..26a2839 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1178,7 +1178,7 @@ public: //Muchnick's Adv. Compiler Design & Implemmntation Fig 7.15 void find_ipostdominators( ); void print_ipostdominators(); - + void do_pdom(); //function to call pdom analysis unsigned get_num_reconvergence_pairs(); @@ -1274,6 +1274,8 @@ public: m_local_mem_framesize = sz; } bool is_entry_point() const { return m_entry_point; } + bool is_pdom_set() const { return pdom_done; } //return pdom flag + void set_pdom() { pdom_done = true; } //set pdom flag private: unsigned m_uid; @@ -1281,6 +1283,7 @@ private: bool m_entry_point; bool m_extern; bool m_assembled; + bool pdom_done; //flag to check whether pdom is completed or not std::string m_name; ptx_instruction **m_instr_mem; unsigned m_start_PC; -- cgit v1.3 From deee9038d3d67e60f106776be3dd0a846dd11df9 Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Sun, 1 Apr 2018 17:31:14 -0700 Subject: fix regressions -- move call to pre_decode into do_pdom --- src/cuda-sim/cuda-sim.cc | 4 ++-- src/cuda-sim/ptx_ir.cc | 57 ++++++++++++++++++++++++++++-------------------- src/cuda-sim/ptx_ir.h | 2 ++ 3 files changed, 37 insertions(+), 26 deletions(-) (limited to 'src/cuda-sim/ptx_ir.h') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 987e3f2..dce35ca 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -252,7 +252,7 @@ void function_info::ptx_assemble() target.set_type(label_t); } } - + m_n = n; printf(" done.\n"); fflush(stdout); @@ -282,7 +282,6 @@ void function_info::ptx_assemble() print_postdominators(); print_ipostdominators(); } -#endif printf("GPGPU-Sim PTX: pre-decoding instructions for \'%s\'...\n", m_name.c_str() ); for ( unsigned ii=0; ii < n; ii += m_instr_mem[ii]->inst_size() ) { // handle branch instructions @@ -293,6 +292,7 @@ void function_info::ptx_assemble() fflush(stdout); m_assembled = true; +#endif } addr_t shared_to_generic( unsigned smid, addr_t addr ) diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 17e91df..be25dbe 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -577,30 +577,39 @@ bool function_info::connect_break_targets() //connecting break instructions with return modified; } -void function_info::do_pdom() { - create_basic_blocks(); - connect_basic_blocks(); - bool modified = false; - do { - find_dominators(); - find_idominators(); - modified = connect_break_targets(); - } while (modified == true); - - if ( g_debug_execution>=50 ) { - print_basic_blocks(); - print_basic_block_links(); - print_basic_block_dot(); - } - if ( g_debug_execution>=2 ) { - print_dominators(); - } - find_postdominators(); - find_ipostdominators(); - if ( g_debug_execution>=50 ) { - print_postdominators(); - print_ipostdominators(); - } +void function_info::do_pdom() +{ + create_basic_blocks(); + connect_basic_blocks(); + bool modified = false; + do { + find_dominators(); + find_idominators(); + modified = connect_break_targets(); + } while (modified == true); + + if ( g_debug_execution>=50 ) { + print_basic_blocks(); + print_basic_block_links(); + print_basic_block_dot(); + } + if ( g_debug_execution>=2 ) { + print_dominators(); + } + find_postdominators(); + find_ipostdominators(); + if ( g_debug_execution>=50 ) { + print_postdominators(); + print_ipostdominators(); + } + printf("GPGPU-Sim PTX: pre-decoding instructions for \'%s\'...\n", m_name.c_str() ); + for ( unsigned ii=0; ii < m_n; ii += m_instr_mem[ii]->inst_size() ) { // handle branch instructions + ptx_instruction *pI = m_instr_mem[ii]; + pI->pre_decode(); + } + printf("GPGPU-Sim PTX: ... done pre-decoding instructions for \'%s\'.\n", m_name.c_str() ); + fflush(stdout); + m_assembled = true; } void intersect( std::set &A, const std::set &B ) { diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 26a2839..85b2a3b 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1308,6 +1308,8 @@ private: //parameter size for device kernels int m_args_aligned_size; + + addr_t m_n; // offset in m_instr_mem (used in do_pdom) }; class arg_buffer_t { -- cgit v1.3 From 8ea33c977b26cfe96beb98cdda289b81b8fda899 Mon Sep 17 00:00:00 2001 From: Amruth Date: Sat, 14 Apr 2018 17:17:11 -0700 Subject: solving alignment issue --- libcuda/cuda_runtime_api.cc | 1 + src/cuda-sim/cuda-sim.cc | 7 ++++++- src/cuda-sim/ptx.l | 5 +++-- src/cuda-sim/ptx_ir.h | 1 + 4 files changed, 11 insertions(+), 3 deletions(-) (limited to 'src/cuda-sim/ptx_ir.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index d6ce4e0..ef46f00 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1923,6 +1923,7 @@ void cuobjdumpParseBinary(unsigned int handle){ if (capability > max_capability) max_capability = capability; } if (max_capability > 20) printf("WARNING: No guarantee that PTX will be parsed for SM version %u\n", max_capability); + if (max_capability == 0) max_capability=context->get_device()->get_gpgpu()->get_config().get_forced_max_capability(); cuobjdumpPTXSection* ptx = NULL; const char* pre_load = getenv("CUOBJDUMP_SIM_FILE"); diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index dce35ca..2c87031 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1139,8 +1139,13 @@ void function_info::finalize( memory_space *param_mem ) } // copy the parameter over word-by-word so that parameter that crosses a memory page can be copied over //Jin: copy parameter using aligned rules + const type_info *paramtype = param->type(); + int align_amount = paramtype->get_key().get_alignment_spec(); + align_amount = (align_amount == -1) ? size : align_amount; + param_address = (param_address + align_amount - 1) / align_amount * align_amount; //aligned + const size_t word_size = 4; - param_address = (param_address + size - 1) / size * size; //aligned with size + //param_address = (param_address + size - 1) / size * size; //aligned with size for (size_t idx = 0; idx < size; idx += word_size) { const char *pdata = reinterpret_cast(param_value.pdata) + idx; // cast to char * for ptr arithmetic param_mem->write(param_address + idx, word_size, pdata,NULL,NULL); diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index 1b5d7f6..908c5be 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -36,7 +36,8 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include "ptx.tab.h" #include -char linebuf[1024]; +#define LINEBUF_SIZE (64*1024) +char linebuf[LINEBUF_SIZE]; unsigned col = 0; #define TC col+=strlen(ptx_text); #define CHECK_UNSIGNED \ @@ -384,7 +385,7 @@ breakaddr TC; ptx_lval.int_value = BREAKADDR_OP; return OPCODE; "//"[^\n]* TC; // eat single -\n.* col=0; strncpy(linebuf, yytext + 1, 1024); yyless( 1 ); +\n.* col=0; strncpy(linebuf, yytext + 1, LINEBUF_SIZE); yyless( 1 ); " " TC; "\t" TC; diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 85b2a3b..6731763 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -91,6 +91,7 @@ public: bool is_tex() const { return m_space_spec == tex_space;} bool is_func_addr() const { return m_is_function?true:false; } int scalar_type() const { return m_scalar_type_spec;} + int get_alignment_spec() const { return m_alignment_spec;} unsigned type_decode( size_t &size, int &t ) const; static unsigned type_decode( int type, size_t &size, int &t ); memory_space_t get_memory_space() const { return m_space_spec; } -- cgit v1.3 From 60017ca1ddbe844a93f631fe2b86bc4101850037 Mon Sep 17 00:00:00 2001 From: Amruth Date: Thu, 19 Apr 2018 18:13:44 -0700 Subject: Crash when array pointers are passed --- README | 27 ++++++++++++++++++++++++++- src/cuda-sim/instructions.cc | 4 +++- src/cuda-sim/ptx_ir.h | 13 +++++++++++++ 3 files changed, 42 insertions(+), 2 deletions(-) (limited to 'src/cuda-sim/ptx_ir.h') diff --git a/README b/README index 543177c..bf5aa62 100644 --- a/README +++ b/README @@ -235,6 +235,14 @@ The documentation resides at doc/doxygen/html. Step 3: Run ============ +Before we run, we need to make sure the application's executable file is dynamically linked to CUDA runtime library. This can be done during compilation of your program by introducing the nvcc flag "--cudart shared" in makefile (quotes should be excluded). + +To confirm the same, type the follwoing command: + +ldd + +You should see that your application is using libcudart.so file in GPGPUSim directory. + Copy the contents of configs/QuadroFX5800/ or configs/GTX480/ to your application's working directory. These files configure the microarchitecture models to resemble the respective GPGPU architectures. @@ -348,7 +356,24 @@ identify any compile time or runtime errors that occur due to the code merging process. -** Debugging failing GPGPU-Sim Regressions ** +4. MISCELLANEOUS + +4.1 Speeding up the execution + +Some applications take several hours to execute on GPGPUSim. This is because the simulator has to dump the PTX, analyze them and get resource usage statistics. This can be avoided everytime we execute the program in the following way: + +Step 1: Execute the program by enabling “-save_embedded_ptx 1” in config file, execute the code and let cuobjdump command dump all necessary files. After this process, you will get 2 new files namely: _cuobjdump_complete_output_ and _1.ptx + +Step 2: Create new environment variables or include the below in your .bashrc file: + a. export PTX_SIM_USE_PTX_FILE=_1.ptx + b. export PTX_SIM_KERNELFILE=_1.ptx + c. export CUOBJDUMP_SIM_FILE=_cuobjdump_complete_output_ + +Step 3: Disable -save_embedded_ptx flag, execute the code again. This will skip the dumping by cuobjdump and directly goes to executing the program thus saving time. + + +4.2 Debugging failing GPGPU-Sim Regressions + Credits: Tor M Aamodt To debug failing GPGPU-Sim regression tests you need to run them locally. The fastest way to do this, assuming you are working with GPGPU-Sim versions more recent than the GPGPU-Sim dev branch circa March 28, 2018 (commit hash 2221d208a745a098a60b0d24c05007e92aaba092), is to install Docker. The instructions below were tested with Docker CE version 18.03 on Ubuntu and Mac OS. Docker will enable you to run the same set of regressions used by GPGPU-Sim when submitting a pull request to https://github.com/gpgpu-sim/gpgpu-sim_distribution and also allow you to log in and launch GPGPU-Sim in gdb so you can inspect failures. diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 0025c52..e53aaab 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -154,7 +154,9 @@ 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 = (size_t)op.get_symbol()->get_pc(); + result.u64 = (size_t)op.get_symbol()->get_pc(); + } else if ( op.is_param_kernel()) { + result.u64 = op.get_symbol()->get_address(); } else { const char *name = op.name().c_str(); printf("GPGPU-Sim PTX: ERROR ** get_operand_value : unknown operand type for %s\n", name ); diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 6731763..58d5f49 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -164,6 +164,7 @@ public: m_is_global = false; m_is_local = false; m_is_param_local = false; + m_is_param_kernel = false; m_is_tex = false; m_is_func_addr = false; m_reg_num_valid = false; @@ -177,6 +178,7 @@ public: if ( type ) m_is_global = type->get_key().is_global(); if ( type ) m_is_local = type->get_key().is_local(); if ( type ) m_is_param_local = type->get_key().is_param_local(); + if ( type ) m_is_param_kernel = type->get_key().is_param_kernel(); if ( type ) m_is_tex = type->get_key().is_tex(); if ( type ) m_is_func_addr = type->get_key().is_func_addr(); } @@ -227,6 +229,7 @@ public: bool is_global() const { return m_is_global;} bool is_local() const { return m_is_local;} bool is_param_local() const { return m_is_param_local; } + bool is_param_kernel() const { return m_is_param_kernel; } bool is_tex() const { return m_is_tex;} bool is_func_addr() const { return m_is_func_addr; } bool is_reg() const @@ -284,6 +287,7 @@ private: bool m_is_global; bool m_is_local; bool m_is_param_local; + bool m_is_param_kernel; bool m_is_tex; bool m_is_func_addr; unsigned m_reg_num; @@ -400,6 +404,8 @@ public: m_type = symbolic_t; } else if ( addr->is_param_local() ) { m_type = symbolic_t; + } else if ( addr->is_param_kernel() ) { + m_type = symbolic_t; } else if ( addr->is_tex() ) { m_type = symbolic_t; } else if ( addr->is_func_addr() ) { @@ -676,6 +682,13 @@ public: return m_value.m_symbolic->type()->get_key().is_param_local(); } + bool is_param_kernel() const + { + if ( m_type != symbolic_t ) + return false; + return m_value.m_symbolic->type()->get_key().is_param_kernel(); + } + bool is_vector() const { if ( m_vector) return true; -- cgit v1.3 From 37dc5311a75548b848a33a3b7369ce4bee64b444 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Thu, 14 Jun 2018 20:15:52 -0700 Subject: bfe bug fix --- .gitignore | 5 +++++ src/cuda-sim/instructions.cc | 30 +++++++++++++++--------------- src/cuda-sim/ptx_ir.h | 2 +- 3 files changed, 21 insertions(+), 16 deletions(-) (limited to 'src/cuda-sim/ptx_ir.h') diff --git a/.gitignore b/.gitignore index 0e2a898..e4e1631 100644 --- a/.gitignore +++ b/.gitignore @@ -34,3 +34,8 @@ cscope* tags regression.sh +#gcov +*.gcov +*.gcda +*.gcno + diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 08bf528..c77e4da 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -1353,45 +1353,45 @@ void bfe_impl( const ptx_instruction *pI, ptx_thread_info *thread ) 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 src = 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); + ptx_reg_t data; 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; + data.u32 = src.u32 >> pos; mask = 0xFFFFFFFF >> (32 - len); - d &= mask; + data.u32 &= mask; break; } case U64_TYPE: { unsigned long mask; - d = a.u64 >> pos; + data.u64 = src.u64 >> pos; mask = 0xFFFFFFFFFFFFFFFF >> (64 - len); - d &= mask; + data.u64 &= 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; + unsigned sbit = len == 0 ? 0 : (src.s32 >> min) & 0x1; + data.s32 = src.s32 >> pos; if (sbit > 0) { mask = 0xFFFFFFFF << len; - d |= mask; + data.s32 |= mask; } else { mask = 0xFFFFFFFF >> (32 - len); - d &= mask; + data.s32 &= mask; } break; } @@ -1399,17 +1399,17 @@ void bfe_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { 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; + unsigned sbit = len == 0 ? 0 : (src.s64 >> min) & 0x1; + data.s64 = src.s64 >> pos; if (sbit > 0) { mask = 0xFFFFFFFFFFFFFFFF << len; - d |= mask; + data.s64 |= mask; } else { mask = 0xFFFFFFFFFFFFFFFF >> (64 - len); - d &= mask; + data.s64 &= mask; } break; } @@ -1418,7 +1418,7 @@ void bfe_impl( const ptx_instruction *pI, ptx_thread_info *thread ) abort(); return; } - thread->set_operand_value(dst,d, i_type, thread, pI); + thread->set_operand_value(dst, data, i_type, thread, pI); } void bfi_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); } diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 58d5f49..5b68fcf 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -745,7 +745,7 @@ public: { ptx_reg_t result; switch ( m_type ) { - case int_t: result.s32 = m_value.m_int; break; + case int_t: result.s64 = m_value.m_int; break; case float_op_t: result.f32 = m_value.m_float; break; case double_op_t: result.f64 = m_value.m_double; break; case unsigned_t: result.u32 = m_value.m_unsigned; break; -- cgit v1.3 From 76a124e9186b9574858238d423b9c5ce715f2c32 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Thu, 21 Jun 2018 17:35:07 -0700 Subject: WIP adding support for PTX JIT and dumping params to cudaLaunches --- libcuda/cuda_runtime_api.cc | 168 ++++++++++++++++++++++++++++++++++++++++--- src/cuda-sim/cuda-sim.cc | 81 +++++++++++++++++++++ src/cuda-sim/cuda-sim.h | 1 + src/cuda-sim/instructions.cc | 1 + src/cuda-sim/ptx_ir.cc | 8 +++ src/cuda-sim/ptx_ir.h | 2 + 6 files changed, 252 insertions(+), 9 deletions(-) (limited to 'src/cuda-sim/ptx_ir.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 08ee413..300fa28 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -126,6 +126,9 @@ #include "host_defines.h" #include "builtin_types.h" #include "driver_types.h" +#if (CUDART_VERSION >= 8000) +#include "cuda.h" +#endif #if (CUDART_VERSION < 8000) #include "__cudaFatFormat.h" #endif @@ -287,6 +290,34 @@ struct CUctx_st { } } + void register_hostFun_function( const char*hostFun, function_info* f){ + m_kernel_lookup[hostFun] = f; + } + + dim3 get_blockdim(const char *hostFun) + { + std::map::iterator i=m_hostFun_blockdim.find(hostFun); + assert( i != m_hostFun_blockdim.end() ); + return i->second; + } + + dim3 get_griddim(const char *hostFun) + { + std::map::iterator i=m_hostFun_griddim.find(hostFun); + assert( i != m_hostFun_griddim.end() ); + return i->second; + } + + void set_blockdim(const char *hostFun, dim3 dims) + { + m_hostFun_blockdim[hostFun] = dims; + } + + void set_griddim(const char *hostFun, dim3 dims) + { + m_hostFun_griddim[hostFun] = dims; + } + function_info *get_kernel(const char *hostFun) { std::map::iterator i=m_kernel_lookup.find(hostFun); @@ -298,6 +329,8 @@ private: _cuda_device_id *m_gpu; // selected gpu std::map m_code; // fat binary handle => global symbol table unsigned m_last_fat_cubin_handle; + std::map m_hostFun_blockdim; + std::map m_hostFun_griddim; std::map m_kernel_lookup; // unique id (CUDA app function address) => kernel entry point struct gpgpu_ptx_sim_info m_binary_info; @@ -1300,16 +1333,15 @@ int CUDARTAPI __cudaSynchronizeThreads(void**, void*) * * *******************************************************************************/ -#if (CUDART_VERSION >= 3010) +#if (CUDART_VERSION >= 3010 && CUDART_VERSION < 8000) typedef struct CUuuid_st { /**< CUDA definition of UUID */ char bytes[16]; } CUuuid; -/** - * CUDA UUID types - */ -// typedef __device_builtin__ struct CUuuid_st cudaUUID_t; +#endif + +#if (CUDART_VERSION >= 3010) __host__ cudaError_t CUDARTAPI cudaGetExportTable(const void **ppExportTable, const cudaUUID_t *pExportTableId) { @@ -1958,10 +1990,10 @@ void cuobjdumpParseBinary(unsigned int handle){ symtab = gpgpu_ptx_sim_load_ptx_from_filename( ptx_filename.c_str() ); } } - name_symtab[fname] = symtab; - context->add_binary(symtab, handle); - load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); - load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); + name_symtab[fname] = symtab; + context->add_binary(symtab, handle); + load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); + load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); return; #endif @@ -2602,5 +2634,123 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, g_ptx_kernel_count++; fflush(stdout); + if(g_debug_execution >= 3){ + entry->debug_param(); + } + return result; } + +CUresult CUDAAPI cuLinkCreate(unsigned int numOptions, CUjit_option *options, void **optionValues, CUlinkState *stateOut) +{ + //currently do not support options or multiple CUlinkStates + return CUDA_SUCCESS; +} + +CUresult CUDAAPI cuLinkAddData(CUlinkState state, CUjitInputType type, void *data, size_t size, const char *name, + unsigned int numOptions, CUjit_option *options, void **optionValues) +{ + assert(type==CU_JIT_INPUT_PTX); + cuda_not_implemented(__my_func__,__LINE__); + return CUDA_ERROR_UNKNOWN; +} + +CUresult CUDAAPI cuLinkAddFile(CUlinkState state, CUjitInputType type, const char *path, + unsigned int numOptions, CUjit_option *options, void **optionValues) +{ + static bool addedFile = false; + if (addedFile){ + printf("GPGPU-Sim PTX: ERROR: cuLinkAddFile does not support multiple file"); + abort(); + } + + //blocking + assert(type==CU_JIT_INPUT_PTX); + CUctx_st *context = GPGPUSim_Context(); + char *file = getenv("PTX_JIT_PATH"); + if(file==NULL){ + printf("GPGPU-Sim PTX: ERROR: PTX_JIT_PATH has not been set"); + abort(); + } + strcat(file,path); + symbol_table *symtab = gpgpu_ptx_sim_load_ptx_from_filename( file ); + std::string fname(path); + name_symtab[fname] = symtab; + context->add_binary(symtab, 1); + load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); + load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); + addedFile = true; + return CUDA_SUCCESS; +} + +CUresult CUDAAPI cuLinkComplete(CUlinkState state, void **cubinOut, size_t *sizeOut) +{ + //all cuLink* function are implemented to block until completion so nothing to do here + return CUDA_SUCCESS; +} + +CUresult CUDAAPI cuLinkDestroy(CUlinkState state) +{ + //currently do not support options or multiple CUlinkStates + return CUDA_SUCCESS; +} + +CUresult CUDAAPI cuModuleLoadData(CUmodule *module, const void *image) +{ + //Currently do not support multiple modules + return CUDA_SUCCESS; +} + +CUresult CUDAAPI cuModuleGetFunction(CUfunction *hfunc, CUmodule hmod, const char *name) +{ + CUctx_st* context = GPGPUSim_Context(); + std::string key(name); + //only support one file + assert(name_symtab.size()==1); + symbol_table* symtab = name_symtab.begin()->second; + function_info* f = symtab->lookup_function( std::string(name) ); + //just need to add given pointer to map for cudaLaunch + context->register_hostFun_function( (const char*) hfunc, f); + return CUDA_SUCCESS; +} + +CUresult CUDAAPI cuModuleUnload(CUmodule hmod) +{ + //Currently do not support multiple modules + return CUDA_SUCCESS; +} + +CUresult CUDAAPI cuFuncSetBlockShape(CUfunction hfunc, int x, int y, int z) +{ + CUctx_st* context = GPGPUSim_Context(); + dim3 dims(x,y,z); + context->set_blockdim((const char *)hfunc, dims); + return CUDA_SUCCESS; +} + +CUresult CUDAAPI cuParamSetSize(CUfunction hfunc, unsigned int numbytes) +{ + //Nothing to do + return CUDA_SUCCESS; +} + +CUresult CUDAAPI cuParamSetv(CUfunction hfunc, int offset, void *ptr, unsigned int numbytes) +{ + cuda_not_implemented(__my_func__,__LINE__); + return CUDA_ERROR_UNKNOWN; +} + +CUresult CUDAAPI cuLaunchGrid(CUfunction f, int grid_width, int grid_height) +{ + cuda_not_implemented(__my_func__,__LINE__); + return CUDA_ERROR_UNKNOWN; + + CUctx_st* context = GPGPUSim_Context(); + const char *hostFun = (const char*) f; + dim3 dims(grid_width,grid_height,1); + context->set_griddim((const char *)f, dims); + cudaConfigureCall(context->get_griddim(hostFun), context->get_blockdim(hostFun), 0, NULL); + + cudaLaunch(hostFun); + return CUDA_SUCCESS; +} diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 34368ce..6875edd 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1226,6 +1226,87 @@ void function_info::list_param( FILE *fout ) const fflush(fout); } +void function_info::debug_param( ) const +{ + char filename[] = "params.txt"; + char buff[1024]; + snprintf(buff,1024,"c++filt %s > %s", get_name().c_str(), filename); + system(buff); + FILE *fp = fopen(filename, "r"); + fgets(buff, 1024, fp); + fclose(fp); + + std::string fn(buff); + size_t pos1, pos2; + pos1 = fn.find("("); + pos2 = fn.find(")"); + assert(pos2>pos1&&pos1>0); + strcpy(buff, fn.substr(pos1 + 1, pos2 - pos1 - 1).c_str()); + printf("params: %s\n", buff); + char *tok; + std::vector params; + tok = strtok(buff, ","); + while(tok!=NULL){ + std::string param(tok); + param.erase(0, param.find_first_not_of(" ")); + param.erase(param.find_last_not_of(" ")+1); + params.push_back(param); + tok = strtok(NULL, ","); + } + for (auto const& it : params){ + std::cout<::const_iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) { + const param_info &p = i->second; + std::string name = p.get_name(); + param_t param_value = p.get_value(); + if(params[i->first].find("const")!=std::string::npos){ + fprintf(fout, "Input: "); + } else { + fprintf(fout, "Input/output: "); + } + + symbol *param = m_symtab->lookup(name.c_str()); + addr_t param_addr = param->get_address(); + fprintf(fout, "%s: %#08x, ", name.c_str(), param_addr); + + if(params[i->first].find("int")!=std::string::npos){ + size_t len = param_value.size/sizeof(int); + int val[len]; + memcpy((void*) val, param_value.pdata+param_value.offset, param_value.size); + fprintf(fout, "val (int) = "); + for (unsigned i = 0; ifirst].find("float")!=std::string::npos){ + size_t len = param_value.size/sizeof(float); + float val[len]; + memcpy((void*) val, param_value.pdata+param_value.offset, param_value.size); + fprintf(fout, "val (float) = "); + for (unsigned i = 0; i bool ptx_debug_exec_dump_cond(int thd_uid, addr_t pc) { diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 958daba..9049a84 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -32,6 +32,7 @@ #include"../gpgpu-sim/shader.h" #include #include +#include #include #include"ptx_sim.h" diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index c77e4da..034a7b9 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -2753,6 +2753,7 @@ void mov_impl( const ptx_instruction *pI, ptx_thread_info *thread ) const operand_info &dst = pI->dst(); const operand_info &src1 = pI->src1(); unsigned i_type = pI->get_type(); + assert( src1.is_param_local() == 0 ); if( (src1.is_vector() || dst.is_vector()) && (i_type != BB64_TYPE) && (i_type != BB128_TYPE) && (i_type != FF64_TYPE) ) { // pack or unpack operation diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 016c600..482b9e0 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -257,6 +257,14 @@ bool symbol_table::add_function_decl( const char *name, int entry_point, functio return prior_decl; } +function_info *symbol_table::lookup_function( std::string name ) +{ + std::string key = std::string(name); + std::map::iterator it = m_function_info_lookup.find(key); + assert ( it != m_function_info_lookup.end() ); + return it->second; +} + type_info *symbol_table::add_type( memory_space_t space_spec, int scalar_type_spec, int vector_spec, int alignment_spec, int extern_spec ) { if( space_spec == param_space_unclassified ) diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 5b68fcf..341f9b7 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -313,6 +313,7 @@ public: symbol *add_variable( const char *identifier, const type_info *type, unsigned size, const char *filename, unsigned line ); void add_function( function_info *func, const char *filename, unsigned linenumber ); bool add_function_decl( const char *name, int entry_point, function_info **func_info, symbol_table **symbol_table ); + function_info *lookup_function(std::string name); type_info *add_type( memory_space_t space_spec, int scalar_type_spec, int vector_spec, int alignment_spec, int extern_spec ); type_info *add_type( function_info *func ); type_info *get_array_type( type_info *base_type, unsigned array_dim ); @@ -1256,6 +1257,7 @@ public: void finalize( memory_space *param_mem ); void param_to_shared( memory_space *shared_mem, symbol_table *symtab ); void list_param( FILE *fout ) const; + void debug_param() const; const struct gpgpu_ptx_sim_info* get_kernel_info () const { -- cgit v1.3 From 1f77b0720f69d684db83beb7a5513cd9461e3676 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Wed, 27 Jun 2018 15:26:00 -0700 Subject: WIP dump params --- debug_tools/WatchYourStep/ptxjitplus/ptxjitplus | Bin 139317 -> 0 bytes .../WatchYourStep/ptxjitplus/ptxjitplus.cpp | 112 ++++++++++++++++++--- libcuda/cuda_runtime_api.cc | 12 ++- src/cuda-sim/cuda-sim.cc | 59 +++++++++-- src/cuda-sim/ptx_ir.h | 2 +- 5 files changed, 157 insertions(+), 28 deletions(-) delete mode 100755 debug_tools/WatchYourStep/ptxjitplus/ptxjitplus (limited to 'src/cuda-sim/ptx_ir.h') diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus deleted file mode 100755 index ddc3435..0000000 Binary files a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus and /dev/null differ diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp index 9954e31..8645114 100644 --- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp +++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp @@ -111,6 +111,87 @@ void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUl checkCudaErrors(cuLinkDestroy(*lState)); } +void function_info::debug_param( ) const +{ + char filename[] = "params.txt"; + char buff[1024]; + snprintf(buff,1024,"c++filt %s > %s", get_name().c_str(), filename); + system(buff); + FILE *fp = fopen(filename, "r"); + fgets(buff, 1024, fp); + fclose(fp); + + std::string fn(buff); + size_t pos1, pos2; + pos1 = fn.find("("); + pos2 = fn.find(")"); + assert(pos2>pos1&&pos1>0); + strcpy(buff, fn.substr(pos1 + 1, pos2 - pos1 - 1).c_str()); + printf("params: %s\n", buff); + char *tok; + std::vector params; + tok = strtok(buff, ","); + while(tok!=NULL){ + std::string param(tok); + param.erase(0, param.find_first_not_of(" ")); + param.erase(param.find_last_not_of(" ")+1); + params.push_back(param); + tok = strtok(NULL, ","); + } + for (auto const& it : params){ + std::cout<::const_iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) { + const param_info &p = i->second; + std::string name = p.get_name(); + param_t param_value = p.get_value(); + if(params[i->first].find("const")!=std::string::npos){ + fprintf(fout, "Input: "); + } else { + fprintf(fout, "Input/output: "); + } + + symbol *param = m_symtab->lookup(name.c_str()); + addr_t param_addr = param->get_address(); + fprintf(fout, "%s: %#08x, ", name.c_str(), param_addr); + + if(params[i->first].find("int")!=std::string::npos){ + size_t len = param_value.size/sizeof(int); + int val[len]; + memcpy((void*) val, param_value.pdata+param_value.offset, param_value.size); + fprintf(fout, "val (int) = "); + for (unsigned i = 0; ifirst].find("float")!=std::string::npos){ + size_t len = param_value.size/sizeof(float); + float val[len]; + memcpy((void*) val, param_value.pdata+param_value.offset, param_value.size); + fprintf(fout, "val (float) = "); + for (unsigned i = 0; i >& param_data) { char *wys_exec_path = getenv("WYS_EXEC_PATH"); @@ -160,19 +241,17 @@ int main(int argc, char **argv) { const unsigned int nThreads = 256; const unsigned int nBlocks = 64; - const size_t memSize = nThreads * nBlocks * sizeof(int); CUmodule hModule = 0; CUfunction hKernel = 0; CUlinkState lState; - int *d_data = 0; - int *h_data = 0; int cuda_device = 0; cudaDeviceProp deviceProp; printf("[%s] - Starting...\n", sSDKname); std::vector< std::pair > param_data; + std::vector< std::pair > device_data; void* storedReg = initializeData(param_data); if (checkCmdLineFlag(argc, (const char **)argv, "device")) @@ -213,15 +292,7 @@ int main(int argc, char **argv) exit(EXIT_WAIVED); } - // Allocate memory on host and device (Runtime API) - // NOTE: The runtime API will create the GPU Context implicitly here - if ((h_data = (int *)malloc(memSize)) == NULL) - { - std::cerr << "Could not allocate host memory" << std::endl; - exit(EXIT_FAILURE); - } - checkCudaErrors(cudaMalloc(&d_data, memSize)); // JIT Compile the Kernel from PTX and get the Handles (Driver API) ptxJIT(argc, argv, &hModule, &hKernel, &lState); @@ -229,16 +300,29 @@ int main(int argc, char **argv) // Set the kernel parameters (Driver API) checkCudaErrors(cuFuncSetBlockShape(hKernel, nThreads, 1, 1)); - //param_data + + //Initialize param_data for kernel int paramOffset = 0; - checkCudaErrors(cuParamSetv(hKernel, paramOffset, &d_data, sizeof(d_data))); - paramOffset += sizeof(d_data); + for( std::vector< std::pair >::const_iterator i=param_data.begin(); i!=param_data.end(); i++ ) { + size_t memSize = nThreads * nBlocks * i->first; + unsigned char *d_data = 0; + checkCudaErrors(cudaMalloc((void**)&d_data, memSize)); + checkCudaErrors(cudaMemcpy(d_data,i->first,cudaMemcpyHostToDevice)); + checkCudaErrors(cuParamSetv(hKernel, paramOffset, &d_data, memSize)); + paramOffset += i->first; + } checkCudaErrors(cuParamSetSize(hKernel, paramOffset)); // Launch the kernel (Driver API_) checkCudaErrors(cuLaunchGrid(hKernel, nBlocks, 1)); std::cout << "CUDA kernel launched" << std::endl; + int *h_data = 0; + if ((h_data = (int *)malloc(memSize)) == NULL) + { + std::cerr << "Could not allocate host memory" << std::endl; + exit(EXIT_FAILURE); + } // Copy the result back to the host checkCudaErrors(cudaMemcpy(h_data, d_data, memSize, cudaMemcpyDeviceToHost)); diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index d2b855c..1e4f1df 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -150,6 +150,7 @@ std::map pinned_memory; //support for pinned memories added std::map pinned_memory_size; +std::map g_devPtr_Size; int no_of_ptx=0; std::map > version_filename; @@ -487,8 +488,10 @@ __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) { CUctx_st* context = GPGPUSim_Context(); *devPtr = context->get_device()->get_gpgpu()->gpu_malloc(size); - if(g_debug_execution >= 3) + if(g_debug_execution >= 3){ printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *devPtr); + g_devPtr_Size[*devPtr] = size; + } if ( *devPtr ) { return g_last_cudaError = cudaSuccess; } else { @@ -2374,8 +2377,10 @@ cudaError_t CUDARTAPI cudaHostGetDevicePointer(void **pDevice, void *pHost, unsi assert(i != pinned_memory_size.end()); size_t size = i->second; *pDevice = gpu->gpu_malloc(size); - if(g_debug_execution >= 3) + if(g_debug_execution >= 3){ printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *pDevice); + g_devPtr_Size[*pDevice] = size; + } if ( *pDevice ) { pinned_memory[pHost]=pDevice; //Copy contents in cpu to gpu @@ -2617,8 +2622,9 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, g_ptx_kernel_count++; fflush(stdout); + if(g_debug_execution >= 3){ - entry->debug_param(); + entry->debug_param(g_devPtr_Size, result->get_param_memory()); } return result; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index bdb1e8d..8795e1c 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1226,33 +1226,72 @@ void function_info::list_param( FILE *fout ) const fflush(fout); } -void function_info::debug_param() const +void function_info::debug_param(std::map devPtr_Size, memory_space *param_mem) const { static unsigned long counter = 0; - std::string gpgpusim_path(getenv("GPGPUSIM_ROOT")); - assert(!gpgpusim_path.empty()); - std::string command = "mkdir " + gpgpusim_path + "/debug_tools/WatchYourStep/data"; - system(command.c_str()); - - std::string filename(gpgpusim_path + "/debug_tools/WatchYourStep/data/params.config" + std::to_string(counter++)); std::vector< std::pair > param_data; + std::vector paramIsPointer; + + char * gpgpusim_path = getenv("GPGPUSIM_ROOT"); + assert(gpgpusim_path!=NULL); + std::string command = std::string("mkdir ") + gpgpusim_path + "/debug_tools/WatchYourStep/data"; + system(command.c_str()); + std::string filename(std::string(gpgpusim_path) + "/debug_tools/WatchYourStep/data/params.config" + std::to_string(counter++)); + + //initialize paramList + char buff[1024]; + std::string filename_c(filename+"_c"); + snprintf(buff,1024,"c++filt %s > %s", get_name().c_str(), filename_c.c_str()); + system(buff); + FILE *fp = fopen(filename_c.c_str(), "r"); + fgets(buff, 1024, fp); + fclose(fp); + std::string fn(buff); + size_t pos1, pos2; + pos1 = fn.find("("); + pos2 = fn.find(")"); + assert(pos2>pos1&&pos1>0); + strcpy(buff, fn.substr(pos1 + 1, pos2 - pos1 - 1).c_str()); + char *tok; + tok = strtok(buff, ","); + while(tok!=NULL){ + std::string param(tok); + if(param.find("*")!=std::string::npos){ + paramIsPointer.push_back(true); + }else{ + paramIsPointer.push_back(false); + } + tok = strtok(NULL, ","); + } for( std::map::const_iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) { const param_info &p = i->second; + std::string name = p.get_name(); + symbol *param = m_symtab->lookup(name.c_str()); + addr_t param_addr = param->get_address(); param_t param_value = p.get_value(); - unsigned char val[param_value.size]; - memcpy((void*) val, (void*)((char*)param_value.pdata+param_value.offset), param_value.size); - param_data.push_back(std::pair(param_value.size,val)); +// if (paramIsPointer[i->first]){ +// assert(param_value.size==8); +// }else{ + unsigned char val[param_value.size]; + param_mem->read(param_addr,param_value.size,(void*)val); + param_data.push_back(std::pair(param_value.size,val)); + //} } FILE *fout = fopen (filename.c_str(), "w"); fprintf(fout, "%s\n", get_name().c_str()); + size_t index = 0; for( std::vector< std::pair >::const_iterator i=param_data.begin(); i!=param_data.end(); i++ ) { + if (paramIsPointer[index]){ + fprintf(fout, "*"); + } fprintf(fout, "%lu :", i->first); for (size_t j = 0; jfirst; j++){ fprintf(fout, " %u", i->second[j]); } fprintf(fout, "\n"); + index++; } fflush(fout); fclose(fout); diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 341f9b7..43907d3 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1257,7 +1257,7 @@ public: void finalize( memory_space *param_mem ); void param_to_shared( memory_space *shared_mem, symbol_table *symtab ); void list_param( FILE *fout ) const; - void debug_param() const; + void debug_param(std::map devPtr_Size, memory_space *param_mem) const; const struct gpgpu_ptx_sim_info* get_kernel_info () const { -- cgit v1.3 From 87cc31c53123f918af3250085da72a3ab8361cf5 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Thu, 28 Jun 2018 13:13:59 -0700 Subject: Tests to find conditions that a value is a pointer and new mallocPtr_Size --- libcuda/cuda_runtime_api.cc | 15 +++++++++------ src/cuda-sim/cuda-sim.cc | 23 +++++++++++++++++++---- src/cuda-sim/ptx_ir.h | 2 +- 3 files changed, 29 insertions(+), 11 deletions(-) (limited to 'src/cuda-sim/ptx_ir.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 1e4f1df..28ff1c4 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -150,7 +150,7 @@ std::map pinned_memory; //support for pinned memories added std::map pinned_memory_size; -std::map g_devPtr_Size; +std::map g_mallocPtr_Size; int no_of_ptx=0; std::map > version_filename; @@ -322,8 +322,8 @@ public: } kernel_config() { - m_GridDim=NULL; - m_BlockDim=NULL; + m_GridDim=dim3(-1,-1,-1); + m_BlockDim=dim3(-1,-1,-1); m_sharedMem=0; m_stream =NULL; } @@ -490,7 +490,7 @@ __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) *devPtr = context->get_device()->get_gpgpu()->gpu_malloc(size); if(g_debug_execution >= 3){ printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *devPtr); - g_devPtr_Size[*devPtr] = size; + g_mallocPtr_Size[(unsigned long long)*devPtr] = size; } if ( *devPtr ) { return g_last_cudaError = cudaSuccess; @@ -1075,6 +1075,9 @@ __host__ cudaError_t CUDARTAPI cudaSetupArgument(const void *arg, size_t size, s gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); kernel_config &config = g_cuda_launch_stack.back(); config.set_arg(arg,size,offset); + printf("GPGPU-Sim PTX: Setting up arguments for %zu bytes starting at 0x%llx..\n",size, (unsigned long long) arg); + assert(size!=8||g_mallocPtr_Size.find(*(long long*)arg)!=g_mallocPtr_Size.end()); + //assert((*(long long*)arg)==0||g_mallocPtr_Size.find(*(long long*)arg)!=g_mallocPtr_Size.end()); return g_last_cudaError = cudaSuccess; } @@ -2379,7 +2382,7 @@ cudaError_t CUDARTAPI cudaHostGetDevicePointer(void **pDevice, void *pHost, unsi *pDevice = gpu->gpu_malloc(size); if(g_debug_execution >= 3){ printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *pDevice); - g_devPtr_Size[*pDevice] = size; + g_mallocPtr_Size[(unsigned long long)*pDevice] = size; } if ( *pDevice ) { pinned_memory[pHost]=pDevice; @@ -2624,7 +2627,7 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, if(g_debug_execution >= 3){ - entry->debug_param(g_devPtr_Size, result->get_param_memory()); + entry->debug_param(g_mallocPtr_Size, result->get_param_memory()); } return result; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 8795e1c..ab607e0 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1226,7 +1226,7 @@ void function_info::list_param( FILE *fout ) const fflush(fout); } -void function_info::debug_param(std::map devPtr_Size, memory_space *param_mem) const +void function_info::debug_param(std::map mallocPtr_Size, memory_space *param_mem) const { static unsigned long counter = 0; std::vector< std::pair > param_data; @@ -1248,8 +1248,8 @@ void function_info::debug_param(std::map devPtr_Size, memory_spa fclose(fp); std::string fn(buff); size_t pos1, pos2; - pos1 = fn.find("("); - pos2 = fn.find(")"); + pos1 = fn.find_last_of("("); + pos2 = fn.find(")", pos1); assert(pos2>pos1&&pos1>0); strcpy(buff, fn.substr(pos1 + 1, pos2 - pos1 - 1).c_str()); char *tok; @@ -1270,8 +1270,23 @@ void function_info::debug_param(std::map devPtr_Size, memory_spa symbol *param = m_symtab->lookup(name.c_str()); addr_t param_addr = param->get_address(); param_t param_value = p.get_value(); + + if (paramIsPointer[i->first]){ + assert(param_value.size==8&&mallocPtr_Size.find(*(long long*)param_value.pdata)!=mallocPtr_Size.end()); + }else{ + assert(!(param_value.size==8&&mallocPtr_Size.find(*(long long*)param_value.pdata)!=mallocPtr_Size.end())); + } // if (paramIsPointer[i->first]){ // assert(param_value.size==8); +// void *array_pointer = (void*)*val; +// assert(devPtr_Size.find(array_pointer)!=devPtr_Size.end()); +// size_t array_size = devPtr_Size[array_pointer]; +// unsigned char array_val[array_size]; +// memcpy((void*) array_val, array_pointer, array_size); +// param_data.push_back(std::pair(array_size,array_val)); +// +// assert(size!=8||g_mallocPtr_Size.find(*(long long*)arg)!=g_mallocPtr_Size.end()); +// // }else{ unsigned char val[param_value.size]; param_mem->read(param_addr,param_value.size,(void*)val); @@ -1295,7 +1310,7 @@ void function_info::debug_param(std::map devPtr_Size, memory_spa } fflush(fout); fclose(fout); - exit(0); + //exit(0); } template diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 43907d3..1eb9d65 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1257,7 +1257,7 @@ public: void finalize( memory_space *param_mem ); void param_to_shared( memory_space *shared_mem, symbol_table *symtab ); void list_param( FILE *fout ) const; - void debug_param(std::map devPtr_Size, memory_space *param_mem) const; + void debug_param(std::map mallocPtr_Size, memory_space *param_mem) const; const struct gpgpu_ptx_sim_info* get_kernel_info () const { -- cgit v1.3 From df58abd050457f062a2ecfe6202c85be5b939230 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Thu, 28 Jun 2018 14:59:07 -0700 Subject: dumps pointers by accessing global memory --- libcuda/cuda_runtime_api.cc | 4 +-- src/cuda-sim/cuda-sim.cc | 59 +++++++++------------------------------------ src/cuda-sim/ptx_ir.h | 2 +- 3 files changed, 14 insertions(+), 51 deletions(-) (limited to 'src/cuda-sim/ptx_ir.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 28ff1c4..53b4da8 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1077,7 +1077,6 @@ __host__ cudaError_t CUDARTAPI cudaSetupArgument(const void *arg, size_t size, s config.set_arg(arg,size,offset); printf("GPGPU-Sim PTX: Setting up arguments for %zu bytes starting at 0x%llx..\n",size, (unsigned long long) arg); assert(size!=8||g_mallocPtr_Size.find(*(long long*)arg)!=g_mallocPtr_Size.end()); - //assert((*(long long*)arg)==0||g_mallocPtr_Size.find(*(long long*)arg)!=g_mallocPtr_Size.end()); return g_last_cudaError = cudaSuccess; } @@ -2624,10 +2623,9 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, entry->finalize(result->get_param_memory()); g_ptx_kernel_count++; fflush(stdout); - if(g_debug_execution >= 3){ - entry->debug_param(g_mallocPtr_Size, result->get_param_memory()); + entry->debug_param(g_mallocPtr_Size, result->get_param_memory(), (gpgpu_t *) context->get_device()->get_gpgpu()); } return result; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index ab607e0..bade750 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1226,7 +1226,7 @@ void function_info::list_param( FILE *fout ) const fflush(fout); } -void function_info::debug_param(std::map mallocPtr_Size, memory_space *param_mem) const +void function_info::debug_param(std::map mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu) const { static unsigned long counter = 0; std::vector< std::pair > param_data; @@ -1234,36 +1234,10 @@ void function_info::debug_param(std::map mallocPtr_S char * gpgpusim_path = getenv("GPGPUSIM_ROOT"); assert(gpgpusim_path!=NULL); - std::string command = std::string("mkdir ") + gpgpusim_path + "/debug_tools/WatchYourStep/data"; - system(command.c_str()); +// std::string command = std::string("mkdir ") + gpgpusim_path + "/debug_tools/WatchYourStep/data"; +// system(command.c_str()); std::string filename(std::string(gpgpusim_path) + "/debug_tools/WatchYourStep/data/params.config" + std::to_string(counter++)); - //initialize paramList - char buff[1024]; - std::string filename_c(filename+"_c"); - snprintf(buff,1024,"c++filt %s > %s", get_name().c_str(), filename_c.c_str()); - system(buff); - FILE *fp = fopen(filename_c.c_str(), "r"); - fgets(buff, 1024, fp); - fclose(fp); - std::string fn(buff); - size_t pos1, pos2; - pos1 = fn.find_last_of("("); - pos2 = fn.find(")", pos1); - assert(pos2>pos1&&pos1>0); - strcpy(buff, fn.substr(pos1 + 1, pos2 - pos1 - 1).c_str()); - char *tok; - tok = strtok(buff, ","); - while(tok!=NULL){ - std::string param(tok); - if(param.find("*")!=std::string::npos){ - paramIsPointer.push_back(true); - }else{ - paramIsPointer.push_back(false); - } - tok = strtok(NULL, ","); - } - for( std::map::const_iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) { const param_info &p = i->second; std::string name = p.get_name(); @@ -1271,30 +1245,22 @@ void function_info::debug_param(std::map mallocPtr_S addr_t param_addr = param->get_address(); param_t param_value = p.get_value(); - if (paramIsPointer[i->first]){ - assert(param_value.size==8&&mallocPtr_Size.find(*(long long*)param_value.pdata)!=mallocPtr_Size.end()); + if(param_value.size==8&&mallocPtr_Size.find(*(unsigned long long*)param_value.pdata)!=mallocPtr_Size.end()){ + size_t array_size = mallocPtr_Size[*(unsigned long long*)param_value.pdata]; + unsigned char val[array_size]; + gpu->get_global_memory()->read(param_addr,array_size,(void*)val); + param_data.push_back(std::pair(array_size,val)); + paramIsPointer.push_back(true); }else{ - assert(!(param_value.size==8&&mallocPtr_Size.find(*(long long*)param_value.pdata)!=mallocPtr_Size.end())); - } -// if (paramIsPointer[i->first]){ -// assert(param_value.size==8); -// void *array_pointer = (void*)*val; -// assert(devPtr_Size.find(array_pointer)!=devPtr_Size.end()); -// size_t array_size = devPtr_Size[array_pointer]; -// unsigned char array_val[array_size]; -// memcpy((void*) array_val, array_pointer, array_size); -// param_data.push_back(std::pair(array_size,array_val)); -// -// assert(size!=8||g_mallocPtr_Size.find(*(long long*)arg)!=g_mallocPtr_Size.end()); -// -// }else{ unsigned char val[param_value.size]; param_mem->read(param_addr,param_value.size,(void*)val); param_data.push_back(std::pair(param_value.size,val)); - //} + paramIsPointer.push_back(false); + } } FILE *fout = fopen (filename.c_str(), "w"); + printf("Writing data to %s ...\n", filename.c_str()); fprintf(fout, "%s\n", get_name().c_str()); size_t index = 0; for( std::vector< std::pair >::const_iterator i=param_data.begin(); i!=param_data.end(); i++ ) { @@ -1310,7 +1276,6 @@ void function_info::debug_param(std::map mallocPtr_S } fflush(fout); fclose(fout); - //exit(0); } template diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 1eb9d65..449d615 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1257,7 +1257,7 @@ public: void finalize( memory_space *param_mem ); void param_to_shared( memory_space *shared_mem, symbol_table *symtab ); void list_param( FILE *fout ) const; - void debug_param(std::map mallocPtr_Size, memory_space *param_mem) const; + void debug_param(std::map mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu) const; const struct gpgpu_ptx_sim_info* get_kernel_info () const { -- cgit v1.3 From 1fddb06b153a3e5fbc255e89bb6861cce1319039 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Wed, 4 Jul 2018 12:52:42 -0700 Subject: dump and load block and grid size and launch --- .../WatchYourStep/ptxjitplus/ptxjitplus.cpp | 14 ++++++------- libcuda/cuda_runtime_api.cc | 24 ++++++++++++++++++++-- src/cuda-sim/cuda-sim.cc | 3 ++- src/cuda-sim/ptx_ir.h | 2 +- 4 files changed, 32 insertions(+), 11 deletions(-) (limited to 'src/cuda-sim/ptx_ir.h') diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp index 65dfa84..df0ff8d 100644 --- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp +++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp @@ -40,6 +40,7 @@ const char *sSDKname = "PTX Just In Time (JIT) Compilation (no-qatest)"; char *wys_exec_path; char *wys_exec_name; char *wys_launch_num; +dim3 gridDim, blockDim; std::string kernelName; void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUlinkState *lState) @@ -133,6 +134,7 @@ void initializeData(std::vector& param_data, std::vector< std::p printf("Processing :%s ...\n", buff); fflush(stdout); kernelName = std::string(buff); + fscanf(fin, "%u,%u,%u %u,%u,%u\n", &gridDim.x, &gridDim.y, &gridDim.z, &blockDim.x, &blockDim.y, &blockDim.z); //fill data structure to pass in params later while (!feof(fin)){ std::pair info; @@ -248,12 +250,6 @@ int main(int argc, char **argv) ptxJIT(argc, argv, &hModule, &hKernel, &lState); checkCudaErrors(cudaFree(d_tmp)); - - // Set the kernel parameters (Driver API) - - // TODO: automatically load these values in - checkCudaErrors(cuFuncSetBlockShape(hKernel, 128, 1, 1)); - //maps param number to pointer to device data std::map< size_t, unsigned char* > m_device_data; //Initialize param_data for kernel @@ -275,7 +271,8 @@ int main(int argc, char **argv) // Launch the kernel (Driver API_) // TODO: automatically load these values in - checkCudaErrors(cuLaunchGrid(hKernel, 1, 20)); + CUDAAPI cuLaunchKernel(hKernel, gridDim.x, gridDim.y, gridDim.z, blockDim.x, blockDim.y, blockDim.z, + 0, NULL, NULL, NULL); std::cout << "CUDA kernel launched" << std::endl; //maps param number to pointer to output data @@ -299,6 +296,9 @@ int main(int argc, char **argv) fprintf(fout, "param %zu: size = %zu, data = ", i->first,param_info[i->first].first); for (size_t j = 0; jfirst].first; j++){ fprintf(fout, " %u", i->second[j]); + if (j&&(!(j%20))){ + fprintf(fout, "\n"); + } } fprintf(fout, "\n"); } diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index cad4d87..410f15f 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -2624,7 +2624,7 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, fflush(stdout); if(g_debug_execution >= 3){ - entry->debug_param(g_mallocPtr_Size, result->get_param_memory(), (gpgpu_t *) context->get_device()->get_gpgpu()); + entry->debug_param(g_mallocPtr_Size, result->get_param_memory(), (gpgpu_t *) context->get_device()->get_gpgpu(), gridDim, blockDim); } return result; @@ -2744,7 +2744,6 @@ CUresult CUDAAPI cuParamSetv(CUfunction hfunc, int offset, void *ptr, unsigned i CUresult CUDAAPI cuLaunchGrid(CUfunction f, int grid_width, int grid_height) { - CUctx_st* context = GPGPUSim_Context(); const char *hostFun = (const char*) f; gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); kernel_config &config = g_cuda_launch_stack.back(); @@ -2754,4 +2753,25 @@ CUresult CUDAAPI cuLaunchGrid(CUfunction f, int grid_width, int grid_height) cudaLaunch(hostFun); return CUDA_SUCCESS; } + + +CUresult CUDAAPI cuLaunchKernel(CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, + unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, + unsigned int sharedMemBytes, CUstream hStream, void **kernelParams, void **extra) +{ + if (sharedMemBytes!=0||hStream!=NULL||kernelParams!=NULL||extra!=NULL){ + printf("GPGPU-Sim CUDA DRIVER API: Warning: Currently do not support \nsharedMemBytes, hStream, kernelParams, and extra parameters.\n"); + } + const char *hostFun = (const char*) f; + gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); + kernel_config &config = g_cuda_launch_stack.back(); + dim3 *d_b = new dim3(blockDimX, blockDimY, blockDimZ); + config.set_block_dim(d_b); + dim3 *d_g = new dim3(gridDimX, gridDimY, gridDimZ); + config.set_grid_dim(d_g); + + cudaLaunch(hostFun); + return CUDA_SUCCESS; +} + #endif diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 62aef6d..3003c4c 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1226,7 +1226,7 @@ void function_info::list_param( FILE *fout ) const fflush(fout); } -void function_info::debug_param(std::map mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu) const +void function_info::debug_param(std::map mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) const { static unsigned long counter = 0; std::vector< std::pair > param_data; @@ -1265,6 +1265,7 @@ void function_info::debug_param(std::map mallocPtr_S FILE *fout = fopen (filename.c_str(), "w"); printf("Writing data to %s ...\n", filename.c_str()); fprintf(fout, "%s\n", get_name().c_str()); + fprintf(fout, "%u,%u,%u %u,%u,%u\n", gridDim.x, gridDim.y, gridDim.z, blockDim.x, blockDim.y, blockDim.z); size_t index = 0; for( std::vector< std::pair >::const_iterator i=param_data.begin(); i!=param_data.end(); i++ ) { if (paramIsPointer[index]){ diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 449d615..b29621c 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1257,7 +1257,7 @@ public: void finalize( memory_space *param_mem ); void param_to_shared( memory_space *shared_mem, symbol_table *symtab ); void list_param( FILE *fout ) const; - void debug_param(std::map mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu) const; + void debug_param(std::map mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) const; const struct gpgpu_ptx_sim_info* get_kernel_info () const { -- cgit v1.3 From 7c4f9d7bff7b1725f3da6ddc8abf3f77a1cbe6f5 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Wed, 4 Jul 2018 16:41:20 -0700 Subject: dumps ptx kernels and scripts to launch all kernels --- debug_tools/WatchYourStep/ptxjitplus/launchkernels | 3 ++ .../WatchYourStep/ptxjitplus/ptxjitplus.cpp | 10 +++--- libcuda/cuda_runtime_api.cc | 2 +- src/cuda-sim/cuda-sim.cc | 37 ++++++++++++++++++++-- src/cuda-sim/ptx_ir.h | 2 +- 5 files changed, 44 insertions(+), 10 deletions(-) create mode 100644 debug_tools/WatchYourStep/ptxjitplus/launchkernels (limited to 'src/cuda-sim/ptx_ir.h') diff --git a/debug_tools/WatchYourStep/ptxjitplus/launchkernels b/debug_tools/WatchYourStep/ptxjitplus/launchkernels new file mode 100644 index 0000000..d2fd015 --- /dev/null +++ b/debug_tools/WatchYourStep/ptxjitplus/launchkernels @@ -0,0 +1,3 @@ +#Launches kernels from $1 to $2 +#Note must source this script +for num in $(eval echo {$1..$2}); do export WYS_LAUNCH_NUM=$num; echo Launching kernel $num...; ./ptxjitplus; done diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp index df0ff8d..1f094ba 100644 --- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp +++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp @@ -81,18 +81,16 @@ void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUl if (sizeof(void *)==4) { // Load the PTX from the string myPtx32 - printf("Loading myPtx32[] program\n"); - // PTX May also be loaded from file, as per below. - myErr = cuLinkAddData(*lState, CU_JIT_INPUT_PTX, (void *)myPtx32, strlen(myPtx32)+1, 0, 0, 0, 0); + printf("Loading myPtx32[] program...\n"); + printf("WARNING: 32-bit execution is untested"); } else { // Load the PTX from the string myPtx (64-bit) printf("Loading myPtx[] program\n"); - //myErr = cuLinkAddData(*lState, CU_JIT_INPUT_PTX, (void *)myPtx64, strlen(myPtx64)+1, 0, 0, 0, 0); - // PTX May also be loaded from file, as per below. - myErr = cuLinkAddFile(*lState, CU_JIT_INPUT_PTX, "myPtx64.ptx",0,0,0); } + std::string ptx_file (std::string("../data/ptx.config") + wys_launch_num); + myErr = cuLinkAddFile(*lState, CU_JIT_INPUT_PTX, ptx_file.c_str(),0,0,0); if (myErr != CUDA_SUCCESS) { diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 410f15f..b67ea85 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -2624,7 +2624,7 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, fflush(stdout); if(g_debug_execution >= 3){ - entry->debug_param(g_mallocPtr_Size, result->get_param_memory(), (gpgpu_t *) context->get_device()->get_gpgpu(), gridDim, blockDim); + entry->ptx_jit_config(g_mallocPtr_Size, result->get_param_memory(), (gpgpu_t *) context->get_device()->get_gpgpu(), gridDim, blockDim); } return result; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 3003c4c..fc4f82a 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1226,7 +1226,7 @@ void function_info::list_param( FILE *fout ) const fflush(fout); } -void function_info::debug_param(std::map mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) const +void function_info::ptx_jit_config(std::map mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) const { static unsigned long counter = 0; std::vector< std::pair > param_data; @@ -1236,7 +1236,7 @@ void function_info::debug_param(std::map mallocPtr_S assert(gpgpusim_path!=NULL); std::string command = std::string("mkdir ") + gpgpusim_path + "/debug_tools/WatchYourStep/data"; system(command.c_str()); - std::string filename(std::string(gpgpusim_path) + "/debug_tools/WatchYourStep/data/params.config" + std::to_string(counter++)); + std::string filename(std::string(gpgpusim_path) + "/debug_tools/WatchYourStep/data/params.config" + std::to_string(counter)); for( std::map::const_iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) { const param_info &p = i->second; @@ -1280,6 +1280,39 @@ void function_info::debug_param(std::map mallocPtr_S } fflush(fout); fclose(fout); + + //ptx config + char buff[1024]; + std::string ptx_config_fn(std::string(gpgpusim_path) + "/debug_tools/WatchYourStep/data/ptx.config" + std::to_string(counter)); + snprintf(buff, 1024, "grep -rn \".entry %s\" *.ptx | cut -d \":\" -f 1-2 > %s", get_name().c_str(), ptx_config_fn.c_str()); + system(buff); + FILE *fin = fopen(ptx_config_fn.c_str(), "r"); + char ptx_source[256]; + unsigned line_number; + int numscanned = fscanf(fin, "%[^:]:%u", ptx_source, &line_number); + assert(numscanned == 2); + fclose(fin); + snprintf(buff, 1024, "grep -rn \".version\" %s | cut -d \":\" -f 1 | xargs -I \"{}\" awk \"NR>={}&&NR<={}+2\" %s > %s", ptx_source, ptx_source, ptx_config_fn.c_str()); + system(buff); + fin = fopen(ptx_source, "r"); + assert(fin!=NULL); + fout = fopen(ptx_config_fn.c_str(), "a"); + assert(fout!=NULL); + for (unsigned i = 0; i diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index b29621c..e0b5e96 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1257,7 +1257,7 @@ public: void finalize( memory_space *param_mem ); void param_to_shared( memory_space *shared_mem, symbol_table *symtab ); void list_param( FILE *fout ) const; - void debug_param(std::map mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) const; + void ptx_jit_config(std::map mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) const; const struct gpgpu_ptx_sim_info* get_kernel_info () const { -- cgit v1.3 From 3a09960577b9c9cbcce8fedeef8874ccc533f378 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Wed, 18 Jul 2018 17:14:41 -0700 Subject: added c++filt that works better, param offset dumping, protection for failing system calls --- .../WatchYourStep/ptxjitplus/ptxjitplus.cpp | 2 +- libcuda/cuda_runtime_api.cc | 5 +- src/cuda-sim/cuda-sim.cc | 78 +++++++++++++++++++--- src/cuda-sim/ptx_ir.h | 2 +- 4 files changed, 74 insertions(+), 13 deletions(-) (limited to 'src/cuda-sim/ptx_ir.h') diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp index b7f9f2d..26a6d29 100644 --- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp +++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp @@ -183,6 +183,7 @@ void initializeData(std::vector& param_data, std::vector< std::p assert( err==1 ); params[i] = (unsigned char) val; } + //TODO: parse param offset param_info.push_back(info); param_data.push_back(params); err = fscanf(fin, "\n"); @@ -290,7 +291,6 @@ int main(int argc, char **argv) checkCudaErrors(cuParamSetSize(hKernel, paramOffset)); // Launch the kernel (Driver API_) - // TODO: automatically load these values in CUDAAPI cuLaunchKernel(hKernel, gridDim.x, gridDim.y, gridDim.z, blockDim.x, blockDim.y, blockDim.z, 0, NULL, paramKernels, NULL); std::cout << "CUDA kernel launched" << std::endl; diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 866fa3b..db4f58e 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1797,7 +1797,10 @@ void extract_code_using_cuobjdump(){ //Running cuobjdump using dynamic link to current process snprintf(command,1000,"md5sum %s ", app_binary.c_str()); printf("Running md5sum using \"%s\"\n", command); - system(command); + if(system(command)){ + std::cout << "Failed to execute: " << command << std::endl; + exit(1); + } // Running cuobjdump using dynamic link to current process // Needs the option '-all' to extract PTX from CDP-enabled binary extern bool g_cdp_enabled; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index d2f096f..a499ce9 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1226,10 +1226,11 @@ void function_info::list_param( FILE *fout ) const fflush(fout); } -void function_info::ptx_jit_config(std::map mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) const +void function_info::ptx_jit_config(std::map mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) { static unsigned long counter = 0; std::vector< std::pair > param_data; + std::vector offsets; std::vector paramIsPointer; char * gpgpusim_path = getenv("GPGPUSIM_ROOT"); @@ -1237,18 +1238,67 @@ void function_info::ptx_jit_config(std::map mallocPt char * wys_exec_path = getenv("WYS_EXEC_PATH"); assert(wys_exec_path!=NULL); std::string command = std::string("mkdir ") + gpgpusim_path + "/debug_tools/WatchYourStep/data"; - system(command.c_str()); + if(system(command.c_str())!=0){ + printf("WARNING: Failed to execute mkdir \n"); + printf("Problematic call: %s", command.c_str()); + abort(); + } std::string filename(std::string(gpgpusim_path) + "/debug_tools/WatchYourStep/data/params.config" + std::to_string(counter)); - for( std::map::const_iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) { - const param_info &p = i->second; + //initialize paramList + char buff[1024]; + std::string filename_c(filename+"_c"); + snprintf(buff,1024,"c++filt %s > %s", get_name().c_str(), filename_c.c_str()); + system(buff); + FILE *fp = fopen(filename_c.c_str(), "r"); + fgets(buff, 1024, fp); + fclose(fp); + std::string fn(buff); + size_t pos1, pos2; + pos1 = fn.find_last_of("("); + pos2 = fn.find(")", pos1); + assert(pos2>pos1&&pos1>0); + strcpy(buff, fn.substr(pos1 + 1, pos2 - pos1 - 1).c_str()); + char *tok; + tok = strtok(buff, ","); + std::string tmp; + while(tok!=NULL){ + std::string param(tok); + if(param.find("<")!=std::string::npos){ + assert(param.find(">")==std::string::npos); + assert(param.find("*")==std::string::npos); + tmp = param; + } else { + if (tmp.length()>0){ + tmp = ""; + assert(param.find(">")!=std::string::npos); + assert(param.find("<")==std::string::npos); + assert(param.find("*")==std::string::npos); + } + if(param.find("*")!=std::string::npos){ + paramIsPointer.push_back(true); + }else{ + paramIsPointer.push_back(false); + } + } + tok = strtok(NULL, ","); + } + + + for( std::map::iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) { + param_info &p = i->second; std::string name = p.get_name(); symbol *param = m_symtab->lookup(name.c_str()); addr_t param_addr = param->get_address(); param_t param_value = p.get_value(); + offsets.push_back((unsigned)p.get_offset()); - if(param_value.size==sizeof(void*) && mallocPtr_Size.find(*(unsigned long long*)param_value.pdata)!=mallocPtr_Size.end()){ + if (paramIsPointer[i->first]){ //is pointer + assert(param_value.size==8); + assert(param_value.size==sizeof(void*) && mallocPtr_Size.find(*(unsigned long long*)param_value.pdata)!=mallocPtr_Size.end()); + //TODO: check in middle of malloc'd memory for pointer + size_t array_size = mallocPtr_Size[*(unsigned long long*)param_value.pdata]; unsigned char* val = (unsigned char*) malloc(param_value.size); param_mem->read(param_addr,param_value.size,(void*)val); @@ -1257,7 +1307,7 @@ void function_info::ptx_jit_config(std::map mallocPt param_data.push_back(std::pair(array_size,array_val)); paramIsPointer.push_back(true); }else{ - unsigned char val[param_value.size]; + unsigned char* val = (unsigned char*) malloc(param_value.size); param_mem->read(param_addr,param_value.size,(void*)val); param_data.push_back(std::pair(param_value.size,val)); paramIsPointer.push_back(false); @@ -1277,6 +1327,8 @@ void function_info::ptx_jit_config(std::map mallocPt for (size_t j = 0; jfirst; j++){ fprintf(fout, " %u", i->second[j]); } + fprintf(fout, " : %u", offsets[index]); + free (i->second); fprintf(fout, "\n"); index++; } @@ -1284,10 +1336,13 @@ void function_info::ptx_jit_config(std::map mallocPt fclose(fout); //ptx config - char buff[1024]; std::string ptx_config_fn(std::string(gpgpusim_path) + "/debug_tools/WatchYourStep/data/ptx.config" + std::to_string(counter)); snprintf(buff, 1024, "grep -rn \".entry %s\" %s/*.ptx | cut -d \":\" -f 1-2 > %s", get_name().c_str(), wys_exec_path, ptx_config_fn.c_str()); - system(buff); + if (system(buff)!=0){ + printf("WARNING: Failed to execute grep to find ptx source \n"); + printf("Problematic call: %s", buff); + abort(); + } FILE *fin = fopen(ptx_config_fn.c_str(), "r"); char ptx_source[256]; unsigned line_number; @@ -1295,7 +1350,11 @@ void function_info::ptx_jit_config(std::map mallocPt assert(numscanned == 2); fclose(fin); snprintf(buff, 1024, "grep -rn \".version\" %s | cut -d \":\" -f 1 | xargs -I \"{}\" awk \"NR>={}&&NR<={}+2\" %s > %s", ptx_source, ptx_source, ptx_config_fn.c_str()); - system(buff); + if (system(buff)!=0){ + printf("WARNING: Failed to execute grep to find ptx header \n"); + printf("Problematic call: %s", buff); + abort(); + } fin = fopen(ptx_source, "r"); assert(fin!=NULL); printf("Writing data to %s ...\n", ptx_config_fn.c_str()); @@ -1318,7 +1377,6 @@ void function_info::ptx_jit_config(std::map mallocPt fflush(fout); fclose(fout); counter++; - //TODO: Free param_data } template diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index e0b5e96..ef4cf48 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1257,7 +1257,7 @@ public: void finalize( memory_space *param_mem ); void param_to_shared( memory_space *shared_mem, symbol_table *symtab ); void list_param( FILE *fout ) const; - void ptx_jit_config(std::map mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) const; + void ptx_jit_config(std::map mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) ; const struct gpgpu_ptx_sim_info* get_kernel_info () const { -- cgit v1.3 From a9af79ac84b69fa18dd395349b88f0d984f0a505 Mon Sep 17 00:00:00 2001 From: J Date: Tue, 7 Aug 2018 15:12:52 -0700 Subject: working fix for deadlock due to operand collector, parser changes to support culaunchkernel --- libcuda/cuda_runtime_api.cc | 74 ++++++--------------------------------------- src/cuda-sim/ptx_ir.h | 12 ++++++++ src/cuda-sim/ptx_parser.cc | 34 +++++++++++++++++++++ src/gpgpu-sim/shader.cc | 6 ++++ 4 files changed, 62 insertions(+), 64 deletions(-) (limited to 'src/cuda-sim/ptx_ir.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index c87c6c3..c12aaeb 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -3033,7 +3033,6 @@ CUresult CUDAAPI cuModuleGetFunction(CUfunction *hfunc, CUmodule hmod, const cha function_info* f = symtab->lookup_function( std::string(name) ); //just need to add given pointer to map for cudaLaunch context->register_hostFun_function( (const char*) hfunc, f); - g_cuda_launch_stack.push_back( kernel_config() ); *hfunc = (CUfunction)hfunc; return CUDA_SUCCESS; } @@ -3047,60 +3046,6 @@ CUresult CUDAAPI cuModuleUnload(CUmodule hmod) return CUDA_SUCCESS; } -CUresult CUDAAPI cuFuncSetBlockShape(CUfunction hfunc, int x, int y, int z) -{ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); - kernel_config &config = g_cuda_launch_stack.back(); - dim3 *d = new dim3(x,y,z); - config.set_block_dim(d); - - return CUDA_SUCCESS; -} - -CUresult CUDAAPI cuParamSetSize(CUfunction hfunc, unsigned int numbytes) -{ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - //check if size matches given args - gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); - kernel_config &config = g_cuda_launch_stack.back(); - gpgpu_ptx_sim_arg_list_t args = config.get_args(); - size_t total_size = 0; - for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) { - total_size += a->m_nbytes; - } - return (numbytes==total_size) ? CUDA_SUCCESS : CUDA_ERROR_INVALID_VALUE; -} - -CUresult CUDAAPI cuParamSetv(CUfunction hfunc, int offset, void *ptr, unsigned int numbytes) -{ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - cudaSetupArgument((const void *) ptr, (size_t) numbytes, (size_t) offset); - return CUDA_SUCCESS; -} - -CUresult CUDAAPI cuLaunchGrid(CUfunction f, int grid_width, int grid_height) -{ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - const char *hostFun = (const char*) f; - gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); - kernel_config &config = g_cuda_launch_stack.back(); - dim3 *d = new dim3(grid_width,grid_height,1); - config.set_grid_dim(d); - - cudaLaunch(hostFun); - return CUDA_SUCCESS; -} - - CUresult CUDAAPI cuLaunchKernel(CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, CUstream hStream, void **kernelParams, void **extra) @@ -3108,17 +3053,18 @@ CUresult CUDAAPI cuLaunchKernel(CUfunction f, unsigned int gridDimX, unsigned in if(g_debug_execution >= 3){ announce_call(__my_func__); } - if (sharedMemBytes!=0||hStream!=NULL||kernelParams!=NULL||extra!=NULL){ - printf("GPGPU-Sim CUDA DRIVER API: Warning: Currently do not support \nsharedMemBytes, hStream, kernelParams, and extra parameters.\n"); + if (extra!=NULL){ + printf("GPGPU-Sim CUDA DRIVER API: ERROR: Currently do not support void** extra.\n"); + abort(); } const char *hostFun = (const char*) f; - gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); - kernel_config &config = g_cuda_launch_stack.back(); - dim3 *d_b = new dim3(blockDimX, blockDimY, blockDimZ); - config.set_block_dim(d_b); - dim3 *d_g = new dim3(gridDimX, gridDimY, gridDimZ); - config.set_grid_dim(d_g); - + CUctx_st *context = GPGPUSim_Context(); + function_info *entry = context->get_kernel(hostFun); + cudaConfigureCall(dim3(gridDimX, gridDimY, gridDimZ), dim3(blockDimX, blockDimY, blockDimZ), sharedMemBytes, (cudaStream_t) hStream); + for(unsigned i = 0; i < entry->num_args(); i++){ + std::pair p = entry->get_param_config(i); + cudaSetupArgument(kernelParams[i], p.first, p.second); + } cudaLaunch(hostFun); return CUDA_SUCCESS; } diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index ef4cf48..e726ab9 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1293,6 +1293,17 @@ public: bool is_pdom_set() const { return pdom_done; } //return pdom flag void set_pdom() { pdom_done = true; } //set pdom flag + void add_config_param( size_t size, unsigned alignment ){ + unsigned offset = 0; + if (m_param_configs.size()>0){ + unsigned offset_nom = m_param_configs.back().first + m_param_configs.back().second; + offset = offset_nom%alignment ? (offset_nom/alignment + 1) * alignment : offset_nom; + } + m_param_configs.push_back(std::pair(size, offset)); + } + + std::pair get_param_config(unsigned param_num) const { return m_param_configs[param_num]; } + private: unsigned m_uid; unsigned m_local_mem_framesize; @@ -1306,6 +1317,7 @@ private: unsigned m_instr_mem_size; std::map m_kernel_params; std::map m_ptx_kernel_param_info; + std::vector< std::pair > m_param_configs; const symbol *m_return_var_sym; std::vector m_args; std::list m_instructions; diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index c418fac..e6d6325 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -61,6 +61,7 @@ memory_space_t g_ptr_spec = undefined_space; int g_scalar_type_spec = -1; int g_vector_spec = -1; int g_alignment_spec = -1; +int g_size = -1; int g_extern_spec = 0; // variable declaration stuff: @@ -116,6 +117,7 @@ void init_directive_state() g_vector_spec=-1; g_opcode=-1; g_alignment_spec = -1; + g_size = -1; g_extern_spec = 0; g_scalar_type.clear(); g_operands.clear(); @@ -373,6 +375,9 @@ int pad_address (new_addr_type address, unsigned size, unsigned maxalign) { void add_identifier( const char *identifier, int array_dim, unsigned array_ident ) { + if(array_ident==ARRAY_IDENTIFIER){ + g_size *= array_dim; + } if( g_func_decl && (g_func_info == NULL) ) { // return variable decl... assert( g_add_identifier_cached__identifier == NULL ); @@ -562,10 +567,13 @@ void add_constptr(const char* identifier1, const char* identifier2, int offset) void add_function_arg() { + assert(g_size>0); if( g_func_info ) { PTX_PARSE_DPRINTF("add_function_arg \"%s\"", g_last_symbol->name().c_str() ); g_func_info->add_arg(g_last_symbol); + g_func_info->add_config_param( g_size, g_alignment_spec ); } + } void add_extern_spec() @@ -617,6 +625,32 @@ void add_vector_spec(int spec ) void add_scalar_type_spec( int type_spec ) { + //save size of parameter + switch ( type_spec ) { + case B8_TYPE: + case S8_TYPE: + case U8_TYPE: + g_size = 1; break; + case B16_TYPE: + case S16_TYPE: + case U16_TYPE: + case F16_TYPE: + g_size = 2; break; + case B32_TYPE: + case S32_TYPE: + case U32_TYPE: + case F32_TYPE: + g_size = 4; break; + case B64_TYPE: + case BB64_TYPE: + case S64_TYPE: + case U64_TYPE: + case F64_TYPE: + case FF64_TYPE: + g_size = 8; break; + case BB128_TYPE: + g_size = 16; break; + } PTX_PARSE_DPRINTF("add_scalar_type_spec \"%s\"", g_ptx_token_decode[type_spec].c_str()); g_scalar_type.push_back( type_spec ); if ( g_scalar_type.size() > 1 ) { diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 4640d65..da85bae 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -3034,6 +3034,12 @@ bool opndcoll_rfu_t::writeback( const warp_inst_t &inst ) for( r=regs.begin(); r!=regs.end();r++,n++ ) { unsigned reg = *r; unsigned bank = register_bank(reg,inst.warp_id(),m_num_banks,m_bank_warp_shift); + unsigned count = 0; + while( !m_arbiter.bank_idle(bank) ) { + assert((++count) Date: Thu, 16 Aug 2018 14:10:55 -0700 Subject: fix alignment bug in parser, modified ptxjitconfig, minor fixes --- src/cuda-sim/cuda-sim.cc | 45 ++++++++++++++++++++++++++++++++++++++++++-- src/cuda-sim/instructions.cc | 8 ++++---- src/cuda-sim/ptx_ir.h | 1 + src/cuda-sim/ptx_parser.cc | 2 +- 4 files changed, 49 insertions(+), 7 deletions(-) (limited to 'src/cuda-sim/ptx_ir.h') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 5b80616..b85ba95 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1232,7 +1232,6 @@ void function_info::ptx_jit_config(std::map mallocPt std::vector< std::pair > param_data; std::vector offsets; std::vector paramIsPointer; - char buff[1024]; char * gpgpusim_path = getenv("GPGPUSIM_ROOT"); assert(gpgpusim_path!=NULL); @@ -1241,6 +1240,47 @@ void function_info::ptx_jit_config(std::map mallocPt std::string command = std::string("mkdir ") + gpgpusim_path + "/debug_tools/WatchYourStep/data"; std::string filename(std::string(gpgpusim_path) + "/debug_tools/WatchYourStep/data/params.config" + std::to_string(counter)); + //initialize paramList + char buff[1024]; + std::string filename_c(filename+"_c"); + snprintf(buff,1024,"c++filt %s > %s", get_name().c_str(), filename_c.c_str()); + system(buff); + FILE *fp = fopen(filename_c.c_str(), "r"); + fgets(buff, 1024, fp); + fclose(fp); + std::string fn(buff); + size_t pos1, pos2; + pos1 = fn.find_last_of("("); + pos2 = fn.find(")", pos1); + assert(pos2>pos1&&pos1>0); + strcpy(buff, fn.substr(pos1 + 1, pos2 - pos1 - 1).c_str()); + char *tok; + tok = strtok(buff, ","); + std::string tmp; + while(tok!=NULL){ + std::string param(tok); + if(param.find("<")!=std::string::npos){ + assert(param.find(">")==std::string::npos); + assert(param.find("*")==std::string::npos); + tmp = param; + } else { + if (tmp.length()>0){ + tmp = ""; + assert(param.find(">")!=std::string::npos); + assert(param.find("<")==std::string::npos); + assert(param.find("*")==std::string::npos); + } + printf("%s\n", param.c_str()); + if(param.find("*")!=std::string::npos){ + paramIsPointer.push_back(true); + }else{ + paramIsPointer.push_back(false); + } + } + tok = strtok(NULL, ","); + } + + for( std::map::iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) { param_info &p = i->second; std::string name = p.get_name(); @@ -1249,8 +1289,9 @@ void function_info::ptx_jit_config(std::map mallocPt param_t param_value = p.get_value(); offsets.push_back((unsigned)p.get_offset()); - if(param_value.size==sizeof(void*) && mallocPtr_Size.find(*(unsigned long long*)param_value.pdata)!=mallocPtr_Size.end()){ + if (paramIsPointer[i->first] && (*(unsigned long long*)param_value.pdata != 0)){ //is pointer + assert(param_value.size==sizeof(void*)&&"MisID'd this param as pointer"); size_t array_size = 0; unsigned long long param_pointer = *(unsigned long long*)param_value.pdata; if(mallocPtr_Size.find(param_pointer)!=mallocPtr_Size.end()){ diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 109cfa7..31fc434 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -1471,17 +1471,17 @@ void brev_impl( const ptx_instruction *pI, ptx_thread_info *thread ) switch(i_type){ case B32_TYPE: msb = 31; - for (unsigned i=0; i<=msb; i++) { + for (unsigned i=0; i<=msb; i++) { if((src1_data.u32 & (1 << i))) data.u32 |= 1 << (msb - i); - } + } break; case B64_TYPE: msb = 63; - for (unsigned i=0; i<=msb; i++) { + for (unsigned i=0; i<=msb; i++) { if((src1_data.u64 & (1 << i))) data.u64 |= 1 << (msb - i); - } + } break; default: assert(0); } diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index e726ab9..768b3e7 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1297,6 +1297,7 @@ public: unsigned offset = 0; if (m_param_configs.size()>0){ unsigned offset_nom = m_param_configs.back().first + m_param_configs.back().second; + //ensure offset matches alignment requirements offset = offset_nom%alignment ? (offset_nom/alignment + 1) * alignment : offset_nom; } m_param_configs.push_back(std::pair(size, offset)); diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 3f3485c..0ba9a5b 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -572,7 +572,7 @@ void add_function_arg() PTX_PARSE_DPRINTF("add_function_arg \"%s\"", g_last_symbol->name().c_str() ); g_func_info->add_arg(g_last_symbol); unsigned alignment = (g_alignment_spec==-1) ? g_size : g_alignment_spec; - assert(alignment<=8); + assert(alignment==1||alignment==2||alignment==4||alignment==8||alignment==16);//known valid alignment values g_func_info->add_config_param( g_size, alignment); } -- cgit v1.3