diff options
| -rw-r--r-- | debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp | 11 | ||||
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 2 | ||||
| -rw-r--r-- | src/cuda-sim/cuda-sim.cc | 5 |
3 files changed, 11 insertions, 7 deletions
diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp index 1f094ba..68964d8 100644 --- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp +++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp @@ -254,11 +254,12 @@ int main(int argc, char **argv) int paramOffset = 0; for( size_t i = 0; i<param_data.size(); i++){ if(param_info[i].second){ - unsigned char *d_data = 0; - checkCudaErrors(cudaMalloc((void**)&d_data, param_info[i].first)); - checkCudaErrors(cudaMemcpy((void*)d_data,(void*)param_data[i],param_info[i].first,cudaMemcpyHostToDevice)); - checkCudaErrors(cuParamSetv(hKernel, paramOffset, &d_data, sizeof(d_data))); - m_device_data[i]=d_data; + unsigned char **d_data = (unsigned char **) malloc(sizeof(unsigned char **)); + *d_data = (unsigned char *) malloc(sizeof(unsigned char *)); + checkCudaErrors(cudaMalloc((void**)d_data, param_info[i].first)); + checkCudaErrors(cudaMemcpy((void*)*d_data,(void*)param_data[i],param_info[i].first,cudaMemcpyHostToDevice)); + checkCudaErrors(cuParamSetv(hKernel, paramOffset, d_data, sizeof(*d_data))); + m_device_data[i]=*d_data; paramOffset += 8; }else{ checkCudaErrors(cuParamSetv(hKernel, paramOffset, param_data[i], param_info[i].first)); diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 5bacb8e..68b4017 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -2883,7 +2883,7 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, g_ptx_kernel_count++; fflush(stdout); - if(g_debug_execution >= 3){ + if(g_debug_execution >= 4){ entry->ptx_jit_config(g_mallocPtr_Size, result->get_param_memory(), (gpgpu_t *) context->get_device()->get_gpgpu(), gridDim, blockDim); } diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index ae2f539..df8d806 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1234,6 +1234,8 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt char * gpgpusim_path = getenv("GPGPUSIM_ROOT"); assert(gpgpusim_path!=NULL); + 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()); std::string filename(std::string(gpgpusim_path) + "/debug_tools/WatchYourStep/data/params.config" + std::to_string(counter)); @@ -1284,7 +1286,7 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt //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()); + 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); FILE *fin = fopen(ptx_config_fn.c_str(), "r"); char ptx_source[256]; @@ -1314,6 +1316,7 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt fflush(fout); fclose(fout); counter++; + //TODO: Free param_data } template<int activate_level> |