3 files changed, 5 insertions, 55 deletions

diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
index 1065870..554e831 100644
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
+++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
@@ -61,7 +61,6 @@ const char *sSDKname = "PTX Just In Time (JIT) Compilation Plus";
 char *wys_exec_path;
 char *wys_exec_name;
 char *wys_launch_num;
-bool gpgpusim = false;
 dim3 gridDim, blockDim;
 std::string kernelName;
 
@@ -138,10 +137,6 @@ void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUl
 
 void initializeData(std::vector<param>& v_params)
 {
-    char *gpgpusim_env = getenv("GPGPUSIM_SETUP_ENVIRONMENT_WAS_RUN");
-    if (gpgpusim_env!=NULL&&gpgpusim_env[0] == '1'){
-        gpgpusim=true;
-    }
     wys_exec_path = getenv("WYS_EXEC_PATH");
     assert(wys_exec_path!=NULL);
     wys_exec_name = getenv("WYS_EXEC_NAME");
@@ -272,23 +267,16 @@ int main(int argc, char **argv)
             unsigned char **d_data = (unsigned char **) malloc(sizeof(unsigned char **));
             checkCudaErrors(cudaMalloc((void**)d_data, p->size));
             checkCudaErrors(cudaMemcpy((void*)*d_data,(void*)p->data,p->size,cudaMemcpyHostToDevice));
-            if (gpgpusim){
-                checkCudaErrors(cuParamSetv(hKernel, p->offset, d_data, sizeof(*d_data)));
-            }
             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, p->offset, p->data, p->size));
-            }
             paramKernels[index] = (void*)p->data;
             paramOffset = p->offset + p->size;
         }
         index ++;
     }
-    checkCudaErrors(cuParamSetSize(hKernel, paramOffset));
 
     // Launch the kernel (Driver API_)
     CUDAAPI cuLaunchKernel(hKernel, gridDim.x, gridDim.y, gridDim.z, blockDim.x, blockDim.y, blockDim.z, 
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index e0cd3f9..5b80616 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -1232,6 +1232,7 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt
     std::vector< std::pair<size_t, unsigned char*> > param_data;
     std::vector<unsigned> offsets;
     std::vector<bool> paramIsPointer;
+    char buff[1024];
 
     char * gpgpusim_path = getenv("GPGPUSIM_ROOT");
     assert(gpgpusim_path!=NULL);
@@ -1240,46 +1241,6 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt
     std::string command = std::string("mkdir ") + gpgpusim_path + "/debug_tools/WatchYourStep/data";
     std::string filename(std::string(gpgpusim_path) + "/debug_tools/WatchYourStep/data/params.config" + std::to_string(counter));
 
-    //initialize paramList
-    char buff[1024];
-    std::string filename_c(filename+"_c");
-    snprintf(buff,1024,"c++filt %s > %s", get_name().c_str(), filename_c.c_str());
-    system(buff);
-    FILE *fp = fopen(filename_c.c_str(), "r");
-    fgets(buff, 1024, fp);
-    fclose(fp);
-    std::string fn(buff);
-    size_t pos1, pos2;
-    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;
-    tok = strtok(buff, ",");
-    std::string tmp;
-    while(tok!=NULL){
-        std::string param(tok);
-        if(param.find("<")!=std::string::npos){
-            assert(param.find(">")==std::string::npos);
-            assert(param.find("*")==std::string::npos);
-            tmp = param;
-        } else {
-            if (tmp.length()>0){
-                tmp = ""; 
-                assert(param.find(">")!=std::string::npos);
-                assert(param.find("<")==std::string::npos);
-                assert(param.find("*")==std::string::npos);
-            }   
-            if(param.find("*")!=std::string::npos){
-                paramIsPointer.push_back(true);
-            }else{                                                                                 
-                paramIsPointer.push_back(false);
-            }   
-        }
-        tok = strtok(NULL, ",");
-    }
-
-
     for( std::map<unsigned,param_info>::iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) {
         param_info &p = i->second;
         std::string name = p.get_name();
@@ -1288,9 +1249,8 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt
         param_t param_value = p.get_value();
         offsets.push_back((unsigned)p.get_offset());
 
-        if (paramIsPointer[i->first]){
+        if(param_value.size==sizeof(void*) && mallocPtr_Size.find(*(unsigned long long*)param_value.pdata)!=mallocPtr_Size.end()){
             //is 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()){
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index e6d6325..3f3485c 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -571,7 +571,9 @@ void add_function_arg()
    if( g_func_info ) {
       PTX_PARSE_DPRINTF("add_function_arg \"%s\"", g_last_symbol->name().c_str() );
       g_func_info->add_arg(g_last_symbol);
-      g_func_info->add_config_param( g_size, g_alignment_spec );
+      unsigned alignment = (g_alignment_spec==-1) ? g_size : g_alignment_spec;
+      assert(alignment<=8);
+      g_func_info->add_config_param( g_size,  alignment);
    }
 
 }