dump and load block and grid size and launch

4 files changed, 32 insertions, 11 deletions

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
    {