From 584ebaa74a838680e6ed1fa13ac266e88c30c071 Mon Sep 17 00:00:00 2001
From: Jonathan <jonathanslew@outlook.com>
Date: Tue, 26 Jun 2018 13:20:39 -0700
Subject: exports and imports param data in new debug tool: WatchYourStep

---
 .../WatchYourStep/ptxjitplus/ptxjitplus.cpp        | 277 +++++++++++++++++++++
 1 file changed, 277 insertions(+)
 create mode 100644 debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp

(limited to 'debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp')

diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
new file mode 100644
index 0000000..9954e31
--- /dev/null
+++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
@@ -0,0 +1,277 @@
+/**
+ * Copyright 1993-2015 NVIDIA Corporation.  All rights reserved.
+ *
+ * Please refer to the NVIDIA end user license agreement (EULA) associated
+ * with this source code for terms and conditions that govern your use of
+ * this software. Any use, reproduction, disclosure, or distribution of
+ * this software and related documentation outside the terms of the EULA
+ * is strictly prohibited.
+ *
+ */
+
+/*
+ * This sample uses the Driver API to just-in-time compile (JIT) a Kernel from PTX code.
+ * Additionally, this sample demonstrates the seamless interoperability capability of CUDA runtime
+ * Runtime and CUDA Driver API calls.
+ * This sample requires Compute Capability 2.0 and higher.
+ *
+ */
+
+// System includes
+#include <iostream>
+#include <math.h>
+#include <string.h>
+#include <stdio.h>
+
+// CUDA driver & runtime
+#include <cuda.h>
+#include <cuda_runtime.h>
+
+// helper functions and utilities to work with CUDA
+#include <helper_cuda.h>
+#include <helper_functions.h>  // helper for shared that are common to CUDA Samples
+
+// sample include
+#include "ptxjitplus.h"
+#include "ptxinst.h"
+
+const char *sSDKname = "PTX Just In Time (JIT) Compilation (no-qatest)";
+
+
+void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUlinkState *lState)
+{
+    CUjit_option options[6];
+    void *optionVals[6];
+    float walltime;
+    char error_log[8192],
+         info_log[8192];
+    unsigned int logSize = 8192;
+    void *cuOut;
+    size_t outSize;
+    int myErr = 0;
+
+    // Setup linker options
+    // Return walltime from JIT compilation
+    options[0] = CU_JIT_WALL_TIME;
+    optionVals[0] = (void *) &walltime;
+    // Pass a buffer for info messages
+    options[1] = CU_JIT_INFO_LOG_BUFFER;
+    optionVals[1] = (void *) info_log;
+    // Pass the size of the info buffer
+    options[2] = CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES;
+    optionVals[2] = (void *) (long)logSize;
+    // Pass a buffer for error message
+    options[3] = CU_JIT_ERROR_LOG_BUFFER;
+    optionVals[3] = (void *) error_log;
+    // Pass the size of the error buffer
+    options[4] = CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES;
+    optionVals[4] = (void *) (long) logSize;
+    // Make the linker verbose
+    options[5] = CU_JIT_LOG_VERBOSE;
+    optionVals[5] = (void *) 1;
+
+    // Create a pending linker invocation
+    checkCudaErrors(cuLinkCreate(6,options, optionVals, lState));
+
+    if (sizeof(void *)==4)
+    {
+        // Load the PTX from the string myPtx32
+        printf("Loading myPtx32[] program\n");
+        // PTX May also be loaded from file, as per below.
+        myErr = cuLinkAddData(*lState, CU_JIT_INPUT_PTX, (void *)myPtx32, strlen(myPtx32)+1, 0, 0, 0, 0);
+    }
+    else
+    {
+        // Load the PTX from the string myPtx (64-bit)
+        printf("Loading myPtx[] program\n");
+        //myErr = cuLinkAddData(*lState, CU_JIT_INPUT_PTX, (void *)myPtx64, strlen(myPtx64)+1, 0, 0, 0, 0);
+        // PTX May also be loaded from file, as per below.
+        myErr = cuLinkAddFile(*lState, CU_JIT_INPUT_PTX, "myPtx64.ptx",0,0,0);
+    }
+
+    if (myErr != CUDA_SUCCESS)
+    {
+        // Errors will be put in error_log, per CU_JIT_ERROR_LOG_BUFFER option above.
+        fprintf(stderr,"PTX Linker Error:\n%s\n",error_log);
+    }
+
+    // Complete the linker step
+    checkCudaErrors(cuLinkComplete(*lState, &cuOut, &outSize));
+
+    // Linker walltime and info_log were requested in options above.
+    printf("CUDA Link Completed in %fms. Linker Output:\n%s\n",walltime,info_log);
+
+    // Load resulting cuBin into module
+    checkCudaErrors(cuModuleLoadData(phModule, cuOut));
+
+    // Locate the kernel entry poin
+    checkCudaErrors(cuModuleGetFunction(phKernel, *phModule, "_Z8myKernelPi"));
+
+    // Destroy the linker invocation
+    checkCudaErrors(cuLinkDestroy(*lState));
+}
+
+void* initializeData(std::vector< std::pair<size_t, unsigned char*> >& param_data)
+{
+    char *wys_exec_path = getenv("WYS_EXEC_PATH");
+    assert(wys_exec_path!=NULL);
+    char *wys_exec_name = getenv("WYS_EXEC_NAME");
+    assert(wys_exec_name!=NULL);
+    std::string path_to_search = std::string(wys_exec_path) + "/" + wys_exec_name + ".*.ptx";
+    char* wys_launch_num(getenv("WYS_LAUNCH_NUM"));
+    assert(wys_launch_num!=NULL);
+    std::string filename = std::string("../data/params.config") + wys_launch_num;
+
+    //instrument ptx
+    FILE *fin = fopen(filename.c_str(), "r");
+    assert(fin);
+    char buff[1024];
+    fscanf(fin, "%s\n", buff);
+    //void *retval = instrument_ptx_from_function(std::string(buff), path_to_search);
+    void *retval = NULL;
+    printf("%s\n", buff);
+    //fill data structure to pass in params later
+    while (!feof(fin)){
+        int err;
+        size_t len;
+        unsigned val;
+        err = fscanf(fin, "%lu : ", &len);
+        assert( err==1 );
+        printf("%lu : ", len);
+        unsigned char params[len];
+        for (size_t i=0; i<len; i++)
+        {
+            err = fscanf(fin, "%u ", &val);
+            assert( err==1 );
+            printf("%u ", val);
+            params[i] = (unsigned char) val;
+        }
+        param_data.push_back(std::pair<size_t, unsigned char*>(len, params));
+        err = fscanf(fin, "\n");
+        assert(err==0);
+        printf("\n");
+    }
+    
+    fclose(fin);
+    return retval;
+}
+
+int main(int argc, char **argv)
+{
+    const unsigned int nThreads = 256;
+    const unsigned int nBlocks  = 64;
+    const size_t memSize = nThreads * nBlocks * sizeof(int);
+
+    CUmodule     hModule  = 0;
+    CUfunction   hKernel  = 0;
+    CUlinkState  lState;
+    int         *d_data   = 0;
+    int         *h_data   = 0;
+
+    int cuda_device = 0;
+    cudaDeviceProp deviceProp;
+
+    printf("[%s] - Starting...\n", sSDKname);
+    std::vector< std::pair<size_t, unsigned char*> > param_data;
+    void* storedReg = initializeData(param_data);
+
+    if (checkCmdLineFlag(argc, (const char **)argv, "device"))
+    {
+        cuda_device = getCmdLineArgumentInt(argc, (const char **)argv, "device=");
+
+        if (cuda_device < 0)
+        {
+            printf("Invalid command line parameters\n");
+            exit(EXIT_FAILURE);
+        }
+        else
+        {
+            printf("cuda_device = %d\n", cuda_device);
+            cuda_device = gpuDeviceInit(cuda_device);
+
+            if (cuda_device < 0)
+            {
+                printf("No CUDA Capable devices found, exiting...\n");
+                exit(EXIT_FAILURE);
+            }
+        }
+    }
+    else
+    {
+        // Otherwise pick the device with the highest Gflops/s
+        cuda_device = gpuGetMaxGflopsDeviceId();
+    }
+
+    checkCudaErrors(cudaSetDevice(cuda_device));
+    checkCudaErrors(cudaGetDeviceProperties(&deviceProp, cuda_device));
+    printf("> Using CUDA device [%d]: %s\n", cuda_device, deviceProp.name);
+
+    if (deviceProp.major < 2)
+    {
+        fprintf(stderr, "Compute Capability 2.0 or greater required for this sample.\n");
+        fprintf(stderr, "Maximum Compute Capability of device[%d] is %d.%d.\n", cuda_device,deviceProp.major,deviceProp.minor);
+        exit(EXIT_WAIVED);
+    }
+
+    // Allocate memory on host and device (Runtime API)
+    // NOTE: The runtime API will create the GPU Context implicitly here
+    if ((h_data = (int *)malloc(memSize)) == NULL)
+    {
+        std::cerr << "Could not allocate host memory" << std::endl;
+        exit(EXIT_FAILURE);
+    }
+
+    checkCudaErrors(cudaMalloc(&d_data, memSize));
+
+    // JIT Compile the Kernel from PTX and get the Handles (Driver API)
+    ptxJIT(argc, argv, &hModule, &hKernel, &lState);
+
+    // Set the kernel parameters (Driver API)
+
+    checkCudaErrors(cuFuncSetBlockShape(hKernel, nThreads, 1, 1));
+    //param_data
+    int paramOffset = 0;
+    checkCudaErrors(cuParamSetv(hKernel, paramOffset, &d_data, sizeof(d_data)));
+    paramOffset += sizeof(d_data);
+    checkCudaErrors(cuParamSetSize(hKernel, paramOffset));
+
+    // Launch the kernel (Driver API_)
+    checkCudaErrors(cuLaunchGrid(hKernel, nBlocks, 1));
+    std::cout << "CUDA kernel launched" << std::endl;
+
+    // Copy the result back to the host
+    checkCudaErrors(cudaMemcpy(h_data, d_data, memSize, cudaMemcpyDeviceToHost));
+
+    // Check the result
+    bool dataGood = true;
+
+    for (unsigned int i = 0 ; dataGood && i < nBlocks * nThreads ; i++)
+    {
+        if (h_data[i] != (int)i)
+        {
+            std::cerr << "Error at " << i << std::endl;
+            dataGood = false;
+        }
+    }
+
+    // Cleanup
+    if (d_data)
+    {
+        checkCudaErrors(cudaFree(d_data));
+        d_data = 0;
+    }
+
+    if (h_data)
+    {
+        free(h_data);
+        h_data = 0;
+    }
+
+    if (hModule)
+    {
+        checkCudaErrors(cuModuleUnload(hModule));
+        hModule = 0;
+    }
+
+    return dataGood ? EXIT_SUCCESS : EXIT_FAILURE;
+}
-- 
cgit v1.3


