From d553124832aca461dda4dd7d503748d22b44bbe2 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 5 Jun 2019 23:28:12 -0400 Subject: Add cuda_api_object and move some var Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 19 +++++++++++++++++-- 1 file changed, 17 insertions(+), 2 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index d6e564b..576ec67 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -1,4 +1,9 @@ +#ifndef __gpgpu_context_h__ +#define __gpgpu_context_h__ #include +#include +#include +#include "cuda_api_object.h" class cuobjdumpSection; class cuobjdumpELFSection; @@ -6,13 +11,23 @@ class cuobjdumpPTXSection; class gpgpu_context { public: + gpgpu_context() { + api = new cuda_runtime_api(); + } + // global list std::list cuobjdumpSectionList; + //maps sm version number to set of filenames + std::map > version_filename; + cuda_runtime_api* api; + // member function list void cuobjdumpInit(); void cuobjdumpParseBinary(unsigned int handle); void extract_code_using_cuobjdump(); std::list pruneSectionList(CUctx_st *context); std::list mergeMatchingSections(std::string identifier); std::list mergeSections(); - cuobjdumpELFSection* findELFSection(const std::string identifier, std::list &libSectionList); - cuobjdumpPTXSection* findPTXSection(const std::string identifier, std::list &libSectionList); + cuobjdumpELFSection* findELFSection(const std::string identifier); + cuobjdumpPTXSection* findPTXSection(const std::string identifier); + void extract_ptx_files_using_cuobjdump(CUctx_st *context); }; +#endif /* __gpgpu_context_h__ */ -- cgit v1.3 From 570ce0a3d6e049ec3ee42329fbf25019fd4eb2e5 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Thu, 6 Jun 2019 01:57:40 -0400 Subject: Move 3 more var Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 92 +++++++++++++++++++++++++-------------------- libcuda/gpgpu_context.h | 6 +++ 2 files changed, 58 insertions(+), 40 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index ea96570..855f11d 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -313,10 +313,7 @@ struct CUctx_st { return i->second; } - std::mapfatbin_registered; - std::map fatbinmap; std::map g_mallocPtr_Size; - std::map name_symtab; std::map pinned_memory; //support for pinned memories added std::map pinned_memory_size; int no_of_ptx; @@ -647,8 +644,8 @@ static int get_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; +void gpgpu_context::cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context){ + fatbinmap[handle] = filename; } /******************************************************************************* @@ -725,7 +722,7 @@ void** cudaRegisterFatBinaryInternal( void *fatCubin, gpgpu_context* gpgpu_ctx = */ assert(fat_cubin_handle >= 1); if (fat_cubin_handle==1) ctx->cuobjdumpInit(); - cuobjdumpRegisterFatBinary(fat_cubin_handle, filename, context); + ctx->cuobjdumpRegisterFatBinary(fat_cubin_handle, filename, context); return (void**)fat_cubin_handle; } @@ -969,6 +966,46 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx = return g_last_cudaError = cudaSuccess; } +#if CUDART_VERSION >= 6050 +CUresult +cuLinkAddFileInternal(CUlinkState state, CUjitInputType type, const char *path, + unsigned int numOptions, CUjit_option *options, void **optionValues, 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__); + } + static bool addedFile = false; + if (addedFile){ + printf("GPGPU-Sim PTX: ERROR: cuLinkAddFile does not support multiple files\n"); + abort(); + } + + //blocking + assert(type==CU_JIT_INPUT_PTX); + CUctx_st *context = GPGPUSim_Context(); + char *file = getenv("PTX_JIT_PATH"); + if(file==NULL){ + printf("GPGPU-Sim PTX: ERROR: PTX_JIT_PATH has not been set\n"); + abort(); + } + strcat(file,"/"); + strcat(file,path); + symbol_table *symtab = gpgpu_ptx_sim_load_ptx_from_filename( file ); + std::string fname(path); + ctx->name_symtab[fname] = symtab; + context->add_binary(symtab, 1); + load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); + load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); + addedFile = true; + return CUDA_SUCCESS; +} +#endif /******************************************************************************* * * @@ -2715,12 +2752,12 @@ void gpgpu_context::cuobjdumpInit(){ void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){ CUctx_st *context = GPGPUSim_Context(); - if(context->fatbin_registered[handle]) return; - context->fatbin_registered[handle] = true; - std::string fname = context->fatbinmap[handle]; + if(fatbin_registered[handle]) return; + fatbin_registered[handle] = true; + std::string fname = fatbinmap[handle]; - if (context->name_symtab.find(fname) != context->name_symtab.end()) { - symbol_table *symtab = context->name_symtab[fname]; + if (name_symtab.find(fname) != name_symtab.end()) { + symbol_table *symtab = name_symtab[fname]; context->add_binary(symtab, handle); return; } @@ -2737,7 +2774,7 @@ void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){ symtab = gpgpu_ptx_sim_load_ptx_from_filename( ptx_filename.c_str() ); } } - context->name_symtab[fname] = symtab; + name_symtab[fname] = symtab; context->add_binary(symtab, handle); load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); @@ -2795,7 +2832,7 @@ void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){ } load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); - context->name_symtab[fname] = symtab; + name_symtab[fname] = symtab; //TODO: Remove temporarily files as per configurations } @@ -3909,33 +3946,8 @@ CUresult CUDAAPI cuLinkAddFile(CUlinkState state, CUjitInputType type, const char *path, unsigned int numOptions, CUjit_option *options, void **optionValues) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - static bool addedFile = false; - if (addedFile){ - printf("GPGPU-Sim PTX: ERROR: cuLinkAddFile does not support multiple files\n"); - abort(); - } - - //blocking - assert(type==CU_JIT_INPUT_PTX); - CUctx_st *context = GPGPUSim_Context(); - char *file = getenv("PTX_JIT_PATH"); - if(file==NULL){ - printf("GPGPU-Sim PTX: ERROR: PTX_JIT_PATH has not been set\n"); - abort(); - } - strcat(file,"/"); - strcat(file,path); - symbol_table *symtab = gpgpu_ptx_sim_load_ptx_from_filename( file ); - std::string fname(path); - context->name_symtab[fname] = symtab; - context->add_binary(symtab, 1); - load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); - load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); - addedFile = true; - return CUDA_SUCCESS; + return cuLinkAddFileInternal(state, type, path, + numOptions, options, optionValues) } #endif diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 576ec67..0543ff8 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -3,11 +3,13 @@ #include #include #include +#include #include "cuda_api_object.h" class cuobjdumpSection; class cuobjdumpELFSection; class cuobjdumpPTXSection; +class symbol_table; class gpgpu_context { public: @@ -16,6 +18,9 @@ class gpgpu_context { } // global list std::list cuobjdumpSectionList; + std::mapfatbin_registered; + std::map fatbinmap; + std::map name_symtab; //maps sm version number to set of filenames std::map > version_filename; cuda_runtime_api* api; @@ -29,5 +34,6 @@ class gpgpu_context { cuobjdumpELFSection* findELFSection(const std::string identifier); cuobjdumpPTXSection* findPTXSection(const std::string identifier); void extract_ptx_files_using_cuobjdump(CUctx_st *context); + void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context); }; #endif /* __gpgpu_context_h__ */ -- cgit v1.3 From bc8d1e10507f043c916ee051dc9a687adc6d9b4b Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Thu, 6 Jun 2019 02:47:07 -0400 Subject: Move more vars Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 117 ++++++++++++++++++++++++++------------------ libcuda/gpgpu_context.h | 11 +++++ 2 files changed, 81 insertions(+), 47 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 66562aa..02e2b2e 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -158,11 +158,11 @@ extern void exit_simulation(); static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr, gpgpu_t *gpu ); static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu ); -static kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, - gpgpu_ptx_sim_arg_list_t args, - struct dim3 gridDim, - struct dim3 blockDim, - struct CUctx_st* context ); +//static kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, +// gpgpu_ptx_sim_arg_list_t args, +// struct dim3 gridDim, +// struct dim3 blockDim, +// struct CUctx_st* context ); /*DEVICE_BUILTIN*/ struct cudaArray @@ -313,7 +313,6 @@ struct CUctx_st { return i->second; } - std::map g_mallocPtr_Size; std::map pinned_memory; //support for pinned memories added std::map pinned_memory_size; int no_of_ptx; @@ -910,7 +909,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx = 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); + kernel_info_t *grid = ctx->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(); @@ -966,6 +965,66 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx = return g_last_cudaError = cudaSuccess; } +cudaError_t cudaMallocInternal(void **devPtr, size_t size, 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(); + *devPtr = context->get_device()->get_gpgpu()->gpu_malloc(size); + if(g_debug_execution >= 3){ + printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *devPtr); + ctx->g_mallocPtr_Size[(unsigned long long)*devPtr] = size; + } + if ( *devPtr ) { + return g_last_cudaError = cudaSuccess; + } else { + return g_last_cudaError = cudaErrorMemoryAllocation; + } +} + +cudaError_t cudaHostGetDevicePointerInternal(void **pDevice, void *pHost, unsigned int flags, 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(g_debug_execution >= 3){ + announce_call(__my_func__); + } + //only cpu memory allocation happens in cudaHostAlloc. Linking with device pointer to pinned memory happens here. + //TODO: once kernel is executed, the contents in global pointer of GPU must be copied back to CPU host pointer! + flags=0; + CUctx_st* context = GPGPUSim_Context(); + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + std::map::const_iterator i = context->pinned_memory_size.find(pHost); + assert(i != context->pinned_memory_size.end()); + size_t size = i->second; + *pDevice = gpu->gpu_malloc(size); + if(g_debug_execution >= 3){ + printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *pDevice); + ctx->g_mallocPtr_Size[(unsigned long long)*pDevice] = size; + } + if ( *pDevice ) { + context->pinned_memory[pHost]=pDevice; + //Copy contents in cpu to gpu + gpu->memcpy_to_gpu((size_t)*pDevice,pHost,size); + return g_last_cudaError = cudaSuccess; + } else { + return g_last_cudaError = cudaErrorMemoryAllocation; + } +} #if CUDART_VERSION >= 6050 CUresult cuLinkAddFileInternal(CUlinkState state, CUjitInputType type, const char *path, @@ -1027,20 +1086,7 @@ cudaError_t cudaPeekAtLastError(void) __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - CUctx_st* context = GPGPUSim_Context(); - *devPtr = context->get_device()->get_gpgpu()->gpu_malloc(size); - if(g_debug_execution >= 3){ - printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *devPtr); - context->g_mallocPtr_Size[(unsigned long long)*devPtr] = size; - } - if ( *devPtr ) { - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorMemoryAllocation; - } + return cudaMallocInternal(devPtr, size); } __host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size) @@ -3111,30 +3157,7 @@ cudaError_t CUDARTAPI cudaHostAlloc(void **pHost, size_t bytes, unsigned int fl cudaError_t CUDARTAPI cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - //only cpu memory allocation happens in cudaHostAlloc. Linking with device pointer to pinned memory happens here. - //TODO: once kernel is executed, the contents in global pointer of GPU must be copied back to CPU host pointer! - flags=0; - CUctx_st* context = GPGPUSim_Context(); - gpgpu_t *gpu = context->get_device()->get_gpgpu(); - std::map::const_iterator i = context->pinned_memory_size.find(pHost); - assert(i != context->pinned_memory_size.end()); - size_t size = i->second; - *pDevice = gpu->gpu_malloc(size); - if(g_debug_execution >= 3){ - printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *pDevice); - context->g_mallocPtr_Size[(unsigned long long)*pDevice] = size; - } - if ( *pDevice ) { - context->pinned_memory[pHost]=pDevice; - //Copy contents in cpu to gpu - gpu->memcpy_to_gpu((size_t)*pDevice,pHost,size); - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorMemoryAllocation; - } + return cudaHostGetDevicePointerInternal(pDevice, pHost, flags); } __host__ cudaError_t CUDARTAPI cudaPointerGetAttributes( @@ -3450,7 +3473,7 @@ static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu return nc_bytes; } -kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, +kernel_info_t * gpgpu_context::gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, gpgpu_ptx_sim_arg_list_t args, struct dim3 gridDim, struct dim3 blockDim, @@ -3482,7 +3505,7 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, fflush(stdout); if(g_debug_execution >= 4){ - entry->ptx_jit_config(context->g_mallocPtr_Size, result->get_param_memory(), (gpgpu_t *) context->get_device()->get_gpgpu(), gridDim, blockDim); + entry->ptx_jit_config(g_mallocPtr_Size, result->get_param_memory(), (gpgpu_t *) context->get_device()->get_gpgpu(), gridDim, blockDim); } return result; diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 0543ff8..f29e2e0 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -10,6 +10,10 @@ class cuobjdumpSection; class cuobjdumpELFSection; class cuobjdumpPTXSection; class symbol_table; +class gpgpu_ptx_sim_arg; +class kernel_info_t; + +typedef std::list gpgpu_ptx_sim_arg_list_t; class gpgpu_context { public: @@ -21,6 +25,7 @@ class gpgpu_context { std::mapfatbin_registered; std::map fatbinmap; std::map name_symtab; + std::map g_mallocPtr_Size; //maps sm version number to set of filenames std::map > version_filename; cuda_runtime_api* api; @@ -35,5 +40,11 @@ class gpgpu_context { cuobjdumpPTXSection* findPTXSection(const std::string identifier); void extract_ptx_files_using_cuobjdump(CUctx_st *context); void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context); + kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, + gpgpu_ptx_sim_arg_list_t args, + struct dim3 gridDim, + struct dim3 blockDim, + struct CUctx_st* context ); + }; #endif /* __gpgpu_context_h__ */ -- cgit v1.3