From 555cce395bb8457a209b26f6f7ff62cc09839fdf Mon Sep 17 00:00:00 2001 From: bftf Date: Mon, 11 Nov 2019 12:29:15 -0800 Subject: Enabled CUDA 10. Added implementations for __cudaPushCallConfiguration and __cudaPopCallConfiguration, neither of which are documented in the CUDA API but they preceed any cudaLaunchKernel() call and push the call configuration on a stack which is then popped when cudaLaunchKernel() is called. Currently the max stack size is 1, an assert will be triggered when this is violated. --- libcuda/cuda_api.h | 3 ++- libcuda/cuda_runtime_api.cc | 54 +++++++++++++++++++++++++++++++++++++++++++-- 2 files changed, 54 insertions(+), 3 deletions(-) (limited to 'libcuda') diff --git a/libcuda/cuda_api.h b/libcuda/cuda_api.h index 3808e8a..b304a59 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 CUDA_VERSION < 1000 typedef struct CUuuid_st { /**< CUDA definition of UUID */ char bytes[16]; } CUuuid; - +#endif #if __CUDA_API_VERSION >= 4010 diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 44f0f4e..fd092e5 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 CUDA_VERSION >= 1000 +/* +* 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 CUDA_VERSION >= 1000 + + __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 CUDA_VERSION >= 1000 + 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 CUDA_VERSION >= 1000 + for(unsigned i = 0; i < entry->num_args(); i++){ std::pair p = entry->get_param_config(i); cudaSetupArgument(args[i], p.first, p.second); -- cgit v1.3