From 1f77b0720f69d684db83beb7a5513cd9461e3676 Mon Sep 17 00:00:00 2001
From: Jonathan <jonathanslew@outlook.com>
Date: Wed, 27 Jun 2018 15:26:00 -0700
Subject: WIP dump params

---
 debug_tools/WatchYourStep/ptxjitplus/ptxjitplus    | Bin 139317 -> 0 bytes
 .../WatchYourStep/ptxjitplus/ptxjitplus.cpp        | 112 ++++++++++++++++++---
 libcuda/cuda_runtime_api.cc                        |  12 ++-
 src/cuda-sim/cuda-sim.cc                           |  59 +++++++++--
 src/cuda-sim/ptx_ir.h                              |   2 +-
 5 files changed, 157 insertions(+), 28 deletions(-)
 delete mode 100755 debug_tools/WatchYourStep/ptxjitplus/ptxjitplus

(limited to 'debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp')

diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus
deleted file mode 100755
index ddc3435..0000000
Binary files a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus and /dev/null differ
diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
index 9954e31..8645114 100644
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
+++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
@@ -111,6 +111,87 @@ void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUl
     checkCudaErrors(cuLinkDestroy(*lState));
 }
 
+void function_info::debug_param( ) const
+{
+   char filename[] = "params.txt";
+   char buff[1024];
+   snprintf(buff,1024,"c++filt %s > %s", get_name().c_str(), filename);
+   system(buff);
+   FILE *fp = fopen(filename, "r");
+   fgets(buff, 1024, fp);
+   fclose(fp);
+
+   std::string fn(buff);
+   size_t pos1, pos2;
+   pos1 = fn.find("(");
+   pos2 = fn.find(")");
+   assert(pos2>pos1&&pos1>0);
+   strcpy(buff, fn.substr(pos1 + 1, pos2 - pos1 - 1).c_str());
+   printf("params: %s\n", buff);
+   char *tok;
+   std::vector<std::string> params;
+   tok = strtok(buff, ",");
+   while(tok!=NULL){
+       std::string param(tok);
+       param.erase(0, param.find_first_not_of(" "));
+       param.erase(param.find_last_not_of(" ")+1);
+       params.push_back(param);
+       tok = strtok(NULL, ",");
+   }
+   for (auto const& it : params){
+      std::cout<<it<<std::endl;
+   }
+
+   FILE *fout  = fopen (filename, "w");
+   fprintf(fout, "Name of function:%s\n", fn.c_str());
+
+   for( std::map<unsigned,param_info>::const_iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) {
+      const param_info &p = i->second;
+      std::string name = p.get_name();
+      param_t param_value = p.get_value();
+      if(params[i->first].find("const")!=std::string::npos){
+         fprintf(fout, "Input: ");
+      } else {
+         fprintf(fout, "Input/output: ");
+      }
+
+      symbol *param = m_symtab->lookup(name.c_str());
+      addr_t param_addr = param->get_address();
+      fprintf(fout, "%s: %#08x, ", name.c_str(), param_addr);
+
+      if(params[i->first].find("int")!=std::string::npos){
+         size_t len = param_value.size/sizeof(int);
+         int val[len];
+         memcpy((void*) val, param_value.pdata+param_value.offset, param_value.size);
+         fprintf(fout, "val (int) = ");
+         for (unsigned i = 0; i<len; i++){
+             fprintf(fout, "%d ", val[i]);
+         }
+         fprintf(fout, "\n");
+      } else if(params[i->first].find("float")!=std::string::npos){
+         size_t len = param_value.size/sizeof(float);
+         float val[len];
+         memcpy((void*) val, param_value.pdata+param_value.offset, param_value.size);
+         fprintf(fout, "val (float) = ");
+         for (unsigned i = 0; i<len; i++){
+             fprintf(fout, "%f ", val[i]);
+         }
+         fprintf(fout, "\n");
+      }else{
+         size_t len = param_value.size/sizeof(char);
+         char val[len];
+         memcpy((void*) val, param_value.pdata+param_value.offset, param_value.size);
+         fprintf(fout, "val (char) = ");
+         for (unsigned i = 0; i<len; i++){
+             fprintf(fout, "%c ", val[i]);
+         }
+         fprintf(fout, "\n");
+      }
+   }
+   fflush(fout);
+   fclose(fout);
+}
+
 void* initializeData(std::vector< std::pair<size_t, unsigned char*> >& param_data)
 {
     char *wys_exec_path = getenv("WYS_EXEC_PATH");
@@ -160,19 +241,17 @@ int main(int argc, char **argv)
 {
     const unsigned int nThreads = 256;
     const unsigned int nBlocks  = 64;
-    const size_t memSize = nThreads * nBlocks * sizeof(int);
 
     CUmodule     hModule  = 0;
     CUfunction   hKernel  = 0;
     CUlinkState  lState;
-    int         *d_data   = 0;
-    int         *h_data   = 0;
 
     int cuda_device = 0;
     cudaDeviceProp deviceProp;
 
     printf("[%s] - Starting...\n", sSDKname);
     std::vector< std::pair<size_t, unsigned char*> > param_data;
+    std::vector< std::pair<size_t, unsigned char*> > device_data;
     void* storedReg = initializeData(param_data);
 
     if (checkCmdLineFlag(argc, (const char **)argv, "device"))
@@ -213,15 +292,7 @@ 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
-    if ((h_data = (int *)malloc(memSize)) == NULL)
-    {
-        std::cerr << "Could not allocate host memory" << std::endl;
-        exit(EXIT_FAILURE);
-    }
 
-    checkCudaErrors(cudaMalloc(&d_data, memSize));
 
     // JIT Compile the Kernel from PTX and get the Handles (Driver API)
     ptxJIT(argc, argv, &hModule, &hKernel, &lState);
@@ -229,16 +300,29 @@ int main(int argc, char **argv)
     // Set the kernel parameters (Driver API)
 
     checkCudaErrors(cuFuncSetBlockShape(hKernel, nThreads, 1, 1));
-    //param_data
+
+    //Initialize param_data for kernel
     int paramOffset = 0;
-    checkCudaErrors(cuParamSetv(hKernel, paramOffset, &d_data, sizeof(d_data)));
-    paramOffset += sizeof(d_data);
+    for( std::vector< std::pair<size_t,unsigned char*> >::const_iterator i=param_data.begin(); i!=param_data.end(); i++ ) {
+        size_t memSize = nThreads * nBlocks * i->first;
+        unsigned char *d_data   = 0;
+        checkCudaErrors(cudaMalloc((void**)&d_data, memSize));
+        checkCudaErrors(cudaMemcpy(d_data,i->first,cudaMemcpyHostToDevice));
+        checkCudaErrors(cuParamSetv(hKernel, paramOffset, &d_data, memSize));
+        paramOffset += i->first;
+    }
     checkCudaErrors(cuParamSetSize(hKernel, paramOffset));
 
     // Launch the kernel (Driver API_)
     checkCudaErrors(cuLaunchGrid(hKernel, nBlocks, 1));
     std::cout << "CUDA kernel launched" << std::endl;
 
+    int         *h_data   = 0;
+    if ((h_data = (int *)malloc(memSize)) == NULL)
+    {
+        std::cerr << "Could not allocate host memory" << std::endl;
+        exit(EXIT_FAILURE);
+    }
     // Copy the result back to the host
     checkCudaErrors(cudaMemcpy(h_data, d_data, memSize, cudaMemcpyDeviceToHost));
 
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index d2b855c..1e4f1df 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -150,6 +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;
 int no_of_ptx=0;
 std::map<int, std::set<std::string> > version_filename;
 
@@ -487,8 +488,10 @@ __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size)
 {
 	CUctx_st* context = GPGPUSim_Context();
 	*devPtr = context->get_device()->get_gpgpu()->gpu_malloc(size);
-	if(g_debug_execution >= 3)
+	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;
+    }
 	if ( *devPtr  ) {
 		return g_last_cudaError = cudaSuccess;
 	} else {
@@ -2374,8 +2377,10 @@ cudaError_t CUDARTAPI cudaHostGetDevicePointer(void **pDevice, void *pHost, unsi
 	assert(i != pinned_memory_size.end());
 	size_t size = i->second;
 	*pDevice = gpu->gpu_malloc(size);
-	if(g_debug_execution >= 3)
+	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;
+    }
 	if ( *pDevice  ) {
 		pinned_memory[pHost]=pDevice;
 		//Copy contents in cpu to gpu
@@ -2617,8 +2622,9 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
 	g_ptx_kernel_count++;
 	fflush(stdout);
 
+	
 	if(g_debug_execution >= 3){
-        entry->debug_param();
+        entry->debug_param(g_devPtr_Size, result->get_param_memory());
     }
 
 	return result;
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index bdb1e8d..8795e1c 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -1226,33 +1226,72 @@ void function_info::list_param( FILE *fout ) const
    fflush(fout);
 }
 
-void function_info::debug_param() const
+void function_info::debug_param(std::map<void *, size_t> devPtr_Size, memory_space *param_mem) const
 {
     static unsigned long counter = 0;
-    std::string gpgpusim_path(getenv("GPGPUSIM_ROOT"));
-    assert(!gpgpusim_path.empty());
-    std::string command = "mkdir " + gpgpusim_path + "/debug_tools/WatchYourStep/data";
-    system(command.c_str());
-
-    std::string filename(gpgpusim_path + "/debug_tools/WatchYourStep/data/params.config" + std::to_string(counter++));
     std::vector< std::pair<size_t, unsigned char*> > param_data;
+    std::vector<bool> paramIsPointer;
+
+    char * gpgpusim_path = getenv("GPGPUSIM_ROOT");
+    assert(gpgpusim_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++));
+
+    //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("(");
+    pos2 = fn.find(")");
+    assert(pos2>pos1&&pos1>0);
+    strcpy(buff, fn.substr(pos1 + 1, pos2 - pos1 - 1).c_str());
+    char *tok;
+    tok = strtok(buff, ",");
+    while(tok!=NULL){
+        std::string param(tok);
+        if(param.find("*")!=std::string::npos){
+            paramIsPointer.push_back(true);
+        }else{
+            paramIsPointer.push_back(false);
+        }
+        tok = strtok(NULL, ",");
+    }
 
     for( std::map<unsigned,param_info>::const_iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) {
         const param_info &p = i->second;
+        std::string name = p.get_name();
+        symbol *param = m_symtab->lookup(name.c_str());
+        addr_t param_addr = param->get_address();
         param_t param_value = p.get_value();
-        unsigned char val[param_value.size];
-        memcpy((void*) val, (void*)((char*)param_value.pdata+param_value.offset), param_value.size);
-        param_data.push_back(std::pair<size_t, unsigned char*>(param_value.size,val));
+//        if (paramIsPointer[i->first]){
+//            assert(param_value.size==8);
+//        }else{
+            unsigned char val[param_value.size];
+            param_mem->read(param_addr,param_value.size,(void*)val);
+            param_data.push_back(std::pair<size_t, unsigned char*>(param_value.size,val));
+        //}
     }
 
     FILE *fout  = fopen (filename.c_str(), "w");
     fprintf(fout, "%s\n", get_name().c_str());
+    size_t index = 0;
     for( std::vector< std::pair<size_t,unsigned char*> >::const_iterator i=param_data.begin(); i!=param_data.end(); i++ ) {
+        if (paramIsPointer[index]){
+            fprintf(fout, "*");
+        }
         fprintf(fout, "%lu :", i->first);
         for (size_t j = 0; j<i->first; j++){
             fprintf(fout, " %u", i->second[j]);
         }
         fprintf(fout, "\n");
+        index++;
     }
     fflush(fout);
     fclose(fout);
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index 341f9b7..43907d3 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() const;
+   void debug_param(std::map<void *, size_t> devPtr_Size, memory_space *param_mem) const;
 
    const struct gpgpu_ptx_sim_info* get_kernel_info () const
    {
-- 
cgit v1.3


From 6ca8b878e29d5e21895a69de9e2240f3c578249c Mon Sep 17 00:00:00 2001
From: Jonathan <jonathanslew@outlook.com>
Date: Thu, 28 Jun 2018 16:46:45 -0700
Subject: loads params

---
 .../WatchYourStep/ptxjitplus/ptxjitplus.cpp        | 180 ++++++++-------------
 1 file changed, 63 insertions(+), 117 deletions(-)

(limited to 'debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp')

diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
index 8645114..c23d7f6 100644
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
+++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
@@ -22,6 +22,7 @@
 #include <math.h>
 #include <string.h>
 #include <stdio.h>
+#include <map>
 
 // CUDA driver & runtime
 #include <cuda.h>
@@ -111,88 +112,7 @@ void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUl
     checkCudaErrors(cuLinkDestroy(*lState));
 }
 
-void function_info::debug_param( ) const
-{
-   char filename[] = "params.txt";
-   char buff[1024];
-   snprintf(buff,1024,"c++filt %s > %s", get_name().c_str(), filename);
-   system(buff);
-   FILE *fp = fopen(filename, "r");
-   fgets(buff, 1024, fp);
-   fclose(fp);
-
-   std::string fn(buff);
-   size_t pos1, pos2;
-   pos1 = fn.find("(");
-   pos2 = fn.find(")");
-   assert(pos2>pos1&&pos1>0);
-   strcpy(buff, fn.substr(pos1 + 1, pos2 - pos1 - 1).c_str());
-   printf("params: %s\n", buff);
-   char *tok;
-   std::vector<std::string> params;
-   tok = strtok(buff, ",");
-   while(tok!=NULL){
-       std::string param(tok);
-       param.erase(0, param.find_first_not_of(" "));
-       param.erase(param.find_last_not_of(" ")+1);
-       params.push_back(param);
-       tok = strtok(NULL, ",");
-   }
-   for (auto const& it : params){
-      std::cout<<it<<std::endl;
-   }
-
-   FILE *fout  = fopen (filename, "w");
-   fprintf(fout, "Name of function:%s\n", fn.c_str());
-
-   for( std::map<unsigned,param_info>::const_iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) {
-      const param_info &p = i->second;
-      std::string name = p.get_name();
-      param_t param_value = p.get_value();
-      if(params[i->first].find("const")!=std::string::npos){
-         fprintf(fout, "Input: ");
-      } else {
-         fprintf(fout, "Input/output: ");
-      }
-
-      symbol *param = m_symtab->lookup(name.c_str());
-      addr_t param_addr = param->get_address();
-      fprintf(fout, "%s: %#08x, ", name.c_str(), param_addr);
-
-      if(params[i->first].find("int")!=std::string::npos){
-         size_t len = param_value.size/sizeof(int);
-         int val[len];
-         memcpy((void*) val, param_value.pdata+param_value.offset, param_value.size);
-         fprintf(fout, "val (int) = ");
-         for (unsigned i = 0; i<len; i++){
-             fprintf(fout, "%d ", val[i]);
-         }
-         fprintf(fout, "\n");
-      } else if(params[i->first].find("float")!=std::string::npos){
-         size_t len = param_value.size/sizeof(float);
-         float val[len];
-         memcpy((void*) val, param_value.pdata+param_value.offset, param_value.size);
-         fprintf(fout, "val (float) = ");
-         for (unsigned i = 0; i<len; i++){
-             fprintf(fout, "%f ", val[i]);
-         }
-         fprintf(fout, "\n");
-      }else{
-         size_t len = param_value.size/sizeof(char);
-         char val[len];
-         memcpy((void*) val, param_value.pdata+param_value.offset, param_value.size);
-         fprintf(fout, "val (char) = ");
-         for (unsigned i = 0; i<len; i++){
-             fprintf(fout, "%c ", val[i]);
-         }
-         fprintf(fout, "\n");
-      }
-   }
-   fflush(fout);
-   fclose(fout);
-}
-
-void* initializeData(std::vector< std::pair<size_t, unsigned char*> >& param_data)
+void initializeData(std::vector<unsigned char*>& param_data, std::vector< std::pair<size_t, bool> >& param_info)
 {
     char *wys_exec_path = getenv("WYS_EXEC_PATH");
     assert(wys_exec_path!=NULL);
@@ -208,33 +128,41 @@ void* initializeData(std::vector< std::pair<size_t, unsigned char*> >& param_dat
     assert(fin);
     char buff[1024];
     fscanf(fin, "%s\n", buff);
-    //void *retval = instrument_ptx_from_function(std::string(buff), path_to_search);
-    void *retval = NULL;
-    printf("%s\n", buff);
+    printf("Processing :%s ...\n", buff);
+    fflush(stdout);
     //fill data structure to pass in params later
     while (!feof(fin)){
+        std::pair<size_t, bool> info;
         int err;
         size_t len;
         unsigned val;
+        int start = fgetc(fin);
+        if (start == '*'){
+            info.second=true;
+        }else{
+            info.second=false;
+            int c = ungetc(start,fin);
+            assert(c==start&&"Couldn't ungetc\n");
+        }
         err = fscanf(fin, "%lu : ", &len);
+        info.first = len;
         assert( err==1 );
-        printf("%lu : ", len);
+        //printf("%lu : ", len);
         unsigned char params[len];
         for (size_t i=0; i<len; i++)
         {
             err = fscanf(fin, "%u ", &val);
             assert( err==1 );
-            printf("%u ", val);
+            //printf("%u ", val);
             params[i] = (unsigned char) val;
         }
-        param_data.push_back(std::pair<size_t, unsigned char*>(len, params));
+        param_info.push_back(info);
+        param_data.push_back(params);
         err = fscanf(fin, "\n");
         assert(err==0);
-        printf("\n");
+        //printf("\n");
     }
-    
     fclose(fin);
-    return retval;
 }
 
 int main(int argc, char **argv)
@@ -250,9 +178,9 @@ int main(int argc, char **argv)
     cudaDeviceProp deviceProp;
 
     printf("[%s] - Starting...\n", sSDKname);
-    std::vector< std::pair<size_t, unsigned char*> > param_data;
-    std::vector< std::pair<size_t, unsigned char*> > device_data;
-    void* storedReg = initializeData(param_data);
+    std::vector<unsigned char*> param_data;
+    std::vector< std::pair<size_t, bool> > param_info;
+    initializeData(param_data,param_info);
 
     if (checkCmdLineFlag(argc, (const char **)argv, "device"))
     {
@@ -302,14 +230,20 @@ int main(int argc, char **argv)
     checkCudaErrors(cuFuncSetBlockShape(hKernel, nThreads, 1, 1));
 
     //Initialize param_data for kernel
+    std::map< size_t, unsigned char* > m_device_data;
     int paramOffset = 0;
-    for( std::vector< std::pair<size_t,unsigned char*> >::const_iterator i=param_data.begin(); i!=param_data.end(); i++ ) {
-        size_t memSize = nThreads * nBlocks * i->first;
-        unsigned char *d_data   = 0;
-        checkCudaErrors(cudaMalloc((void**)&d_data, memSize));
-        checkCudaErrors(cudaMemcpy(d_data,i->first,cudaMemcpyHostToDevice));
-        checkCudaErrors(cuParamSetv(hKernel, paramOffset, &d_data, memSize));
-        paramOffset += i->first;
+    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;
+            paramOffset += 8;
+        }else{
+            checkCudaErrors(cuParamSetv(hKernel, paramOffset, param_data[i], param_info[i].first));
+            paramOffset += param_info[i].first;
+        }
     }
     checkCudaErrors(cuParamSetSize(hKernel, paramOffset));
 
@@ -317,15 +251,24 @@ int main(int argc, char **argv)
     checkCudaErrors(cuLaunchGrid(hKernel, nBlocks, 1));
     std::cout << "CUDA kernel launched" << std::endl;
 
-    int         *h_data   = 0;
-    if ((h_data = (int *)malloc(memSize)) == NULL)
-    {
-        std::cerr << "Could not allocate host memory" << std::endl;
-        exit(EXIT_FAILURE);
+    std::map< size_t, unsigned char* > m_output_data;
+    for(std::map< size_t, unsigned char* >::iterator i = m_device_data.begin(); i!=m_device_data.begin(); i++){
+        unsigned char *h_data   = 0;
+        if ((h_data = (unsigned char *)malloc(param_info[i->first].first)) == 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));
+        m_output_data[i->first] = h_data;
+    }
+
+    for(std::map< size_t, unsigned char* >::iterator i = m_output_data.begin(); i!=m_output_data.begin(); i++){
+        //print data out to file
     }
-    // Copy the result back to the host
-    checkCudaErrors(cudaMemcpy(h_data, d_data, memSize, cudaMemcpyDeviceToHost));
 
+    int* h_data;
     // Check the result
     bool dataGood = true;
 
@@ -338,17 +281,20 @@ int main(int argc, char **argv)
         }
     }
 
-    // Cleanup
-    if (d_data)
-    {
-        checkCudaErrors(cudaFree(d_data));
-        d_data = 0;
+    //Cleanup
+    for(std::map< size_t, unsigned char* >::iterator i = m_device_data.begin(); i!=m_device_data.begin(); i++){
+        if (i->second){
+            checkCudaErrors(cudaFree(i->second));
+            i->second = 0;
+        }
     }
 
-    if (h_data)
-    {
-        free(h_data);
-        h_data = 0;
+    for(std::map< size_t, unsigned char* >::iterator i = m_output_data.begin(); i!=m_output_data.begin(); i++){
+        if (i->second)
+        {
+            free(i->second);
+            i->second = 0;
+        }
     }
 
     if (hModule)
-- 
cgit v1.3


From 2e8ecbcb1a30768ddcac9a1c6e386b873bbc8e5e Mon Sep 17 00:00:00 2001
From: Jonathan <jonathanslew@outlook.com>
Date: Fri, 29 Jun 2018 16:55:52 -0700
Subject: dumps global using pointer, dump output from ptxjit to file

---
 .../WatchYourStep/ptxjitplus/ptxjitplus.cpp        | 25 ++++++++++++++++------
 src/cuda-sim/cuda-sim.cc                           | 15 +++++++------
 2 files changed, 28 insertions(+), 12 deletions(-)

(limited to 'debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp')

diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
index c23d7f6..e539dbd 100644
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
+++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
@@ -37,7 +37,9 @@
 #include "ptxinst.h"
 
 const char *sSDKname = "PTX Just In Time (JIT) Compilation (no-qatest)";
-
+char *wys_exec_path;
+char *wys_exec_name;
+char *wys_launch_num;
 
 void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUlinkState *lState)
 {
@@ -114,12 +116,12 @@ void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUl
 
 void initializeData(std::vector<unsigned char*>& param_data, std::vector< std::pair<size_t, bool> >& param_info)
 {
-    char *wys_exec_path = getenv("WYS_EXEC_PATH");
+    wys_exec_path = getenv("WYS_EXEC_PATH");
     assert(wys_exec_path!=NULL);
-    char *wys_exec_name = getenv("WYS_EXEC_NAME");
+    wys_exec_name = getenv("WYS_EXEC_NAME");
     assert(wys_exec_name!=NULL);
     std::string path_to_search = std::string(wys_exec_path) + "/" + wys_exec_name + ".*.ptx";
-    char* wys_launch_num(getenv("WYS_LAUNCH_NUM"));
+    wys_launch_num = getenv("WYS_LAUNCH_NUM");
     assert(wys_launch_num!=NULL);
     std::string filename = std::string("../data/params.config") + wys_launch_num;
 
@@ -178,7 +180,9 @@ int main(int argc, char **argv)
     cudaDeviceProp deviceProp;
 
     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);
 
@@ -229,8 +233,9 @@ int main(int argc, char **argv)
 
     checkCudaErrors(cuFuncSetBlockShape(hKernel, nThreads, 1, 1));
 
-    //Initialize param_data for kernel
+    //maps param number to pointer to device data
     std::map< size_t, unsigned char* > m_device_data;
+    //Initialize param_data for kernel
     int paramOffset = 0;
     for( size_t i = 0; i<param_data.size(); i++){
         if(param_info[i].second){
@@ -251,6 +256,7 @@ int main(int argc, char **argv)
     checkCudaErrors(cuLaunchGrid(hKernel, nBlocks, 1));
     std::cout << "CUDA kernel launched" << std::endl;
 
+    //maps param number to pointer to output data
     std::map< size_t, unsigned char* > m_output_data;
     for(std::map< size_t, unsigned char* >::iterator i = m_device_data.begin(); i!=m_device_data.begin(); i++){
         unsigned char *h_data   = 0;
@@ -264,8 +270,15 @@ int main(int argc, char **argv)
         m_output_data[i->first] = h_data;
     }
 
+    std::string filename = std::string("../data/wys.out") + wys_launch_num;
+    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.begin(); i++){
-        //print data out to file
+        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, " %u", i->second[j]);
+        }
+        fprintf(fout, "\n");
     }
 
     int* h_data;
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index bade750..8f06a12 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -1234,8 +1234,8 @@ void function_info::debug_param(std::map<unsigned long long, size_t> mallocPtr_S
 
     char * gpgpusim_path = getenv("GPGPUSIM_ROOT");
     assert(gpgpusim_path!=NULL);
-//    std::string command = std::string("mkdir ") + gpgpusim_path + "/debug_tools/WatchYourStep/data";
-//    system(command.c_str());
+    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++));
 
     for( std::map<unsigned,param_info>::const_iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) {
@@ -1245,11 +1245,14 @@ void function_info::debug_param(std::map<unsigned long long, size_t> mallocPtr_S
         addr_t param_addr = param->get_address();
         param_t param_value = p.get_value();
 
-        if(param_value.size==8&&mallocPtr_Size.find(*(unsigned long long*)param_value.pdata)!=mallocPtr_Size.end()){
+        if(param_value.size==sizeof(void*) && mallocPtr_Size.find(*(unsigned long long*)param_value.pdata)!=mallocPtr_Size.end()){
+            //is pointer
             size_t array_size = mallocPtr_Size[*(unsigned long long*)param_value.pdata];
-            unsigned char val[array_size];
-            gpu->get_global_memory()->read(param_addr,array_size,(void*)val);
-            param_data.push_back(std::pair<size_t, unsigned char*>(array_size,val));
+            unsigned char val[param_value.size];
+            param_mem->read(param_addr,param_value.size,(void*)val);
+            unsigned char array_val[array_size];
+            gpu->get_global_memory()->read(*(unsigned*)((void*)val),array_size,(void*)array_val);
+            param_data.push_back(std::pair<size_t, unsigned char*>(array_size,array_val));
             paramIsPointer.push_back(true);
         }else{
             unsigned char val[param_value.size];
-- 
cgit v1.3


From b1ce8d27bae0bbf48d97bfc83b001b8d95209f99 Mon Sep 17 00:00:00 2001
From: Jonathan <jonathanslew@outlook.com>
Date: Tue, 3 Jul 2018 09:40:33 -0700
Subject: minor bug fix

---
 debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp | 13 +++++++++----
 1 file changed, 9 insertions(+), 4 deletions(-)

(limited to 'debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp')

diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
index e539dbd..df034aa 100644
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
+++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
@@ -258,7 +258,7 @@ int main(int argc, char **argv)
 
     //maps param number to pointer to output data
     std::map< size_t, unsigned char* > m_output_data;
-    for(std::map< size_t, unsigned char* >::iterator i = m_device_data.begin(); i!=m_device_data.begin(); i++){
+    for(std::map< size_t, unsigned char* >::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)
         {
@@ -273,15 +273,17 @@ int main(int argc, char **argv)
     std::string filename = std::string("../data/wys.out") + wys_launch_num;
     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.begin(); i++){
+    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, " %u", i->second[j]);
         }
         fprintf(fout, "\n");
     }
+    fflush(fout);
+    fclose(fout);
 
-    int* h_data;
+    int* h_data = (int*) m_output_data[0];
     // Check the result
     bool dataGood = true;
 
@@ -293,9 +295,12 @@ int main(int argc, char **argv)
             dataGood = false;
         }
     }
+    if(dataGood){
+        std::cout<<"OK!"<<std::endl;
+    }
 
     //Cleanup
-    for(std::map< size_t, unsigned char* >::iterator i = m_device_data.begin(); i!=m_device_data.begin(); i++){
+    for(std::map< size_t, unsigned char* >::iterator i = m_device_data.begin(); i!=m_device_data.end(); i++){
         if (i->second){
             checkCudaErrors(cudaFree(i->second));
             i->second = 0;
-- 
cgit v1.3


From 330cca35c59490bd8b96bfaa6688b5ea5163a735 Mon Sep 17 00:00:00 2001
From: Jonathan <jonathanslew@outlook.com>
Date: Tue, 3 Jul 2018 11:08:13 -0700
Subject: correctly load params

---
 debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp | 14 +++++++++++++-
 1 file changed, 13 insertions(+), 1 deletion(-)

(limited to 'debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp')

diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
index df034aa..e9b16f1 100644
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
+++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
@@ -150,7 +150,7 @@ void initializeData(std::vector<unsigned char*>& param_data, std::vector< std::p
         info.first = len;
         assert( err==1 );
         //printf("%lu : ", len);
-        unsigned char params[len];
+        unsigned char* params = (unsigned char*) malloc(len*sizeof(unsigned char));
         for (size_t i=0; i<len; i++)
         {
             err = fscanf(fin, "%u ", &val);
@@ -165,6 +165,16 @@ void initializeData(std::vector<unsigned char*>& param_data, std::vector< std::p
         //printf("\n");
     }
     fclose(fin);
+    //filename = std::string("../data/wys.out") + wys_launch_num + "_param";
+    //fout = fopen(filename.c_str(), "w");
+    //assert(fout);
+    //fprintf(fout, "param %zu: size = %zu, data = ", 0,param_info[0].first);
+    //for (size_t j = 0; j<param_info[0].first; j++){
+    //    fprintf(fout, " %u", i->second[j]);
+    //}
+    //fprintf(fout, "\n");
+    //fflush(fout);
+    //fclose(fout);
 }
 
 int main(int argc, char **argv)
@@ -186,6 +196,8 @@ int main(int argc, char **argv)
     std::vector< std::pair<size_t, bool> > param_info;
     initializeData(param_data,param_info);
 
+
+
     if (checkCmdLineFlag(argc, (const char **)argv, "device"))
     {
         cuda_device = getCmdLineArgumentInt(argc, (const char **)argv, "device=");
-- 
cgit v1.3


From 2dfdeafc0c4dbc82cfaa8d22c5a5dcc303568a34 Mon Sep 17 00:00:00 2001
From: Jonathan <jonathanslew@outlook.com>
Date: Tue, 3 Jul 2018 16:56:08 -0700
Subject: loading params works now

---
 .gitignore                                         |  1 +
 debug_tools/WatchYourStep/ptxjitplus/myPtx32.ptx   | 33 -------------
 debug_tools/WatchYourStep/ptxjitplus/myPtx64.ptx   | 34 --------------
 .../WatchYourStep/ptxjitplus/ptxjitplus.cpp        | 54 +++++++++++++---------
 libcuda/cuda_runtime_api.cc                        |  3 +-
 src/cuda-sim/cuda-sim.cc                           |  4 +-
 6 files changed, 37 insertions(+), 92 deletions(-)
 delete mode 100644 debug_tools/WatchYourStep/ptxjitplus/myPtx32.ptx
 delete mode 100644 debug_tools/WatchYourStep/ptxjitplus/myPtx64.ptx

(limited to 'debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp')

diff --git a/.gitignore b/.gitignore
index 0d43bd1..150efbf 100644
--- a/.gitignore
+++ b/.gitignore
@@ -50,3 +50,4 @@ debug_tools/WatchYourStep/ptxjitplus/*.icnt
 debug_tools/WatchYourStep/ptxjitplus/gpgpu*
 debug_tools/WatchYourStep/ptxjitplus/*.old
 debug_tools/WatchYourStep/ptxjitplus/ptxjitplus
+debug_tools/WatchYourStep/ptxjitplus/*.ptx
diff --git a/debug_tools/WatchYourStep/ptxjitplus/myPtx32.ptx b/debug_tools/WatchYourStep/ptxjitplus/myPtx32.ptx
deleted file mode 100644
index 28e8e4d..0000000
--- a/debug_tools/WatchYourStep/ptxjitplus/myPtx32.ptx
+++ /dev/null
@@ -1,33 +0,0 @@
-//
-// Generated by NVIDIA NVVM Compiler
-// Compiler built on Sat Mar 30 20:51:33 2013 (1364691093)
-// Cuda compilation tools, release 5.5, V5.5.0
-//
-
-.version 3.2
-.target sm_20
-.address_size 32
-
-.visible .entry _Z8myKernelPi(
-	.param .u32 _Z8myKernelPi_param_0
-)
-{
-	.reg .s32 	%r<9>;
-
-
-	ld.param.u32 	%r1, [_Z8myKernelPi_param_0];
-	cvta.to.global.u32 	%r2, %r1;
-	.loc 1 3 1
-	mov.u32 	%r3, %ntid.x;
-	mov.u32 	%r4, %ctaid.x;
-	mov.u32 	%r5, %tid.x;
-	mad.lo.s32 	%r6, %r3, %r4, %r5;
-	.loc 1 4 1
-	shl.b32 	%r7, %r6, 2;
-	add.s32 	%r8, %r2, %r7;
-	st.global.u32 	[%r8], %r6;
-	.loc 1 5 2
-	ret;
-}
-
-
diff --git a/debug_tools/WatchYourStep/ptxjitplus/myPtx64.ptx b/debug_tools/WatchYourStep/ptxjitplus/myPtx64.ptx
deleted file mode 100644
index 892aaa5..0000000
--- a/debug_tools/WatchYourStep/ptxjitplus/myPtx64.ptx
+++ /dev/null
@@ -1,34 +0,0 @@
-//
-// Generated by NVIDIA NVVM Compiler
-// Compiler built on Sat Mar 30 20:51:33 2013 (1364691093)
-// Cuda compilation tools, release 5.5, V5.5.0
-//
-
-.version 3.2
-.target sm_20
-.address_size 64
-
-.visible .entry _Z8myKernelPi(
-	.param .u64 _Z8myKernelPi_param_0
-)
-{
-	.reg .s32 	%r<5>;
-	.reg .s64 	%rd<5>;
-
-
-	ld.param.u64 	%rd1, [_Z8myKernelPi_param_0];
-	cvta.to.global.u64 	%rd2, %rd1;
-	.loc 1 3 1
-	mov.u32 	%r1, %ntid.x;
-	mov.u32 	%r2, %ctaid.x;
-	mov.u32 	%r3, %tid.x;
-	mad.lo.s32 	%r4, %r1, %r2, %r3;
-	mul.wide.s32 	%rd3, %r4, 4;
-	add.s64 	%rd4, %rd2, %rd3;
-	.loc 1 4 1
-	st.global.u32 	[%rd4], %r4;
-	.loc 1 5 2
-	ret;
-}
-
-
diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
index e9b16f1..65dfa84 100644
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
+++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
@@ -40,6 +40,7 @@ const char *sSDKname = "PTX Just In Time (JIT) Compilation (no-qatest)";
 char *wys_exec_path;
 char *wys_exec_name;
 char *wys_launch_num;
+std::string kernelName;
 
 void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUlinkState *lState)
 {
@@ -108,7 +109,7 @@ void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUl
     checkCudaErrors(cuModuleLoadData(phModule, cuOut));
 
     // Locate the kernel entry poin
-    checkCudaErrors(cuModuleGetFunction(phKernel, *phModule, "_Z8myKernelPi"));
+    checkCudaErrors(cuModuleGetFunction(phKernel, *phModule, kernelName.c_str()));
 
     // Destroy the linker invocation
     checkCudaErrors(cuLinkDestroy(*lState));
@@ -125,13 +126,13 @@ void initializeData(std::vector<unsigned char*>& param_data, std::vector< std::p
     assert(wys_launch_num!=NULL);
     std::string filename = std::string("../data/params.config") + wys_launch_num;
 
-    //instrument ptx
     FILE *fin = fopen(filename.c_str(), "r");
     assert(fin);
     char buff[1024];
     fscanf(fin, "%s\n", buff);
     printf("Processing :%s ...\n", buff);
     fflush(stdout);
+    kernelName = std::string(buff);
     //fill data structure to pass in params later
     while (!feof(fin)){
         std::pair<size_t, bool> info;
@@ -179,8 +180,8 @@ void initializeData(std::vector<unsigned char*>& param_data, std::vector< std::p
 
 int main(int argc, char **argv)
 {
-    const unsigned int nThreads = 256;
-    const unsigned int nBlocks  = 64;
+    const unsigned int nThreads = 2;
+    const unsigned int nBlocks  = 2;
 
     CUmodule     hModule  = 0;
     CUfunction   hKernel  = 0;
@@ -238,12 +239,20 @@ int main(int argc, char **argv)
 
 
 
+    // Allocate memory on host and device (Runtime API)
+    // NOTE: The runtime API will create the GPU Context implicitly here
+    int         *d_tmp   = 0;
+    checkCudaErrors(cudaMalloc(&d_tmp, 1));
+
     // JIT Compile the Kernel from PTX and get the Handles (Driver API)
     ptxJIT(argc, argv, &hModule, &hKernel, &lState);
+    checkCudaErrors(cudaFree(d_tmp));
+
 
     // Set the kernel parameters (Driver API)
 
-    checkCudaErrors(cuFuncSetBlockShape(hKernel, nThreads, 1, 1));
+    // TODO: automatically load these values in
+    checkCudaErrors(cuFuncSetBlockShape(hKernel, 128, 1, 1));
 
     //maps param number to pointer to device data
     std::map< size_t, unsigned char* > m_device_data;
@@ -265,7 +274,8 @@ int main(int argc, char **argv)
     checkCudaErrors(cuParamSetSize(hKernel, paramOffset));
 
     // Launch the kernel (Driver API_)
-    checkCudaErrors(cuLaunchGrid(hKernel, nBlocks, 1));
+    // TODO: automatically load these values in
+    checkCudaErrors(cuLaunchGrid(hKernel, 1, 20));
     std::cout << "CUDA kernel launched" << std::endl;
 
     //maps param number to pointer to output data
@@ -295,21 +305,21 @@ int main(int argc, char **argv)
     fflush(fout);
     fclose(fout);
 
-    int* h_data = (int*) m_output_data[0];
-    // Check the result
-    bool dataGood = true;
-
-    for (unsigned int i = 0 ; dataGood && i < nBlocks * nThreads ; i++)
-    {
-        if (h_data[i] != (int)i)
-        {
-            std::cerr << "Error at " << i << std::endl;
-            dataGood = false;
-        }
-    }
-    if(dataGood){
-        std::cout<<"OK!"<<std::endl;
-    }
+//    int* h_data = (int*) m_output_data[0];
+//    // Check the result
+//    bool dataGood = true;
+//
+//    for (unsigned int i = 0 ; dataGood && i < nBlocks * nThreads ; i++)
+//    {
+//        if (h_data[i] != (int)i)
+//        {
+//            std::cerr << "Error at " << i << std::endl;
+//            dataGood = false;
+//        }
+//    }
+//    if(dataGood){
+//        std::cout<<"OK!"<<std::endl;
+//    }
 
     //Cleanup
     for(std::map< size_t, unsigned char* >::iterator i = m_device_data.begin(); i!=m_device_data.end(); i++){
@@ -333,5 +343,5 @@ int main(int argc, char **argv)
         hModule = 0;
     }
 
-    return dataGood ? EXIT_SUCCESS : EXIT_FAILURE;
+    //return dataGood ? EXIT_SUCCESS : EXIT_FAILURE;
 }
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 53b4da8..cad4d87 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -1076,7 +1076,6 @@ __host__ cudaError_t CUDARTAPI cudaSetupArgument(const void *arg, size_t size, s
 	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());
 
 	return g_last_cudaError = cudaSuccess;
 }
@@ -2631,6 +2630,7 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
 	return result;
 }
 
+#if (CUDART_VERSION >= 8000)
 CUresult CUDAAPI cuLinkCreate(unsigned int numOptions, CUjit_option *options, void **optionValues, CUlinkState *stateOut)
 {
     //currently do not support options or multiple CUlinkStates
@@ -2754,3 +2754,4 @@ CUresult CUDAAPI cuLaunchGrid(CUfunction f, int grid_width, int grid_height)
     cudaLaunch(hostFun);
 	return CUDA_SUCCESS;
 }
+#endif
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 8f06a12..62aef6d 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -1248,9 +1248,9 @@ void function_info::debug_param(std::map<unsigned long long, size_t> mallocPtr_S
         if(param_value.size==sizeof(void*) && mallocPtr_Size.find(*(unsigned long long*)param_value.pdata)!=mallocPtr_Size.end()){
             //is pointer
             size_t array_size = mallocPtr_Size[*(unsigned long long*)param_value.pdata];
-            unsigned char val[param_value.size];
+            unsigned char* val = (unsigned char*) malloc(param_value.size);
             param_mem->read(param_addr,param_value.size,(void*)val);
-            unsigned char array_val[array_size];
+            unsigned char* array_val = (unsigned char*) malloc(array_size);
             gpu->get_global_memory()->read(*(unsigned*)((void*)val),array_size,(void*)array_val);
             param_data.push_back(std::pair<size_t, unsigned char*>(array_size,array_val));
             paramIsPointer.push_back(true);
-- 
cgit v1.3


From 1fddb06b153a3e5fbc255e89bb6861cce1319039 Mon Sep 17 00:00:00 2001
From: Jonathan <jonathanslew@outlook.com>
Date: Wed, 4 Jul 2018 12:52:42 -0700
Subject: dump and load block and grid size and launch

---
 .../WatchYourStep/ptxjitplus/ptxjitplus.cpp        | 14 ++++++-------
 libcuda/cuda_runtime_api.cc                        | 24 ++++++++++++++++++++--
 src/cuda-sim/cuda-sim.cc                           |  3 ++-
 src/cuda-sim/ptx_ir.h                              |  2 +-
 4 files changed, 32 insertions(+), 11 deletions(-)

(limited to 'debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp')

diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
index 65dfa84..df0ff8d 100644
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
+++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
@@ -40,6 +40,7 @@ const char *sSDKname = "PTX Just In Time (JIT) Compilation (no-qatest)";
 char *wys_exec_path;
 char *wys_exec_name;
 char *wys_launch_num;
+dim3 gridDim, blockDim;
 std::string kernelName;
 
 void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUlinkState *lState)
@@ -133,6 +134,7 @@ void initializeData(std::vector<unsigned char*>& param_data, std::vector< std::p
     printf("Processing :%s ...\n", buff);
     fflush(stdout);
     kernelName = std::string(buff);
+    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;
@@ -248,12 +250,6 @@ int main(int argc, char **argv)
     ptxJIT(argc, argv, &hModule, &hKernel, &lState);
     checkCudaErrors(cudaFree(d_tmp));
 
-
-    // Set the kernel parameters (Driver API)
-
-    // TODO: automatically load these values in
-    checkCudaErrors(cuFuncSetBlockShape(hKernel, 128, 1, 1));
-
     //maps param number to pointer to device data
     std::map< size_t, unsigned char* > m_device_data;
     //Initialize param_data for kernel
@@ -275,7 +271,8 @@ int main(int argc, char **argv)
 
     // Launch the kernel (Driver API_)
     // TODO: automatically load these values in
-    checkCudaErrors(cuLaunchGrid(hKernel, 1, 20));
+    CUDAAPI cuLaunchKernel(hKernel, gridDim.x, gridDim.y, gridDim.z, blockDim.x, blockDim.y, blockDim.z, 
+        0, NULL, NULL, NULL);
     std::cout << "CUDA kernel launched" << std::endl;
 
     //maps param number to pointer to output data
@@ -299,6 +296,9 @@ int main(int argc, char **argv)
         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, " %u", i->second[j]);
+            if (j&&(!(j%20))){
+                fprintf(fout, "\n");
+            }
         }
         fprintf(fout, "\n");
     }
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index cad4d87..410f15f 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -2624,7 +2624,7 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
 	fflush(stdout);
 	
 	if(g_debug_execution >= 3){
-        entry->debug_param(g_mallocPtr_Size, result->get_param_memory(), (gpgpu_t *) context->get_device()->get_gpgpu());
+        entry->debug_param(g_mallocPtr_Size, result->get_param_memory(), (gpgpu_t *) context->get_device()->get_gpgpu(), gridDim, blockDim);
     }
 
 	return result;
@@ -2744,7 +2744,6 @@ CUresult CUDAAPI cuParamSetv(CUfunction hfunc, int offset, void *ptr, unsigned i
 
 CUresult CUDAAPI cuLaunchGrid(CUfunction f, int grid_width, int grid_height)
 {
-	CUctx_st* context = GPGPUSim_Context();
     const char *hostFun = (const char*) f;
 	gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
 	kernel_config &config = g_cuda_launch_stack.back();
@@ -2754,4 +2753,25 @@ CUresult CUDAAPI cuLaunchGrid(CUfunction f, int grid_width, int grid_height)
     cudaLaunch(hostFun);
 	return CUDA_SUCCESS;
 }
+
+
+CUresult CUDAAPI cuLaunchKernel(CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, 
+        unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, 
+        unsigned int sharedMemBytes, CUstream hStream, void **kernelParams, void **extra)
+{
+    if (sharedMemBytes!=0||hStream!=NULL||kernelParams!=NULL||extra!=NULL){
+	    printf("GPGPU-Sim CUDA DRIVER API: Warning: Currently do not support \nsharedMemBytes, hStream, kernelParams, and extra parameters.\n");
+    }
+    const char *hostFun = (const char*) f;
+	gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
+	kernel_config &config = g_cuda_launch_stack.back();
+    dim3 *d_b = new dim3(blockDimX, blockDimY, blockDimZ);
+	config.set_block_dim(d_b);
+    dim3 *d_g = new dim3(gridDimX, gridDimY, gridDimZ);
+	config.set_grid_dim(d_g);
+
+    cudaLaunch(hostFun);
+	return CUDA_SUCCESS;
+}
+
 #endif
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 62aef6d..3003c4c 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<unsigned long long, size_t> mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu) const
+void function_info::debug_param(std::map<unsigned long long, size_t> mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) const
 {
     static unsigned long counter = 0;
     std::vector< std::pair<size_t, unsigned char*> > param_data;
@@ -1265,6 +1265,7 @@ void function_info::debug_param(std::map<unsigned long long, size_t> mallocPtr_S
     FILE *fout  = fopen (filename.c_str(), "w");
     printf("Writing data to %s ...\n", filename.c_str());
     fprintf(fout, "%s\n", get_name().c_str());
+    fprintf(fout, "%u,%u,%u %u,%u,%u\n", gridDim.x, gridDim.y, gridDim.z, blockDim.x, blockDim.y, blockDim.z);
     size_t index = 0;
     for( std::vector< std::pair<size_t,unsigned char*> >::const_iterator i=param_data.begin(); i!=param_data.end(); i++ ) {
         if (paramIsPointer[index]){
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index 449d615..b29621c 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<unsigned long long, size_t> mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu) const;
+   void debug_param(std::map<unsigned long long, size_t> mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) const;
 
    const struct gpgpu_ptx_sim_info* get_kernel_info () const
    {
-- 
cgit v1.3


From 7c4f9d7bff7b1725f3da6ddc8abf3f77a1cbe6f5 Mon Sep 17 00:00:00 2001
From: Jonathan <jonathanslew@outlook.com>
Date: Wed, 4 Jul 2018 16:41:20 -0700
Subject: dumps ptx kernels and scripts to launch all kernels

---
 debug_tools/WatchYourStep/ptxjitplus/launchkernels |  3 ++
 .../WatchYourStep/ptxjitplus/ptxjitplus.cpp        | 10 +++---
 libcuda/cuda_runtime_api.cc                        |  2 +-
 src/cuda-sim/cuda-sim.cc                           | 37 ++++++++++++++++++++--
 src/cuda-sim/ptx_ir.h                              |  2 +-
 5 files changed, 44 insertions(+), 10 deletions(-)
 create mode 100644 debug_tools/WatchYourStep/ptxjitplus/launchkernels

(limited to 'debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp')

diff --git a/debug_tools/WatchYourStep/ptxjitplus/launchkernels b/debug_tools/WatchYourStep/ptxjitplus/launchkernels
new file mode 100644
index 0000000..d2fd015
--- /dev/null
+++ b/debug_tools/WatchYourStep/ptxjitplus/launchkernels
@@ -0,0 +1,3 @@
+#Launches kernels from $1 to $2
+#Note must source this script
+for num in $(eval echo {$1..$2}); do export WYS_LAUNCH_NUM=$num; echo Launching kernel $num...; ./ptxjitplus; done
diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
index df0ff8d..1f094ba 100644
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
+++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
@@ -81,18 +81,16 @@ void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUl
     if (sizeof(void *)==4)
     {
         // Load the PTX from the string myPtx32
-        printf("Loading myPtx32[] program\n");
-        // PTX May also be loaded from file, as per below.
-        myErr = cuLinkAddData(*lState, CU_JIT_INPUT_PTX, (void *)myPtx32, strlen(myPtx32)+1, 0, 0, 0, 0);
+        printf("Loading myPtx32[] program...\n");
+        printf("WARNING: 32-bit execution is untested");
     }
     else
     {
         // Load the PTX from the string myPtx (64-bit)
         printf("Loading myPtx[] program\n");
-        //myErr = cuLinkAddData(*lState, CU_JIT_INPUT_PTX, (void *)myPtx64, strlen(myPtx64)+1, 0, 0, 0, 0);
-        // PTX May also be loaded from file, as per below.
-        myErr = cuLinkAddFile(*lState, CU_JIT_INPUT_PTX, "myPtx64.ptx",0,0,0);
     }
+    std::string ptx_file (std::string("../data/ptx.config") + wys_launch_num);
+    myErr = cuLinkAddFile(*lState, CU_JIT_INPUT_PTX, ptx_file.c_str(),0,0,0);
 
     if (myErr != CUDA_SUCCESS)
     {
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 410f15f..b67ea85 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -2624,7 +2624,7 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
 	fflush(stdout);
 	
 	if(g_debug_execution >= 3){
-        entry->debug_param(g_mallocPtr_Size, result->get_param_memory(), (gpgpu_t *) context->get_device()->get_gpgpu(), gridDim, blockDim);
+        entry->ptx_jit_config(g_mallocPtr_Size, result->get_param_memory(), (gpgpu_t *) context->get_device()->get_gpgpu(), gridDim, blockDim);
     }
 
 	return result;
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 3003c4c..fc4f82a 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<unsigned long long, size_t> mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) const
+void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) const
 {
     static unsigned long counter = 0;
     std::vector< std::pair<size_t, unsigned char*> > param_data;
@@ -1236,7 +1236,7 @@ void function_info::debug_param(std::map<unsigned long long, size_t> mallocPtr_S
     assert(gpgpusim_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++));
+    std::string filename(std::string(gpgpusim_path) + "/debug_tools/WatchYourStep/data/params.config" + std::to_string(counter));
 
     for( std::map<unsigned,param_info>::const_iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) {
         const param_info &p = i->second;
@@ -1280,6 +1280,39 @@ void function_info::debug_param(std::map<unsigned long long, size_t> mallocPtr_S
     }
     fflush(fout);
     fclose(fout);
+    
+    //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());
+    system(buff);
+    FILE *fin = fopen(ptx_config_fn.c_str(), "r");
+    char ptx_source[256];
+    unsigned line_number;
+    int numscanned = fscanf(fin, "%[^:]:%u", ptx_source, &line_number);
+    assert(numscanned == 2);
+    fclose(fin);
+    snprintf(buff, 1024, "grep -rn \".version\" %s | cut -d \":\" -f 1 | xargs -I \"{}\" awk \"NR>={}&&NR<={}+2\" %s > %s", ptx_source, ptx_source, ptx_config_fn.c_str());
+    system(buff);
+    fin = fopen(ptx_source, "r");
+    assert(fin!=NULL);
+    fout = fopen(ptx_config_fn.c_str(), "a");
+    assert(fout!=NULL);
+    for (unsigned i = 0; i<line_number; i++){
+        fgets(buff, 1024, fin);
+        assert(!feof(fin));
+    }
+    fprintf(fout, "\n\n");
+    do{
+        fprintf(fout, "%s", buff);
+        fgets(buff, 1024, fin);
+        assert(!feof(fin));
+    } while(strstr(buff, "entry")==NULL);
+
+    fclose(fin);
+    fflush(fout);
+    fclose(fout);
+    counter++;
 }
 
 template<int activate_level> 
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index b29621c..e0b5e96 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<unsigned long long, size_t> mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) const;
+   void ptx_jit_config(std::map<unsigned long long, size_t> mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) const;
 
    const struct gpgpu_ptx_sim_info* get_kernel_info () const
    {
-- 
cgit v1.3


From dd4728921136fa5f3a2e6c9591eed852ba871cb2 Mon Sep 17 00:00:00 2001
From: Jonathan <jonathanslew@outlook.com>
Date: Mon, 9 Jul 2018 15:02:24 -0700
Subject: fixed pointer bug in ptx jit and added support for executing
 ptxjitconfig when running ptxjit

---
 debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp | 11 ++++++-----
 libcuda/cuda_runtime_api.cc                         |  2 +-
 src/cuda-sim/cuda-sim.cc                            |  5 ++++-
 3 files changed, 11 insertions(+), 7 deletions(-)

(limited to 'debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp')

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> 
-- 
cgit v1.3


From d072fa72c6ffce7f34520d24c8cb285e0d5b92b2 Mon Sep 17 00:00:00 2001
From: Jonathan <jonathanslew@outlook.com>
Date: Fri, 13 Jul 2018 10:46:47 -0700
Subject: Add more print statements to debug api calls Fixed bug with 32bit rem
 Fixed bug with launching kernels in ptxjitplus Removed extraneous files
 Cleaned up more pointers Added documentation for ptxjitplus

---
 debug_tools/WatchYourStep/ptxjitplus/Makefile      |   5 +-
 debug_tools/WatchYourStep/ptxjitplus/ptxinst.cpp   |  14 ---
 debug_tools/WatchYourStep/ptxjitplus/ptxinst.h     |  14 ---
 .../WatchYourStep/ptxjitplus/ptxjitplus.cpp        | 103 ++++++++++++---------
 libcuda/cuda_runtime_api.cc                        |  36 +++++++
 src/cuda-sim/cuda-sim.cc                           |   4 +-
 src/cuda-sim/instructions.cc                       |  16 +++-
 7 files changed, 113 insertions(+), 79 deletions(-)
 delete mode 100644 debug_tools/WatchYourStep/ptxjitplus/ptxinst.cpp
 delete mode 100644 debug_tools/WatchYourStep/ptxjitplus/ptxinst.h

(limited to 'debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp')

diff --git a/debug_tools/WatchYourStep/ptxjitplus/Makefile b/debug_tools/WatchYourStep/ptxjitplus/Makefile
index d571567..b273ac2 100644
--- a/debug_tools/WatchYourStep/ptxjitplus/Makefile
+++ b/debug_tools/WatchYourStep/ptxjitplus/Makefile
@@ -315,14 +315,11 @@ endif
 ptxjitplus.o:ptxjitplus.cpp
 	$(EXEC) $(NVCC) $(INCLUDES)  --cudart shared $(ALL_CCFLAGS) $(GENCODE_FLAGS) -g -o $@ -c $<
 
-ptxjitplus: ptxjitplus.o ptxinst.o
+ptxjitplus: ptxjitplus.o 
 	$(EXEC) $(NVCC) $(ALL_LDFLAGS) --cudart shared $(GENCODE_FLAGS) -g -o $@ $+ $(LIBRARIES)
 	$(EXEC) mkdir -p ./bin/$(TARGET_ARCH)/$(TARGET_OS)/$(BUILD_TYPE)
 	$(EXEC) cp $@ ./bin/$(TARGET_ARCH)/$(TARGET_OS)/$(BUILD_TYPE)
 
-ptxinst.o:ptxinst.cpp
-	$(EXEC) $(NVCC) $(INCLUDES)  --cudart shared $(ALL_CCFLAGS) $(GENCODE_FLAGS) -g -o $@ -c $<
-
 run: build
 	$(EXEC) ./ptxjitplus
 
diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxinst.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxinst.cpp
deleted file mode 100644
index 6b39eb1..0000000
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxinst.cpp
+++ /dev/null
@@ -1,14 +0,0 @@
-/* ptxinst.cpp 
- * Jonathan Lew 
- * University of British Columbia
- */
-#include "ptxinst.h"
-
-void* instrument_ptx_from_function(std::string function, std::string path)
-{
-    return NULL;
-}
-void* instrument_ptx_from_string(std::string ptxcode)
-{
-    return NULL;
-}
diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxinst.h b/debug_tools/WatchYourStep/ptxjitplus/ptxinst.h
deleted file mode 100644
index de2595e..0000000
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxinst.h
+++ /dev/null
@@ -1,14 +0,0 @@
-/* ptxinst.h 
- * Jonathan Lew 
- * University of British Columbia
- */
- 
-#ifndef _PTXINST_H_
-#define _PTXINST_H_
-
-#include <string>
-
-void* instrument_ptx_from_function(std::string function, std::string path);
-void* instrument_ptx_from_string(std::string ptxcode);
-
-#endif
diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
index 68964d8..b7f9f2d 100644
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
+++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
@@ -1,5 +1,4 @@
-/**
- * Copyright 1993-2015 NVIDIA Corporation.  All rights reserved.
+/** Copyright 1993-2015 NVIDIA Corporation.  All rights reserved.
  *
  * Please refer to the NVIDIA end user license agreement (EULA) associated
  * with this source code for terms and conditions that govern your use of
@@ -7,9 +6,6 @@
  * this software and related documentation outside the terms of the EULA
  * is strictly prohibited.
  *
- */
-
-/*
  * This sample uses the Driver API to just-in-time compile (JIT) a Kernel from PTX code.
  * Additionally, this sample demonstrates the seamless interoperability capability of CUDA runtime
  * Runtime and CUDA Driver API calls.
@@ -17,6 +13,32 @@
  *
  */
 
+/**
+ * Modified by: Jonathan Lew
+ * PTX JIT PLUS
+ * 
+ * **********
+ * User Guide
+ * **********
+ *
+ * Welcome to WatchYourStep, a debugging tool that allows you launch individual
+ * kernels using parameters captured from cudaLaunch and outputs the values in 
+ * the arrays from the kernel. It allows you to watch each step you program takes, 
+ * kernel by kernel. 
+ *
+ * 1. Set environment variables to create params.config* and ptx.config* files. 
+ *      a)export PTX_SIM_DEBUG=4
+ *      b)export PTX_JIT_PATH=[path to this file]
+ *      c)export WYS_EXEC_PATH=[path to executable (program to debug)]
+ *      d)export WYS_EXEC_NAME=[name of executable (program to debug)]
+ *      e)Make sure all GPGPU-Sim path variables are set (see GPGPU-Sim documentation)
+ * 2. Run executable (program to debug) using GPGPU-Sim 
+ * 3. export PTX_SIM_DEBUG=[less than 4 to not dump config files again]
+ * 4-1. Run one kernel at a time: export WYS_LAUNCH_NUM=[kernel to launch] and compile ptxjitplus and run ptxjitplus
+ * 4-2. Run all kernels: compile and run ". launchkernels 0 [max number of kernels]" in terminal
+ * 5. Find output in ../data/wys.out* where * is the launch number
+ */
+
 // System includes
 #include <iostream>
 #include <math.h>
@@ -34,12 +56,12 @@
 
 // sample include
 #include "ptxjitplus.h"
-#include "ptxinst.h"
 
-const char *sSDKname = "PTX Just In Time (JIT) Compilation (no-qatest)";
+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;
 
@@ -116,6 +138,10 @@ void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUl
 
 void initializeData(std::vector<unsigned char*>& param_data, std::vector< std::pair<size_t, bool> >& param_info)
 {
+    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");
@@ -150,32 +176,19 @@ void initializeData(std::vector<unsigned char*>& param_data, std::vector< std::p
         err = fscanf(fin, "%lu : ", &len);
         info.first = len;
         assert( err==1 );
-        //printf("%lu : ", len);
         unsigned char* params = (unsigned char*) malloc(len*sizeof(unsigned char));
         for (size_t i=0; i<len; i++)
         {
             err = fscanf(fin, "%u ", &val);
             assert( err==1 );
-            //printf("%u ", val);
             params[i] = (unsigned char) val;
         }
         param_info.push_back(info);
         param_data.push_back(params);
         err = fscanf(fin, "\n");
         assert(err==0);
-        //printf("\n");
     }
     fclose(fin);
-    //filename = std::string("../data/wys.out") + wys_launch_num + "_param";
-    //fout = fopen(filename.c_str(), "w");
-    //assert(fout);
-    //fprintf(fout, "param %zu: size = %zu, data = ", 0,param_info[0].first);
-    //for (size_t j = 0; j<param_info[0].first; j++){
-    //    fprintf(fout, " %u", i->second[j]);
-    //}
-    //fprintf(fout, "\n");
-    //fflush(fout);
-    //fclose(fout);
 }
 
 int main(int argc, char **argv)
@@ -249,20 +262,28 @@ int main(int argc, char **argv)
     checkCudaErrors(cudaFree(d_tmp));
 
     //maps param number to pointer to device data
-    std::map< size_t, unsigned char* > m_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
     int paramOffset = 0;
     for( size_t i = 0; i<param_data.size(); i++){
         if(param_info[i].second){
             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)));
+            if (gpgpusim){
+                checkCudaErrors(cuParamSetv(hKernel, paramOffset, d_data, sizeof(*d_data)));
+            }
+            paramKernels[i] = (void*)d_data;
             m_device_data[i]=*d_data;
+            m_cleanup[i]=d_data;
             paramOffset += 8;
         }else{
-            checkCudaErrors(cuParamSetv(hKernel, paramOffset, param_data[i], param_info[i].first));
+            if (gpgpusim){
+                checkCudaErrors(cuParamSetv(hKernel, paramOffset, param_data[i], param_info[i].first));
+            }
+            paramKernels[i] = (void*)param_data[i];
             paramOffset += param_info[i].first;
         }
     }
@@ -271,12 +292,12 @@ int main(int argc, char **argv)
     // Launch the kernel (Driver API_)
     // TODO: automatically load these values in
     CUDAAPI cuLaunchKernel(hKernel, gridDim.x, gridDim.y, gridDim.z, blockDim.x, blockDim.y, blockDim.z, 
-        0, NULL, NULL, NULL);
+        0, NULL, paramKernels, NULL);
     std::cout << "CUDA kernel launched" << std::endl;
 
     //maps param number to pointer to output data
     std::map< size_t, unsigned char* > m_output_data;
-    for(std::map< size_t, unsigned char* >::iterator i = m_device_data.begin(); i!=m_device_data.end(); i++){
+    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)
         {
@@ -294,40 +315,32 @@ int main(int argc, char **argv)
     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, " %u", i->second[j]);
-            if (j&&(!(j%20))){
+            if (!(j%24)){ 
                 fprintf(fout, "\n");
             }
+            fprintf(fout, " %u", i->second[j]);
         }
         fprintf(fout, "\n");
     }
     fflush(fout);
     fclose(fout);
 
-//    int* h_data = (int*) m_output_data[0];
-//    // Check the result
-//    bool dataGood = true;
-//
-//    for (unsigned int i = 0 ; dataGood && i < nBlocks * nThreads ; i++)
-//    {
-//        if (h_data[i] != (int)i)
-//        {
-//            std::cerr << "Error at " << i << std::endl;
-//            dataGood = false;
-//        }
-//    }
-//    if(dataGood){
-//        std::cout<<"OK!"<<std::endl;
-//    }
 
     //Cleanup
-    for(std::map< size_t, unsigned char* >::iterator i = m_device_data.begin(); i!=m_device_data.end(); i++){
+    for(std::map< size_t, void* >::iterator i = m_device_data.begin(); i!=m_device_data.end(); i++){
         if (i->second){
             checkCudaErrors(cudaFree(i->second));
             i->second = 0;
         }
     }
 
+    for(std::map< size_t, void* >::iterator i = m_cleanup.begin(); i!=m_cleanup.end(); i++){
+        if (i->second){
+            free(i->second);
+            i->second = 0;
+        }
+    }
+
     for(std::map< size_t, unsigned char* >::iterator i = m_output_data.begin(); i!=m_output_data.begin(); i++){
         if (i->second)
         {
@@ -342,5 +355,5 @@ int main(int argc, char **argv)
         hModule = 0;
     }
 
-    //return dataGood ? EXIT_SUCCESS : EXIT_FAILURE;
+    return 0;
 }
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 68b4017..866fa3b 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -2347,11 +2347,17 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
 
 void __cudaUnregisterFatBinary(void **fatCubinHandle)
 {
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
 	;
 }
 
 cudaError_t cudaDeviceReset ( void ) {
 	// Should reset the simulated GPU
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
 	return g_last_cudaError = cudaSuccess;
 }
 cudaError_t CUDARTAPI cudaDeviceSynchronize(void){
@@ -2397,6 +2403,9 @@ extern void __cudaRegisterVar(
 		int constant,
 		int global )
 {
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
 	printf("GPGPU-Sim PTX: __cudaRegisterVar: hostVar = %p; deviceAddress = %s; deviceName = %s\n", hostVar, deviceAddress, deviceName);
 	printf("GPGPU-Sim PTX: __cudaRegisterVar: Registering const memory space of %d bytes\n", size);
 	if(GPGPUSim_Context()->get_device()->get_gpgpu()->get_config().use_cuobjdump())
@@ -2415,6 +2424,9 @@ void __cudaRegisterShared(
 		void **devicePtr
 )
 {
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
 	// we don't do anything here
 	printf("GPGPU-Sim PTX: __cudaRegisterShared\n" );
 }
@@ -2444,6 +2456,9 @@ void __cudaRegisterTexture(
 		int ext
 ) //passes in a newly created textureReference
 {
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
 	std::string devStr (deviceName);
 	#if (CUDART_VERSION > 4020)
 	if (devStr.size() > 2 && devStr.data()[0] == ':' && devStr.data()[1] == ':')
@@ -2465,6 +2480,9 @@ typedef unsigned long GLuint;
 
 cudaError_t cudaGLRegisterBufferObject(GLuint bufferObj)
 {
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
 	printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ );
 	return g_last_cudaError = cudaSuccess;
 }
@@ -2481,6 +2499,9 @@ glbmap_entry_t* g_glbmap = NULL;
 
 cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj) 
 {
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
 #ifdef OPENGL_SUPPORT
 	GLint buffer_size=0;
 	CUctx_st* ctx = GPGPUSim_Context();
@@ -2534,6 +2555,9 @@ cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj)
 
 cudaError_t cudaGLUnmapBufferObject(GLuint bufferObj)
 {
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
 #ifdef OPENGL_SUPPORT
 	glbmap_entry_t *p = g_glbmap;
 	while ( p && p->m_bufferObj != bufferObj )
@@ -2558,6 +2582,9 @@ cudaError_t cudaGLUnmapBufferObject(GLuint bufferObj)
 
 cudaError_t cudaGLUnregisterBufferObject(GLuint bufferObj) 
 {
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
 	printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ );
 	return g_last_cudaError = cudaSuccess;
 }
@@ -2783,6 +2810,9 @@ extern FILE *ptxinfo_in;
 
 static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr, gpgpu_t *gpu ) 
 {
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
 	printf( "GPGPU-Sim PTX: loading globals with explicit initializers... \n" );
 	fflush(stdout);
 	int ng_bytes=0;
@@ -2819,6 +2849,9 @@ static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsign
 
 static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu ) 
 {
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
 	printf( "GPGPU-Sim PTX: loading constants with explicit initializers... " );
 	fflush(stdout);
 	int nc_bytes = 0;
@@ -2866,6 +2899,9 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
 		struct dim3 blockDim,
 		CUctx_st* context )
 {
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
 	function_info *entry = context->get_kernel(hostFun);
 	kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry);
 	if( entry == NULL ) {
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index df8d806..d2f096f 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -1309,7 +1309,9 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt
     do{
         fprintf(fout, "%s", buff);
         fgets(buff, 1024, fin);
-        assert(!feof(fin));
+        if(feof(fin)){
+            break;
+        }
     } while(strstr(buff, "entry")==NULL);
 
     fclose(fin);
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index 034a7b9..9e18772 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -3190,7 +3190,21 @@ void rem_impl( const ptx_instruction *pI, ptx_thread_info *thread )
    src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1);
    src2_data = thread->get_operand_value(src2, dst, i_type, thread, 1);
 
-   data.u64 = src1_data.u64 % src2_data.u64;
+   switch ( i_type ) {
+   case S32_TYPE:
+      data.s32 = src1_data.s32 % src2_data.s32;
+      break;
+   case S64_TYPE:
+      data.s64 = src1_data.s64 % src2_data.s64;
+      break;
+   case U32_TYPE:
+      data.u32 = src1_data.u32 % src2_data.u32;
+      break;
+   case U64_TYPE:
+      data.u64 = src1_data.u64 % src2_data.u64;
+      break;
+   default: assert(0); break;
+   }
 
    thread->set_operand_value(dst,data, i_type, thread, pI);
 }
-- 
cgit v1.3


From 3a09960577b9c9cbcce8fedeef8874ccc533f378 Mon Sep 17 00:00:00 2001
From: Jonathan <jonathanslew@outlook.com>
Date: Wed, 18 Jul 2018 17:14:41 -0700
Subject: added c++filt that works better, param offset dumping, protection for
 failing system calls

---
 .../WatchYourStep/ptxjitplus/ptxjitplus.cpp        |  2 +-
 libcuda/cuda_runtime_api.cc                        |  5 +-
 src/cuda-sim/cuda-sim.cc                           | 78 +++++++++++++++++++---
 src/cuda-sim/ptx_ir.h                              |  2 +-
 4 files changed, 74 insertions(+), 13 deletions(-)

(limited to 'debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp')

diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
index b7f9f2d..26a6d29 100644
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
+++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
@@ -183,6 +183,7 @@ 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);
         err = fscanf(fin, "\n");
@@ -290,7 +291,6 @@ int main(int argc, char **argv)
     checkCudaErrors(cuParamSetSize(hKernel, paramOffset));
 
     // Launch the kernel (Driver API_)
-    // TODO: automatically load these values in
     CUDAAPI cuLaunchKernel(hKernel, gridDim.x, gridDim.y, gridDim.z, blockDim.x, blockDim.y, blockDim.z, 
         0, NULL, paramKernels, NULL);
     std::cout << "CUDA kernel launched" << std::endl;
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 866fa3b..db4f58e 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -1797,7 +1797,10 @@ void extract_code_using_cuobjdump(){
     //Running cuobjdump using dynamic link to current process
     snprintf(command,1000,"md5sum %s ", app_binary.c_str());
     printf("Running md5sum using \"%s\"\n", command);
-    system(command);
+    if(system(command)){
+        std::cout << "Failed to execute: " << command << std::endl;
+        exit(1);
+    }
     // Running cuobjdump using dynamic link to current process
     // Needs the option '-all' to extract PTX from CDP-enabled binary 
     extern bool g_cdp_enabled;
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index d2f096f..a499ce9 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -1226,10 +1226,11 @@ void function_info::list_param( FILE *fout ) const
    fflush(fout);
 }
 
-void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) const
+void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) 
 {
     static unsigned long counter = 0;
     std::vector< std::pair<size_t, unsigned char*> > param_data;
+    std::vector<unsigned> offsets;
     std::vector<bool> paramIsPointer;
 
     char * gpgpusim_path = getenv("GPGPUSIM_ROOT");
@@ -1237,18 +1238,67 @@ 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";
-    system(command.c_str());
+    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));
 
