Tests to find conditions that a value is a pointer and new mallocPtr_Size

3 files changed, 29 insertions, 11 deletions

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 1e4f1df..28ff1c4 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -150,7 +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;
+std::map<unsigned long long, size_t> g_mallocPtr_Size;
 int no_of_ptx=0;
 std::map<int, std::set<std::string> > version_filename;
 
@@ -322,8 +322,8 @@ public:
 	}
 	kernel_config()
 	{
-		m_GridDim=NULL;
-		m_BlockDim=NULL;
+		m_GridDim=dim3(-1,-1,-1);
+		m_BlockDim=dim3(-1,-1,-1);
 		m_sharedMem=0;
 		m_stream =NULL;
 	}
@@ -490,7 +490,7 @@ __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size)
 	*devPtr = context->get_device()->get_gpgpu()->gpu_malloc(size);
 	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;
+        g_mallocPtr_Size[(unsigned long long)*devPtr] = size;
     }
 	if ( *devPtr  ) {
 		return g_last_cudaError = cudaSuccess;
@@ -1075,6 +1075,9 @@ __host__ cudaError_t CUDARTAPI cudaSetupArgument(const void *arg, size_t size, s
 	gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
 	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());
+    //assert((*(long long*)arg)==0||g_mallocPtr_Size.find(*(long long*)arg)!=g_mallocPtr_Size.end());
 
 	return g_last_cudaError = cudaSuccess;
 }
@@ -2379,7 +2382,7 @@ cudaError_t CUDARTAPI cudaHostGetDevicePointer(void **pDevice, void *pHost, unsi
 	*pDevice = gpu->gpu_malloc(size);
 	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;
+        g_mallocPtr_Size[(unsigned long long)*pDevice] = size;
     }
 	if ( *pDevice  ) {
 		pinned_memory[pHost]=pDevice;
@@ -2624,7 +2627,7 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
 
 	
 	if(g_debug_execution >= 3){
-        entry->debug_param(g_devPtr_Size, result->get_param_memory());
+        entry->debug_param(g_mallocPtr_Size, result->get_param_memory());
     }
 
 	return result;
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 8795e1c..ab607e0 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<void *, size_t> devPtr_Size, memory_space *param_mem) const
+void function_info::debug_param(std::map<unsigned long long, size_t> mallocPtr_Size, memory_space *param_mem) const
 {
     static unsigned long counter = 0;
     std::vector< std::pair<size_t, unsigned char*> > param_data;
@@ -1248,8 +1248,8 @@ void function_info::debug_param(std::map<void *, size_t> devPtr_Size, memory_spa
     fclose(fp);
     std::string fn(buff);
     size_t pos1, pos2;
-    pos1 = fn.find("(");
-    pos2 = fn.find(")");
+    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;
@@ -1270,8 +1270,23 @@ void function_info::debug_param(std::map<void *, size_t> devPtr_Size, memory_spa
         symbol *param = m_symtab->lookup(name.c_str());
         addr_t param_addr = param->get_address();
         param_t param_value = p.get_value();
+
+        if (paramIsPointer[i->first]){
+            assert(param_value.size==8&&mallocPtr_Size.find(*(long long*)param_value.pdata)!=mallocPtr_Size.end());
+        }else{
+            assert(!(param_value.size==8&&mallocPtr_Size.find(*(long long*)param_value.pdata)!=mallocPtr_Size.end()));
+        }
 //        if (paramIsPointer[i->first]){
 //            assert(param_value.size==8);
+//            void *array_pointer = (void*)*val;
+//            assert(devPtr_Size.find(array_pointer)!=devPtr_Size.end());
+//            size_t array_size = devPtr_Size[array_pointer];
+//            unsigned char array_val[array_size];
+//            memcpy((void*) array_val, array_pointer, array_size);
+//            param_data.push_back(std::pair<size_t, unsigned char*>(array_size,array_val));
+//
+//            assert(size!=8||g_mallocPtr_Size.find(*(long long*)arg)!=g_mallocPtr_Size.end());
+//
 //        }else{
             unsigned char val[param_value.size];
             param_mem->read(param_addr,param_value.size,(void*)val);
@@ -1295,7 +1310,7 @@ void function_info::debug_param(std::map<void *, size_t> devPtr_Size, memory_spa
     }
     fflush(fout);
     fclose(fout);
-    exit(0);
+    //exit(0);
 }
 
 template<int activate_level> 
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index 43907d3..1eb9d65 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<void *, size_t> devPtr_Size, memory_space *param_mem) const;
+   void debug_param(std::map<unsigned long long, size_t> mallocPtr_Size, memory_space *param_mem) const;
 
    const struct gpgpu_ptx_sim_info* get_kernel_info () const
    {