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.cc | 26 ++++++++++++++++++++++++++ 1 file changed, 26 insertions(+) (limited to 'src/cuda-sim/ptx_ir.cc') diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 8ebdcf8..6a17eaf 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -575,6 +575,31 @@ 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 intersect( std::set &A, const std::set &B ) { // return intersection of A and B in A @@ -1305,6 +1330,7 @@ function_info::function_info(int entry_point ) m_kernel_info.smem = 0; m_local_mem_framesize = 0; m_args_aligned_size = -1; + pdom_done = false; //initialize it to false } unsigned function_info::print_insn( unsigned pc, FILE * fp ) const -- cgit v1.3 From 5e8b10d9a0ea518dcb6c3c88b0a9cfd143363bcd Mon Sep 17 00:00:00 2001 From: Amruth Date: Tue, 27 Mar 2018 14:12:05 -0700 Subject: support for pinned memories - temporary fix --- libcuda/cuda_runtime_api.cc | 53 +++++++++++++++++++++++++++++++++++++++++-- src/abstract_hardware_model.h | 3 +++ src/cuda-sim/ptx_ir.cc | 6 +++-- 3 files changed, 58 insertions(+), 4 deletions(-) (limited to 'src/cuda-sim/ptx_ir.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index f74c4eb..97d702c 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -145,6 +145,8 @@ #include #endif +std::map pinned_memory; //support for pinned memories added +std::map pinned_memory_size; int no_of_ptx=0; extern void synchronize(); @@ -476,6 +478,8 @@ __host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size) GPGPUSim_Context(); *ptr = malloc(size); if ( *ptr ) { + //track pinned memory size allocated in the host so that same amount of memory is also allocated in GPU. + pinned_memory_size[*ptr]=size; return g_last_cudaError = cudaSuccess; } else { return g_last_cudaError = cudaErrorMemoryAllocation; @@ -766,6 +770,16 @@ __host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count) return g_last_cudaError = cudaSuccess; } +//memset operation is done but i think its not async? +__host__ cudaError_t CUDARTAPI cudaMemsetAsync(void *mem, int c, size_t count, cudaStream_t stream=0) +{ + printf("GPGPU-Sim PTX: WARNING: Asynchronous memset not supported (%s)\n", __my_func__); + CUctx_st *context = GPGPUSim_Context(); + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + gpu->gpu_memset((size_t)mem, c, count); + return g_last_cudaError = cudaSuccess; +} + __host__ cudaError_t CUDARTAPI cudaMemset2D(void *mem, size_t pitch, int c, size_t width, size_t height) { cuda_not_implemented(__my_func__,__LINE__); @@ -855,6 +869,12 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic case 76: *value= 3 ; break; + case 78: + *value= 0 ; //TODO: as of now, we dont support stream priorities. + break; + default: + printf("ERROR: implement the attribute numbered %d \n",attr); + abort(); } return g_last_cudaError = cudaSuccess; } else { @@ -1054,6 +1074,15 @@ __host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream) return g_last_cudaError = cudaSuccess; } +//TODO: introduce priorities +__host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithPriority(cudaStream_t *stream, unsigned int flags, int priority) { + return cudaStreamCreate(stream); +} + +__host__ __device__ cudaError_t CUDARTAPI cudaDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPriority) { + return cudaSuccess; +} + __host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t *stream, unsigned int flags) { return cudaStreamCreate(stream); } @@ -2206,6 +2235,9 @@ cudaError_t cudaGLUnregisterBufferObject(GLuint bufferObj) cudaError_t CUDARTAPI cudaHostAlloc(void **pHost, size_t bytes, unsigned int flags) { *pHost = malloc(bytes); + //need to track the size allocated so that cudaHostGetDevicePointer() can function properly. + //TODO: vary this function behavior based on flags value (following nvidia documentation) + pinned_memory_size[*pHost]=bytes; if( *pHost ) return g_last_cudaError = cudaSuccess; else @@ -2214,8 +2246,25 @@ cudaError_t CUDARTAPI cudaHostAlloc(void **pHost, size_t bytes, unsigned int fl cudaError_t CUDARTAPI cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags) { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + //only cpu memory allocation happens in cudaHostAlloc. Linking with device pointer to pinned memory happens here. + //TODO: once kernel is executed, the contents in global pointer of GPU must be copied back to CPU host pointer! + flags=0; + CUctx_st* context = GPGPUSim_Context(); + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + std::map::const_iterator i = pinned_memory_size.find(pHost); + assert(i != pinned_memory_size.end()); + size_t size = i->second; + *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); + if ( *pDevice ) { + pinned_memory[pHost]=pDevice; + //Copy contents in cpu to gpu + gpu->memcpy_to_gpu((size_t)*pDevice,pHost,size); + return g_last_cudaError = cudaSuccess; + } else { + return g_last_cudaError = cudaErrorMemoryAllocation; + } } cudaError_t CUDARTAPI cudaSetValidDevices(int *device_arr, int len) diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index aaa4b00..67b36c7 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -182,6 +182,9 @@ void increment_x_then_y_then_z( dim3 &i, const dim3 &bound); class stream_manager; struct CUstream_st; extern stream_manager * g_stream_manager; +//support for pinned memories added +extern std::map pinned_memory; +extern std::map pinned_memory_size; class kernel_info_t { public: diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 6a17eaf..17e91df 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -280,8 +280,10 @@ type_info *symbol_table::get_array_type( type_info *base_type, unsigned array_di { type_info_key t = base_type->get_key(); t.set_array_dim(array_dim); - type_info *pt; - pt = m_types[t] = new type_info(this,t); + type_info *pt = new type_info(this,t); + //Where else is m_types being used? As of now, I dont find any use of it and causing seg fault. So disabling m_types. + //TODO: find where m_types can be used in future and solve the seg fault. + //pt = m_types[t] = new type_info(this,t); return pt; } -- 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.cc') 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 f7b0d64c68f12d604e09aec8dbba569df354faf6 Mon Sep 17 00:00:00 2001 From: Jonathan Date: Tue, 5 Jun 2018 12:58:08 -0700 Subject: parse all ptx and add to symbol table --- src/cuda-sim/instructions.cc | 2 +- src/cuda-sim/ptx_ir.cc | 2 +- src/cuda-sim/ptx_parser.cc | 69 ++++++++++++++++++++++---------------------- 3 files changed, 37 insertions(+), 36 deletions(-) (limited to 'src/cuda-sim/ptx_ir.cc') diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 37438fa..08bf528 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -2291,7 +2291,7 @@ void div_impl( const ptx_instruction *pI, ptx_thread_info *thread ) void dp4a_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { - printf("instruction not implemented yet"); + printf("DP4A instruction not implemented yet"); assert(0); } diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index be25dbe..016c600 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1031,7 +1031,7 @@ static std::list check_operands( int opcode, const std::list &operands ) { static int g_warn_literal_operands_two_type_inst; - if( (opcode == CVT_OP) || (opcode == SET_OP) || (opcode == SLCT_OP) || (opcode == TEX_OP) ) { + if( (opcode == CVT_OP) || (opcode == SET_OP) || (opcode == SLCT_OP) || (opcode == TEX_OP) || (opcode == DP4A_OP) ) { // just make sure these do not have have const operands... if( !g_warn_literal_operands_two_type_inst ) { std::list::const_iterator o; diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index a51799a..06ca870 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -107,6 +107,34 @@ void read_parser_environment_variables() } } +void init_directive_state() +{ + PTX_PARSE_DPRINTF("init_directive_state"); + g_space_spec=undefined_space; + g_ptr_spec=undefined_space; + g_scalar_type_spec=-1; + g_vector_spec=-1; + g_opcode=-1; + g_alignment_spec = -1; + g_extern_spec = 0; + g_scalar_type.clear(); + g_operands.clear(); + g_last_symbol = NULL; +} + +void init_instruction_state() +{ + PTX_PARSE_DPRINTF("init_instruction_state"); + g_pred = NULL; + g_neg_pred = 0; + g_pred_mod = -1; + g_label = NULL; + g_opcode = -1; + g_options.clear(); + g_return_var = operand_info(); + init_directive_state(); +} + symbol_table *init_parser( const char *ptx_filename ) { g_filename = strdup(ptx_filename); @@ -114,9 +142,9 @@ symbol_table *init_parser( const char *ptx_filename ) g_global_allfiles_symbol_table = new symbol_table("global_allfiles", 0, NULL); g_global_symbol_table = g_current_symbol_table = g_global_allfiles_symbol_table; } - else { - g_global_symbol_table = g_current_symbol_table = new symbol_table("global",0,g_global_allfiles_symbol_table); - } +// else { +// g_global_symbol_table = g_current_symbol_table = new symbol_table("global",0,g_global_allfiles_symbol_table); +// } ptx_lineno = 1; #define DEF(X,Y) g_ptx_token_decode[X] = Y; @@ -136,7 +164,8 @@ symbol_table *init_parser( const char *ptx_filename ) g_ptx_token_decode[global_space] = "global_space"; g_ptx_token_decode[generic_space] = "generic_space"; g_ptx_token_decode[instruction_space] = "instruction_space"; - + init_directive_state(); + init_instruction_state(); ptx_in = fopen(ptx_filename, "r"); ptx_parse(); @@ -144,34 +173,6 @@ symbol_table *init_parser( const char *ptx_filename ) return g_global_symbol_table; } -void init_directive_state() -{ - PTX_PARSE_DPRINTF("init_directive_state"); - g_space_spec=undefined_space; - g_ptr_spec=undefined_space; - g_scalar_type_spec=-1; - g_vector_spec=-1; - g_opcode=-1; - g_alignment_spec = -1; - g_extern_spec = 0; - g_scalar_type.clear(); - g_operands.clear(); - g_last_symbol = NULL; -} - -void init_instruction_state() -{ - PTX_PARSE_DPRINTF("init_instruction_state"); - g_pred = NULL; - g_neg_pred = 0; - g_pred_mod = -1; - g_label = NULL; - g_opcode = -1; - g_options.clear(); - g_return_var = operand_info(); - init_directive_state(); -} - static int g_entry_point; void start_function( int entry_point ) @@ -621,8 +622,8 @@ void add_scalar_type_spec( int type_spec ) g_scalar_type.push_back( type_spec ); if ( g_scalar_type.size() > 1 ) { parse_assert( (g_opcode == -1) || (g_opcode == CVT_OP) || (g_opcode == SET_OP) || (g_opcode == SLCT_OP) - || (g_opcode == TEX_OP), - "only cvt, set, slct, and tex can have more than one type specifier."); + || (g_opcode == TEX_OP)|| (g_opcode == DP4A_OP), + "only cvt, set, slct, tex, and dp4a can have more than one type specifier."); } g_scalar_type_spec = type_spec; } -- 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.cc') 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