-    for( std::map<unsigned,param_info>::const_iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) {
-        const param_info &p = i->second;
+    //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();
         symbol *param = m_symtab->lookup(name.c_str());
         addr_t param_addr = param->get_address();
         param_t param_value = p.get_value();
+        offsets.push_back((unsigned)p.get_offset());
 
-        if(param_value.size==sizeof(void*) && mallocPtr_Size.find(*(unsigned long long*)param_value.pdata)!=mallocPtr_Size.end()){
+        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
+
             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);
@@ -1257,7 +1307,7 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt
             param_data.push_back(std::pair<size_t, unsigned char*>(array_size,array_val));
             paramIsPointer.push_back(true);
         }else{
-            unsigned char val[param_value.size];
+            unsigned char* val = (unsigned char*) malloc(param_value.size);
             param_mem->read(param_addr,param_value.size,(void*)val);
             param_data.push_back(std::pair<size_t, unsigned char*>(param_value.size,val));
             paramIsPointer.push_back(false);
@@ -1277,6 +1327,8 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt
         for (size_t j = 0; j<i->first; j++){
             fprintf(fout, " %u", i->second[j]);
         }
+        fprintf(fout, " : %u", offsets[index]);
+        free (i->second);
         fprintf(fout, "\n");
         index++;
     }
