Merge branch 'dev' of https://github.com/gpgpu-sim/gpgpu-sim_distribution into dev

1 files changed, 116 insertions, 40 deletions

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 95a3c24..3a9d613 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -354,8 +354,8 @@ struct _cuda_device_id *GPGPUSim_Init()
 
 		cudaDeviceProp *prop = (cudaDeviceProp *) calloc(sizeof(cudaDeviceProp),1);
 		snprintf(prop->name,256,"GPGPU-Sim_v%s", g_gpgpusim_version_string );
-		prop->major = 5;
-		prop->minor = 2;
+		prop->major = the_gpu->compute_capability_major();
+		prop->minor = the_gpu->compute_capability_minor();
 		prop->totalGlobalMem = 0x80000000 /* 2 GB */;
 		prop->memPitch = 0;
 		if(prop->major >= 2) {
@@ -1120,11 +1120,11 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic
                 case 41:
                         *value= 0;
                         break;
-                case 75:
-                        *value= 9 ;
+                case 75://cudaDevAttrComputeCapabilityMajor
+                        *value= prop->major ;
                         break;
-                case 76:
-                        *value= 3 ;
+                case 76://cudaDevAttrComputeCapabilityMinor
+                        *value= prop->minor ;
                         break;
                 case 78:
                         *value= 0 ; //TODO: as of now, we dont support stream priorities.
@@ -1200,6 +1200,59 @@ __host__ cudaError_t CUDARTAPI cudaGetDevice(int *device)
 	return g_last_cudaError = cudaSuccess;
 }
 
+__host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit limit )
+{
+	if(g_debug_execution >= 3){
+            announce_call(__my_func__);
+   	 }
+        _cuda_device_id *dev = GPGPUSim_Init();
+	const struct cudaDeviceProp *prop = dev->get_prop();
+	const gpgpu_sim_config& config=dev->get_gpgpu()->get_config();
+        switch(limit) {
+        case 0:  // cudaLimitStackSize
+        	*pValue=config.stack_limit();
+                break;
+        case 2:  // cudaLimitMallocHeapSize
+                *pValue=config.heap_limit();
+                break;
+#if (CUDART_VERSION > 5050)
+        case 3: // cudaLimitDevRuntimeSyncDepth
+		if(prop->major > 2){
+			*pValue=config.sync_depth_limit();
+	                break;
+		}
+		else{
+			printf("ERROR:Limit %s is not supported on this architecture \n",limit);
+			abort();
+		}
+        case 4: // cudaLimitDevRuntimePendingLaunchCount
+		if(prop->major > 2){
+     	        	*pValue=config.pending_launch_count_limit();
+	                break;
+		}
+		else{
+			printf("ERROR:Limit %s is not supported on this architecture \n",limit);
+			abort();
+		}
+#endif
+        default:
+                        printf("ERROR:Limit %s unimplemented \n",limit);
+                        abort();
+        }
+	return g_last_cudaError = cudaSuccess;
+
+}
+
+__host__ cudaError_t CUDARTAPI cudaStreamGetPriority ( cudaStream_t hStream, int* priority )
+{
+        if(g_debug_execution >= 3){
+            announce_call(__my_func__);
+   	 }
+        cuda_not_implemented(__my_func__,__LINE__);
+        return g_last_cudaError = cudaSuccess;
+
+}
+
 __host__ cudaError_t CUDARTAPI cudaDeviceGetPCIBusId (
 		char *pciBusId,
 		int len,
@@ -1235,6 +1288,16 @@ __host__ cudaError_t cudaIpcOpenMemHandle(
 	return g_last_cudaError = cudaErrorUnknown;
 }
 
+__host__ cudaError_t CUDARTAPI cudaDestroyTextureObject(cudaTextureObject_t texObject)
+{
+        if(g_debug_execution >= 3){
+            announce_call(__my_func__);
+   	 }
+        cuda_not_implemented(__my_func__,__LINE__);
+        return g_last_cudaError = cudaErrorUnknown;
+}
+
+
 /*******************************************************************************
  *                                                                              *
  *                                                                              *
@@ -1469,42 +1532,26 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun )
 	return g_last_cudaError = cudaSuccess;
 }
 
-
 __host__ cudaError_t CUDARTAPI cudaLaunchKernel ( const char* hostFun, dim3 gridDim, dim3 blockDim, const void** args, size_t sharedMem, cudaStream_t stream )
 {
-	struct CUstream_st *s = (struct CUstream_st *)stream;
-	g_cuda_launch_stack.push_back( kernel_config(gridDim,blockDim,sharedMem,s) );
-
-	//printf("cudaLaunchKernel:sizeof(Arg[0])=%d)\n ",sizeof(args[0]));
-	kernel_config &config = g_cuda_launch_stack.back();
-	config.set_arg(args[0],432,0);//standard interface for cutlass library #TODO Implementing a generalized kernel
-
-	CUctx_st* context = GPGPUSim_Context();
-	char *mode = getenv("PTX_SIM_MODE_FUNC");
-	if( mode )
-		sscanf(mode,"%u", &g_ptx_sim_mode);
-	gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
-	kernel_config config1 = g_cuda_launch_stack.back();
-	struct CUstream_st *stream1 = config1.get_stream();
-	printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun,
-			g_ptx_sim_mode?"functional simulation":"performance simulation", stream1?stream1->get_uid():0 );
-	kernel_info_t *grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config1.get_args(),config1.grid_dim(),config1.block_dim(),context);
-	std::string kname = grid->name();
-	dim3 gridDim1 = config1.grid_dim();
-	dim3 blockDim1 = config1.block_dim();
-	printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n",
-			kname.c_str(), stream1?stream1->get_uid():0, gridDim1.x,gridDim1.y,gridDim1.z,blockDim1.x,blockDim1.y,blockDim1.z );
 
-	/*Kernel is hardcoded to enable the cutlass library*/
-	std::string cutlass("cutlass");
-	assert(kname.find(cutlass) != std::string::npos);
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    	}
+        CUctx_st *context = GPGPUSim_Context();
+        function_info *entry = context->get_kernel(hostFun);
+    
+	cudaConfigureCall(gridDim, blockDim, sharedMem, stream);
+    	for(unsigned i = 0; i < entry->num_args(); i++){
+        	std::pair<size_t, unsigned> p = entry->get_param_config(i);
+        	cudaSetupArgument(args[i], p.first, p.second);
+    	}  
 
