fixed pointer bug in ptx jit and added support for executing ptxjitconfig when running ptxjit

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>