Merge pull request #4 from gpgpu-sim/dev

Dev

2 files changed, 25 insertions, 36 deletions

diff --git a/configs/deprecated-cfgs/SM6_GTX1080/gpgpusim.config b/configs/deprecated-cfgs/SM6_GTX1080/gpgpusim.config
index fb044c6..3261d5a 100644
--- a/configs/deprecated-cfgs/SM6_GTX1080/gpgpusim.config
+++ b/configs/deprecated-cfgs/SM6_GTX1080/gpgpusim.config
@@ -31,12 +31,14 @@
 -gpgpu_simd_model 1 
 
 # Pipeline widths and number of FUs
-# ID_OC_SP,ID_OC_SFU,ID_OC_MEM,OC_EX_SP,OC_EX_SFU,OC_EX_MEM,EX_WB
+# ID_OC_SP, ID_OC_DP, ID_OC_INT, ID_OC_SFU, ID_OC_MEM, OC_EX_SP, OC_EX_DP, OC_EX_INT, OC_EX_SFU, OC_EX_MEM, EX_WB
 ## Pascal GP102 has 4 SP SIMD units and 1 SFU unit
 ## we need to scale the number of pipeline registers to be equal to the number of SP units
--gpgpu_pipeline_widths 4,1,1,4,1,1,6
+-gpgpu_pipeline_widths 4,0,0,1,1,4,0,0,1,1,6
 -gpgpu_num_sp_units 4
 -gpgpu_num_sfu_units 1
+-gpgpu_tensor_core_avail 0
+-gpgpu_num_tensor_core_units 0
 
 # Instruction latencies and initiation intervals
 # "ADD,MAX,MUL,MAD,DIV"
@@ -54,20 +56,20 @@
 # Pascal GP102 has 96KB Shared memory
 # Pascal GP102 has 64KB L1 cache
 # The default is to disable the L1 cache, unless cache modifieres is used
--gpgpu_cache:dl1  64:128:6,L:L:m:N:H,A:128:8,8
+-gpgpu_cache:dl1  N:64:128:6,L:L:m:N:H,A:128:8,8
 -gpgpu_shmem_size 98304
 -gmem_skip_L1D 1
 
 # 64 sets, each 128 bytes 16-way for each memory sub partition (128 KB per memory sub partition). This gives 3MB L2 cache
--gpgpu_cache:dl2 64:128:16,L:B:m:W:L,A:1024:1024,4:0,32 # used to be 128:4
+-gpgpu_cache:dl2 N:64:128:16,L:B:m:W:L,A:1024:1024,4:0,32 # used to be 128:4
 -gpgpu_cache:dl2_texture_only 0 
 
 # 4 KB Inst.
--gpgpu_cache:il1 8:128:4,L:R:f:N:L,A:2:48,4
+-gpgpu_cache:il1 N:8:128:4,L:R:f:N:L,A:2:48,4
 # 48 KB Tex 
--gpgpu_tex_cache:l1 16:128:24,L:R:m:N:L,F:128:4,128:2
+-gpgpu_tex_cache:l1 N:16:128:24,L:R:m:N:L,F:128:4,128:2
 # 12 KB Const
--gpgpu_const_cache:l1 128:64:2,L:R:f:N:L,A:2:64,4
+-gpgpu_const_cache:l1 N:128:64:2,L:R:f:N:L,A:2:64,4
 
 # enable operand collector 
 ## larger operand collectors and reg_banks are needed for the 4 warp schedulers and 4 SIMD units
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index e6b53c3..cf16d9b 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -1499,42 +1499,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;
 }
 
+
 /*******************************************************************************
  *                                                                              *
  *                                                                              *
@@ -1585,6 +1569,9 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream)
 	    announce_call(__my_func__);
     }
 #if (CUDART_VERSION >= 3000)
+	//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. 
+	synchronize();
 	g_stream_manager->destroy_stream(stream);
 #endif
 	return g_last_cudaError = cudaSuccess;