Merge pull request #156 from bftf/dev

Enabled CUDA 10. Added implementations for __cudaPushCallConfiguratio…

6 files changed, 61 insertions, 5 deletions

diff --git a/CHANGES b/CHANGES
index 2a4222b..31ae672 100644
--- a/CHANGES
+++ b/CHANGES
@@ -37,6 +37,7 @@ Version 4.0.0 (development branch) versus 3.2.3
 3- Addig new system stats: gpu occupancy, L2BW, etc
 -Library:
 1  Enabled CUTLASS Library on GPGPU-Sim
+2  Enabled CUDA 10
 -Regression:
 1- Added TensorCore Regression Kernel
 -Configs:
diff --git a/Makefile b/Makefile
index 3db8ce8..8df3b57 100644
--- a/Makefile
+++ b/Makefile
@@ -164,6 +164,7 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.0; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.1; fi
 	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.9.2 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.9.2; fi
+	if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.0; fi
 
 $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib
 	g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\
diff --git a/libcuda/cuda_api.h b/libcuda/cuda_api.h
index 3808e8a..0ded242 100644
--- a/libcuda/cuda_api.h
+++ b/libcuda/cuda_api.h
@@ -234,10 +234,11 @@ typedef struct CUgraphicsResource_st *CUgraphicsResource; /**< CUDA graphics int
 typedef unsigned long long CUtexObject;                   /**< An opaque value that represents a CUDA texture object */
 typedef unsigned long long CUsurfObject;                  /**< An opaque value that represents a CUDA surface object */
 
+#if CUDART_VERSION < 10000
 typedef struct CUuuid_st {                                /**< CUDA definition of UUID */
     char bytes[16];
 } CUuuid;
-
+#endif // #if CUDART_VERSION < 10000
 
 #if __CUDA_API_VERSION >= 4010
 
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 44f0f4e..71b8722 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -1448,6 +1448,50 @@ __host__ cudaError_t CUDARTAPI cudaConfigureCall(dim3 gridDim, dim3 blockDim, si
 	return g_last_cudaError = cudaSuccess;
 }
 
+
+#if CUDART_VERSION >= 10000
+/*
+* CUDA 10 requires a new CUDA kernel launch sequence
+* A call to __cudaPushCallConfiguration() preceeds any call to cudaLaunchKernel()
+* __cudaPushCallConfiguration is undocumented in the API but it simply sets up a buffer with the arguments which is accessed in cudaLaunchKernel()
+* __cudaPopCallConfiguration is undocumented in the API but it simply pops the configuration set in cudaLaunchKernel()
+*
+* pushing more than 1 configuration without popping is currently not implemented in GPGPU-Sim and will result in an assert error
+*/
+namespace g_cudaPushArgsBuffer
+{
+  bool g_is_initialized = false;
+  dim3 g_gridDim;
+  dim3 g_blockDim;
+  size_t g_sharedMem;
+  cudaStream_t g_stream;
+}
+
+__host__ cudaError_t CUDARTAPI __cudaPushCallConfiguration(dim3 gridDim, dim3 blockDim, size_t sharedMem, cudaStream_t stream)
+{
+  assert(g_cudaPushArgsBuffer::g_is_initialized == false);
+  printf("Pushing cuda call configuration \n");
+  g_cudaPushArgsBuffer::g_is_initialized = true;
+  g_cudaPushArgsBuffer::g_gridDim = gridDim;
+  g_cudaPushArgsBuffer::g_blockDim = blockDim;
+  g_cudaPushArgsBuffer::g_sharedMem = sharedMem;
+  g_cudaPushArgsBuffer::g_stream = stream;
+
+  return cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI __cudaPopCallConfiguration()
+{
+  printf("Inside __cudaPopCallConfiguration\n");
+  assert(g_cudaPushArgsBuffer::g_is_initialized == true);
+  printf("Poping cuda call configuration \n");
+  g_cudaPushArgsBuffer::g_is_initialized = false;
+  return cudaSuccess;
+}
+
+#endif // #if CUDART_VERSION >= 10000
+
+
 __host__ cudaError_t CUDARTAPI cudaSetupArgument(const void *arg, size_t size, size_t offset)
 {
 	if(g_debug_execution >= 3){
@@ -1551,8 +1595,14 @@ __host__ cudaError_t CUDARTAPI cudaLaunchKernel ( const char* hostFun, dim3 grid
     	}
         CUctx_st *context = GPGPUSim_Context();
         function_info *entry = context->get_kernel(hostFun);
-    
-	cudaConfigureCall(gridDim, blockDim, sharedMem, stream);
+
+#if CUDART_VERSION >= 10000
+  assert(g_cudaPushArgsBuffer::g_is_initialized == false);
+  cudaConfigureCall(g_cudaPushArgsBuffer::g_gridDim, g_cudaPushArgsBuffer::g_blockDim, g_cudaPushArgsBuffer::g_sharedMem, g_cudaPushArgsBuffer::g_stream);
+#else
+  cudaConfigureCall(gridDim, blockDim, sharedMem, stream);
+#endif // #if CUDART_VERSION >= 10000
+  
     	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);
diff --git a/linux-so-version.txt b/linux-so-version.txt
index a7c2d3c..357c6cd 100644
--- a/linux-so-version.txt
+++ b/linux-so-version.txt
@@ -6,3 +6,5 @@ libcudart.so.9.2{
 };
 libcuda.so.1{
 };
+libcudart.so.10.0{
+};
\ No newline at end of file
diff --git a/setup_environment b/setup_environment
index b420584..154dea0 100644
--- a/setup_environment
+++ b/setup_environment
@@ -51,9 +51,10 @@ CC_VERSION=`gcc --version | head -1 | awk '{for(i=1;i<=NF;i++){ if(match($i,/^[0
 
 CUDA_VERSION_STRING=`$CUDA_INSTALL_PATH/bin/nvcc --version | awk '/release/ {print $5;}' | sed 's/,//'`;
 export CUDA_VERSION_NUMBER=`echo $CUDA_VERSION_STRING | sed 's/\./ /' | awk '{printf("%02u%02u", 10*int($1), 10*$2);}'`
-if [ $CUDA_VERSION_NUMBER -gt 9100 -o $CUDA_VERSION_NUMBER -lt 2030  ]; then
+if [ $CUDA_VERSION_NUMBER -gt 10000 -o $CUDA_VERSION_NUMBER -lt 2030  ]; then
 	echo "ERROR ** GPGPU-Sim version $GPGPUSIM_VERSION_STRING not tested with CUDA version $CUDA_VERSION_STRING (please see README)";
-	return
+	echo $CUDA_VERSION_NUMBER
+  return
 fi
 
 if [ $CUDA_VERSION_NUMBER -ge 6000 ]; then