From dab24b95a31bf1401e42237ae2b46ab3df22058c Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 22 Apr 2019 23:27:52 -0400 Subject: Move ptxinfo to reentrant Signed-off-by: Mengchi Zhang --- src/cuda-sim/ptx_loader.cc | 34 +++++++++++++++++++++++++++------- 1 file changed, 27 insertions(+), 7 deletions(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 8c6f361..da7c32d 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -50,9 +50,12 @@ extern int ptx__scan_string(const char*); extern std::map get_duplicate(); const char *g_ptxinfo_filename; -extern int ptxinfo_parse(); -extern int ptxinfo_debug; -extern FILE *ptxinfo_in; + +typedef void * yyscan_t; +extern int ptxinfo_lex_init(yyscan_t* scanner); +extern void ptxinfo_set_in (FILE * _in_str ,yyscan_t yyscanner ); +extern int ptxinfo_parse(yyscan_t scanner); +extern int ptxinfo_lex_destroy(yyscan_t scanner); static bool g_save_embedded_ptx; static int g_occupancy_sm_number; @@ -347,8 +350,13 @@ void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_versio } g_ptxinfo_filename = strdup(ptxas_filename.c_str()); + FILE *ptxinfo_in; ptxinfo_in = fopen(g_ptxinfo_filename,"r"); - ptxinfo_parse(); + yyscan_t scanner; + ptxinfo_lex_init(&scanner); + ptxinfo_set_in(ptxinfo_in, scanner); + ptxinfo_parse(scanner); + ptxinfo_lex_destroy(scanner); fclose(ptxinfo_in); } @@ -410,9 +418,15 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num if( result != 0 ) { // 65280 = duplicate errors if (result == 65280) { + FILE *ptxinfo_in; ptxinfo_in = fopen(tempfile_ptxinfo,"r"); g_ptxinfo_filename = tempfile_ptxinfo; - ptxinfo_parse(); + yyscan_t scanner; + ptxinfo_lex_init(&scanner); + ptxinfo_set_in(ptxinfo_in, scanner); + ptxinfo_parse(scanner); + ptxinfo_lex_destroy(scanner); + fclose(ptxinfo_in); fix_duplicate_errors(fname2); snprintf(commandline,1024,"$CUDA_INSTALL_PATH/bin/ptxas %s -v %s --output-file /dev/null 2> %s", @@ -495,9 +509,15 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num g_ptxinfo_filename = final_tempfile_ptxinfo; else g_ptxinfo_filename = tempfile_ptxinfo; + FILE *ptxinfo_in; ptxinfo_in = fopen(g_ptxinfo_filename,"r"); - ptxinfo_parse(); + yyscan_t scanner; + ptxinfo_lex_init(&scanner); + ptxinfo_set_in(ptxinfo_in, scanner); + ptxinfo_parse(scanner); + ptxinfo_lex_destroy(scanner); + fclose(ptxinfo_in); snprintf(commandline,1024,"rm -f *info"); if( system(commandline) != 0 ) { @@ -515,4 +535,4 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num exit(1); } } -} \ No newline at end of file +} -- cgit v1.3 From 434925ba00c0ad680d1fe75b68dcb46d8fe44e24 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Sat, 18 May 2019 21:36:11 -0400 Subject: Fix parser for loader Signed-off-by: Mengchi Zhang --- src/cuda-sim/ptx_loader.cc | 15 +++++++++++---- 1 file changed, 11 insertions(+), 4 deletions(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index da7c32d..7eb18ee 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -44,8 +44,12 @@ bool g_override_embedded_ptx = false; /// extern prototypes -extern int ptx_parse(); -extern int ptx__scan_string(const char*); +extern int ptx_error( yyscan_t yyscanner, const char *s ); +extern int ptx_lex_init(yyscan_t* scanner); +extern void ptx_set_in(FILE * _in_str ,yyscan_t yyscanner ); +extern int ptx_parse(yyscan_t scanner, ptx_recognizer* recognizer); +extern int ptx_lex_destroy(yyscan_t scanner); +extern int ptx__scan_string(const char*, yyscan_t scanner); extern std::map get_duplicate(); @@ -171,8 +175,10 @@ symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source fclose(fp); } symbol_table *symtab=init_parser(buf); - ptx__scan_string(p); - int errors = ptx_parse (); + ptx_recognizer recognizer; + ptx_lex_init(&(recognizer.scanner)); + ptx__scan_string(p, recognizer.scanner); + int errors = ptx_parse (recognizer.scanner, &recognizer); if ( errors ) { char fname[1024]; snprintf(fname,1024,"_ptx_errors_XXXXXX"); @@ -185,6 +191,7 @@ symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source abort(); exit(40); } + ptx_lex_destroy(recognizer.scanner); if ( g_debug_execution >= 100 ) print_ptx_file(p,source_num,buf); -- cgit v1.3 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.cc | 32 ++++++++++++++++---------------- src/cuda-sim/ptx_loader.h | 10 ++++++++++ src/cuda-sim/ptxinfo.l | 19 +++++++++++-------- src/cuda-sim/ptxinfo.y | 10 +++++----- 4 files changed, 42 insertions(+), 29 deletions(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 7eb18ee..735ff84 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -58,7 +58,7 @@ 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 ); -extern int ptxinfo_parse(yyscan_t scanner); +extern int ptxinfo_parse(yyscan_t scanner, ptxinfo_data* ptxinfo); extern int ptxinfo_lex_destroy(yyscan_t scanner); static bool g_save_embedded_ptx; @@ -359,11 +359,11 @@ void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_versio g_ptxinfo_filename = strdup(ptxas_filename.c_str()); FILE *ptxinfo_in; ptxinfo_in = fopen(g_ptxinfo_filename,"r"); - yyscan_t scanner; - ptxinfo_lex_init(&scanner); - ptxinfo_set_in(ptxinfo_in, scanner); - ptxinfo_parse(scanner); - ptxinfo_lex_destroy(scanner); + ptxinfo_data 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); } @@ -428,11 +428,11 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num FILE *ptxinfo_in; ptxinfo_in = fopen(tempfile_ptxinfo,"r"); g_ptxinfo_filename = tempfile_ptxinfo; - yyscan_t scanner; - ptxinfo_lex_init(&scanner); - ptxinfo_set_in(ptxinfo_in, scanner); - ptxinfo_parse(scanner); - ptxinfo_lex_destroy(scanner); + ptxinfo_data 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); @@ -519,11 +519,11 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num FILE *ptxinfo_in; ptxinfo_in = fopen(g_ptxinfo_filename,"r"); - yyscan_t scanner; - ptxinfo_lex_init(&scanner); - ptxinfo_set_in(ptxinfo_in, scanner); - ptxinfo_parse(scanner); - ptxinfo_lex_destroy(scanner); + ptxinfo_data 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); snprintf(commandline,1024,"rm -f *info"); 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. diff --git a/src/cuda-sim/ptxinfo.l b/src/cuda-sim/ptxinfo.l index a190e6d..92f7a30 100644 --- a/src/cuda-sim/ptxinfo.l +++ b/src/cuda-sim/ptxinfo.l @@ -36,14 +36,17 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. %option reentrant %{ +#include "ptx_loader.h" #include "ptxinfo.tab.h" #include #define LINEBUF_SIZE 1024 -#define TC if( (yylineno == 1) && (yylval->col + strlen(yytext) < LINEBUF_SIZE) ) { \ - strncpy(yylval->linebuf+yylval->col,yytext,strlen(yytext)); \ +#define TC if( (yylineno == 1) && (ptxinfo->col + strlen(yytext) < LINEBUF_SIZE) ) { \ + strncpy(ptxinfo->linebuf+ptxinfo->col,yytext,strlen(yytext)); \ } \ - yylval->col+=strlen(yytext); + ptxinfo->col+=strlen(yytext); +#define YY_DECL int ptxinfo_lex \ + (YYSTYPE * yylval_param , yyscan_t yyscanner, ptxinfo_data* ptxinfo) %} %% @@ -80,7 +83,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. " " TC; "\t" TC; -\n.* yylval->col=0; strncpy(yylval->linebuf, yytext + 1, 1024); yyless( 1 ); +\n.* ptxinfo->col=0; strncpy(ptxinfo->linebuf, yytext + 1, 1024); yyless( 1 ); %% @@ -88,7 +91,7 @@ extern int g_ptxinfo_error_detected; extern const char *g_filename; extern const char *g_ptxinfo_filename; -int ptxinfo_error(yyscan_t yyscanner, const char* msg) +int ptxinfo_error(yyscan_t yyscanner, ptxinfo_data* ptxinfo, const char* msg) { struct yyguts_t * yyg = (struct yyguts_t*)yyscanner; int i; @@ -97,10 +100,10 @@ int ptxinfo_error(yyscan_t yyscanner, const char* msg) 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(" %s\n", yylval->linebuf ); + printf(" %s\n", ptxinfo->linebuf ); printf(" "); - for( i=0; i < yylval->col-1; i++ ) { - if( yylval->linebuf[i] == '\t' ) printf("\t"); + for( i=0; i < ptxinfo->col-1; i++ ) { + if( ptxinfo->linebuf[i] == '\t' ) printf("\t"); else printf(" "); } diff --git a/src/cuda-sim/ptxinfo.y b/src/cuda-sim/ptxinfo.y index eed304c..00c81e0 100644 --- a/src/cuda-sim/ptxinfo.y +++ b/src/cuda-sim/ptxinfo.y @@ -29,18 +29,18 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. %{ typedef void * yyscan_t; +class ptxinfo_data; %} %define api.pure full %parse-param {yyscan_t scanner} +%parse-param {ptxinfo_data* ptxinfo} %lex-param {yyscan_t scanner} +%lex-param {ptxinfo_data* ptxinfo} %union { -#define LINEBUF_SIZE 1024 int int_value; char * string_value; - char linebuf[LINEBUF_SIZE]; - unsigned col; } %token INT_OPERAND @@ -77,8 +77,8 @@ typedef void * yyscan_t; static unsigned g_declared; static unsigned g_system; - int ptxinfo_lex(YYSTYPE * yylval_param, yyscan_t yyscanner); - void yyerror(yyscan_t yyscanner, const char* msg); + 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 ); -- 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.cc') 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.cc') 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 f257d5fdec428b6627426b0231a5afe977f4ce38 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 12 Jun 2019 10:27:54 -0400 Subject: Remove unused vars Signed-off-by: Mengchi Zhang --- src/cuda-sim/ptx_loader.cc | 4 ---- 1 file changed, 4 deletions(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index d7e9b71..2d931f1 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -37,10 +37,6 @@ /// globals -memory_space *g_global_mem; -memory_space *g_tex_mem; -memory_space *g_surf_mem; -memory_space *g_param_mem; bool g_override_embedded_ptx = false; /// extern prototypes -- 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.cc') 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.cc') 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 7ac6034a99b52d09db7ef07bc008b8f6b039f929 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 12 Jun 2019 18:03:35 -0400 Subject: Move print_ptx_file Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 1 + src/cuda-sim/ptx_loader.cc | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index d3db1ad..dbc3140 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -19,6 +19,7 @@ class gpgpu_context { 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 ); + void print_ptx_file( const char *p, unsigned source_num, const char *filename ); }; gpgpu_context* GPGPU_Context(); diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 2a6d930..38597e0 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -81,7 +81,7 @@ void ptx_reg_options(option_parser_t opp) "0"); } -void print_ptx_file( const char *p, unsigned source_num, const char *filename ) +void gpgpu_context::print_ptx_file( const char *p, unsigned source_num, const char *filename ) { printf("\nGPGPU-Sim PTX: file _%u.ptx contents:\n\n", source_num ); char *s = strdup(p); -- cgit v1.3 From c146d3707216f235a1dde3bd978239fd55b111ed Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 12 Jun 2019 18:17:45 -0400 Subject: Move all ptx_parser inside gpgpu_context Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 4 ++++ src/cuda-sim/ptx_loader.cc | 9 ++++----- src/cuda-sim/ptx_parser.cc | 18 +++++++++--------- src/cuda-sim/ptx_parser.h | 1 - 4 files changed, 17 insertions(+), 15 deletions(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index dbc3140..b586e47 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -2,17 +2,20 @@ #define __gpgpu_context_h__ #include "cuda_api_object.h" #include "../src/cuda-sim/ptx_loader.h" +#include "../src/cuda-sim/ptx_parser.h" class gpgpu_context { public: gpgpu_context() { api = new cuda_runtime_api(); ptxinfo = new ptxinfo_data(); + ptx_parser = new ptx_recognizer(); } // global list // objects pointers for each file cuda_runtime_api* api; ptxinfo_data* ptxinfo; + ptx_recognizer* ptx_parser; // member function list void cuobjdumpParseBinary(unsigned int handle); class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); @@ -20,6 +23,7 @@ class gpgpu_context { 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 ); void print_ptx_file( const char *p, unsigned source_num, const char *filename ); + class symbol_table* init_parser(const char*); }; gpgpu_context* GPGPU_Context(); diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 38597e0..921d1e6 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -170,10 +170,9 @@ symbol_table *gpgpu_context::gpgpu_ptx_sim_load_ptx_from_string( const char *p, fclose(fp); } symbol_table *symtab=init_parser(buf); - ptx_recognizer recognizer; - ptx_lex_init(&(recognizer.scanner)); - ptx__scan_string(p, recognizer.scanner); - int errors = ptx_parse (recognizer.scanner, &recognizer); + ptx_lex_init(&(ptx_parser->scanner)); + ptx__scan_string(p, ptx_parser->scanner); + int errors = ptx_parse (ptx_parser->scanner, ptx_parser); if ( errors ) { char fname[1024]; snprintf(fname,1024,"_ptx_errors_XXXXXX"); @@ -186,7 +185,7 @@ symbol_table *gpgpu_context::gpgpu_ptx_sim_load_ptx_from_string( const char *p, abort(); exit(40); } - ptx_lex_destroy(recognizer.scanner); + ptx_lex_destroy(ptx_parser->scanner); if ( g_debug_execution >= 100 ) print_ptx_file(p,source_num,buf); diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 897beaf..7080ee9 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -27,6 +27,7 @@ #include "ptx_parser.h" #include "ptx_ir.h" +#include "../../libcuda/gpgpu_context.h" typedef void * yyscan_t; #include "ptx.tab.h" @@ -122,7 +123,7 @@ void ptx_recognizer::init_instruction_state() init_directive_state(); } -symbol_table *init_parser( const char *ptx_filename ) +symbol_table * gpgpu_context::init_parser( const char *ptx_filename ) { g_filename = strdup(ptx_filename); if (g_global_allfiles_symbol_table == NULL) { @@ -151,17 +152,16 @@ symbol_table *init_parser( const char *ptx_filename ) g_ptx_token_decode[generic_space] = "generic_space"; g_ptx_token_decode[instruction_space] = "instruction_space"; - ptx_recognizer recognizer; - ptx_lex_init(&(recognizer.scanner)); - recognizer.init_directive_state(); - recognizer.init_instruction_state(); + ptx_lex_init(&(ptx_parser->scanner)); + ptx_parser->init_directive_state(); + ptx_parser->init_instruction_state(); FILE *ptx_in; ptx_in = fopen(ptx_filename, "r"); - ptx_set_in(ptx_in, recognizer.scanner); - ptx_parse(recognizer.scanner, &recognizer); - ptx_in = ptx_get_in(recognizer.scanner); - ptx_lex_destroy(recognizer.scanner); + ptx_set_in(ptx_in, ptx_parser->scanner); + ptx_parse(ptx_parser->scanner, ptx_parser); + ptx_in = ptx_get_in(ptx_parser->scanner); + ptx_lex_destroy(ptx_parser->scanner); fclose(ptx_in); return g_global_symbol_table; } diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h index c52fe20..5293c2e 100644 --- a/src/cuda-sim/ptx_parser.h +++ b/src/cuda-sim/ptx_parser.h @@ -34,7 +34,6 @@ extern const char *g_filename; extern int g_error_detected; #ifdef __cplusplus -class symbol_table* init_parser(const char*); const class ptx_instruction *ptx_instruction_lookup( const char *filename, unsigned linenumber ); #endif -- 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.cc') 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.cc') 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 e6f475eaa9c27cd601033c37321466010299dabf Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Thu, 11 Jul 2019 01:47:04 -0400 Subject: Move g_inst_lookup Signed-off-by: Mengchi Zhang --- src/cuda-sim/ptx_loader.cc | 2 +- src/cuda-sim/ptx_parser.cc | 3 +-- src/cuda-sim/ptx_parser.h | 6 ++---- 3 files changed, 4 insertions(+), 7 deletions(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 2bf464c..891828a 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -90,7 +90,7 @@ void gpgpu_context::print_ptx_file( const char *p, unsigned source_num, const ch while ( (*u != '\n') && (*u != '\0') ) u++; unsigned last = (*u == '\0'); *u = '\0'; - const ptx_instruction *pI = ptx_instruction_lookup(filename,n); + const ptx_instruction *pI = ptx_parser->ptx_instruction_lookup(filename,n); char pc[64]; if( pI && pI->get_PC() ) snprintf(pc,64,"%4u", pI->get_PC() ); diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 5a94679..8b2b72e 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -254,9 +254,8 @@ void ptx_recognizer::set_return() g_return_var = g_operands.front(); } -std::map > g_inst_lookup; -const ptx_instruction *ptx_instruction_lookup( const char *filename, unsigned linenumber ) +const ptx_instruction *ptx_recognizer::ptx_instruction_lookup( const char *filename, unsigned linenumber ) { std::map >::iterator f=g_inst_lookup.find(filename); if( f == g_inst_lookup.end() ) diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h index ef422e5..7e87ec8 100644 --- a/src/cuda-sim/ptx_parser.h +++ b/src/cuda-sim/ptx_parser.h @@ -31,10 +31,6 @@ #include "../abstract_hardware_model.h" #include "ptx_ir.h" -#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 { @@ -107,6 +103,7 @@ class ptx_recognizer { bool g_debug_ir_generation; int g_entry_point; const struct core_config *g_shader_core_config; + std::map > g_inst_lookup; // backward pointer class gpgpu_context* gpgpu_ctx; @@ -177,6 +174,7 @@ class ptx_recognizer { bool check_for_duplicates( const char *identifier ); void read_parser_environment_variables(); void set_ptx_warp_size(const struct core_config * warp_size); + const class ptx_instruction *ptx_instruction_lookup( const char *filename, unsigned linenumber ); }; -- 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.cc') 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 4671280cfe7252bf875d071e667f57064250a1b7 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Fri, 12 Jul 2019 00:09:29 -0400 Subject: Move g_cdp_enabled Signed-off-by: Mengchi Zhang --- libcuda/cuda_api_object.h | 1 + libcuda/cuda_runtime_api.cc | 11 ++++------- src/abstract_hardware_model.cc | 3 --- src/cuda-sim/cuda_device_runtime.h | 2 ++ src/cuda-sim/ptx_loader.cc | 9 +++------ src/gpgpu-sim/gpu-sim.cc | 3 +-- 6 files changed, 11 insertions(+), 18 deletions(-) (limited to 'src/cuda-sim/ptx_loader.cc') diff --git a/libcuda/cuda_api_object.h b/libcuda/cuda_api_object.h index db5e6a4..51416f2 100644 --- a/libcuda/cuda_api_object.h +++ b/libcuda/cuda_api_object.h @@ -199,6 +199,7 @@ class cuda_runtime_api { std::list mergeSections(); cuobjdumpELFSection* findELFSection(const std::string identifier); cuobjdumpPTXSection* findPTXSection(const std::string identifier); + cuobjdumpPTXSection* findPTXSectionInList(std::list §ionlist, const std::string identifier); 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, diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 59d2a60..10a651a 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -2421,7 +2421,6 @@ __host__ cudaError_t CUDARTAPI cudaGetExportTable(const void **ppExportTable, co //extracts all ptx files from binary and dumps into prog_name.unique_no.sm_<>.ptx files void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context){ - extern bool g_cdp_enabled; char command[1000]; char *pytorch_bin = getenv("PYTORCH_BIN"); std::string app_binary = get_app_binary(); @@ -2442,7 +2441,7 @@ void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context){ printf("WARNING: Failed to execute cuobjdump to get list of ptx files \n"); exit(0); } - if(!g_cdp_enabled) { + if(!gpgpu_ctx->device_runtime->g_cdp_enabled) { //based on the list above, dump ptx files individually. Format of dumped ptx file is prog_name.unique_no.sm_<>.ptx std::ifstream infile(ptx_list_file_name); @@ -2515,7 +2514,6 @@ void cuda_runtime_api::extract_code_using_cuobjdump(){ } // Running cuobjdump using dynamic link to current process // Needs the option '-all' to extract PTX from CDP-enabled binary - extern bool g_cdp_enabled; //dump ptx for all individial ptx files into sepearte files which is later used by ptxas. int result=0; @@ -2530,7 +2528,7 @@ void cuda_runtime_api::extract_code_using_cuobjdump(){ snprintf(fname,1024,"_cuobjdump_complete_output_XXXXXX"); int fd=mkstemp(fname); close(fd); - if(!g_cdp_enabled) + if(!gpgpu_ctx->device_runtime->g_cdp_enabled) snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass %s > %s", app_binary.c_str(), fname); else snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass -all %s > %s", app_binary.c_str(), fname); @@ -2822,7 +2820,7 @@ cuobjdumpELFSection* cuda_runtime_api::findELFSection(const std::string identifi } //! Within the section list, find the PTX section corresponding to a given identifier -cuobjdumpPTXSection* findPTXSectionInList(std::list §ionlist, const std::string identifier){ +cuobjdumpPTXSection* cuda_runtime_api::findPTXSectionInList(std::list §ionlist, const std::string identifier){ std::list::iterator iter; for ( iter = sectionlist.begin(); iter != sectionlist.end(); @@ -2833,8 +2831,7 @@ cuobjdumpPTXSection* findPTXSectionInList(std::list §ionl if(ptxsection->getIdentifier() == identifier) return ptxsection; else { - extern bool g_cdp_enabled; - if(g_cdp_enabled) { + if(gpgpu_ctx->device_runtime->g_cdp_enabled) { printf("Warning: __cudaRegisterFatBinary needs %s, but find PTX section with %s\n", identifier.c_str(), ptxsection->getIdentifier().c_str()); return ptxsection; diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 450087b..fde7874 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -691,9 +691,6 @@ void warp_inst_t::completed( unsigned long long cycle ) const ptx_file_line_stats_add_latency(pc, latency * active_count()); } -//Jin: CDP support -bool g_cdp_enabled; - unsigned kernel_info_t::m_next_uid = 1; /*A snapshot of the texture mappings needs to be stored in the kernel's info as diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h index f37849b..7f7a0ca 100644 --- a/src/cuda-sim/cuda_device_runtime.h +++ b/src/cuda-sim/cuda_device_runtime.h @@ -51,6 +51,8 @@ class cuda_device_runtime { std::list g_cuda_device_launch_op; unsigned g_kernel_launch_latency; unsigned long long g_max_total_param_size; + bool g_cdp_enabled; + // backward pointer class gpgpu_context* gpgpu_ctx; #if (CUDART_VERSION >= 5000) diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 6e36a62..dca3cec 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -330,8 +330,7 @@ void gpgpu_context::gpgpu_ptx_info_load_from_filename( const char *filename, uns std::string ptxas_filename(std::string(filename) + "as"); char buff[1024], extra_flags[1024]; extra_flags[0]=0; - extern bool g_cdp_enabled; - if(!g_cdp_enabled) + if(!device_runtime->g_cdp_enabled) snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version); else snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",sm_version); @@ -398,8 +397,7 @@ void gpgpu_context::gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsi "A register size/SM mismatch may result in occupancy differences." ); exit(1); } - extern bool g_cdp_enabled; - if(!g_cdp_enabled) + if(!device_runtime->g_cdp_enabled) snprintf(extra_flags,1024,"--gpu-name=sm_%u", g_occupancy_sm_number); else snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",g_occupancy_sm_number); @@ -467,8 +465,7 @@ void gpgpu_context::gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsi #if CUDART_VERSION >= 3000 if (sm_version == 0) sm_version = 20; - extern bool g_cdp_enabled; - if(!g_cdp_enabled) + if(!device_runtime->g_cdp_enabled) snprintf(extra_flags,1024,"--gpu-name=sm_%u",sm_version); else snprintf(extra_flags,1024,"--compile-only --gpu-name=sm_%u",sm_version); diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index a3d6a8a..0481259 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -551,9 +551,8 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) option_parser_register(opp, "-gpgpu_kernel_launch_latency", OPT_INT32, &(gpgpu_ctx->device_runtime->g_kernel_launch_latency), "Kernel launch latency in cycles. Default: 0", "0"); - extern bool g_cdp_enabled; option_parser_register(opp, "-gpgpu_cdp_enabled", OPT_BOOL, - &g_cdp_enabled, "Turn on CDP", + &(gpgpu_ctx->device_runtime->g_cdp_enabled), "Turn on CDP", "0"); } -- cgit v1.3