load param offsets with refactored code and look in middle of malloc'd area for pointer

3 files changed, 55 insertions, 44 deletions

diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
index 26a6d29..1065870 100644
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
+++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
@@ -136,7 +136,7 @@ void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUl
     checkCudaErrors(cuLinkDestroy(*lState));
 }
 
-void initializeData(std::vector<unsigned char*>& param_data, std::vector< std::pair<size_t, bool> >& param_info)
+void initializeData(std::vector<param>& v_params)
 {
     char *gpgpusim_env = getenv("GPGPUSIM_SETUP_ENVIRONMENT_WAS_RUN");
     if (gpgpusim_env!=NULL&&gpgpusim_env[0] == '1'){
@@ -161,21 +161,21 @@ void initializeData(std::vector<unsigned char*>& param_data, std::vector< std::p
     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<size_t, bool> info;
+        param p;
         int err;
         size_t len;
         unsigned val;
         int start = fgetc(fin);
         if (start == '*'){
-            info.second=true;
+            p.isPointer = true;
         }else{
-            info.second=false;
+            p.isPointer = false;
             int c = ungetc(start,fin);
             assert(c==start&&"Couldn't ungetc\n");
         }
         err = fscanf(fin, "%lu : ", &len);
-        info.first = len;
         assert( err==1 );
+        p.size = len;
         unsigned char* params = (unsigned char*) malloc(len*sizeof(unsigned char));
         for (size_t i=0; i<len; i++)
         {
@@ -183,9 +183,12 @@ void initializeData(std::vector<unsigned char*>& 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);
+        p.data = params;
+        unsigned offset;
+        err = fscanf(fin, " : %u", &offset);
+        assert( err==1 );
+        p.offset = offset;
+        v_params.push_back(p);
         err = fscanf(fin, "\n");
         assert(err==0);
     }
@@ -206,12 +209,8 @@ int main(int argc, char **argv)
 
     printf("[%s] - Starting...\n", sSDKname);
     //parameter data
-    std::vector<unsigned char*> param_data;
-    //parameter data size and isPointer
-    std::vector< std::pair<size_t, bool> > param_info;
-    initializeData(param_data,param_info);
-
-
+    std::vector<param> v_params;
+    initializeData(v_params);
 
     if (checkCmdLineFlag(argc, (const char **)argv, "device"))
     {
@@ -251,8 +250,6 @@ 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
     int         *d_tmp   = 0;
@@ -265,28 +262,31 @@ int main(int argc, char **argv)
     //maps param number to pointer to device data
     std::map< size_t, void* > m_device_data;
     std::map< size_t, void* > m_cleanup;
-    void * paramKernels[param_data.size()];
-    //Initialize param_data for kernel
+    void * paramKernels[v_params.size()];
+    //Initialize param data for kernel
     int paramOffset = 0;
-    for( size_t i = 0; i<param_data.size(); i++){
-        if(param_info[i].second){
+
+    int index = 0;
+    for(std::vector<param>::iterator p = v_params.begin(); p!=v_params.end(); p++){
+        if(p->isPointer){
             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(cudaMalloc((void**)d_data, p->size));
+            checkCudaErrors(cudaMemcpy((void*)*d_data,(void*)p->data,p->size,cudaMemcpyHostToDevice));
             if (gpgpusim){
-                checkCudaErrors(cuParamSetv(hKernel, paramOffset, d_data, sizeof(*d_data)));
+                checkCudaErrors(cuParamSetv(hKernel, p->offset, d_data, sizeof(*d_data)));
             }
-            paramKernels[i] = (void*)d_data;
-            m_device_data[i]=*d_data;
-            m_cleanup[i]=d_data;
-            paramOffset += 8;
+            paramKernels[index] = (void*)d_data;
+            m_device_data[index]=*d_data;
+            m_cleanup[index]=d_data;
+            paramOffset = p->offset + 8;
         }else{
             if (gpgpusim){
-                checkCudaErrors(cuParamSetv(hKernel, paramOffset, param_data[i], param_info[i].first));
+                checkCudaErrors(cuParamSetv(hKernel, p->offset, p->data, p->size));
             }
-            paramKernels[i] = (void*)param_data[i];
-            paramOffset += param_info[i].first;
+            paramKernels[index] = (void*)p->data;
+            paramOffset = p->offset + p->size;
         }
+        index ++;
     }
     checkCudaErrors(cuParamSetSize(hKernel, paramOffset));
 
@@ -299,13 +299,13 @@ int main(int argc, char **argv)
     std::map< size_t, unsigned char* > m_output_data;
     for(std::map< size_t, void* >::iterator i = m_device_data.begin(); i!=m_device_data.end(); i++){
         unsigned char *h_data   = 0;
-        if ((h_data = (unsigned char *)malloc(param_info[i->first].first)) == NULL)
+        if ((h_data = (unsigned char *)malloc(v_params[i->first].size)) == NULL)
         {
             std::cerr << "Could not allocate host memory" << std::endl;
             exit(EXIT_FAILURE);
         }
         // Copy the result back to the host
-        checkCudaErrors(cudaMemcpy(h_data, i->second, param_info[i->first].first, cudaMemcpyDeviceToHost));
+        checkCudaErrors(cudaMemcpy(h_data, i->second, v_params[i->first].size, cudaMemcpyDeviceToHost));
         m_output_data[i->first] = h_data;
     }
 
@@ -313,8 +313,8 @@ int main(int argc, char **argv)
     FILE *fout = fopen(filename.c_str(), "w");
     assert(fout);
     for(std::map< size_t, unsigned char* >::iterator i = m_output_data.begin(); i!=m_output_data.end(); i++){
-        fprintf(fout, "param %zu: size = %zu, data = ", i->first,param_info[i->first].first);
-        for (size_t j = 0; j<param_info[i->first].first; j++){
+        fprintf(fout, "param %zu: size = %zu, data = ", i->first, v_params[i->first].size);
+        for (size_t j = 0; j<v_params[i->first].size; j++){
             if (!(j%24)){ 
                 fprintf(fout, "\n");
             }
@@ -325,7 +325,6 @@ int main(int argc, char **argv)
     fflush(fout);
     fclose(fout);
 
-
     //Cleanup
     for(std::map< size_t, void* >::iterator i = m_device_data.begin(); i!=m_device_data.end(); i++){
         if (i->second){
diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.h b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.h
index 8fb5611..3c6e181 100644
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.h
+++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.h
@@ -12,6 +12,13 @@
 #ifndef _PTXJIT_H_
 #define _PTXJIT_H_
 
+struct param{
+    bool isPointer;
+    size_t size;
+    unsigned char *data;
+    unsigned offset;
+};
+
 /*
  * PTX is equivalent to the following kernel:
  *
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index a499ce9..e0cd3f9 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -1238,11 +1238,6 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> 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";
-    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));
 
     //initialize paramList
@@ -1295,11 +1290,21 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt
 
         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
+            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()){
+                array_size = mallocPtr_Size[param_pointer];
+            }else{
+                for( std::map<unsigned long long, size_t>::iterator j=mallocPtr_Size.begin(); j!=mallocPtr_Size.end(); j++ ) {
+                    if(param_pointer>j->first&&param_pointer<j->first + j->second){
+                        array_size = j->first + j->second - param_pointer;
+                        break;
+                    }
+                }
+                assert(array_size>0&&"pointer was not previously malloc'd");
+            }
 
-            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);
             unsigned char* array_val = (unsigned char*) malloc(array_size);