diff options
Diffstat (limited to 'libcuda/cuda_runtime_api.cc')
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 182 |
1 files changed, 160 insertions, 22 deletions
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 1c80941..4b50e34 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -346,7 +346,7 @@ private: gpgpu_ptx_sim_arg_list_t m_args; }; -class _cuda_device_id *GPGPUSim_Init() +struct _cuda_device_id *GPGPUSim_Init() { static _cuda_device_id *the_device = NULL; if( !the_device ) { @@ -358,10 +358,19 @@ class _cuda_device_id *GPGPUSim_Init() prop->minor = 2; prop->totalGlobalMem = 0x80000000 /* 2 GB */; prop->memPitch = 0; - prop->maxThreadsPerBlock = 512; - prop->maxThreadsDim[0] = 512; - prop->maxThreadsDim[1] = 512; - prop->maxThreadsDim[2] = 512; + if(prop->major >= 2) { + prop->maxThreadsPerBlock = 1024; + prop->maxThreadsDim[0] = 1024; + prop->maxThreadsDim[1] = 1024; + } + else + { + prop->maxThreadsPerBlock = 512; + prop->maxThreadsDim[0] = 512; + prop->maxThreadsDim[1] = 512; + } + + prop->maxThreadsDim[2] = 64; prop->maxGridSize[0] = 0x40000000; prop->maxGridSize[1] = 0x40000000; prop->maxGridSize[2] = 0x40000000; @@ -380,6 +389,9 @@ class _cuda_device_id *GPGPUSim_Init() #if (CUDART_VERSION >= 2010) prop->multiProcessorCount = the_gpu->get_config().num_shader(); #endif +#if (CUDART_VERSION >= 4000) + prop->maxThreadsPerMultiProcessor = the_gpu->threads_per_core(); +#endif the_gpu->set_prop(prop); the_device = new _cuda_device_id(the_gpu); } @@ -1412,6 +1424,41 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) } dim3 gridDim = config.grid_dim(); dim3 blockDim = config.block_dim(); + + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + checkpoint *g_checkpoint; + g_checkpoint = new checkpoint(); + class memory_space *global_mem; + global_mem = gpu->get_global_memory(); + + if(gpu->resume_option ==1 && (grid->get_uid()==gpu->resume_kernel)) + { + + char f1name[2048]; + snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", grid->get_uid()); + + g_checkpoint->load_global_mem(global_mem, f1name); + for (int i=0;i<gpu->resume_CTA;i++) + grid->increment_cta_id(); + } + if(gpu->resume_option==1 && (grid->get_uid()<gpu->resume_kernel)) + { + char f1name[2048]; + snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", grid->get_uid()); + + g_checkpoint->load_global_mem(global_mem, f1name); + printf("Skipping kernel %d as resuming from kernel %d\n",grid->get_uid(),gpu->resume_kernel ); + g_cuda_launch_stack.pop_back(); + return g_last_cudaError = cudaSuccess; + + } + if(gpu->checkpoint_option==1 && (grid->get_uid()>gpu->checkpoint_kernel)) + { + printf("Skipping kernel %d as checkpoint from kernel %d\n",grid->get_uid(),gpu->checkpoint_kernel ); + g_cuda_launch_stack.pop_back(); + return g_last_cudaError = cudaSuccess; + + } printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n", kname.c_str(), stream?stream->get_uid():0, gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z ); stream_operation op(grid,g_ptx_sim_mode,stream); @@ -1420,6 +1467,42 @@ __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); + + stream_operation op(grid,g_ptx_sim_mode,stream1); + g_stream_manager->push(op); + g_cuda_launch_stack.pop_back(); + return g_last_cudaError = cudaSuccess; +} + /******************************************************************************* * * * * @@ -1968,6 +2051,29 @@ void extract_ptx_files_using_cuobjdump(){ } +static int get_app_cuda_version() { + int app_cuda_version = 0; + char fname[1024]; + snprintf(fname,1024,"_app_cuda_version_XXXXXX"); + int fd=mkstemp(fname); + close(fd); + std::string app_cuda_version_command = "ldd " + get_app_binary() + " | grep libcudart.so | sed 's/.*libcudart.so.\\(.*\\) =>.*/\\1/' > " + fname; + system(app_cuda_version_command.c_str()); + FILE * cmd = fopen(fname, "r"); + char buf[256]; + while (fgets(buf, sizeof(buf), cmd) != 0) { + std::cout << buf; + app_cuda_version = atoi(buf); + } + fclose(cmd); + if ( app_cuda_version == 0 ) { + printf( "Error - Cannot detect the app's CUDA version.\n" ); + exit(1); + } + return app_cuda_version; +} + + //! Call cuobjdump to extract everything (-elf -sass -ptx) /*! * This Function extract the whole PTX (for all the files) using cuobjdump @@ -2447,21 +2553,30 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) if (sizeof(void*) == 4) printf("GPGPU-Sim PTX: FatBin file name extraction has not been tested on 32-bit system.\n"); - #if (CUDART_VERSION <= 6000) - // FatBin handle from the .fatbin.c file (one of the intermediate files generated by NVCC) - typedef struct {int m; int v; const unsigned long long* d; char* f;} __fatDeviceText __attribute__ ((aligned (8))); - __fatDeviceText * fatDeviceText = (__fatDeviceText *) fatCubin; + // This code will get the CUDA version the app was compiled with. + // We need this to determine how to handle the parsing of the binary. + // 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." ); + const char* filename; +#if CUDART_VERSION < 6000 + // FatBin handle from the .fatbin.c file (one of the intermediate files generated by NVCC) + typedef struct {int m; int v; const unsigned long long* d; char* f;} __fatDeviceText __attribute__ ((aligned (8))); + __fatDeviceText * fatDeviceText = (__fatDeviceText *) fatCubin; + + // Extract the source code file name that generate the given FatBin. + // - Obtains the pointer to the actual fatbin structure from the FatBin handle (fatCubin). + // - An integer inside the fatbin structure contains the relative offset to the source code file name. + // - This offset differs among different CUDA and GCC versions. + char * pfatbin = (char*) fatDeviceText->d; + int offset = *((int*)(pfatbin+48)); + filename = (pfatbin+16+offset); +#else + filename = "default"; +#endif - // Extract the source code file name that generate the given FatBin. - // - Obtains the pointer to the actual fatbin structure from the FatBin handle (fatCubin). - // - An integer inside the fatbin structure contains the relative offset to the source code file name. - // - This offset differs among different CUDA and GCC versions. - char * pfatbin = (char*) fatDeviceText->d; - int offset = *((int*)(pfatbin+48)); - char * filename = (pfatbin+16+offset); - #else - const char * filename = "default"; - #endif // The extracted file name is associated with a fat_cubin_handle passed // into cudaLaunch(). Inside cudaLaunch(), the associated file name is // used to find the PTX/SASS section from cuobjdump, which contains the @@ -2900,8 +3015,28 @@ cudaError_t CUDARTAPI cudaSetDeviceFlags( int flags ) if(g_debug_execution >= 3){ announce_call(__my_func__); } - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + // This flag is implicitly always on (unless you are using the driver API). It is safe for GPGPU-Sim to + // just ignore it. + if ( cudaDeviceMapHost == flags ) { + return g_last_cudaError = cudaSuccess; + } else { + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; + } +} + +size_t getMaxThreadsPerBlock(struct cudaFuncAttributes *attr) { + _cuda_device_id *dev = GPGPUSim_Init(); + struct cudaDeviceProp prop; + + prop = *dev->get_prop(); + + size_t max = prop.maxThreadsPerBlock; + + if ((prop.regsPerBlock / attr->numRegs) < max) + max = prop.regsPerBlock / attr->numRegs; + + return max; } cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *attr, const char *hostFun ) @@ -2917,7 +3052,10 @@ cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *attr, con attr->constSizeBytes = kinfo->cmem; attr->localSizeBytes = kinfo->lmem; attr->numRegs = kinfo->regs; - attr->maxThreadsPerBlock = 0; // from pragmas? + if(kinfo->maxthreads > 0) + attr->maxThreadsPerBlock = kinfo->maxthreads; + else + attr->maxThreadsPerBlock = getMaxThreadsPerBlock(attr); #if CUDART_VERSION >= 3000 attr->ptxVersion = kinfo->ptx_version; attr->binaryVersion = kinfo->sm_target; |