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