@@ -1284,10 +1336,13 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt
     fclose(fout);
     
     //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\" %s/*.ptx | cut -d \":\" -f 1-2 > %s", get_name().c_str(), wys_exec_path, ptx_config_fn.c_str());
-    system(buff);
+    if (system(buff)!=0){
+        printf("WARNING: Failed to execute grep to find ptx source \n");
+        printf("Problematic call: %s", buff);
+        abort();
+    }
     FILE *fin = fopen(ptx_config_fn.c_str(), "r");
     char ptx_source[256];
     unsigned line_number;
@@ -1295,7 +1350,11 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt
     assert(numscanned == 2);
     fclose(fin);
     snprintf(buff, 1024, "grep -rn \".version\" %s | cut -d \":\" -f 1 | xargs -I \"{}\" awk \"NR>={}&&NR<={}+2\" %s > %s", ptx_source, ptx_source, ptx_config_fn.c_str());
-    system(buff);
+    if (system(buff)!=0){
+        printf("WARNING: Failed to execute grep to find ptx header \n");
+        printf("Problematic call: %s", buff);
+        abort();
+    }
     fin = fopen(ptx_source, "r");
     assert(fin!=NULL);
     printf("Writing data to %s ...\n", ptx_config_fn.c_str());
@@ -1318,7 +1377,6 @@ 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> 
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index e0b5e96..ef4cf48 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 ptx_jit_config(std::map<unsigned long long, size_t> mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) const;
+   void ptx_jit_config(std::map<unsigned long long, size_t> mallocPtr_Size, memory_space *param_mem, gpgpu_t* gpu, dim3 gridDim, dim3 blockDim) ;
 
    const struct gpgpu_ptx_sim_info* get_kernel_info () const
    {
-- 
cgit v1.3


From e22b3179ccdee8bfbd1e74f86d93e0d7fc9ae6dc Mon Sep 17 00:00:00 2001
From: J <jonathanslew@outlook.com>
Date: Fri, 20 Jul 2018 11:40:44 -0700
Subject: load param offsets with refactored code and look in middle of
 malloc'd area for pointer

---
 .../WatchYourStep/ptxjitplus/ptxjitplus.cpp        | 69 +++++++++++-----------
 debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.h  |  7 +++
 src/cuda-sim/cuda-sim.cc                           | 23 +++++---
 3 files changed, 55 insertions(+), 44 deletions(-)

(limited to 'debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp')

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);
-- 
cgit v1.3


From 1a3de9dba3ddf4adab4256fadf5e815b410862ae Mon Sep 17 00:00:00 2001
From: Jonathan <jonathanslew@outlook.com>
Date: Fri, 10 Aug 2018 16:27:26 -0700
Subject: fixes ptxjitplus and cuLaunchKernel

---
 .../WatchYourStep/ptxjitplus/ptxjitplus.cpp        | 12 ------
 src/cuda-sim/cuda-sim.cc                           | 44 +---------------------
 src/cuda-sim/ptx_parser.cc                         |  4 +-
 3 files changed, 5 insertions(+), 55 deletions(-)

(limited to 'debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp')

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);
    }
 
 }
-- 
cgit v1.3