loading params works now

6 files changed, 37 insertions, 92 deletions

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);