diff options
Diffstat (limited to 'libcuda')
| -rw-r--r-- | libcuda/cuda_api_object.h | 3 | ||||
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 792 |
2 files changed, 416 insertions, 379 deletions
diff --git a/libcuda/cuda_api_object.h b/libcuda/cuda_api_object.h index 86ffa98..2001f91 100644 --- a/libcuda/cuda_api_object.h +++ b/libcuda/cuda_api_object.h @@ -1,11 +1,14 @@ #ifndef __cuda_api_object_h__ #define __cuda_api_object_h__ class cuobjdumpSection; +class kernel_config; + class cuda_runtime_api { public: // global list std::list<cuobjdumpSection*> libSectionList; + std::list<kernel_config> g_cuda_launch_stack; // member function list }; #endif /* __cuda_api_object_h__ */ diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 4221634..ea96570 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -237,8 +237,6 @@ private: struct _cuda_device_id *m_next; }; -class kernel_config; - #ifndef OPENGL_SUPPORT typedef unsigned long GLuint; #endif @@ -315,7 +313,6 @@ struct CUctx_st { return i->second; } - std::list<kernel_config> g_cuda_launch_stack; std::map<int, bool>fatbin_registered; std::map<int, std::string> fatbinmap; std::map<unsigned long long, size_t> g_mallocPtr_Size; @@ -576,6 +573,403 @@ void setCuobjdumpsassfilename(const char* filename, std::list<cuobjdumpSection*> (dynamic_cast<cuobjdumpELFSection*>(cuobjdumpSectionList.front()))->setSASSfilename(filename); } +//! Return the executable file of the process containing the PTX/SASS code +//! +//! This Function returns the executable file ran by the process. This +//! executable is supposed to contain the PTX/SASS code. It provides workaround +//! for processes running on valgrind by dereferencing /proc/<pid>/exe within the +//! GPGPU-Sim process before calling cuobjdump to extract PTX/SASS. This is +//! needed because valgrind uses x86 emulation to detect memory leak. Other +//! processes (e.g. cuobjdump) reading /proc/<pid>/exe will see the emulator +//! executable instead of the application binary. +//! +std::string get_app_binary(){ + char self_exe_path[1025]; +#ifdef __APPLE__ + uint32_t size = sizeof(self_exe_path); + if( _NSGetExecutablePath(self_exe_path,&size) != 0 ) { + printf("GPGPU-Sim ** ERROR: _NSGetExecutablePath input buffer too small\n"); + exit(1); + } +#else + std::stringstream exec_link; + exec_link << "/proc/self/exe"; + + ssize_t path_length = readlink(exec_link.str().c_str(), self_exe_path, 1024); + assert(path_length != -1); + self_exe_path[path_length] = '\0'; +#endif + + printf("self exe links to: %s\n", self_exe_path); + return self_exe_path; +} + +//above func gives abs path whereas this give just the name of application. +char* get_app_binary_name(std::string abs_path){ + char *self_exe_path; +#ifdef __APPLE__ + //TODO: get apple device and check the result. + printf("WARNING: not tested for Apple-mac devices \n"); + abort(); +#else + char* buf = strdup(abs_path.c_str()); + char *token = strtok(buf, "/"); + while(token !=NULL){ + self_exe_path = token; + token = strtok(NULL,"/"); + } +#endif + self_exe_path = strtok(self_exe_path, "."); + printf("self exe links to: %s\n", self_exe_path); + return self_exe_path; +} + +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; +} + +//! Keep track of the association between filename and cubin handle +void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context){ + context->fatbinmap[handle] = filename; +} + +/******************************************************************************* + * Add internal cuda runtime API call to accept gpgpu_context * + *******************************************************************************/ + +void** cudaRegisterFatBinaryInternal( void *fatCubin, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } +#if (CUDART_VERSION < 2010) + printf("GPGPU-Sim PTX: ERROR ** this version of GPGPU-Sim requires CUDA 2.1 or higher\n"); + exit(1); +#endif + CUctx_st *context = GPGPUSim_Context(); + static unsigned next_fat_bin_handle = 1; + if(context->get_device()->get_gpgpu()->get_config().use_cuobjdump()) { + // The following workaround has only been verified on 64-bit systems. + if (sizeof(void*) == 4) + printf("GPGPU-Sim PTX: FatBin file name extraction has not been tested on 32-bit system.\n"); + + // 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. + //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) + 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 + + // 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 + // PTX/SASS code for the launched kernel function. + // This allows us to work around the fact that cuobjdump only outputs the + // file name associated with each section. + unsigned long long fat_cubin_handle = next_fat_bin_handle; + next_fat_bin_handle++; + printf("GPGPU-Sim PTX: __cudaRegisterFatBinary, fat_cubin_handle = %llu, filename=%s\n", fat_cubin_handle, filename); + /*! + * This function extracts all data from all files in first call + * then for next calls, only returns the appropriate number + */ + assert(fat_cubin_handle >= 1); + if (fat_cubin_handle==1) ctx->cuobjdumpInit(); + cuobjdumpRegisterFatBinary(fat_cubin_handle, filename, context); + + return (void**)fat_cubin_handle; + } +#if (CUDART_VERSION < 8000) + else { + static unsigned source_num=1; + unsigned long long fat_cubin_handle = next_fat_bin_handle++; + __cudaFatCudaBinary *info = (__cudaFatCudaBinary *)fatCubin; + assert( info->version >= 3 ); + unsigned num_ptx_versions=0; + unsigned max_capability=0; + unsigned selected_capability=0; + bool found=false; + unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability(); + if (!info->ptx){ + printf("ERROR: Cannot find ptx code in cubin file\n" + "\tIf you are using CUDA 4.0 or higher, please enable -gpgpu_ptx_use_cuobjdump or downgrade to CUDA 3.1\n"); + exit(1); + } + while( info->ptx[num_ptx_versions].gpuProfileName != NULL ) { + unsigned capability=0; + sscanf(info->ptx[num_ptx_versions].gpuProfileName,"compute_%u",&capability); + printf("GPGPU-Sim PTX: __cudaRegisterFatBinary found PTX versions for '%s', ", info->ident); + printf("capability = %s\n", info->ptx[num_ptx_versions].gpuProfileName ); + if( forced_max_capability ) { + if( capability > max_capability && capability <= forced_max_capability ) { + found = true; + max_capability=capability; + selected_capability = num_ptx_versions; + } + } else { + if( capability > max_capability ) { + found = true; + max_capability=capability; + selected_capability = num_ptx_versions; + } + } + num_ptx_versions++; + } + if( found ) { + printf("GPGPU-Sim PTX: Loading PTX for %s, capability = %s\n", + info->ident, info->ptx[selected_capability].gpuProfileName ); + symbol_table *symtab; + const char *ptx = info->ptx[selected_capability].ptx; + if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) { + printf("GPGPU-Sim PTX: ERROR ** PTXPlus is only supported through cuobjdump\n" + "\tEither enable cuobjdump or disable PTXPlus in your configuration file\n"); + exit(1); + } else { + symtab=gpgpu_ptx_sim_load_ptx_from_string(ptx,source_num); + context->add_binary(symtab,fat_cubin_handle); + gpgpu_ptxinfo_load_from_string( ptx, source_num, max_capability, context->no_of_ptx ); + } + source_num++; + load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); + load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); + } else { + printf("GPGPU-Sim PTX: warning -- did not find an appropriate PTX in cubin\n"); + } + return (void**)fat_cubin_handle; + } +#else + else { + printf("ERROR ** __cudaRegisterFatBinary() needs to be updated\n"); + abort(); + } +#endif +} + +void cudaRegisterFunctionInternal( + void **fatCubinHandle, + const char *hostFun, + char *deviceFun, + const char *deviceName, + int thread_limit, + uint3 *tid, + uint3 *bid, + dim3 *bDim, + dim3 *gDim, + gpgpu_context *gpgpu_ctx = NULL +) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + CUctx_st *context = GPGPUSim_Context(); + unsigned fat_cubin_handle = (unsigned)(unsigned long long)fatCubinHandle; + printf("GPGPU-Sim PTX: __cudaRegisterFunction %s : hostFun 0x%p, fat_cubin_handle = %u\n", + deviceFun, hostFun, fat_cubin_handle); + if(context->get_device()->get_gpgpu()->get_config().use_cuobjdump()) + ctx->cuobjdumpParseBinary(fat_cubin_handle); + context->register_function( fat_cubin_handle, hostFun, deviceFun ); +} + +void cudaRegisterVarInternal( + void **fatCubinHandle, + char *hostVar, //pointer to...something + char *deviceAddress, //name of variable + const char *deviceName, //name of variable (same as above) + int ext, + int size, + int constant, + int global, + gpgpu_context *gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + printf("GPGPU-Sim PTX: __cudaRegisterVar: hostVar = %p; deviceAddress = %s; deviceName = %s\n", hostVar, deviceAddress, deviceName); + printf("GPGPU-Sim PTX: __cudaRegisterVar: Registering const memory space of %d bytes\n", size); + if(GPGPUSim_Context()->get_device()->get_gpgpu()->get_config().use_cuobjdump()) + ctx->cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle); + fflush(stdout); + if ( constant && !global && !ext ) { + gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size); + } else if ( !constant && !global && !ext ) { + gpgpu_ptx_sim_register_global_variable(hostVar,deviceName,size); + } else cuda_not_implemented(__my_func__,__LINE__); +} + +cudaError_t cudaConfigureCallInternal(dim3 gridDim, dim3 blockDim, size_t sharedMem, cudaStream_t stream, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + struct CUstream_st *s = (struct CUstream_st *)stream; + ctx->api->g_cuda_launch_stack.push_back( kernel_config(gridDim,blockDim,sharedMem,s) ); + return g_last_cudaError = cudaSuccess; +} + +cudaError_t cudaSetupArgumentInternal(const void *arg, size_t size, size_t offset, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + gpgpusim_ptx_assert( !ctx->api->g_cuda_launch_stack.empty(), "empty launch stack" ); + kernel_config &config = ctx->api->g_cuda_launch_stack.back(); + config.set_arg(arg,size,offset); + printf("GPGPU-Sim PTX: Setting up arguments for %zu bytes starting at 0x%llx..\n",size, (unsigned long long) arg); + + return g_last_cudaError = cudaSuccess; +} + +cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx = NULL ) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + CUctx_st* context = GPGPUSim_Context(); + char *mode = getenv("PTX_SIM_MODE_FUNC"); + if( mode ) + sscanf(mode,"%u", &g_ptx_sim_mode); + gpgpusim_ptx_assert( !ctx->api->g_cuda_launch_stack.empty(), "empty launch stack" ); + kernel_config config = ctx->api->g_cuda_launch_stack.back(); + struct CUstream_st *stream = config.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", stream?stream->get_uid():0 ); + kernel_info_t *grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context); + //do dynamic PDOM analysis for performance simulation scenario + std::string kname = grid->name(); + function_info *kernel_func_info = grid->entry(); + if (kernel_func_info->is_pdom_set()) { + printf("GPGPU-Sim PTX: PDOM analysis already done for %s \n", kname.c_str() ); + } else { + printf("GPGPU-Sim PTX: finding reconvergence points for \'%s\'...\n", kname.c_str() ); + kernel_func_info->do_pdom(); + kernel_func_info->set_pdom(); + } + 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 ); + ctx->api->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 ); + ctx->api->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); + g_stream_manager()->push(op); + ctx->api->g_cuda_launch_stack.pop_back(); + return g_last_cudaError = cudaSuccess; +} + + /******************************************************************************* * * * * @@ -1531,100 +1925,15 @@ __host__ const char* CUDARTAPI cudaGetErrorString(cudaError_t error) return strdup(buf); } -__host__ cudaError_t CUDARTAPI cudaConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, cudaStream_t stream) -{ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - struct CUstream_st *s = (struct CUstream_st *)stream; - CUctx_st *context = GPGPUSim_Context(); - context->g_cuda_launch_stack.push_back( kernel_config(gridDim,blockDim,sharedMem,s) ); - return g_last_cudaError = cudaSuccess; -} - __host__ cudaError_t CUDARTAPI cudaSetupArgument(const void *arg, size_t size, size_t offset) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - CUctx_st *context = GPGPUSim_Context(); - gpgpusim_ptx_assert( !context->g_cuda_launch_stack.empty(), "empty launch stack" ); - kernel_config &config = context->g_cuda_launch_stack.back(); - config.set_arg(arg,size,offset); - printf("GPGPU-Sim PTX: Setting up arguments for %zu bytes starting at 0x%llx..\n",size, (unsigned long long) arg); - - return g_last_cudaError = cudaSuccess; + return cudaSetupArgumentInternal(arg, size, offset); } __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - CUctx_st* context = GPGPUSim_Context(); - char *mode = getenv("PTX_SIM_MODE_FUNC"); - if( mode ) - sscanf(mode,"%u", &g_ptx_sim_mode); - gpgpusim_ptx_assert( !context->g_cuda_launch_stack.empty(), "empty launch stack" ); - kernel_config config = context->g_cuda_launch_stack.back(); - struct CUstream_st *stream = config.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", stream?stream->get_uid():0 ); - kernel_info_t *grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context); - //do dynamic PDOM analysis for performance simulation scenario - std::string kname = grid->name(); - function_info *kernel_func_info = grid->entry(); - if (kernel_func_info->is_pdom_set()) { - printf("GPGPU-Sim PTX: PDOM analysis already done for %s \n", kname.c_str() ); - } else { - printf("GPGPU-Sim PTX: finding reconvergence points for \'%s\'...\n", kname.c_str() ); - kernel_func_info->do_pdom(); - kernel_func_info->set_pdom(); - } - 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 ); - context->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 ); - context->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); - g_stream_manager()->push(op); - context->g_cuda_launch_stack.pop_back(); - return g_last_cudaError = cudaSuccess; + return cudaLaunchInternal( hostFun ); } __host__ cudaError_t CUDARTAPI cudaLaunchKernel ( const char* hostFun, dim3 gridDim, dim3 blockDim, const void** args, size_t sharedMem, cudaStream_t stream ) @@ -1636,13 +1945,13 @@ __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); + cudaConfigureCallInternal(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); + cudaSetupArgumentInternal(args[i], p.first, p.second); } - cudaLaunch(hostFun); + cudaLaunchInternal(hostFun); return g_last_cudaError = cudaSuccess; } @@ -1951,57 +2260,6 @@ __host__ cudaError_t CUDARTAPI cudaGetExportTable(const void **ppExportTable, co //#include "../../cuobjdump_to_ptxplus/cuobjdump_parser.h" -//! Return the executable file of the process containing the PTX/SASS code -//! -//! This Function returns the executable file ran by the process. This -//! executable is supposed to contain the PTX/SASS code. It provides workaround -//! for processes running on valgrind by dereferencing /proc/<pid>/exe within the -//! GPGPU-Sim process before calling cuobjdump to extract PTX/SASS. This is -//! needed because valgrind uses x86 emulation to detect memory leak. Other -//! processes (e.g. cuobjdump) reading /proc/<pid>/exe will see the emulator -//! executable instead of the application binary. -//! -std::string get_app_binary(){ - char self_exe_path[1025]; -#ifdef __APPLE__ - uint32_t size = sizeof(self_exe_path); - if( _NSGetExecutablePath(self_exe_path,&size) != 0 ) { - printf("GPGPU-Sim ** ERROR: _NSGetExecutablePath input buffer too small\n"); - exit(1); - } -#else - std::stringstream exec_link; - exec_link << "/proc/self/exe"; - - ssize_t path_length = readlink(exec_link.str().c_str(), self_exe_path, 1024); - assert(path_length != -1); - self_exe_path[path_length] = '\0'; -#endif - - printf("self exe links to: %s\n", self_exe_path); - return self_exe_path; -} - -//above func gives abs path whereas this give just the name of application. -char* get_app_binary_name(std::string abs_path){ - char *self_exe_path; -#ifdef __APPLE__ - //TODO: get apple device and check the result. - printf("WARNING: not tested for Apple-mac devices \n"); - abort(); -#else - char* buf = strdup(abs_path.c_str()); - char *token = strtok(buf, "/"); - while(token !=NULL){ - self_exe_path = token; - token = strtok(NULL,"/"); - } -#endif - self_exe_path = strtok(self_exe_path, "."); - printf("self exe links to: %s\n", self_exe_path); - return self_exe_path; -} - //extracts all ptx files from binary and dumps into prog_name.unique_no.sm_<>.ptx files void gpgpu_context::extract_ptx_files_using_cuobjdump(CUctx_st *context){ extern bool g_cdp_enabled; @@ -2072,28 +2330,6 @@ void gpgpu_context::extract_ptx_files_using_cuobjdump(CUctx_st *context){ } -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) /*! @@ -2475,11 +2711,6 @@ void gpgpu_context::cuobjdumpInit(){ } -//! Keep track of the association between filename and cubin handle -void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context){ - context->fatbinmap[handle] = filename; -} - //! Either submit PTX for simulation or convert SASS to PTXPlus and submit it void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){ @@ -2570,214 +2801,11 @@ void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){ } } -void** cudaRegisterFatBinary( void *fatCubin, gpgpu_context* gpgpu_ctx = NULL) -{ - gpgpu_context *ctx; - if (gpgpu_ctx){ - ctx = gpgpu_ctx; - } else { - ctx = GPGPU_Context(); - } - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } -#if (CUDART_VERSION < 2010) - printf("GPGPU-Sim PTX: ERROR ** this version of GPGPU-Sim requires CUDA 2.1 or higher\n"); - exit(1); -#endif - CUctx_st *context = GPGPUSim_Context(); - static unsigned next_fat_bin_handle = 1; - if(context->get_device()->get_gpgpu()->get_config().use_cuobjdump()) { - // The following workaround has only been verified on 64-bit systems. - if (sizeof(void*) == 4) - printf("GPGPU-Sim PTX: FatBin file name extraction has not been tested on 32-bit system.\n"); - - // 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. - //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) - 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 - - // 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 - // PTX/SASS code for the launched kernel function. - // This allows us to work around the fact that cuobjdump only outputs the - // file name associated with each section. - unsigned long long fat_cubin_handle = next_fat_bin_handle; - next_fat_bin_handle++; - printf("GPGPU-Sim PTX: __cudaRegisterFatBinary, fat_cubin_handle = %llu, filename=%s\n", fat_cubin_handle, filename); - /*! - * This function extracts all data from all files in first call - * then for next calls, only returns the appropriate number - */ - assert(fat_cubin_handle >= 1); - if (fat_cubin_handle==1) ctx->cuobjdumpInit(); - cuobjdumpRegisterFatBinary(fat_cubin_handle, filename, context); - - return (void**)fat_cubin_handle; - } -#if (CUDART_VERSION < 8000) - else { - static unsigned source_num=1; - unsigned long long fat_cubin_handle = next_fat_bin_handle++; - __cudaFatCudaBinary *info = (__cudaFatCudaBinary *)fatCubin; - assert( info->version >= 3 ); - unsigned num_ptx_versions=0; - unsigned max_capability=0; - unsigned selected_capability=0; - bool found=false; - unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability(); - if (!info->ptx){ - printf("ERROR: Cannot find ptx code in cubin file\n" - "\tIf you are using CUDA 4.0 or higher, please enable -gpgpu_ptx_use_cuobjdump or downgrade to CUDA 3.1\n"); - exit(1); - } - while( info->ptx[num_ptx_versions].gpuProfileName != NULL ) { - unsigned capability=0; - sscanf(info->ptx[num_ptx_versions].gpuProfileName,"compute_%u",&capability); - printf("GPGPU-Sim PTX: __cudaRegisterFatBinary found PTX versions for '%s', ", info->ident); - printf("capability = %s\n", info->ptx[num_ptx_versions].gpuProfileName ); - if( forced_max_capability ) { - if( capability > max_capability && capability <= forced_max_capability ) { - found = true; - max_capability=capability; - selected_capability = num_ptx_versions; - } - } else { - if( capability > max_capability ) { - found = true; - max_capability=capability; - selected_capability = num_ptx_versions; - } - } - num_ptx_versions++; - } - if( found ) { - printf("GPGPU-Sim PTX: Loading PTX for %s, capability = %s\n", - info->ident, info->ptx[selected_capability].gpuProfileName ); - symbol_table *symtab; - const char *ptx = info->ptx[selected_capability].ptx; - if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) { - printf("GPGPU-Sim PTX: ERROR ** PTXPlus is only supported through cuobjdump\n" - "\tEither enable cuobjdump or disable PTXPlus in your configuration file\n"); - exit(1); - } else { - symtab=gpgpu_ptx_sim_load_ptx_from_string(ptx,source_num); - context->add_binary(symtab,fat_cubin_handle); - gpgpu_ptxinfo_load_from_string( ptx, source_num, max_capability, context->no_of_ptx ); - } - source_num++; - load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); - load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); - } else { - printf("GPGPU-Sim PTX: warning -- did not find an appropriate PTX in cubin\n"); - } - return (void**)fat_cubin_handle; - } -#else - else { - printf("ERROR ** __cudaRegisterFatBinary() needs to be updated\n"); - abort(); - } -#endif -} - -void cudaRegisterFunction( - void **fatCubinHandle, - const char *hostFun, - char *deviceFun, - const char *deviceName, - int thread_limit, - uint3 *tid, - uint3 *bid, - dim3 *bDim, - dim3 *gDim, - gpgpu_context *gpgpu_ctx = NULL -) -{ - gpgpu_context *ctx; - if (gpgpu_ctx){ - ctx = gpgpu_ctx; - } else { - ctx = GPGPU_Context(); - } - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - CUctx_st *context = GPGPUSim_Context(); - unsigned fat_cubin_handle = (unsigned)(unsigned long long)fatCubinHandle; - printf("GPGPU-Sim PTX: __cudaRegisterFunction %s : hostFun 0x%p, fat_cubin_handle = %u\n", - deviceFun, hostFun, fat_cubin_handle); - if(context->get_device()->get_gpgpu()->get_config().use_cuobjdump()) - ctx->cuobjdumpParseBinary(fat_cubin_handle); - context->register_function( fat_cubin_handle, hostFun, deviceFun ); -} - -void cudaRegisterVar( - void **fatCubinHandle, - char *hostVar, //pointer to...something - char *deviceAddress, //name of variable - const char *deviceName, //name of variable (same as above) - int ext, - int size, - int constant, - int global, - gpgpu_context *gpgpu_ctx = NULL) -{ - gpgpu_context *ctx; - if (gpgpu_ctx){ - ctx = gpgpu_ctx; - } else { - ctx = GPGPU_Context(); - } - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - printf("GPGPU-Sim PTX: __cudaRegisterVar: hostVar = %p; deviceAddress = %s; deviceName = %s\n", hostVar, deviceAddress, deviceName); - printf("GPGPU-Sim PTX: __cudaRegisterVar: Registering const memory space of %d bytes\n", size); - if(GPGPUSim_Context()->get_device()->get_gpgpu()->get_config().use_cuobjdump()) - ctx->cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle); - fflush(stdout); - if ( constant && !global && !ext ) { - gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size); - } else if ( !constant && !global && !ext ) { - gpgpu_ptx_sim_register_global_variable(hostVar,deviceName,size); - } else cuda_not_implemented(__my_func__,__LINE__); -} - extern "C" { void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) { - return cudaRegisterFatBinary(fatCubin); + return cudaRegisterFatBinaryInternal(fatCubin); } void CUDARTAPI __cudaRegisterFunction( @@ -2791,7 +2819,7 @@ void CUDARTAPI __cudaRegisterFunction( dim3 *bDim, dim3 *gDim ) { - cudaRegisterFunction( + cudaRegisterFunctionInternal( fatCubinHandle, hostFun, deviceFun, @@ -2815,7 +2843,7 @@ extern void __cudaRegisterVar( int constant, int global ) { - cudaRegisterVar( + cudaRegisterVarInternal( fatCubinHandle, hostVar, deviceAddress, @@ -2825,6 +2853,12 @@ extern void __cudaRegisterVar( constant, global ); } + +__host__ cudaError_t CUDARTAPI cudaConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, cudaStream_t stream) +{ + return cudaConfigureCallInternal(gridDim, blockDim, sharedMem, stream); +} + void __cudaUnregisterFatBinary(void **fatCubinHandle) { if(g_debug_execution >= 3){ @@ -4868,12 +4902,12 @@ CUresult CUDAAPI cuLaunchKernel(CUfunction f, const char *hostFun = (const char*) f; CUctx_st *context = GPGPUSim_Context(); function_info *entry = context->get_kernel(hostFun); - cudaConfigureCall(dim3(gridDimX, gridDimY, gridDimZ), dim3(blockDimX, blockDimY, blockDimZ), sharedMemBytes, (cudaStream_t) hStream); + cudaConfigureCallInternal(dim3(gridDimX, gridDimY, gridDimZ), dim3(blockDimX, blockDimY, blockDimZ), sharedMemBytes, (cudaStream_t) hStream); for(unsigned i = 0; i < entry->num_args(); i++){ std::pair<size_t, unsigned> p = entry->get_param_config(i); cudaSetupArgument(kernelParams[i], p.first, p.second); } - cudaLaunch(hostFun); + cudaLaunchInternal(hostFun); return CUDA_SUCCESS; } #endif /* CUDART_VERSION >= 4000 */ |