diff options
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 15 | ||||
| -rw-r--r-- | src/cuda-sim/cuda-sim.cc | 23 | ||||
| -rw-r--r-- | src/cuda-sim/ptx_ir.h | 2 |
3 files changed, 29 insertions, 11 deletions
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<void *,void **> pinned_memory; //support for pinned memories added std::map<void *, size_t> pinned_memory_size; -std::map<void *, size_t> g_devPtr_Size; +std::map<unsigned long long, size_t> g_mallocPtr_Size; int no_of_ptx=0; std::map<int, std::set<std::string> > 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<void *, size_t> devPtr_Size, memory_space *param_mem) const +void function_info::debug_param(std::map<unsigned long long, size_t> mallocPtr_Size, memory_space *param_mem) const { static unsigned long counter = 0; std::vector< std::pair<size_t, unsigned char*> > param_data; @@ -1248,8 +1248,8 @@ void function_info::debug_param(std::map<void *, size_t> 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<void *, size_t> 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<size_t, unsigned char*>(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<void *, size_t> devPtr_Size, memory_spa } fflush(fout); fclose(fout); - exit(0); + //exit(0); } template<int activate_level> 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<void *, size_t> devPtr_Size, memory_space *param_mem) const; + void debug_param(std::map<unsigned long long, size_t> mallocPtr_Size, memory_space *param_mem) const; const struct gpgpu_ptx_sim_info* get_kernel_info () const { |