-	stream_operation op(grid,g_ptx_sim_mode,stream1);
-	g_stream_manager->push(op);
-	g_cuda_launch_stack.pop_back();
+	cudaLaunch(hostFun);
 	return g_last_cudaError = cudaSuccess;
 }
 
+
 /*******************************************************************************
  *                                                                              *
  *                                                                              *
@@ -1555,6 +1602,9 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream)
 	    announce_call(__my_func__);
     }
 #if (CUDART_VERSION >= 3000)
+	//per-stream synchronization required for application using external libraries without explicit synchronization in the code to 
+	//avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. 
+	stream->synchronize();
 	g_stream_manager->destroy_stream(stream);
 #endif
 	return g_last_cudaError = cudaSuccess;
@@ -2552,8 +2602,17 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
         // Making this a runtime variable based on the app, enables GPGPU-Sim compiled
         // with a newer version of CUDA to run apps compiled with older versions of
         // CUDA. This is especially useful for PTXPLUS execution.
-        int app_cuda_version = get_app_cuda_version();
-        assert( app_cuda_version == CUDART_VERSION / 1000  && "The app must be compiled with same major version as the simulator." );
+        //Skip cuda version check for pytorch application
+	std::string app_binary_path =  get_app_binary();
+        int pos = app_binary_path.find("python");
+        if (pos==std::string::npos){
+        	// Not pytorch app : checking cuda version
+		int app_cuda_version = get_app_cuda_version();
+        	assert( app_cuda_version == CUDART_VERSION / 1000  && "The app must be compiled with same major version as the simulator." );
+        }
+
+	//int app_cuda_version = get_app_cuda_version();
+        //assert( app_cuda_version == CUDART_VERSION / 1000  && "The app must be compiled with same major version as the simulator." );
         const char* filename;
 #if CUDART_VERSION < 6000
             // FatBin handle from the .fatbin.c file (one of the intermediate files generated by NVCC)
@@ -2672,14 +2731,14 @@ cudaError_t cudaDeviceReset ( void ) {
 	return g_last_cudaError = cudaSuccess;
 }
 cudaError_t CUDARTAPI cudaDeviceSynchronize(void){
-	// I don't know what this should do
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
+	//Blocks until the device has completed all preceding requested tasks
+	synchronize();
 	return g_last_cudaError = cudaSuccess;
 }
 
-
 void CUDARTAPI __cudaRegisterFunction(
 		void   **fatCubinHandle,
 		const char    *hostFun,
@@ -3109,6 +3168,8 @@ __host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t v
     }
     return g_last_cudaError = cudaSuccess;
 }
+
+
 #endif
 
 #endif
@@ -3287,7 +3348,12 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
 	    announce_call(__my_func__);
     }
 	function_info *entry = context->get_kernel(hostFun);
-	kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry);
+	gpgpu_t* gpu= context->get_device()->get_gpgpu();
+	/*
+	Passing a snapshot of the GPU's current texture mapping to the kernel's info
+	as kernels should use texture bindings present at the time of their launch.
+	*/
+	kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry,gpu->getNameArrayMapping(),gpu->getNameInfoMapping());
 	if( entry == NULL ) {
 		printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found for %p\n", hostFun);
 		abort();
@@ -4447,6 +4513,16 @@ CUresult CUDAAPI cuPointerGetAttribute(void *data, CUpointer_attribute attribute
 #endif /* CUDART_VERSION >= 4000 */
 
 #if CUDART_VERSION >= 8000
+__host__ cudaError_t CUDARTAPI cudaCreateTextureObject ( cudaTextureObject_t* pTexObject, const cudaResourceDesc* pResDesc, const cudaTextureDesc* pTexDesc, const cudaResourceViewDesc* pResViewDesc )
+{
+         if(g_debug_execution >= 3){
+            announce_call(__my_func__);
+        }
+        cuda_not_implemented(__my_func__,__LINE__);
+        return g_last_cudaError = cudaSuccess;
+
+}
+
 CUresult CUDAAPI cuMemPrefetchAsync(CUdeviceptr devPtr, size_t count, CUdevice dstDevice, CUstream hStream)
 {
 	if(g_debug_execution >= 3){