From dc835911d5320008b2c227722a90240a8f6b0f3a Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Thu, 23 May 2019 22:44:37 -0400 Subject: Fix linebuf for ptxinfo Signed-off-by: Mengchi Zhang --- src/cuda-sim/ptx_loader.h | 10 ++++++++++ 1 file changed, 10 insertions(+) (limited to 'src/cuda-sim/ptx_loader.h') diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index e5df6a9..36e439e 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -29,6 +29,16 @@ #define PTX_LOADER_H_INCLUDED #include +#define LINEBUF_SIZE 1024 +typedef void * yyscan_t; +class ptxinfo_data{ + public: + yyscan_t scanner; + char linebuf[LINEBUF_SIZE]; + unsigned col; +}; + + extern bool g_override_embedded_ptx; extern int no_of_ptx; //counter to track number of ptx files to be extracted in an application. -- cgit v1.3 From 9e3a9ac5ed0a70ec9b048bd3cb4df781687e85f8 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 29 May 2019 19:23:40 -0400 Subject: Move fatbin etc globals Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 111 +++++++++++++++++++++----------------------- src/cuda-sim/ptx_loader.cc | 2 +- src/cuda-sim/ptx_loader.h | 2 +- 3 files changed, 56 insertions(+), 59 deletions(-) (limited to 'src/cuda-sim/ptx_loader.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index f3c827c..fcd5b07 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -149,10 +149,6 @@ typedef void * yyscan_t; #include #endif -std::map pinned_memory; //support for pinned memories added -std::map pinned_memory_size; -std::map g_mallocPtr_Size; -int no_of_ptx=0; extern void synchronize(); extern void exit_simulation(); @@ -241,12 +237,25 @@ private: class kernel_config; +#ifndef OPENGL_SUPPORT +typedef unsigned long GLuint; +#endif + +struct glbmap_entry { + GLuint m_bufferObj; + void *m_devPtr; + size_t m_size; + struct glbmap_entry *m_next; +}; + struct CUctx_st { CUctx_st( _cuda_device_id *gpu ) { m_gpu = gpu; m_binary_info.cmem = 0; m_binary_info.gmem = 0; + no_of_ptx=0; + g_glbmap = NULL; } _cuda_device_id *get_device() { return m_gpu; } @@ -309,6 +318,16 @@ struct CUctx_st { //maps sm version number to set of filenames std::map > version_filename; std::list g_cuda_launch_stack; + 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; + typedef struct glbmap_entry glbmap_entry_t; + + glbmap_entry_t* g_glbmap; private: _cuda_device_id *m_gpu; // selected gpu @@ -573,7 +592,7 @@ __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) *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); - g_mallocPtr_Size[(unsigned long long)*devPtr] = size; + context->g_mallocPtr_Size[(unsigned long long)*devPtr] = size; } if ( *devPtr ) { return g_last_cudaError = cudaSuccess; @@ -587,11 +606,11 @@ __host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size) if(g_debug_execution >= 3){ announce_call(__my_func__); } - GPGPUSim_Context(); + CUctx_st* context = GPGPUSim_Context(); *ptr = malloc(size); if ( *ptr ) { //track pinned memory size allocated in the host so that same amount of memory is also allocated in GPU. - pinned_memory_size[*ptr]=size; + context->pinned_memory_size[*ptr]=size; return g_last_cudaError = cudaSuccess; } else { return g_last_cudaError = cudaErrorMemoryAllocation; @@ -2010,11 +2029,11 @@ void extract_ptx_files_using_cuobjdump(CUctx_st *context){ printf("ERROR: command: %s failed \n",command); exit(0); } - no_of_ptx++; + context->no_of_ptx++; } } - if(!no_of_ptx){ + if(!context->no_of_ptx){ printf("WARNING: Number of ptx in the executable file are 0. One of the reasons might be\n"); printf("\t1. CDP is enabled\n"); printf("\t2. When using PyTorch, PYTORCH_BIN is not set correctly\n"); @@ -2444,25 +2463,22 @@ void cuobjdumpInit(std::list &cuobjdumpSectionList){ } } -std::map fatbinmap; -std::mapfatbin_registered; -std::map name_symtab; //! Keep track of the association between filename and cubin handle -void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename){ - fatbinmap[handle] = filename; +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 cuobjdumpParseBinary(unsigned int handle){ - if(fatbin_registered[handle]) return; - fatbin_registered[handle] = true; CUctx_st *context = GPGPUSim_Context(); - std::string fname = fatbinmap[handle]; + if(context->fatbin_registered[handle]) return; + context->fatbin_registered[handle] = true; + std::string fname = context->fatbinmap[handle]; - if (name_symtab.find(fname) != name_symtab.end()) { - symbol_table *symtab = name_symtab[fname]; + if (context->name_symtab.find(fname) != context->name_symtab.end()) { + symbol_table *symtab = context->name_symtab[fname]; context->add_binary(symtab, handle); return; } @@ -2479,7 +2495,7 @@ void cuobjdumpParseBinary(unsigned int handle){ symtab = gpgpu_ptx_sim_load_ptx_from_filename( ptx_filename.c_str() ); } } - name_symtab[fname] = symtab; + context->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()); @@ -2526,18 +2542,18 @@ void cuobjdumpParseBinary(unsigned int handle){ symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxplus_str, handle); printf("Adding %s with cubin handle %u\n", ptx->getPTXfilename().c_str(), handle); context->add_binary(symtab, handle); - gpgpu_ptxinfo_load_from_string( ptxcode, handle, max_capability ); + gpgpu_ptxinfo_load_from_string( ptxcode, handle, max_capability, context->no_of_ptx ); delete[] ptxplus_str; } else { symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxcode, handle); //if CUOBJDUMP_SIM_FILE is not set, ptx is NULL. So comment below. //printf("Adding %s with cubin handle %u\n", ptx->getPTXfilename().c_str(), handle); context->add_binary(symtab, handle); - gpgpu_ptxinfo_load_from_string( ptxcode, handle, max_capability ); + gpgpu_ptxinfo_load_from_string( ptxcode, handle, max_capability, context->no_of_ptx ); } load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); - name_symtab[fname] = symtab; + context->name_symtab[fname] = symtab; //TODO: Remove temporarily files as per configurations } @@ -2606,7 +2622,7 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) */ assert(fat_cubin_handle >= 1); if (fat_cubin_handle==1) cuobjdumpInit(context->cuobjdumpSectionList); - cuobjdumpRegisterFatBinary(fat_cubin_handle, filename); + cuobjdumpRegisterFatBinary(fat_cubin_handle, filename, context); return (void**)fat_cubin_handle; } @@ -2658,7 +2674,7 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) } 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 ); + 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()); @@ -2818,10 +2834,6 @@ char __cudaInitModule( } -#ifndef OPENGL_SUPPORT -typedef unsigned long GLuint; -#endif - cudaError_t cudaGLRegisterBufferObject(GLuint bufferObj) { if(g_debug_execution >= 3){ @@ -2831,16 +2843,6 @@ cudaError_t cudaGLRegisterBufferObject(GLuint bufferObj) return g_last_cudaError = cudaSuccess; } -struct glbmap_entry { - GLuint m_bufferObj; - void *m_devPtr; - size_t m_size; - struct glbmap_entry *m_next; -}; -typedef struct glbmap_entry glbmap_entry_t; - -glbmap_entry_t* g_glbmap = NULL; - cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj) { if(g_debug_execution >= 3){ @@ -2850,7 +2852,7 @@ cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj) GLint buffer_size=0; CUctx_st* ctx = GPGPUSim_Context(); - glbmap_entry_t *p = g_glbmap; + glbmap_entry_t *p = ctx->g_glbmap; while ( p && p->m_bufferObj != bufferObj ) p = p->m_next; if ( p == NULL ) { @@ -2861,8 +2863,8 @@ cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj) // create entry and insert to front of list glbmap_entry_t *n = (glbmap_entry_t *) calloc(1,sizeof(glbmap_entry_t)); - n->m_next = g_glbmap; - g_glbmap = n; + n->m_next = ctx->g_glbmap; + ctx->g_glbmap = n; // initialize entry n->m_bufferObj = bufferObj; @@ -2903,7 +2905,8 @@ cudaError_t cudaGLUnmapBufferObject(GLuint bufferObj) announce_call(__my_func__); } #ifdef OPENGL_SUPPORT - glbmap_entry_t *p = g_glbmap; + CUctx_st* ctx = GPGPUSim_Context(); + glbmap_entry_t *p = ctx->g_glbmap; while ( p && p->m_bufferObj != bufferObj ) p = p->m_next; if ( p == NULL ) @@ -2943,7 +2946,8 @@ cudaError_t CUDARTAPI cudaHostAlloc(void **pHost, size_t bytes, unsigned int fl *pHost = malloc(bytes); //need to track the size allocated so that cudaHostGetDevicePointer() can function properly. //TODO: vary this function behavior based on flags value (following nvidia documentation) - pinned_memory_size[*pHost]=bytes; + CUctx_st* context = GPGPUSim_Context(); + context->pinned_memory_size[*pHost]=bytes; if( *pHost ) return g_last_cudaError = cudaSuccess; else @@ -2960,16 +2964,16 @@ cudaError_t CUDARTAPI cudaHostGetDevicePointer(void **pDevice, void *pHost, unsi flags=0; CUctx_st* context = GPGPUSim_Context(); gpgpu_t *gpu = context->get_device()->get_gpgpu(); - std::map::const_iterator i = pinned_memory_size.find(pHost); - assert(i != pinned_memory_size.end()); + 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); - g_mallocPtr_Size[(unsigned long long)*pDevice] = size; + context->g_mallocPtr_Size[(unsigned long long)*pDevice] = size; } if ( *pDevice ) { - pinned_memory[pHost]=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; @@ -3204,13 +3208,6 @@ int CUDARTAPI __cudaSynchronizeThreads(void**, void*) //////// -extern int ptx_parse(); -extern int ptx__scan_string(const char*); -extern FILE *ptx_in; - -extern int ptxinfo_parse(); -extern FILE *ptxinfo_in; - /// static functions static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr, gpgpu_t *gpu ) @@ -3330,7 +3327,7 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, fflush(stdout); if(g_debug_execution >= 4){ - entry->ptx_jit_config(g_mallocPtr_Size, result->get_param_memory(), (gpgpu_t *) context->get_device()->get_gpgpu(), gridDim, blockDim); + entry->ptx_jit_config(context->g_mallocPtr_Size, result->get_param_memory(), (gpgpu_t *) context->get_device()->get_gpgpu(), gridDim, blockDim); } return result; @@ -3815,7 +3812,7 @@ cuLinkAddFile(CUlinkState state, CUjitInputType type, const char *path, strcat(file,path); symbol_table *symtab = gpgpu_ptx_sim_load_ptx_from_filename( file ); std::string fname(path); - name_symtab[fname] = symtab; + 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()); diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 735ff84..f037c34 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -367,7 +367,7 @@ void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_versio fclose(ptxinfo_in); } -void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version ) +void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version, int no_of_ptx ) { //do ptxas for individual files instead of one big embedded ptx. This prevents the duplicate defs and declarations. char ptx_file[1000]; diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index 36e439e..c4d8292 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -44,7 +44,7 @@ extern int no_of_ptx; //counter to track number of ptx files to be extracted in class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); class symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename ); -void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version=20 ); +void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version=20, int no_of_ptx=0 ); void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version ); char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str); bool keep_intermediate_files(); -- cgit v1.3 From 9e52c143a883f682c02d81149748cdf8aa5508f7 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 12 Jun 2019 01:42:52 -0400 Subject: Move some function from ptx_loader to gpgpu_context Signed-off-by: Mengchi Zhang --- libcuda/cuda_api_object.h | 1 - libcuda/cuda_runtime_api.cc | 42 ++++++++++++++++++++--------------------- libcuda/gpgpu_context.h | 3 +++ libopencl/opencl_runtime_api.cc | 5 ++++- src/cuda-sim/ptx_loader.cc | 5 +++-- src/cuda-sim/ptx_loader.h | 2 -- 6 files changed, 31 insertions(+), 27 deletions(-) (limited to 'src/cuda-sim/ptx_loader.h') diff --git a/libcuda/cuda_api_object.h b/libcuda/cuda_api_object.h index f9a4fde..0054697 100644 --- a/libcuda/cuda_api_object.h +++ b/libcuda/cuda_api_object.h @@ -191,7 +191,6 @@ class cuda_runtime_api { void cuobjdumpInit(); void extract_code_using_cuobjdump(); void extract_ptx_files_using_cuobjdump(CUctx_st *context); - void cuobjdumpParseBinary(unsigned int handle); std::list pruneSectionList(CUctx_st *context); std::list mergeMatchingSections(std::string identifier); std::list mergeSections(); diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 09a13a7..25642a7 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -685,7 +685,7 @@ void cudaRegisterFunctionInternal( 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->api->cuobjdumpParseBinary(fat_cubin_handle); + ctx->cuobjdumpParseBinary(fat_cubin_handle); context->register_function( fat_cubin_handle, hostFun, deviceFun ); } @@ -712,7 +712,7 @@ void cudaRegisterVarInternal( 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->api->cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle); + ctx->cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle); fflush(stdout); if ( constant && !global && !ext ) { gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size); @@ -1009,7 +1009,7 @@ cuLinkAddFileInternal(CUlinkState state, CUjitInputType type, const char *path, } strcat(file,"/"); strcat(file,path); - symbol_table *symtab = gpgpu_ptx_sim_load_ptx_from_filename( file ); + symbol_table *symtab = ctx->gpgpu_ptx_sim_load_ptx_from_filename( file ); std::string fname(path); ctx->api->name_symtab[fname] = symtab; context->add_binary(symtab, 1); @@ -2750,15 +2750,15 @@ void cuda_runtime_api::cuobjdumpInit(){ //! Either submit PTX for simulation or convert SASS to PTXPlus and submit it -void cuda_runtime_api::cuobjdumpParseBinary(unsigned int handle){ +void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){ CUctx_st *context = GPGPUSim_Context(); - if(fatbin_registered[handle]) return; - fatbin_registered[handle] = true; - std::string fname = fatbinmap[handle]; + if(api->fatbin_registered[handle]) return; + api->fatbin_registered[handle] = true; + std::string fname = api->fatbinmap[handle]; - if (name_symtab.find(fname) != name_symtab.end()) { - symbol_table *symtab = name_symtab[fname]; + if (api->name_symtab.find(fname) != api->name_symtab.end()) { + symbol_table *symtab = api->name_symtab[fname]; context->add_binary(symtab, handle); return; } @@ -2767,7 +2767,7 @@ void cuda_runtime_api::cuobjdumpParseBinary(unsigned int handle){ #if (CUDART_VERSION >= 6000) //loops through all ptx files from smallest sm version to largest std::map >::iterator itr_m; - for (itr_m = version_filename.begin(); itr_m!=version_filename.end(); itr_m++){ + for (itr_m = api->version_filename.begin(); itr_m!=api->version_filename.end(); itr_m++){ std::set::iterator itr_s; for (itr_s = itr_m->second.begin(); itr_s!=itr_m->second.end(); itr_s++){ std::string ptx_filename = *itr_s; @@ -2775,11 +2775,11 @@ void cuda_runtime_api::cuobjdumpParseBinary(unsigned int handle){ symtab = gpgpu_ptx_sim_load_ptx_from_filename( ptx_filename.c_str() ); } } - name_symtab[fname] = symtab; + api->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()); - for (itr_m = version_filename.begin(); itr_m!=version_filename.end(); itr_m++){ + api->load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); + api->load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); + for (itr_m = api->version_filename.begin(); itr_m!=api->version_filename.end(); itr_m++){ std::set::iterator itr_s; for (itr_s = itr_m->second.begin(); itr_s!=itr_m->second.end(); itr_s++){ std::string ptx_filename = *itr_s; @@ -2791,8 +2791,8 @@ void cuda_runtime_api::cuobjdumpParseBinary(unsigned int handle){ #endif unsigned max_capability = 0; - for ( std::list::iterator iter = cuobjdumpSectionList.begin(); - iter != cuobjdumpSectionList.end(); + for ( std::list::iterator iter = api->cuobjdumpSectionList.begin(); + iter != api->cuobjdumpSectionList.end(); iter++){ unsigned capability = (*iter)->getArch(); if (capability > max_capability) max_capability = capability; @@ -2803,7 +2803,7 @@ void cuda_runtime_api::cuobjdumpParseBinary(unsigned int handle){ cuobjdumpPTXSection* ptx = NULL; const char* pre_load = getenv("CUOBJDUMP_SIM_FILE"); if(pre_load==NULL || strlen(pre_load)==0) - ptx = findPTXSection(fname); + ptx = api->findPTXSection(fname); char *ptxcode; const char *override_ptx_name = getenv("PTX_SIM_KERNELFILE"); if (override_ptx_name == NULL or getenv("PTX_SIM_USE_PTX_FILE") == NULL or strlen(getenv("PTX_SIM_USE_PTX_FILE"))==0) { @@ -2813,7 +2813,7 @@ void cuda_runtime_api::cuobjdumpParseBinary(unsigned int handle){ ptxcode = readfile(override_ptx_name); } if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) { - cuobjdumpELFSection* elfsection = findELFSection(ptx->getIdentifier()); + cuobjdumpELFSection* elfsection = api->findELFSection(ptx->getIdentifier()); assert (elfsection!= NULL); char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus( ptx->getPTXfilename(), @@ -2831,9 +2831,9 @@ void cuda_runtime_api::cuobjdumpParseBinary(unsigned int handle){ context->add_binary(symtab, handle); gpgpu_ptxinfo_load_from_string( ptxcode, handle, max_capability, context->no_of_ptx ); } - load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); - load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); - name_symtab[fname] = symtab; + api->load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); + api->load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); + api->name_symtab[fname] = symtab; //TODO: Remove temporarily files as per configurations } diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 6c8a293..16626eb 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -11,6 +11,9 @@ class gpgpu_context { // objects pointers for each file cuda_runtime_api* api; // member function list + void cuobjdumpParseBinary(unsigned int handle); + class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); + class symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename ); }; gpgpu_context* GPGPU_Context(); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 97a54d8..0ec635e 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -84,6 +84,7 @@ #include "../src/gpgpusim_entrypoint.h" #include "../src/gpgpu-sim/gpu-sim.h" #include "../src/gpgpu-sim/shader.h" +#include "../libcuda/gpgpu_context.h" //# define __my_func__ __PRETTY_FUNCTION__ # if defined __cplusplus ? __GNUC_PREREQ (2, 6) : __GNUC_PREREQ (2, 4) @@ -424,6 +425,8 @@ void register_ptx_function( const char *name, function_info *impl ) void _cl_program::Build(const char *options) { + gpgpu_context *ctx; + ctx = GPGPU_Context(); printf("GPGPU-Sim OpenCL API: compiling OpenCL kernels...\n"); std::map::iterator i; for( i = m_pgm.begin(); i!= m_pgm.end(); i++ ) { @@ -576,7 +579,7 @@ void _cl_program::Build(const char *options) } } info.m_asm = tmp; - info.m_symtab = gpgpu_ptx_sim_load_ptx_from_string( tmp, source_num ); + info.m_symtab = ctx->gpgpu_ptx_sim_load_ptx_from_string( tmp, source_num ); gpgpu_ptxinfo_load_from_string( tmp, source_num ); free(tmp); } diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index f037c34..d7e9b71 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -33,6 +33,7 @@ #include #include #include +#include "../../libcuda/gpgpu_context.h" /// globals @@ -165,7 +166,7 @@ char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptxfilenam } -symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ) +symbol_table *gpgpu_context::gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ) { char buf[1024]; snprintf(buf,1024,"_%u.ptx", source_num ); @@ -200,7 +201,7 @@ symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source return symtab; } -symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename ) +symbol_table *gpgpu_context::gpgpu_ptx_sim_load_ptx_from_filename( const char *filename ) { symbol_table *symtab=init_parser(filename); printf("GPGPU-Sim PTX: finished parsing EMBEDDED .ptx file %s\n",filename); diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index c4d8292..2af611a 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -42,8 +42,6 @@ class ptxinfo_data{ extern bool g_override_embedded_ptx; extern int no_of_ptx; //counter to track number of ptx files to be extracted in an application. -class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); -class symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename ); void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version=20, int no_of_ptx=0 ); void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version ); char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str); -- cgit v1.3 From 7a6ff7084c1c65450775b4d58e80f8c789888815 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 12 Jun 2019 11:09:11 -0400 Subject: Move g_ptxinfo_filename Signed-off-by: Mengchi Zhang --- src/cuda-sim/ptx_loader.cc | 18 ++++++++---------- src/cuda-sim/ptx_loader.h | 2 ++ src/cuda-sim/ptxinfo.l | 3 +-- 3 files changed, 11 insertions(+), 12 deletions(-) (limited to 'src/cuda-sim/ptx_loader.h') diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 2d931f1..9857741 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -50,8 +50,6 @@ extern int ptx__scan_string(const char*, yyscan_t scanner); extern std::map get_duplicate(); -const char *g_ptxinfo_filename; - typedef void * yyscan_t; extern int ptxinfo_lex_init(yyscan_t* scanner); extern void ptxinfo_set_in (FILE * _in_str ,yyscan_t yyscanner ); @@ -353,10 +351,10 @@ void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_versio exit(1); } - g_ptxinfo_filename = strdup(ptxas_filename.c_str()); FILE *ptxinfo_in; - ptxinfo_in = fopen(g_ptxinfo_filename,"r"); ptxinfo_data ptxinfo; + ptxinfo.g_ptxinfo_filename = strdup(ptxas_filename.c_str()); + ptxinfo_in = fopen(ptxinfo.g_ptxinfo_filename,"r"); ptxinfo_lex_init(&(ptxinfo.scanner)); ptxinfo_set_in(ptxinfo_in, ptxinfo.scanner); ptxinfo_parse(ptxinfo.scanner, &ptxinfo); @@ -422,10 +420,10 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num if( result != 0 ) { // 65280 = duplicate errors if (result == 65280) { + ptxinfo_data ptxinfo; FILE *ptxinfo_in; ptxinfo_in = fopen(tempfile_ptxinfo,"r"); - g_ptxinfo_filename = tempfile_ptxinfo; - ptxinfo_data ptxinfo; + ptxinfo.g_ptxinfo_filename = tempfile_ptxinfo; ptxinfo_lex_init(&(ptxinfo.scanner)); ptxinfo_set_in(ptxinfo_in, ptxinfo.scanner); ptxinfo_parse(ptxinfo.scanner, &ptxinfo); @@ -509,14 +507,14 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num } } + ptxinfo_data ptxinfo; if(no_of_ptx>0) - g_ptxinfo_filename = final_tempfile_ptxinfo; + ptxinfo.g_ptxinfo_filename = final_tempfile_ptxinfo; else - g_ptxinfo_filename = tempfile_ptxinfo; + ptxinfo.g_ptxinfo_filename = tempfile_ptxinfo; FILE *ptxinfo_in; - ptxinfo_in = fopen(g_ptxinfo_filename,"r"); + ptxinfo_in = fopen(ptxinfo.g_ptxinfo_filename,"r"); - ptxinfo_data ptxinfo; ptxinfo_lex_init(&(ptxinfo.scanner)); ptxinfo_set_in(ptxinfo_in, ptxinfo.scanner); ptxinfo_parse(ptxinfo.scanner, &ptxinfo); diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index 2af611a..c3ce888 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -36,6 +36,8 @@ class ptxinfo_data{ yyscan_t scanner; char linebuf[LINEBUF_SIZE]; unsigned col; + const char *g_ptxinfo_filename; + }; diff --git a/src/cuda-sim/ptxinfo.l b/src/cuda-sim/ptxinfo.l index 92f7a30..b0ada09 100644 --- a/src/cuda-sim/ptxinfo.l +++ b/src/cuda-sim/ptxinfo.l @@ -89,7 +89,6 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. extern int g_ptxinfo_error_detected; extern const char *g_filename; -extern const char *g_ptxinfo_filename; int ptxinfo_error(yyscan_t yyscanner, ptxinfo_data* ptxinfo, const char* msg) { @@ -99,7 +98,7 @@ int ptxinfo_error(yyscan_t yyscanner, ptxinfo_data* ptxinfo, const char* msg) fflush(stdout); printf("GPGPU-Sim: ERROR while parsing output of ptxas (used to capture resource usage information)\n"); if( msg != NULL ) - printf("GPGPU-Sim: %s (%s:%u) Syntax error:\n\n", g_filename, g_ptxinfo_filename, yylineno ); + printf("GPGPU-Sim: %s (%s:%u) Syntax error:\n\n", g_filename, ptxinfo->g_ptxinfo_filename, yylineno ); printf(" %s\n", ptxinfo->linebuf ); printf(" "); for( i=0; i < ptxinfo->col-1; i++ ) { -- cgit v1.3 From 7d02cbb061485db38ed8e5f6bf06c9b2fa40eed2 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 12 Jun 2019 13:19:08 -0400 Subject: Integrate ptxinfo into gpgpu_context Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 2 +- libcuda/gpgpu_context.h | 5 +++++ libopencl/opencl_runtime_api.cc | 2 +- src/cuda-sim/ptx_loader.cc | 43 +++++++++++++++++++---------------------- src/cuda-sim/ptx_loader.h | 2 -- 5 files changed, 27 insertions(+), 27 deletions(-) (limited to 'src/cuda-sim/ptx_loader.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 14e3329..0a6eabb 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -640,7 +640,7 @@ void** cudaRegisterFatBinaryInternal( void *fatCubin, gpgpu_context* gpgpu_ctx = } else { symtab=ctx->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 ); + ctx->gpgpu_ptxinfo_load_from_string( ptx, source_num, max_capability, context->no_of_ptx ); } source_num++; ctx->api->load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 16626eb..d3db1ad 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -1,19 +1,24 @@ #ifndef __gpgpu_context_h__ #define __gpgpu_context_h__ #include "cuda_api_object.h" +#include "../src/cuda-sim/ptx_loader.h" class gpgpu_context { public: gpgpu_context() { api = new cuda_runtime_api(); + ptxinfo = new ptxinfo_data(); } // global list // objects pointers for each file cuda_runtime_api* api; + ptxinfo_data* ptxinfo; // member function list void cuobjdumpParseBinary(unsigned int handle); class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); class symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename ); + void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version); + void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version=20, int no_of_ptx=0 ); }; gpgpu_context* GPGPU_Context(); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 0ec635e..50a02fa 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -580,7 +580,7 @@ void _cl_program::Build(const char *options) } info.m_asm = tmp; info.m_symtab = ctx->gpgpu_ptx_sim_load_ptx_from_string( tmp, source_num ); - gpgpu_ptxinfo_load_from_string( tmp, source_num ); + ctx->gpgpu_ptxinfo_load_from_string( tmp, source_num ); free(tmp); } printf("GPGPU-Sim OpenCL API: finished compiling OpenCL kernels.\n"); diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 9857741..2a6d930 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -332,7 +332,7 @@ char* get_app_binary_name(){ return self_exe_path; } -void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version) +void gpgpu_context::gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version) { std::string ptxas_filename(std::string(filename) + "as"); char buff[1024], extra_flags[1024]; @@ -352,17 +352,16 @@ void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_versio } FILE *ptxinfo_in; - ptxinfo_data ptxinfo; - ptxinfo.g_ptxinfo_filename = strdup(ptxas_filename.c_str()); - ptxinfo_in = fopen(ptxinfo.g_ptxinfo_filename,"r"); - ptxinfo_lex_init(&(ptxinfo.scanner)); - ptxinfo_set_in(ptxinfo_in, ptxinfo.scanner); - ptxinfo_parse(ptxinfo.scanner, &ptxinfo); - ptxinfo_lex_destroy(ptxinfo.scanner); + ptxinfo->g_ptxinfo_filename = strdup(ptxas_filename.c_str()); + ptxinfo_in = fopen(ptxinfo->g_ptxinfo_filename,"r"); + ptxinfo_lex_init(&(ptxinfo->scanner)); + ptxinfo_set_in(ptxinfo_in, ptxinfo->scanner); + ptxinfo_parse(ptxinfo->scanner, ptxinfo); + ptxinfo_lex_destroy(ptxinfo->scanner); fclose(ptxinfo_in); } -void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version, int no_of_ptx ) +void gpgpu_context::gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version, int no_of_ptx ) { //do ptxas for individual files instead of one big embedded ptx. This prevents the duplicate defs and declarations. char ptx_file[1000]; @@ -420,14 +419,13 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num if( result != 0 ) { // 65280 = duplicate errors if (result == 65280) { - ptxinfo_data ptxinfo; FILE *ptxinfo_in; ptxinfo_in = fopen(tempfile_ptxinfo,"r"); - ptxinfo.g_ptxinfo_filename = tempfile_ptxinfo; - ptxinfo_lex_init(&(ptxinfo.scanner)); - ptxinfo_set_in(ptxinfo_in, ptxinfo.scanner); - ptxinfo_parse(ptxinfo.scanner, &ptxinfo); - ptxinfo_lex_destroy(ptxinfo.scanner); + ptxinfo->g_ptxinfo_filename = tempfile_ptxinfo; + ptxinfo_lex_init(&(ptxinfo->scanner)); + ptxinfo_set_in(ptxinfo_in, ptxinfo->scanner); + ptxinfo_parse(ptxinfo->scanner, ptxinfo); + ptxinfo_lex_destroy(ptxinfo->scanner); fclose(ptxinfo_in); fix_duplicate_errors(fname2); @@ -507,18 +505,17 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num } } - ptxinfo_data ptxinfo; if(no_of_ptx>0) - ptxinfo.g_ptxinfo_filename = final_tempfile_ptxinfo; + ptxinfo->g_ptxinfo_filename = final_tempfile_ptxinfo; else - ptxinfo.g_ptxinfo_filename = tempfile_ptxinfo; + ptxinfo->g_ptxinfo_filename = tempfile_ptxinfo; FILE *ptxinfo_in; - ptxinfo_in = fopen(ptxinfo.g_ptxinfo_filename,"r"); + ptxinfo_in = fopen(ptxinfo->g_ptxinfo_filename,"r"); - ptxinfo_lex_init(&(ptxinfo.scanner)); - ptxinfo_set_in(ptxinfo_in, ptxinfo.scanner); - ptxinfo_parse(ptxinfo.scanner, &ptxinfo); - ptxinfo_lex_destroy(ptxinfo.scanner); + ptxinfo_lex_init(&(ptxinfo->scanner)); + ptxinfo_set_in(ptxinfo_in, ptxinfo->scanner); + ptxinfo_parse(ptxinfo->scanner, ptxinfo); + ptxinfo_lex_destroy(ptxinfo->scanner); fclose(ptxinfo_in); snprintf(commandline,1024,"rm -f *info"); diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index c3ce888..3637bff 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -44,8 +44,6 @@ class ptxinfo_data{ extern bool g_override_embedded_ptx; extern int no_of_ptx; //counter to track number of ptx files to be extracted in an application. -void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version=20, int no_of_ptx=0 ); -void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version ); char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str); bool keep_intermediate_files(); -- cgit v1.3 From 632bdaab79b9c1fd51ef8152aa00209f76c03101 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 12 Jun 2019 19:55:35 -0400 Subject: Fix some compiler warning Signed-off-by: Mengchi Zhang --- src/cuda-sim/ptx_loader.h | 4 ++-- src/cuda-sim/ptx_parser.cc | 2 -- src/cuda-sim/ptx_parser.h | 9 ++++++--- 3 files changed, 8 insertions(+), 7 deletions(-) (limited to 'src/cuda-sim/ptx_loader.h') diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index 3637bff..ee09e16 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -29,12 +29,12 @@ #define PTX_LOADER_H_INCLUDED #include -#define LINEBUF_SIZE 1024 +#define PTXINFO_LINEBUF_SIZE 1024 typedef void * yyscan_t; class ptxinfo_data{ public: yyscan_t scanner; - char linebuf[LINEBUF_SIZE]; + char linebuf[PTXINFO_LINEBUF_SIZE]; unsigned col; const char *g_ptxinfo_filename; diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index a01d1e3..2872b84 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -63,9 +63,7 @@ std::map g_sym_name_to_symbol_table; fflush(stdout); \ } -static function_info *g_func_info = NULL; static std::map g_ptx_token_decode; -static operand_info g_return_var; const char *decode_token( int type ) { diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h index c4e0b89..25c01fe 100644 --- a/src/cuda-sim/ptx_parser.h +++ b/src/cuda-sim/ptx_parser.h @@ -58,14 +58,15 @@ class ptx_recognizer { g_max_regs_per_thread = 0; g_global_symbol_table = NULL; g_current_symbol_table = NULL; - symbol *g_last_symbol = NULL; + g_last_symbol = NULL; g_error_detected = 0; g_entry_func_param_index=0; + g_func_info = NULL; } // global list yyscan_t scanner; -#define LINEBUF_SIZE (4*1024) - char linebuf[LINEBUF_SIZE]; +#define PTX_LINEBUF_SIZE (4*1024) + char linebuf[PTX_LINEBUF_SIZE]; unsigned col; int g_size; char *g_add_identifier_cached__identifier; @@ -100,6 +101,8 @@ class ptx_recognizer { std::list g_instructions; int g_error_detected; unsigned g_entry_func_param_index; + function_info *g_func_info; + operand_info g_return_var; // member function list void init_directive_state(); -- cgit v1.3 From 3600ed5c59adafe40840524d19f622aa25e60dd6 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 2 Jul 2019 00:22:15 -0400 Subject: Move ptxinfo_addinfo Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 2 +- src/cuda-sim/ptx_loader.h | 2 +- src/cuda-sim/ptxinfo.y | 5 ++--- 3 files changed, 4 insertions(+), 5 deletions(-) (limited to 'src/cuda-sim/ptx_loader.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index a34727f..eb645ce 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -274,7 +274,7 @@ gpgpu_context* GPGPU_Context() return gpgpu_ctx; } - void ptxinfo_addinfo() + void ptxinfo_data::ptxinfo_addinfo() { if(!get_ptxinfo_kname()){ /* This info is not per kernel (since CUDA 5.0 some info (e.g. gmem, and cmem) is added at the beginning for the whole binary ) */ diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index ee09e16..77f27c8 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -37,7 +37,7 @@ class ptxinfo_data{ char linebuf[PTXINFO_LINEBUF_SIZE]; unsigned col; const char *g_ptxinfo_filename; - + void ptxinfo_addinfo(); }; diff --git a/src/cuda-sim/ptxinfo.y b/src/cuda-sim/ptxinfo.y index 00c81e0..b303958 100644 --- a/src/cuda-sim/ptxinfo.y +++ b/src/cuda-sim/ptxinfo.y @@ -29,7 +29,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. %{ typedef void * yyscan_t; -class ptxinfo_data; +#include "ptx_loader.h" %} %define api.pure full @@ -79,7 +79,6 @@ class ptxinfo_data; static unsigned g_system; int ptxinfo_lex(YYSTYPE * yylval_param, yyscan_t yyscanner, ptxinfo_data* ptxinfo); void yyerror(yyscan_t yyscanner, ptxinfo_data* ptxinfo, const char* msg); - void ptxinfo_addinfo(); void ptxinfo_function(const char *fname ); void ptxinfo_regs( unsigned nregs ); void ptxinfo_lmem( unsigned declared, unsigned system ); @@ -104,7 +103,7 @@ line: HEADER INFO COLON line_info ; line_info: function_name - | function_info { ptxinfo_addinfo(); } + | function_info { ptxinfo->ptxinfo_addinfo(); } | gmem_info ; -- cgit v1.3 From 60c38b0f77378eb111e88f632702d19dc3746cc7 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 3 Jul 2019 14:54:16 -0400 Subject: Remove g_filename Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 5 +++-- src/cuda-sim/ptx.l | 5 ++--- src/cuda-sim/ptx.y | 5 ++--- src/cuda-sim/ptx_loader.h | 5 +++++ src/cuda-sim/ptx_parser.cc | 26 ++++++++++++-------------- src/cuda-sim/ptx_parser.h | 7 +++++-- src/cuda-sim/ptxinfo.l | 4 ++-- src/gpgpu-sim/gpu-sim.h | 1 + 8 files changed, 32 insertions(+), 26 deletions(-) (limited to 'src/cuda-sim/ptx_loader.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 69341be..d819559 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -10,12 +10,13 @@ class gpgpu_context { gpgpu_context() { g_global_allfiles_symbol_table = NULL; api = new cuda_runtime_api(); - ptxinfo = new ptxinfo_data(); - ptx_parser = new ptx_recognizer(); + ptxinfo = new ptxinfo_data(this); + ptx_parser = new ptx_recognizer(this); the_gpgpusim = new GPGPUsim_ctx(this); } // global list symbol_table *g_global_allfiles_symbol_table; + const char *g_filename; // objects pointers for each file cuda_runtime_api* api; ptxinfo_data* ptxinfo; diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index 18952a9..3a2a839 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -40,6 +40,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include "ptx_parser.h" #include "ptx.tab.h" #include +#include "../../libcuda/gpgpu_context.h" #define LINEBUF_SIZE (4*1024) #define TC recognizer->col+=strlen(yytext); @@ -452,8 +453,6 @@ breakaddr TC; yylval->int_value = BREAKADDR_OP; return OPCODE; . TC; ptx_error(yyscanner, recognizer, (const char*)NULL); %% -extern const char *g_filename; - int ptx_error( yyscan_t yyscanner, ptx_recognizer* recognizer, const char *s ) { struct yyguts_t * yyg = (struct yyguts_t*)yyscanner; @@ -461,7 +460,7 @@ int ptx_error( yyscan_t yyscanner, ptx_recognizer* recognizer, const char *s ) recognizer->g_error_detected = 1; fflush(stdout); if( s != NULL ) - printf("%s:%u Syntax error:\n\n", g_filename, yylineno ); + printf("%s:%u Syntax error:\n\n", recognizer->gpgpu_ctx->g_filename, yylineno ); printf(" %s\n", recognizer->linebuf ); printf(" "); for( i=0; i < recognizer->col-1; i++ ) { diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y index 837fbe9..a01c3c6 100644 --- a/src/cuda-sim/ptx.y +++ b/src/cuda-sim/ptx.y @@ -30,6 +30,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. %{ typedef void * yyscan_t; class ptx_recognizer; +#include "../../libcuda/gpgpu_context.h" %} %define api.pure full @@ -628,11 +629,9 @@ address_expression: IDENTIFIER { recognizer->add_address_operand($1,0); } %% -extern const char *g_filename; - void syntax_not_implemented(yyscan_t yyscanner, ptx_recognizer* recognizer) { - printf("Parse error (%s): this syntax is not (yet) implemented:\n",g_filename); + printf("Parse error (%s): this syntax is not (yet) implemented:\n", recognizer->gpgpu_ctx->g_filename); ptx_error(yyscanner, recognizer, NULL); abort(); } diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index 77f27c8..c214b95 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -30,13 +30,18 @@ #include #define PTXINFO_LINEBUF_SIZE 1024 +class gpgpu_context; typedef void * yyscan_t; class ptxinfo_data{ public: + ptxinfo_data(gpgpu_context* ctx) { + gpgpu_ctx = ctx; + } yyscan_t scanner; char linebuf[PTXINFO_LINEBUF_SIZE]; unsigned col; const char *g_ptxinfo_filename; + class gpgpu_context* gpgpu_ctx; void ptxinfo_addinfo(); }; diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 09bc15d..05fc618 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -47,14 +47,12 @@ void ptx_recognizer::set_ptx_warp_size(const struct core_config * warp_size) g_shader_core_config=warp_size; } -const char *g_filename; - // the program intermediate representation... std::map g_sym_name_to_symbol_table; #define PTX_PARSE_DPRINTF(...) \ if( g_debug_ir_generation ) { \ - printf(" %s:%u => ",g_filename,ptx_get_lineno(scanner)); \ + printf(" %s:%u => ",gpgpu_ctx->g_filename,ptx_get_lineno(scanner)); \ printf(" (%s:%u) ", __FILE__, __LINE__); \ printf(__VA_ARGS__); \ printf("\n"); \ @@ -70,7 +68,7 @@ const char *decode_token( int type ) void ptx_recognizer::read_parser_environment_variables() { - g_filename = getenv("PTX_SIM_KERNELFILE"); + gpgpu_ctx->g_filename = getenv("PTX_SIM_KERNELFILE"); char *dbg_level = getenv("PTX_SIM_DEBUG"); if ( dbg_level && strlen(dbg_level) ) { int debug_execution=0; @@ -180,7 +178,7 @@ void ptx_recognizer::add_function_name( const char *name ) if( prior_decl ) { g_func_info->remove_args(); } - g_global_symbol_table->add_function( g_func_info, g_filename, ptx_get_lineno(scanner) ); + g_global_symbol_table->add_function( g_func_info, gpgpu_ctx->g_filename, ptx_get_lineno(scanner) ); } //Jin: handle instruction group for cdp @@ -229,7 +227,7 @@ void ptx_recognizer::parse_error_impl( const char *file, unsigned line, const ch va_end(ap); g_error_detected = 1; - printf("%s:%u: Parse error: %s (%s:%u)\n\n", g_filename, ptx_get_lineno(scanner), buf, file, line); + printf("%s:%u: Parse error: %s (%s:%u)\n\n", gpgpu_ctx->g_filename, ptx_get_lineno(scanner), buf, file, line); ptx_error(scanner, NULL); abort(); exit(1); @@ -284,12 +282,12 @@ void ptx_recognizer::add_instruction() g_wmma_options, g_scalar_type, g_space_spec, - g_filename, + gpgpu_ctx->g_filename, ptx_get_lineno(scanner), linebuf, g_shader_core_config ); g_instructions.push_back(i); - g_inst_lookup[g_filename][ptx_get_lineno(scanner)] = i; + g_inst_lookup[gpgpu_ctx->g_filename][ptx_get_lineno(scanner)] = i; init_instruction_state(); } @@ -391,7 +389,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi default: break; } - g_last_symbol = g_current_symbol_table->add_variable(identifier,type,num_bits/8,g_filename,ptx_get_lineno(scanner)); + g_last_symbol = g_current_symbol_table->add_variable(identifier,type,num_bits/8,gpgpu_ctx->g_filename,ptx_get_lineno(scanner)); switch ( ti.get_memory_space().get_type() ) { case reg_space: { regnum = g_current_symbol_table->next_reg_num(); @@ -654,7 +652,7 @@ void ptx_recognizer::add_label( const char *identifier ) if ( s != NULL ) { g_label = s; } else { - g_label = g_current_symbol_table->add_variable(identifier,NULL,0,g_filename,ptx_get_lineno(scanner)); + g_label = g_current_symbol_table->add_variable(identifier,NULL,0,gpgpu_ctx->g_filename,ptx_get_lineno(scanner)); } } @@ -912,7 +910,7 @@ void ptx_recognizer::add_scalar_operand( const char *identifier ) if ( s == NULL ) { if ( g_opcode == BRA_OP || g_opcode == CALLP_OP) { // forward branch target... - s = g_current_symbol_table->add_variable(identifier,NULL,0,g_filename,ptx_get_lineno(scanner)); + s = g_current_symbol_table->add_variable(identifier,NULL,0,gpgpu_ctx->g_filename,ptx_get_lineno(scanner)); } else { std::string msg = std::string("operand \"") + identifier + "\" has no declaration."; parse_error( msg.c_str() ); @@ -926,7 +924,7 @@ void ptx_recognizer::add_neg_pred_operand( const char *identifier ) PTX_PARSE_DPRINTF("add_neg_pred_operand"); const symbol *s = g_current_symbol_table->lookup(identifier); if ( s == NULL ) { - s = g_current_symbol_table->add_variable(identifier,NULL,1,g_filename,ptx_get_lineno(scanner)); + s = g_current_symbol_table->add_variable(identifier,NULL,1,gpgpu_ctx->g_filename,ptx_get_lineno(scanner)); } operand_info op(s); op.set_neg_pred(); @@ -962,7 +960,7 @@ void ptx_recognizer::add_version_info( float ver, unsigned ext ) void ptx_recognizer::add_file( unsigned num, const char *filename ) { - if( g_filename == NULL ) { + if( gpgpu_ctx->g_filename == NULL ) { char *b = strdup(filename); char *l=b; char *n=b; @@ -978,7 +976,7 @@ void ptx_recognizer::add_file( unsigned num, const char *filename ) char *q = strtok(NULL,"."); if( q && !strcmp(q,"cu") ) { - g_filename = strdup(buf); + gpgpu_ctx->g_filename = strdup(buf); } free( b ); diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h index ec300a1..e09aab1 100644 --- a/src/cuda-sim/ptx_parser.h +++ b/src/cuda-sim/ptx_parser.h @@ -30,17 +30,17 @@ #include "../abstract_hardware_model.h" #include "ptx_ir.h" -extern const char *g_filename; extern int g_error_detected; #ifdef __cplusplus const class ptx_instruction *ptx_instruction_lookup( const char *filename, unsigned linenumber ); #endif +class gpgpu_context; typedef void * yyscan_t; class ptx_recognizer { public: - ptx_recognizer() { + ptx_recognizer( gpgpu_context* ctx ) { scanner = NULL; g_size = -1; g_add_identifier_cached__identifier = NULL; @@ -63,6 +63,7 @@ class ptx_recognizer { g_entry_func_param_index=0; g_func_info = NULL; g_debug_ir_generation=false; + gpgpu_ctx = ctx; } // global list yyscan_t scanner; @@ -107,6 +108,8 @@ class ptx_recognizer { bool g_debug_ir_generation; int g_entry_point; const struct core_config *g_shader_core_config; + // backward pointer + class gpgpu_context* gpgpu_ctx; // member function list void init_directive_state(); diff --git a/src/cuda-sim/ptxinfo.l b/src/cuda-sim/ptxinfo.l index b0ada09..3a152b0 100644 --- a/src/cuda-sim/ptxinfo.l +++ b/src/cuda-sim/ptxinfo.l @@ -39,6 +39,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include "ptx_loader.h" #include "ptxinfo.tab.h" #include +#include "../../libcuda/gpgpu_context.h" #define LINEBUF_SIZE 1024 #define TC if( (yylineno == 1) && (ptxinfo->col + strlen(yytext) < LINEBUF_SIZE) ) { \ @@ -88,7 +89,6 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. %% extern int g_ptxinfo_error_detected; -extern const char *g_filename; int ptxinfo_error(yyscan_t yyscanner, ptxinfo_data* ptxinfo, const char* msg) { @@ -98,7 +98,7 @@ int ptxinfo_error(yyscan_t yyscanner, ptxinfo_data* ptxinfo, const char* msg) fflush(stdout); printf("GPGPU-Sim: ERROR while parsing output of ptxas (used to capture resource usage information)\n"); if( msg != NULL ) - printf("GPGPU-Sim: %s (%s:%u) Syntax error:\n\n", g_filename, ptxinfo->g_ptxinfo_filename, yylineno ); + printf("GPGPU-Sim: %s (%s:%u) Syntax error:\n\n", ptxinfo->gpgpu_ctx->g_filename, ptxinfo->g_ptxinfo_filename, yylineno ); printf(" %s\n", ptxinfo->linebuf ); printf(" "); for( i=0; i < ptxinfo->col-1; i++ ) { diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index b1d2e45..e2c913a 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -506,6 +506,7 @@ private: void gpgpu_debug(); ///// data ///// +// backward pointer class gpgpu_context* gpgpu_ctx; class simt_core_cluster **m_cluster; -- cgit v1.3 From 7e9591a24ffbe645cfaa7a84b33a82edb9fca0bc Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Sun, 7 Jul 2019 23:56:21 -0400 Subject: g_keep_intermediate_files Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 1 + libopencl/opencl_runtime_api.cc | 2 +- src/cuda-sim/ptx_ir.h | 3 --- src/cuda-sim/ptx_loader.cc | 7 +++---- src/cuda-sim/ptx_loader.h | 3 ++- 5 files changed, 7 insertions(+), 9 deletions(-) (limited to 'src/cuda-sim/ptx_loader.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index d819559..a8b60f4 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -32,6 +32,7 @@ class gpgpu_context { class symbol_table* init_parser(const char*); class gpgpu_sim *gpgpu_ptx_sim_init_perf(); struct _cuda_device_id *GPGPUSim_Init(); + void ptx_reg_options(option_parser_t opp); }; gpgpu_context* GPGPU_Context(); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index e91e2e0..03ec80c 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -537,7 +537,7 @@ void _cl_program::Build(const char *options) exit(1); } } - if( !g_keep_intermediate_files ) { + if( !ctx->ptxinfo->g_keep_intermediate_files ) { // clean up files... snprintf(commandline,1024,"rm -f %s", cl_fname ); int result = system(commandline); diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 1af85de..fd869c6 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1575,11 +1575,8 @@ struct textureInfo { extern std::map g_sym_name_to_symbol_table; -extern bool g_keep_intermediate_files; - void gpgpu_ptx_assemble( std::string kname, void *kinfo ); #include "../option_parser.h" -void ptx_reg_options(option_parser_t opp); unsigned ptx_kernel_shmem_size( void *kernel_impl ); unsigned ptx_kernel_nregs( void *kernel_impl ); diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 921d1e6..e0d9a11 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -58,17 +58,16 @@ extern int ptxinfo_lex_destroy(yyscan_t scanner); static bool g_save_embedded_ptx; static int g_occupancy_sm_number; -bool g_keep_intermediate_files; bool m_ptx_save_converted_ptxplus; -bool keep_intermediate_files() {return g_keep_intermediate_files;} +bool ptxinfo_data::keep_intermediate_files() {return g_keep_intermediate_files;} -void ptx_reg_options(option_parser_t opp) +void gpgpu_context::ptx_reg_options(option_parser_t opp) { option_parser_register(opp, "-save_embedded_ptx", OPT_BOOL, &g_save_embedded_ptx, "saves ptx files embedded in binary as .ptx", "0"); - option_parser_register(opp, "-keep", OPT_BOOL, &g_keep_intermediate_files, + option_parser_register(opp, "-keep", OPT_BOOL, &(ptxinfo->g_keep_intermediate_files), "keep intermediate files created by GPGPU-Sim when interfacing with external programs", "0"); option_parser_register(opp, "-gpgpu_ptx_save_converted_ptxplus", OPT_BOOL, diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index c214b95..aa3e2f3 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -42,7 +42,9 @@ class ptxinfo_data{ unsigned col; const char *g_ptxinfo_filename; class gpgpu_context* gpgpu_ctx; + bool g_keep_intermediate_files; void ptxinfo_addinfo(); + bool keep_intermediate_files(); }; @@ -50,6 +52,5 @@ extern bool g_override_embedded_ptx; extern int no_of_ptx; //counter to track number of ptx files to be extracted in an application. char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str); -bool keep_intermediate_files(); #endif -- cgit v1.3 From 6e051c42b8c9c820af50a10025d182a96fdd12a6 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 9 Jul 2019 10:06:31 -0400 Subject: Move m_ptx_save_converted_ptxplus Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 2 +- src/cuda-sim/ptx_loader.cc | 5 ++--- src/cuda-sim/ptx_loader.h | 4 ++-- 3 files changed, 5 insertions(+), 6 deletions(-) (limited to 'src/cuda-sim/ptx_loader.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index eb645ce..43c5e1f 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -2936,7 +2936,7 @@ void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){ if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) { cuobjdumpELFSection* elfsection = api->findELFSection(ptx->getIdentifier()); assert (elfsection!= NULL); - char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus( + char *ptxplus_str = ptxinfo->gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus( ptx->getPTXfilename(), elfsection->getELFfilename(), elfsection->getSASSfilename()); diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index e0d9a11..2bf464c 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -58,7 +58,6 @@ extern int ptxinfo_lex_destroy(yyscan_t scanner); static bool g_save_embedded_ptx; static int g_occupancy_sm_number; -bool m_ptx_save_converted_ptxplus; bool ptxinfo_data::keep_intermediate_files() {return g_keep_intermediate_files;} @@ -71,7 +70,7 @@ void gpgpu_context::ptx_reg_options(option_parser_t opp) "keep intermediate files created by GPGPU-Sim when interfacing with external programs", "0"); option_parser_register(opp, "-gpgpu_ptx_save_converted_ptxplus", OPT_BOOL, - &m_ptx_save_converted_ptxplus, + &(ptxinfo->m_ptx_save_converted_ptxplus), "Saved converted ptxplus to a file", "0"); option_parser_register(opp, "-gpgpu_occupancy_sm_number", OPT_INT32, &g_occupancy_sm_number, @@ -106,7 +105,7 @@ void gpgpu_context::print_ptx_file( const char *p, unsigned source_num, const ch fflush(stdout); } -char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptxfilename, const std::string elffilename, const std::string sassfilename) +char* ptxinfo_data::gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptxfilename, const std::string elffilename, const std::string sassfilename) { printf("GPGPU-Sim PTX: converting EMBEDDED .ptx file to ptxplus \n"); diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index aa3e2f3..decfdc8 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -43,14 +43,14 @@ class ptxinfo_data{ const char *g_ptxinfo_filename; class gpgpu_context* gpgpu_ctx; bool g_keep_intermediate_files; + bool m_ptx_save_converted_ptxplus; void ptxinfo_addinfo(); bool keep_intermediate_files(); + char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str); }; extern bool g_override_embedded_ptx; extern int no_of_ptx; //counter to track number of ptx files to be extracted in an application. -char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str); - #endif -- cgit v1.3 From 5ac9a28382becfdfafba744857f7fb56469440d1 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Thu, 11 Jul 2019 10:40:24 -0400 Subject: Move g_override_embedded_ptx Signed-off-by: Mengchi Zhang --- src/cuda-sim/cuda-sim.h | 2 ++ src/cuda-sim/ptx_loader.cc | 4 ---- src/cuda-sim/ptx_loader.h | 2 -- 3 files changed, 2 insertions(+), 6 deletions(-) (limited to 'src/cuda-sim/ptx_loader.h') diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index e4d34fe..16eca19 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -132,6 +132,7 @@ class cuda_sim { g_inst_op_classification_stat= NULL; g_assemble_code_next_pc=0; g_debug_thread_uid = 0; + g_override_embedded_ptx = false; gpgpu_ctx = ctx; } //global variables @@ -166,6 +167,7 @@ class cuda_sim { unsigned cdp_latency[5]; unsigned g_assemble_code_next_pc; int g_debug_thread_uid; + bool g_override_embedded_ptx; // backward pointer class gpgpu_context* gpgpu_ctx; //global functions diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 891828a..6e36a62 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -35,10 +35,6 @@ #include #include "../../libcuda/gpgpu_context.h" -/// globals - -bool g_override_embedded_ptx = false; - /// extern prototypes extern int ptx_error( yyscan_t yyscanner, const char *s ); diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index decfdc8..5a468da 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -49,8 +49,6 @@ class ptxinfo_data{ char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str); }; - -extern bool g_override_embedded_ptx; extern int no_of_ptx; //counter to track number of ptx files to be extracted in an application. #endif -- cgit v1.3 From bfd493e03d52def57152bc9195c88eef8d0a2580 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Thu, 11 Jul 2019 10:45:22 -0400 Subject: Remove standalone extern no_of_ptx Signed-off-by: Mengchi Zhang --- src/cuda-sim/ptx_loader.h | 2 -- 1 file changed, 2 deletions(-) (limited to 'src/cuda-sim/ptx_loader.h') diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index 5a468da..d8f1cbc 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -49,6 +49,4 @@ class ptxinfo_data{ char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str); }; -extern int no_of_ptx; //counter to track number of ptx files to be extracted in an application. - #endif -- cgit v1.3