From 10156f6728f10fdc33ca8796e9cc6ade306f06ec Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Thu, 16 May 2019 21:08:39 -0400 Subject: Remove unused extern Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 1 - 1 file changed, 1 deletion(-) (limited to 'libcuda/cuda_runtime_api.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 11dcd5a..fc152c2 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -3261,7 +3261,6 @@ extern int ptx__scan_string(const char*); extern FILE *ptx_in; extern int ptxinfo_parse(); -extern int ptxinfo_debug; extern FILE *ptxinfo_in; /// static functions -- cgit v1.3 From 913865e6ac7d0f7c12faeb8430dc5724fdc4be80 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 27 May 2019 17:27:58 -0400 Subject: Move some cuobjdump parser variables Signed-off-by: Mengchi Zhang --- libcuda/Makefile | 4 ++-- libcuda/cuda_runtime_api.cc | 23 ++++++++++++----------- libcuda/cuobjdump.h | 18 ++++++++++++++++++ libcuda/cuobjdump.l | 9 ++++++--- libcuda/cuobjdump.y | 43 ++++++++++++++++++++----------------------- 5 files changed, 58 insertions(+), 39 deletions(-) create mode 100644 libcuda/cuobjdump.h (limited to 'libcuda/cuda_runtime_api.cc') diff --git a/libcuda/Makefile b/libcuda/Makefile index 13932e2..c8ff2e3 100644 --- a/libcuda/Makefile +++ b/libcuda/Makefile @@ -111,10 +111,10 @@ $(OUTPUT_DIR)/%.o: %.cc $(CPP) $(CXXFLAGS) -I./ -I$(OUTPUT_DIR) -I$(CUDA_INSTALL_PATH)/include -c $< -o $@ $(OUTPUT_DIR)/%.o: %.c - $(CC) $(CCFLAGS) -I./ -I$(OUTPUT_DIR) -I$(CUDA_INSTALL_PATH)/include -c $< -o $@ + $(CPP) $(CCFLAGS) -I./ -I$(OUTPUT_DIR) -I$(CUDA_INSTALL_PATH)/include -c $< -o $@ $(OUTPUT_DIR)/%.o: $(OUTPUT_DIR)/%.c - $(CC) $(CCFLAGS) -I./ -I$(OUTPUT_DIR) -I$(CUDA_INSTALL_PATH)/include -c $< -o $@ + $(CPP) $(CCFLAGS) -I./ -I$(OUTPUT_DIR) -I$(CUDA_INSTALL_PATH)/include -c $< -o $@ $(OUTPUT_DIR)/cuobjdump_parser.c: cuobjdump.y $(YACC) $(YFLAGS) -p cuobjdump_ -o$@ $< --file-prefix=$(OUTPUT_DIR)/cuobjdump diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index fc152c2..1d4e870 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1972,9 +1972,10 @@ void setCuobjdumpsassfilename(const char* filename){ (dynamic_cast(cuobjdumpSectionList.front()))->setSASSfilename(filename); } typedef void * yyscan_t; +#include "cuobjdump.h" extern int cuobjdump_lex_init(yyscan_t* scanner); extern void cuobjdump_set_in (FILE * _in_str ,yyscan_t yyscanner ); -extern int cuobjdump_parse(yyscan_t scanner); +extern int cuobjdump_parse(yyscan_t scanner, cuobjdump_parser* parser); extern int cuobjdump_lex_destroy(yyscan_t scanner); //! Return the executable file of the process containing the PTX/SASS code @@ -2187,11 +2188,11 @@ void extract_code_using_cuobjdump(){ FILE *cuobjdump_in; cuobjdump_in = fopen(fname, "r"); - yyscan_t scanner; - cuobjdump_lex_init(&scanner); - cuobjdump_set_in(cuobjdump_in, scanner); - cuobjdump_parse(scanner); - cuobjdump_lex_destroy(scanner); + cuobjdump_parser parser; + cuobjdump_lex_init(&(parser.scanner)); + cuobjdump_set_in(cuobjdump_in, (parser.scanner)); + cuobjdump_parse(parser.scanner, &parser); + cuobjdump_lex_destroy(parser.scanner); fclose(cuobjdump_in); printf("Done parsing!!!\n"); } else { @@ -2241,11 +2242,11 @@ void extract_code_using_cuobjdump(){ std::cout << "Trying to parse " << libcodfn.str() << std::endl; FILE *cuobjdump_in; cuobjdump_in = fopen(libcodfn.str().c_str(), "r"); - yyscan_t scanner; - cuobjdump_lex_init(&scanner); - cuobjdump_set_in(cuobjdump_in, scanner); - cuobjdump_parse(scanner); - cuobjdump_lex_destroy(scanner); + cuobjdump_parser parser; + cuobjdump_lex_init(&(parser.scanner)); + cuobjdump_set_in(cuobjdump_in, (parser.scanner)); + cuobjdump_parse(parser.scanner, &parser); + cuobjdump_lex_destroy(parser.scanner); fclose(cuobjdump_in); std::getline(libsf, line); } diff --git a/libcuda/cuobjdump.h b/libcuda/cuobjdump.h new file mode 100644 index 0000000..61bf806 --- /dev/null +++ b/libcuda/cuobjdump.h @@ -0,0 +1,18 @@ +#ifndef __cuobjdump_h__ +#define __cuobjdump_h__ +class cuobjdump_parser { + + public: + yyscan_t scanner; + int elfserial; + int ptxserial; + FILE *ptxfile; + FILE *elffile; + FILE *sassfile; + char filename [1024]; + cuobjdump_parser() { + int elfserial = 1; + int ptxserial = 1; + } +}; +#endif /* __cuobjdump_h__ */ diff --git a/libcuda/cuobjdump.l b/libcuda/cuobjdump.l index 9359281..26fbb55 100644 --- a/libcuda/cuobjdump.l +++ b/libcuda/cuobjdump.l @@ -30,6 +30,7 @@ %{ #include #include +#include "cuobjdump.h" #include "cuobjdump_parser.h" #define YY_NEVER_INTERACTIVE 1 @@ -38,7 +39,9 @@ #define YYDEBUG 1 -void cuobjdump_error(yyscan_t yyscanner, const char* msg); +void cuobjdump_error(yyscan_t yyscanner, cuobjdump_parser* parser, const char* msg); +#define YY_DECL int cuobjdump_lex \ + (YYSTYPE * yylval_param , yyscan_t yyscanner, cuobjdump_parser* parser) %} %option stack @@ -153,10 +156,10 @@ newlines {newline}+ <> BEGIN(INITIAL);return 0; /*No other rule matched. Throw an error*/ -. cuobjdump_error(yyscanner, "Invalid token"); +. cuobjdump_error(yyscanner, parser, "Invalid token"); %% -void cuobjdump_error(yyscan_t yyscanner, const char* msg) +void cuobjdump_error(yyscan_t yyscanner, cuobjdump_parser* parser, const char* msg) { struct yyguts_t * yyg = (struct yyguts_t*)yyscanner; printf(" %s near \"%s\"",msg, yytext); diff --git a/libcuda/cuobjdump.y b/libcuda/cuobjdump.y index 66cbace..9c0c28d 100644 --- a/libcuda/cuobjdump.y +++ b/libcuda/cuobjdump.y @@ -30,6 +30,7 @@ #include typedef void * yyscan_t; +#include "cuobjdump.h" extern void addCuobjdumpSection(int sectiontype); void setCuobjdumparch(const char* arch); @@ -37,23 +38,19 @@ void setCuobjdumpidentifier(const char* identifier); void setCuobjdumpptxfilename(const char* filename); void setCuobjdumpelffilename(const char* filename); void setCuobjdumpsassfilename(const char* filename); -int elfserial = 1; -int ptxserial = 1; -FILE *ptxfile; -FILE *elffile; -FILE *sassfile; -char filename [1024]; %} %define api.pure full %parse-param {yyscan_t scanner} +%parse-param {cuobjdump_parser* parser} %lex-param {yyscan_t scanner} +%lex-param {cuobjdump_parser* parser} %union { char* string_value; } %{ -int yylex(YYSTYPE * yylval_param, yyscan_t yyscanner); -void yyerror(yyscan_t yyscanner, const char* msg); +int yylex(YYSTYPE * yylval_param, yyscan_t yyscanner, cuobjdump_parser* parser); +void yyerror(yyscan_t yyscanner, cuobjdump_parser* parser, const char* msg); %} %token H_SEPARATOR H_ARCH H_CODEVERSION H_PRODUCER H_HOST H_COMPILESIZE H_IDENTIFIER H_UNKNOWN H_COMPRESSED %token CODEVERSION @@ -79,24 +76,24 @@ emptylines : emptylines NEWLINE section : PTXHEADER { addCuobjdumpSection(0); - snprintf(filename, 1024, "_cuobjdump_%d.ptx", ptxserial++); - ptxfile = fopen(filename, "w"); - setCuobjdumpptxfilename(filename); + snprintf(parser->filename, 1024, "_cuobjdump_%d.ptx", parser->ptxserial++); + parser->ptxfile = fopen(parser->filename, "w"); + setCuobjdumpptxfilename(parser->filename); } headerinfo compressedkeyword identifier ptxcode { - fclose(ptxfile); + fclose(parser->ptxfile); } | ELFHEADER { addCuobjdumpSection(1); - snprintf(filename, 1024, "_cuobjdump_%d.elf", elfserial); - elffile = fopen(filename, "w"); - setCuobjdumpelffilename(filename); + snprintf(parser->filename, 1024, "_cuobjdump_%d.elf", parser->elfserial); + parser->elffile = fopen(parser->filename, "w"); + setCuobjdumpelffilename(parser->filename); } headerinfo compressedkeyword identifier elfcode { - fclose(elffile); - snprintf(filename, 1024, "_cuobjdump_%d.sass", elfserial++); - sassfile = fopen(filename, "w"); - setCuobjdumpsassfilename(filename); + fclose(parser->elffile); + snprintf(parser->filename, 1024, "_cuobjdump_%d.sass", parser->elfserial++); + parser->sassfile = fopen(parser->filename, "w"); + setCuobjdumpsassfilename(parser->filename); } sasscode { - fclose(sassfile); + fclose(parser->sassfile); }; headerinfo : H_SEPARATOR NEWLINE @@ -118,13 +115,13 @@ identifier : H_IDENTIFIER FILENAME emptylines {setCuobjdumpidentifier($2);} compressedkeyword : H_COMPRESSED emptylines | ; -ptxcode : ptxcode PTXLINE {fprintf(ptxfile, "%s", $2);} +ptxcode : ptxcode PTXLINE {fprintf(parser->ptxfile, "%s", $2);} | ; -elfcode : elfcode ELFLINE {fprintf(elffile, "%s", $2);} +elfcode : elfcode ELFLINE {fprintf(parser->elffile, "%s", $2);} | ; -sasscode : sasscode SASSLINE {fprintf(sassfile, "%s", $2);} +sasscode : sasscode SASSLINE {fprintf(parser->sassfile, "%s", $2);} | ; -- cgit v1.3 From 568e3185c58b07ac31fdd59f6c2aa7af7533939e Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 28 May 2019 15:14:52 -0400 Subject: Move some functions to enable C++ Signed-off-by: Mengchi Zhang --- libcuda/Makefile | 4 +- libcuda/cuda_runtime_api.cc | 246 ++++++++++++++++++++++---------------------- 2 files changed, 126 insertions(+), 124 deletions(-) (limited to 'libcuda/cuda_runtime_api.cc') diff --git a/libcuda/Makefile b/libcuda/Makefile index 13932e2..c8ff2e3 100644 --- a/libcuda/Makefile +++ b/libcuda/Makefile @@ -111,10 +111,10 @@ $(OUTPUT_DIR)/%.o: %.cc $(CPP) $(CXXFLAGS) -I./ -I$(OUTPUT_DIR) -I$(CUDA_INSTALL_PATH)/include -c $< -o $@ $(OUTPUT_DIR)/%.o: %.c - $(CC) $(CCFLAGS) -I./ -I$(OUTPUT_DIR) -I$(CUDA_INSTALL_PATH)/include -c $< -o $@ + $(CPP) $(CCFLAGS) -I./ -I$(OUTPUT_DIR) -I$(CUDA_INSTALL_PATH)/include -c $< -o $@ $(OUTPUT_DIR)/%.o: $(OUTPUT_DIR)/%.c - $(CC) $(CCFLAGS) -I./ -I$(OUTPUT_DIR) -I$(CUDA_INSTALL_PATH)/include -c $< -o $@ + $(CPP) $(CCFLAGS) -I./ -I$(OUTPUT_DIR) -I$(CUDA_INSTALL_PATH)/include -c $< -o $@ $(OUTPUT_DIR)/cuobjdump_parser.c: cuobjdump.y $(YACC) $(YFLAGS) -p cuobjdump_ -o$@ $< --file-prefix=$(OUTPUT_DIR)/cuobjdump diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 168c80f..afe47f8 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -483,6 +483,130 @@ event_tracker_t g_timer_events; int g_active_device = 0; //active gpu that runs the code std::list g_cuda_launch_stack; +typedef void * yyscan_t; +#include "cuobjdump.h" +extern int cuobjdump_lex_init(yyscan_t* scanner); +extern void cuobjdump_set_in (FILE * _in_str ,yyscan_t yyscanner ); +extern int cuobjdump_parse(yyscan_t scanner, struct cuobjdump_parser* parser); +extern int cuobjdump_lex_destroy(yyscan_t scanner); + +enum cuobjdumpSectionType { + PTXSECTION=0, + ELFSECTION +}; + + +class cuobjdumpSection { +public: + //Constructor + cuobjdumpSection() { + arch = 0; + identifier = ""; + } + virtual ~cuobjdumpSection() {} + unsigned getArch() {return arch;} + void setArch(unsigned a) {arch = a;} + std::string getIdentifier() {return identifier;} + void setIdentifier(std::string i) {identifier = i;} + virtual void print(){std::cout << "cuobjdump Section: unknown type" << std::endl;} +private: + unsigned arch; + std::string identifier; +}; + +class cuobjdumpELFSection : public cuobjdumpSection +{ +public: + cuobjdumpELFSection() {} + virtual ~cuobjdumpELFSection() { + elffilename = ""; + sassfilename = ""; + } + std::string getELFfilename() {return elffilename;} + void setELFfilename(std::string f) {elffilename = f;} + std::string getSASSfilename() {return sassfilename;} + void setSASSfilename(std::string f) {sassfilename = f;} + virtual void print() { + std::cout << "ELF Section:" << std::endl; + std::cout << "arch: sm_" << getArch() << std::endl; + std::cout << "identifier: " << getIdentifier() << std::endl; + std::cout << "elf filename: " << getELFfilename() << std::endl; + std::cout << "sass filename: " << getSASSfilename() << std::endl; + std::cout << std::endl; + } +private: + std::string elffilename; + std::string sassfilename; +}; + +class cuobjdumpPTXSection : public cuobjdumpSection +{ +public: + cuobjdumpPTXSection(){ + ptxfilename = ""; + } + std::string getPTXfilename() {return ptxfilename;} + void setPTXfilename(std::string f) {ptxfilename = f;} + virtual void print() { + std::cout << "PTX Section:" << std::endl; + std::cout << "arch: sm_" << getArch() << std::endl; + std::cout << "identifier: " << getIdentifier() << std::endl; + std::cout << "ptx filename: " << getPTXfilename() << std::endl; + std::cout << std::endl; + } +private: + std::string ptxfilename; +}; + + +std::list cuobjdumpSectionList; +std::list libSectionList; + +// sectiontype: 0 for ptx, 1 for elf +void addCuobjdumpSection(int sectiontype){ + if (sectiontype) + cuobjdumpSectionList.push_front(new cuobjdumpELFSection()); + else + cuobjdumpSectionList.push_front(new cuobjdumpPTXSection()); + printf("## Adding new section %s\n", sectiontype?"ELF":"PTX"); +} + +void setCuobjdumparch(const char* arch){ + unsigned archnum; + sscanf(arch, "sm_%u", &archnum); + assert (archnum && "cannot have sm_0"); + printf("Adding arch: %s\n", arch); + cuobjdumpSectionList.front()->setArch(archnum); +} + +void setCuobjdumpidentifier(const char* identifier){ + printf("Adding identifier: %s\n", identifier); + cuobjdumpSectionList.front()->setIdentifier(identifier); +} + +void setCuobjdumpptxfilename(const char* filename){ + printf("Adding ptx filename: %s\n", filename); + cuobjdumpSection* x = cuobjdumpSectionList.front(); + if (dynamic_cast(x) == NULL){ + assert (0 && "You shouldn't be trying to add a ptxfilename to an elf section"); + } + (dynamic_cast(x))->setPTXfilename(filename); +} + +void setCuobjdumpelffilename(const char* filename){ + if (dynamic_cast(cuobjdumpSectionList.front()) == NULL){ + assert (0 && "You shouldn't be trying to add a elffilename to an ptx section"); + } + (dynamic_cast(cuobjdumpSectionList.front()))->setELFfilename(filename); +} + +void setCuobjdumpsassfilename(const char* filename){ + if (dynamic_cast(cuobjdumpSectionList.front()) == NULL){ + assert (0 && "You shouldn't be trying to add a sassfilename to an ptx section"); + } + (dynamic_cast(cuobjdumpSectionList.front()))->setSASSfilename(filename); +} + /******************************************************************************* * * * * @@ -1856,128 +1980,6 @@ __host__ cudaError_t CUDARTAPI cudaGetExportTable(const void **ppExportTable, co //#include "../../cuobjdump_to_ptxplus/cuobjdump_parser.h" -enum cuobjdumpSectionType { - PTXSECTION=0, - ELFSECTION -}; - - -class cuobjdumpSection { -public: - //Constructor - cuobjdumpSection() { - arch = 0; - identifier = ""; - } - virtual ~cuobjdumpSection() {} - unsigned getArch() {return arch;} - void setArch(unsigned a) {arch = a;} - std::string getIdentifier() {return identifier;} - void setIdentifier(std::string i) {identifier = i;} - virtual void print(){std::cout << "cuobjdump Section: unknown type" << std::endl;} -private: - unsigned arch; - std::string identifier; -}; - -class cuobjdumpELFSection : public cuobjdumpSection -{ -public: - cuobjdumpELFSection() {} - virtual ~cuobjdumpELFSection() { - elffilename = ""; - sassfilename = ""; - } - std::string getELFfilename() {return elffilename;} - void setELFfilename(std::string f) {elffilename = f;} - std::string getSASSfilename() {return sassfilename;} - void setSASSfilename(std::string f) {sassfilename = f;} - virtual void print() { - std::cout << "ELF Section:" << std::endl; - std::cout << "arch: sm_" << getArch() << std::endl; - std::cout << "identifier: " << getIdentifier() << std::endl; - std::cout << "elf filename: " << getELFfilename() << std::endl; - std::cout << "sass filename: " << getSASSfilename() << std::endl; - std::cout << std::endl; - } -private: - std::string elffilename; - std::string sassfilename; -}; - -class cuobjdumpPTXSection : public cuobjdumpSection -{ -public: - cuobjdumpPTXSection(){ - ptxfilename = ""; - } - std::string getPTXfilename() {return ptxfilename;} - void setPTXfilename(std::string f) {ptxfilename = f;} - virtual void print() { - std::cout << "PTX Section:" << std::endl; - std::cout << "arch: sm_" << getArch() << std::endl; - std::cout << "identifier: " << getIdentifier() << std::endl; - std::cout << "ptx filename: " << getPTXfilename() << std::endl; - std::cout << std::endl; - } -private: - std::string ptxfilename; -}; - -std::list cuobjdumpSectionList; -std::list libSectionList; - -// sectiontype: 0 for ptx, 1 for elf -void addCuobjdumpSection(int sectiontype){ - if (sectiontype) - cuobjdumpSectionList.push_front(new cuobjdumpELFSection()); - else - cuobjdumpSectionList.push_front(new cuobjdumpPTXSection()); - printf("## Adding new section %s\n", sectiontype?"ELF":"PTX"); -} - -void setCuobjdumparch(const char* arch){ - unsigned archnum; - sscanf(arch, "sm_%u", &archnum); - assert (archnum && "cannot have sm_0"); - printf("Adding arch: %s\n", arch); - cuobjdumpSectionList.front()->setArch(archnum); -} - -void setCuobjdumpidentifier(const char* identifier){ - printf("Adding identifier: %s\n", identifier); - cuobjdumpSectionList.front()->setIdentifier(identifier); -} - -void setCuobjdumpptxfilename(const char* filename){ - printf("Adding ptx filename: %s\n", filename); - cuobjdumpSection* x = cuobjdumpSectionList.front(); - if (dynamic_cast(x) == NULL){ - assert (0 && "You shouldn't be trying to add a ptxfilename to an elf section"); - } - (dynamic_cast(x))->setPTXfilename(filename); -} - -void setCuobjdumpelffilename(const char* filename){ - if (dynamic_cast(cuobjdumpSectionList.front()) == NULL){ - assert (0 && "You shouldn't be trying to add a elffilename to an ptx section"); - } - (dynamic_cast(cuobjdumpSectionList.front()))->setELFfilename(filename); -} - -void setCuobjdumpsassfilename(const char* filename){ - if (dynamic_cast(cuobjdumpSectionList.front()) == NULL){ - assert (0 && "You shouldn't be trying to add a sassfilename to an ptx section"); - } - (dynamic_cast(cuobjdumpSectionList.front()))->setSASSfilename(filename); -} -typedef void * yyscan_t; -#include "cuobjdump.h" -extern int cuobjdump_lex_init(yyscan_t* scanner); -extern void cuobjdump_set_in (FILE * _in_str ,yyscan_t yyscanner ); -extern int cuobjdump_parse(yyscan_t scanner, struct cuobjdump_parser* parser); -extern int cuobjdump_lex_destroy(yyscan_t scanner); - //! Return the executable file of the process containing the PTX/SASS code //! //! This Function returns the executable file ran by the process. This -- cgit v1.3 From 25bdc0dd89932f95ace0fc617649a4e041aaadd9 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 29 May 2019 01:28:29 -0400 Subject: Move SectionList to context Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 109 ++++++++++---------------------------------- libcuda/cuobjdump.h | 67 +++++++++++++++++++++++++++ libcuda/cuobjdump.l | 8 ++-- libcuda/cuobjdump.y | 36 ++++++++------- 4 files changed, 113 insertions(+), 107 deletions(-) (limited to 'libcuda/cuda_runtime_api.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index afe47f8..df7ddc7 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -139,6 +139,8 @@ #include "../src/gpgpusim_entrypoint.h" #include "../src/stream_manager.h" #include "../src/abstract_hardware_model.h" +typedef void * yyscan_t; +#include "cuobjdump.h" #include #include @@ -302,6 +304,7 @@ struct CUctx_st { return i->second; } + std::list cuobjdumpSectionList; private: _cuda_device_id *m_gpu; // selected gpu std::map m_code; // fat binary handle => global symbol table @@ -483,11 +486,9 @@ event_tracker_t g_timer_events; int g_active_device = 0; //active gpu that runs the code std::list g_cuda_launch_stack; -typedef void * yyscan_t; -#include "cuobjdump.h" extern int cuobjdump_lex_init(yyscan_t* scanner); extern void cuobjdump_set_in (FILE * _in_str ,yyscan_t yyscanner ); -extern int cuobjdump_parse(yyscan_t scanner, struct cuobjdump_parser* parser); +extern int cuobjdump_parse(yyscan_t scanner, struct cuobjdump_parser* parser, std::list &cuobjdumpSectionList); extern int cuobjdump_lex_destroy(yyscan_t scanner); enum cuobjdumpSectionType { @@ -496,74 +497,10 @@ enum cuobjdumpSectionType { }; -class cuobjdumpSection { -public: - //Constructor - cuobjdumpSection() { - arch = 0; - identifier = ""; - } - virtual ~cuobjdumpSection() {} - unsigned getArch() {return arch;} - void setArch(unsigned a) {arch = a;} - std::string getIdentifier() {return identifier;} - void setIdentifier(std::string i) {identifier = i;} - virtual void print(){std::cout << "cuobjdump Section: unknown type" << std::endl;} -private: - unsigned arch; - std::string identifier; -}; - -class cuobjdumpELFSection : public cuobjdumpSection -{ -public: - cuobjdumpELFSection() {} - virtual ~cuobjdumpELFSection() { - elffilename = ""; - sassfilename = ""; - } - std::string getELFfilename() {return elffilename;} - void setELFfilename(std::string f) {elffilename = f;} - std::string getSASSfilename() {return sassfilename;} - void setSASSfilename(std::string f) {sassfilename = f;} - virtual void print() { - std::cout << "ELF Section:" << std::endl; - std::cout << "arch: sm_" << getArch() << std::endl; - std::cout << "identifier: " << getIdentifier() << std::endl; - std::cout << "elf filename: " << getELFfilename() << std::endl; - std::cout << "sass filename: " << getSASSfilename() << std::endl; - std::cout << std::endl; - } -private: - std::string elffilename; - std::string sassfilename; -}; - -class cuobjdumpPTXSection : public cuobjdumpSection -{ -public: - cuobjdumpPTXSection(){ - ptxfilename = ""; - } - std::string getPTXfilename() {return ptxfilename;} - void setPTXfilename(std::string f) {ptxfilename = f;} - virtual void print() { - std::cout << "PTX Section:" << std::endl; - std::cout << "arch: sm_" << getArch() << std::endl; - std::cout << "identifier: " << getIdentifier() << std::endl; - std::cout << "ptx filename: " << getPTXfilename() << std::endl; - std::cout << std::endl; - } -private: - std::string ptxfilename; -}; - - -std::list cuobjdumpSectionList; std::list libSectionList; // sectiontype: 0 for ptx, 1 for elf -void addCuobjdumpSection(int sectiontype){ +void addCuobjdumpSection(int sectiontype, std::list &cuobjdumpSectionList){ if (sectiontype) cuobjdumpSectionList.push_front(new cuobjdumpELFSection()); else @@ -571,7 +508,7 @@ void addCuobjdumpSection(int sectiontype){ printf("## Adding new section %s\n", sectiontype?"ELF":"PTX"); } -void setCuobjdumparch(const char* arch){ +void setCuobjdumparch(const char* arch, std::list &cuobjdumpSectionList){ unsigned archnum; sscanf(arch, "sm_%u", &archnum); assert (archnum && "cannot have sm_0"); @@ -579,12 +516,12 @@ void setCuobjdumparch(const char* arch){ cuobjdumpSectionList.front()->setArch(archnum); } -void setCuobjdumpidentifier(const char* identifier){ +void setCuobjdumpidentifier(const char* identifier, std::list &cuobjdumpSectionList){ printf("Adding identifier: %s\n", identifier); cuobjdumpSectionList.front()->setIdentifier(identifier); } -void setCuobjdumpptxfilename(const char* filename){ +void setCuobjdumpptxfilename(const char* filename, std::list &cuobjdumpSectionList){ printf("Adding ptx filename: %s\n", filename); cuobjdumpSection* x = cuobjdumpSectionList.front(); if (dynamic_cast(x) == NULL){ @@ -593,14 +530,14 @@ void setCuobjdumpptxfilename(const char* filename){ (dynamic_cast(x))->setPTXfilename(filename); } -void setCuobjdumpelffilename(const char* filename){ +void setCuobjdumpelffilename(const char* filename, std::list &cuobjdumpSectionList){ if (dynamic_cast(cuobjdumpSectionList.front()) == NULL){ assert (0 && "You shouldn't be trying to add a elffilename to an ptx section"); } (dynamic_cast(cuobjdumpSectionList.front()))->setELFfilename(filename); } -void setCuobjdumpsassfilename(const char* filename){ +void setCuobjdumpsassfilename(const char* filename, std::list &cuobjdumpSectionList){ if (dynamic_cast(cuobjdumpSectionList.front()) == NULL){ assert (0 && "You shouldn't be trying to add a sassfilename to an ptx section"); } @@ -2132,7 +2069,7 @@ static int get_app_cuda_version() { * It is also responsible for extracting the libraries linked to the binary if the option is * enabled * */ -void extract_code_using_cuobjdump(){ +void extract_code_using_cuobjdump(std::list &cuobjdumpSectionList){ CUctx_st *context = GPGPUSim_Context(); unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability(); @@ -2195,7 +2132,7 @@ void extract_code_using_cuobjdump(){ parser.ptxserial = 1; cuobjdump_lex_init(&(parser.scanner)); cuobjdump_set_in(cuobjdump_in, (parser.scanner)); - cuobjdump_parse(parser.scanner, &parser); + cuobjdump_parse(parser.scanner, &parser, cuobjdumpSectionList); cuobjdump_lex_destroy(parser.scanner); fclose(cuobjdump_in); printf("Done parsing!!!\n"); @@ -2251,7 +2188,7 @@ void extract_code_using_cuobjdump(){ parser.ptxserial = 1; cuobjdump_lex_init(&(parser.scanner)); cuobjdump_set_in(cuobjdump_in, (parser.scanner)); - cuobjdump_parse(parser.scanner, &parser); + cuobjdump_parse(parser.scanner, &parser, cuobjdumpSectionList); cuobjdump_lex_destroy(parser.scanner); fclose(cuobjdump_in); std::getline(libsf, line); @@ -2445,7 +2382,7 @@ cuobjdumpELFSection* findELFSectionInList(std::list sectionli } //! Find an ELF section in all the known lists -cuobjdumpELFSection* findELFSection(const std::string identifier){ +cuobjdumpELFSection* findELFSection(const std::string identifier, std::list cuobjdumpSectionList){ cuobjdumpELFSection* sec = findELFSectionInList(cuobjdumpSectionList, identifier); if (sec!=NULL)return sec; sec = findELFSectionInList(libSectionList, identifier); @@ -2480,7 +2417,7 @@ cuobjdumpPTXSection* findPTXSectionInList(std::list sectionli } //! Find an PTX section in all the known lists -cuobjdumpPTXSection* findPTXSection(const std::string identifier){ +cuobjdumpPTXSection* findPTXSection(const std::string identifier, std::list cuobjdumpSectionList){ cuobjdumpPTXSection* sec = findPTXSectionInList(cuobjdumpSectionList, identifier); if (sec!=NULL)return sec; sec = findPTXSectionInList(libSectionList, identifier); @@ -2493,9 +2430,9 @@ cuobjdumpPTXSection* findPTXSection(const std::string identifier){ //! Extract the code using cuobjdump and remove unnecessary sections -void cuobjdumpInit(){ +void cuobjdumpInit(std::list &cuobjdumpSectionList){ CUctx_st *context = GPGPUSim_Context(); - extract_code_using_cuobjdump(); //extract all the output of cuobjdump to _cuobjdump_*.* + extract_code_using_cuobjdump(cuobjdumpSectionList); //extract all the output of cuobjdump to _cuobjdump_*.* const char* pre_load = getenv("CUOBJDUMP_SIM_FILE"); if (pre_load ==NULL || strlen(pre_load)==0){ cuobjdumpSectionList = pruneSectionList(cuobjdumpSectionList, context); @@ -2513,7 +2450,7 @@ void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename){ } //! Either submit PTX for simulation or convert SASS to PTXPlus and submit it -void cuobjdumpParseBinary(unsigned int handle){ +void cuobjdumpParseBinary(unsigned int handle, std::list &cuobjdumpSectionList){ if(fatbin_registered[handle]) return; fatbin_registered[handle] = true; @@ -2566,7 +2503,7 @@ void 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 = findPTXSection(fname, context->cuobjdumpSectionList); 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) { @@ -2576,7 +2513,7 @@ void 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 = findELFSection(ptx->getIdentifier(), context->cuobjdumpSectionList); assert (elfsection!= NULL); char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus( ptx->getPTXfilename(), @@ -2664,7 +2601,7 @@ void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin ) * then for next calls, only returns the appropriate number */ assert(fat_cubin_handle >= 1); - if (fat_cubin_handle==1) cuobjdumpInit(); + if (fat_cubin_handle==1) cuobjdumpInit(context->cuobjdumpSectionList); cuobjdumpRegisterFatBinary(fat_cubin_handle, filename); return (void**)fat_cubin_handle; @@ -2779,7 +2716,7 @@ void CUDARTAPI __cudaRegisterFunction( 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()) - cuobjdumpParseBinary(fat_cubin_handle); + cuobjdumpParseBinary(fat_cubin_handle, context->cuobjdumpSectionList); context->register_function( fat_cubin_handle, hostFun, deviceFun ); } @@ -2799,7 +2736,7 @@ extern void __cudaRegisterVar( 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()) - cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle); + cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle, GPGPUSim_Context()->cuobjdumpSectionList ); fflush(stdout); if ( constant && !global && !ext ) { gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size); diff --git a/libcuda/cuobjdump.h b/libcuda/cuobjdump.h index 66cd736..49af3e2 100644 --- a/libcuda/cuobjdump.h +++ b/libcuda/cuobjdump.h @@ -1,5 +1,9 @@ #ifndef __cuobjdump_h__ #define __cuobjdump_h__ +#include +#include +#include + struct cuobjdump_parser { yyscan_t scanner; int elfserial; @@ -9,4 +13,67 @@ struct cuobjdump_parser { FILE *sassfile; char filename [1024]; }; + +class cuobjdumpSection { +public: + //Constructor + cuobjdumpSection() { + arch = 0; + identifier = ""; + } + virtual ~cuobjdumpSection() {} + unsigned getArch() {return arch;} + void setArch(unsigned a) {arch = a;} + std::string getIdentifier() {return identifier;} + void setIdentifier(std::string i) {identifier = i;} + virtual void print(){std::cout << "cuobjdump Section: unknown type" << std::endl;} +private: + unsigned arch; + std::string identifier; +}; + +class cuobjdumpELFSection : public cuobjdumpSection +{ +public: + cuobjdumpELFSection() {} + virtual ~cuobjdumpELFSection() { + elffilename = ""; + sassfilename = ""; + } + std::string getELFfilename() {return elffilename;} + void setELFfilename(std::string f) {elffilename = f;} + std::string getSASSfilename() {return sassfilename;} + void setSASSfilename(std::string f) {sassfilename = f;} + virtual void print() { + std::cout << "ELF Section:" << std::endl; + std::cout << "arch: sm_" << getArch() << std::endl; + std::cout << "identifier: " << getIdentifier() << std::endl; + std::cout << "elf filename: " << getELFfilename() << std::endl; + std::cout << "sass filename: " << getSASSfilename() << std::endl; + std::cout << std::endl; + } +private: + std::string elffilename; + std::string sassfilename; +}; + +class cuobjdumpPTXSection : public cuobjdumpSection +{ +public: + cuobjdumpPTXSection(){ + ptxfilename = ""; + } + std::string getPTXfilename() {return ptxfilename;} + void setPTXfilename(std::string f) {ptxfilename = f;} + virtual void print() { + std::cout << "PTX Section:" << std::endl; + std::cout << "arch: sm_" << getArch() << std::endl; + std::cout << "identifier: " << getIdentifier() << std::endl; + std::cout << "ptx filename: " << getPTXfilename() << std::endl; + std::cout << std::endl; + } +private: + std::string ptxfilename; +}; + #endif /* __cuobjdump_h__ */ diff --git a/libcuda/cuobjdump.l b/libcuda/cuobjdump.l index eccc1f2..5a19d65 100644 --- a/libcuda/cuobjdump.l +++ b/libcuda/cuobjdump.l @@ -39,9 +39,9 @@ #define YYDEBUG 1 -void cuobjdump_error(yyscan_t yyscanner, struct cuobjdump_parser* parser, const char* msg); +void cuobjdump_error(yyscan_t yyscanner, struct cuobjdump_parser* parser, std::list &cuobjdumpSectionList, const char* msg); #define YY_DECL int cuobjdump_lex \ - (YYSTYPE * yylval_param , yyscan_t yyscanner, struct cuobjdump_parser* parser) + (YYSTYPE * yylval_param , yyscan_t yyscanner, struct cuobjdump_parser* parser, std::list &cuobjdumpSectionList) %} %option stack @@ -156,10 +156,10 @@ newlines {newline}+ <> BEGIN(INITIAL);return 0; /*No other rule matched. Throw an error*/ -. cuobjdump_error(yyscanner, parser, "Invalid token"); +. cuobjdump_error(yyscanner, parser, cuobjdumpSectionList, "Invalid token"); %% -void cuobjdump_error(yyscan_t yyscanner, struct cuobjdump_parser* parser, const char* msg) +void cuobjdump_error(yyscan_t yyscanner, struct cuobjdump_parser* parser, std::list &cuobjdumpSectionList, const char* msg) { struct yyguts_t * yyg = (struct yyguts_t*)yyscanner; printf(" %s near \"%s\"",msg, yytext); diff --git a/libcuda/cuobjdump.y b/libcuda/cuobjdump.y index fcc863e..8d1bca6 100644 --- a/libcuda/cuobjdump.y +++ b/libcuda/cuobjdump.y @@ -32,25 +32,27 @@ typedef void * yyscan_t; #include "cuobjdump.h" -extern void addCuobjdumpSection(int sectiontype); -void setCuobjdumparch(const char* arch); -void setCuobjdumpidentifier(const char* identifier); -void setCuobjdumpptxfilename(const char* filename); -void setCuobjdumpelffilename(const char* filename); -void setCuobjdumpsassfilename(const char* filename); +extern void addCuobjdumpSection(int sectiontype, std::list &cuobjdumpSectionList); +void setCuobjdumparch(const char* arch, std::list &cuobjdumpSectionList); +void setCuobjdumpidentifier(const char* identifier, std::list &cuobjdumpSectionList); +void setCuobjdumpptxfilename(const char* filename, std::list &cuobjdumpSectionList); +void setCuobjdumpelffilename(const char* filename, std::list &cuobjdumpSectionList); +void setCuobjdumpsassfilename(const char* filename, std::list &cuobjdumpSectionList); %} %define api.pure full %parse-param {yyscan_t scanner} %parse-param {struct cuobjdump_parser* parser} +%parse-param {std::list &cuobjdumpSectionList} %lex-param {yyscan_t scanner} %lex-param {struct cuobjdump_parser* parser} +%lex-param {std::list &cuobjdumpSectionList} %union { char* string_value; } %{ -int yylex(YYSTYPE * yylval_param, yyscan_t yyscanner, struct cuobjdump_parser* parser); -void yyerror(yyscan_t yyscanner, struct cuobjdump_parser* parser, const char* msg); +int yylex(YYSTYPE * yylval_param, yyscan_t yyscanner, struct cuobjdump_parser* parser, std::list &cuobjdumpSectionList); +void yyerror(yyscan_t yyscanner, struct cuobjdump_parser* parser, std::list &cuobjdumpSectionList, const char* msg); %} %token H_SEPARATOR H_ARCH H_CODEVERSION H_PRODUCER H_HOST H_COMPILESIZE H_IDENTIFIER H_UNKNOWN H_COMPRESSED %token CODEVERSION @@ -75,23 +77,23 @@ emptylines : emptylines NEWLINE | ; section : PTXHEADER { - addCuobjdumpSection(0); + addCuobjdumpSection(0, cuobjdumpSectionList); snprintf(parser->filename, 1024, "_cuobjdump_%d.ptx", parser->ptxserial++); parser->ptxfile = fopen(parser->filename, "w"); - setCuobjdumpptxfilename(parser->filename); + setCuobjdumpptxfilename(parser->filename, cuobjdumpSectionList); } headerinfo compressedkeyword identifier ptxcode { fclose(parser->ptxfile); } | ELFHEADER { - addCuobjdumpSection(1); + addCuobjdumpSection(1, cuobjdumpSectionList); snprintf(parser->filename, 1024, "_cuobjdump_%d.elf", parser->elfserial); parser->elffile = fopen(parser->filename, "w"); - setCuobjdumpelffilename(parser->filename); + setCuobjdumpelffilename(parser->filename, cuobjdumpSectionList); } headerinfo compressedkeyword identifier elfcode { fclose(parser->elffile); snprintf(parser->filename, 1024, "_cuobjdump_%d.sass", parser->elfserial++); parser->sassfile = fopen(parser->filename, "w"); - setCuobjdumpsassfilename(parser->filename); + setCuobjdumpsassfilename(parser->filename, cuobjdumpSectionList); } sasscode { fclose(parser->sassfile); }; @@ -101,16 +103,16 @@ headerinfo : H_SEPARATOR NEWLINE H_CODEVERSION CODEVERSION NEWLINE H_PRODUCER H_UNKNOWN NEWLINE H_HOST IDENTIFIER NEWLINE - H_COMPILESIZE IDENTIFIER {setCuobjdumparch($4);}; + H_COMPILESIZE IDENTIFIER {setCuobjdumparch($4, cuobjdumpSectionList);}; | H_SEPARATOR NEWLINE H_ARCH IDENTIFIER NEWLINE H_CODEVERSION CODEVERSION NEWLINE H_PRODUCER IDENTIFIER NEWLINE H_HOST IDENTIFIER NEWLINE - H_COMPILESIZE IDENTIFIER {setCuobjdumparch($4);}; + H_COMPILESIZE IDENTIFIER {setCuobjdumparch($4, cuobjdumpSectionList);}; -identifier : H_IDENTIFIER FILENAME emptylines {setCuobjdumpidentifier($2);} - | {setCuobjdumpidentifier("default");}; +identifier : H_IDENTIFIER FILENAME emptylines {setCuobjdumpidentifier($2, cuobjdumpSectionList);} + | {setCuobjdumpidentifier("default", cuobjdumpSectionList);}; compressedkeyword : H_COMPRESSED emptylines | ; -- cgit v1.3 From 7864d663823cb5f6af7d9d4eb25c4527a8b4e80f Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 29 May 2019 12:25:48 -0400 Subject: Move libSectionList Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) (limited to 'libcuda/cuda_runtime_api.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index df7ddc7..225e93b 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -305,6 +305,8 @@ struct CUctx_st { } std::list cuobjdumpSectionList; + std::list libSectionList; + private: _cuda_device_id *m_gpu; // selected gpu std::map m_code; // fat binary handle => global symbol table @@ -497,8 +499,6 @@ enum cuobjdumpSectionType { }; -std::list libSectionList; - // sectiontype: 0 for ptx, 1 for elf void addCuobjdumpSection(int sectiontype, std::list &cuobjdumpSectionList){ if (sectiontype) @@ -2193,7 +2193,7 @@ void extract_code_using_cuobjdump(std::list &cuobjdumpSection fclose(cuobjdump_in); std::getline(libsf, line); } - libSectionList = cuobjdumpSectionList; + context->libSectionList = cuobjdumpSectionList; //Restore the original section list cuobjdumpSectionList = tmpsl; @@ -2382,7 +2382,7 @@ cuobjdumpELFSection* findELFSectionInList(std::list sectionli } //! Find an ELF section in all the known lists -cuobjdumpELFSection* findELFSection(const std::string identifier, std::list cuobjdumpSectionList){ +cuobjdumpELFSection* findELFSection(const std::string identifier, std::list cuobjdumpSectionList, std::list &libSectionList){ cuobjdumpELFSection* sec = findELFSectionInList(cuobjdumpSectionList, identifier); if (sec!=NULL)return sec; sec = findELFSectionInList(libSectionList, identifier); @@ -2417,7 +2417,7 @@ cuobjdumpPTXSection* findPTXSectionInList(std::list sectionli } //! Find an PTX section in all the known lists -cuobjdumpPTXSection* findPTXSection(const std::string identifier, std::list cuobjdumpSectionList){ +cuobjdumpPTXSection* findPTXSection(const std::string identifier, std::list cuobjdumpSectionList, std::list &libSectionList){ cuobjdumpPTXSection* sec = findPTXSectionInList(cuobjdumpSectionList, identifier); if (sec!=NULL)return sec; sec = findPTXSectionInList(libSectionList, identifier); @@ -2503,7 +2503,7 @@ void cuobjdumpParseBinary(unsigned int handle, std::list &cuo cuobjdumpPTXSection* ptx = NULL; const char* pre_load = getenv("CUOBJDUMP_SIM_FILE"); if(pre_load==NULL || strlen(pre_load)==0) - ptx = findPTXSection(fname, context->cuobjdumpSectionList); + ptx = findPTXSection(fname, context->cuobjdumpSectionList, context->libSectionList); 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) { @@ -2513,7 +2513,7 @@ void cuobjdumpParseBinary(unsigned int handle, std::list &cuo ptxcode = readfile(override_ptx_name); } if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) { - cuobjdumpELFSection* elfsection = findELFSection(ptx->getIdentifier(), context->cuobjdumpSectionList); + cuobjdumpELFSection* elfsection = findELFSection(ptx->getIdentifier(), context->cuobjdumpSectionList, context->libSectionList); assert (elfsection!= NULL); char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus( ptx->getPTXfilename(), -- cgit v1.3 From d9ca3558c774b8b86bb18024bbbf330df53722f7 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 29 May 2019 14:48:14 -0400 Subject: Move version_filename and g_cuda_launch_stack Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 50 ++++++++++++++++++++++++--------------------- 1 file changed, 27 insertions(+), 23 deletions(-) (limited to 'libcuda/cuda_runtime_api.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 225e93b..f3c827c 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -153,8 +153,6 @@ std::map pinned_memory; //support for pinned memories added std::map pinned_memory_size; std::map g_mallocPtr_Size; int no_of_ptx=0; -//maps sm version number to set of filenames -std::map > version_filename; extern void synchronize(); extern void exit_simulation(); @@ -241,6 +239,8 @@ private: struct _cuda_device_id *m_next; }; +class kernel_config; + struct CUctx_st { CUctx_st( _cuda_device_id *gpu ) { @@ -306,6 +306,9 @@ struct CUctx_st { std::list cuobjdumpSectionList; std::list libSectionList; + //maps sm version number to set of filenames + std::map > version_filename; + std::list g_cuda_launch_stack; private: _cuda_device_id *m_gpu; // selected gpu @@ -486,7 +489,6 @@ typedef std::map event_tracker_t; int CUevent_st::m_next_event_uid; event_tracker_t g_timer_events; int g_active_device = 0; //active gpu that runs the code -std::list g_cuda_launch_stack; extern int cuobjdump_lex_init(yyscan_t* scanner); extern void cuobjdump_set_in (FILE * _in_str ,yyscan_t yyscanner ); @@ -1505,7 +1507,8 @@ __host__ cudaError_t CUDARTAPI cudaConfigureCall(dim3 gridDim, dim3 blockDim, si announce_call(__my_func__); } struct CUstream_st *s = (struct CUstream_st *)stream; - g_cuda_launch_stack.push_back( kernel_config(gridDim,blockDim,sharedMem,s) ); + CUctx_st *context = GPGPUSim_Context(); + context->g_cuda_launch_stack.push_back( kernel_config(gridDim,blockDim,sharedMem,s) ); return g_last_cudaError = cudaSuccess; } @@ -1514,8 +1517,9 @@ __host__ cudaError_t CUDARTAPI cudaSetupArgument(const void *arg, size_t size, s if(g_debug_execution >= 3){ announce_call(__my_func__); } - gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); - kernel_config &config = g_cuda_launch_stack.back(); + CUctx_st *context = GPGPUSim_Context(); + gpgpusim_ptx_assert( !context->g_cuda_launch_stack.empty(), "empty launch stack" ); + kernel_config &config = context->g_cuda_launch_stack.back(); config.set_arg(arg,size,offset); printf("GPGPU-Sim PTX: Setting up arguments for %zu bytes starting at 0x%llx..\n",size, (unsigned long long) arg); @@ -1532,8 +1536,8 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) char *mode = getenv("PTX_SIM_MODE_FUNC"); if( mode ) sscanf(mode,"%u", &g_ptx_sim_mode); - gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); - kernel_config config = g_cuda_launch_stack.back(); + gpgpusim_ptx_assert( !context->g_cuda_launch_stack.empty(), "empty launch stack" ); + kernel_config config = context->g_cuda_launch_stack.back(); struct CUstream_st *stream = config.get_stream(); printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun, g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 ); @@ -1574,14 +1578,14 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) g_checkpoint->load_global_mem(global_mem, f1name); printf("Skipping kernel %d as resuming from kernel %d\n",grid->get_uid(),gpu->resume_kernel ); - g_cuda_launch_stack.pop_back(); + context->g_cuda_launch_stack.pop_back(); return g_last_cudaError = cudaSuccess; } if(gpu->checkpoint_option==1 && (grid->get_uid()>gpu->checkpoint_kernel)) { printf("Skipping kernel %d as checkpoint from kernel %d\n",grid->get_uid(),gpu->checkpoint_kernel ); - g_cuda_launch_stack.pop_back(); + context->g_cuda_launch_stack.pop_back(); return g_last_cudaError = cudaSuccess; } @@ -1589,7 +1593,7 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) kname.c_str(), stream?stream->get_uid():0, gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z ); stream_operation op(grid,g_ptx_sim_mode,stream); g_stream_manager->push(op); - g_cuda_launch_stack.pop_back(); + context->g_cuda_launch_stack.pop_back(); return g_last_cudaError = cudaSuccess; } @@ -1969,7 +1973,7 @@ char* get_app_binary_name(std::string abs_path){ } //extracts all ptx files from binary and dumps into prog_name.unique_no.sm_<>.ptx files -void extract_ptx_files_using_cuobjdump(){ +void extract_ptx_files_using_cuobjdump(CUctx_st *context){ extern bool g_cdp_enabled; char command[1000]; char *pytorch_bin = getenv("PYTORCH_BIN"); @@ -2030,10 +2034,10 @@ void extract_ptx_files_using_cuobjdump(){ } std::string vstr = line.substr(pos1+3,pos2-pos1-3); int version = atoi(vstr.c_str()); - if (version_filename.find(version)==version_filename.end()){ - version_filename[version] = std::set(); + if (context->version_filename.find(version)==context->version_filename.end()){ + context->version_filename[version] = std::set(); } - version_filename[version].insert(line); + context->version_filename[version].insert(line); } } @@ -2091,7 +2095,7 @@ void extract_code_using_cuobjdump(std::list &cuobjdumpSection //dump ptx for all individial ptx files into sepearte files which is later used by ptxas. int result=0; #if (CUDART_VERSION >= 6000) - extract_ptx_files_using_cuobjdump(); + extract_ptx_files_using_cuobjdump(context); return; #endif //TODO: redundant to dump twice. how can it be prevented? @@ -2450,7 +2454,7 @@ void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename){ } //! Either submit PTX for simulation or convert SASS to PTXPlus and submit it -void cuobjdumpParseBinary(unsigned int handle, std::list &cuobjdumpSectionList){ +void cuobjdumpParseBinary(unsigned int handle){ if(fatbin_registered[handle]) return; fatbin_registered[handle] = true; @@ -2467,7 +2471,7 @@ void cuobjdumpParseBinary(unsigned int handle, std::list &cuo #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 = context->version_filename.begin(); itr_m!=context->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; @@ -2479,7 +2483,7 @@ void cuobjdumpParseBinary(unsigned int handle, std::list &cuo 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++){ + for (itr_m = context->version_filename.begin(); itr_m!=context->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; @@ -2491,8 +2495,8 @@ void cuobjdumpParseBinary(unsigned int handle, std::list &cuo #endif unsigned max_capability = 0; - for ( std::list::iterator iter = cuobjdumpSectionList.begin(); - iter != cuobjdumpSectionList.end(); + for ( std::list::iterator iter = context->cuobjdumpSectionList.begin(); + iter != context->cuobjdumpSectionList.end(); iter++){ unsigned capability = (*iter)->getArch(); if (capability > max_capability) max_capability = capability; @@ -2716,7 +2720,7 @@ void CUDARTAPI __cudaRegisterFunction( 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()) - cuobjdumpParseBinary(fat_cubin_handle, context->cuobjdumpSectionList); + cuobjdumpParseBinary(fat_cubin_handle); context->register_function( fat_cubin_handle, hostFun, deviceFun ); } @@ -2736,7 +2740,7 @@ extern void __cudaRegisterVar( 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()) - cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle, GPGPUSim_Context()->cuobjdumpSectionList ); + cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle); fflush(stdout); if ( constant && !global && !ext ) { gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size); -- 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 'libcuda/cuda_runtime_api.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