From 1fb385af6e036588a0b807aa51d74dcf34f841cd Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Fri, 17 May 2019 10:59:27 -0400
Subject: Move ptx parser to reentrant

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 src/cuda-sim/ptx_parser.cc | 168 ++++++++++++++++++++++++---------------------
 1 file changed, 88 insertions(+), 80 deletions(-)

(limited to 'src/cuda-sim/ptx_parser.cc')

diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 25758dd..a1f6e00 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -27,13 +27,17 @@
 
 #include "ptx_parser.h"
 #include "ptx_ir.h"
+
+typedef void * yyscan_t;
 #include "ptx.tab.h"
 #include <stdarg.h>
 
-extern int ptx_error( const char *s );
-extern int ptx_lineno;
-extern int ptx_parse();
-extern FILE *ptx_in;
+extern int ptx_get_lineno (yyscan_t yyscanner );
+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);
 
 static const struct core_config *g_shader_core_config;
 void set_ptx_warp_size(const struct core_config * warp_size)
@@ -80,7 +84,7 @@ std::list<int> g_scalar_type;
 
 #define PTX_PARSE_DPRINTF(...) \
    if( g_debug_ir_generation ) { \
-      printf(" %s:%u => ",g_filename,ptx_lineno); \
+      printf(" %s:%u => ",g_filename,ptx_get_lineno(scanner)); \
       printf("   (%s:%u) ", __FILE__, __LINE__); \
       printf(__VA_ARGS__); \
       printf("\n"); \
@@ -109,7 +113,7 @@ void read_parser_environment_variables()
    }
 }
 
-void init_directive_state()
+void ptx_recognizer::init_directive_state()
 {
    PTX_PARSE_DPRINTF("init_directive_state");
    g_space_spec=undefined_space;
@@ -125,7 +129,7 @@ void init_directive_state()
    g_last_symbol = NULL;
 }
 
-void init_instruction_state()
+void ptx_recognizer::init_instruction_state()
 {
    PTX_PARSE_DPRINTF("init_instruction_state");
    g_pred = NULL;
@@ -149,7 +153,6 @@ symbol_table *init_parser( const char *ptx_filename )
    /*else {
        g_global_symbol_table = g_current_symbol_table = new symbol_table("global",0,g_global_allfiles_symbol_table);
    }*/
-   ptx_lineno = 1;
 
 #define DEF(X,Y) g_ptx_token_decode[X] = Y;
 #include "ptx_parser_decode.def"
@@ -169,18 +172,23 @@ symbol_table *init_parser( const char *ptx_filename )
    g_ptx_token_decode[generic_space] = "generic_space";
    g_ptx_token_decode[instruction_space] = "instruction_space";
 
-   init_directive_state();
-   init_instruction_state();
+   ptx_recognizer recognizer;
+   ptx_lex_init(&(recognizer.scanner));
+   recognizer.init_directive_state();
+   recognizer.init_instruction_state();
 
+   FILE *ptx_in;
    ptx_in = fopen(ptx_filename, "r");
-   ptx_parse();
+   ptx_set_in(ptx_in, recognizer.scanner);
+   ptx_parse(recognizer.scanner, &recognizer);
+   ptx_lex_destroy(recognizer.scanner);
    fclose(ptx_in);
    return g_global_symbol_table;
 }
 
 static int g_entry_point;
 
-void start_function( int entry_point ) 
+void ptx_recognizer::start_function( int entry_point )
 {
    PTX_PARSE_DPRINTF("start_function");
    init_directive_state();
@@ -194,7 +202,7 @@ char *g_add_identifier_cached__identifier = NULL;
 int g_add_identifier_cached__array_dim;
 int g_add_identifier_cached__array_ident;
 
-void add_function_name( const char *name ) 
+void ptx_recognizer::add_function_name( const char *name )
 {
    PTX_PARSE_DPRINTF("add_function_name %s %s", name,  ((g_entry_point==1)?"(entrypoint)":((g_entry_point==2)?"(extern)":"")));
    bool prior_decl = g_global_symbol_table->add_function_decl( name, g_entry_point, &g_func_info, &g_current_symbol_table );
@@ -210,21 +218,21 @@ void add_function_name( const char *name )
    if( prior_decl ) {
       g_func_info->remove_args();
    }
-   g_global_symbol_table->add_function( g_func_info, g_filename, ptx_lineno );
+   g_global_symbol_table->add_function( g_func_info, g_filename, ptx_get_lineno(scanner) );
 }
 
 //Jin: handle instruction group for cdp
-void start_inst_group() {
+void ptx_recognizer::start_inst_group() {
    PTX_PARSE_DPRINTF("start_instruction_group");
    g_current_symbol_table = g_current_symbol_table->start_inst_group();
 }
 
-void end_inst_group() {
+void ptx_recognizer::end_inst_group() {
    PTX_PARSE_DPRINTF("end_instruction_group");
    g_current_symbol_table = g_current_symbol_table->end_inst_group();
 }
 
-void add_directive() 
+void ptx_recognizer::add_directive()
 {
    PTX_PARSE_DPRINTF("add_directive");
    init_directive_state();
@@ -232,7 +240,7 @@ void add_directive()
 
 #define mymax(a,b) ((a)>(b)?(a):(b))
 
-void end_function() 
+void ptx_recognizer::end_function()
 {
    PTX_PARSE_DPRINTF("end_function");
 
@@ -250,7 +258,7 @@ void end_function()
 #define parse_error(msg, ...) parse_error_impl(__FILE__,__LINE__, msg, ##__VA_ARGS__)
 #define parse_assert(cond,msg, ...) parse_assert_impl((cond),__FILE__,__LINE__, msg, ##__VA_ARGS__)
 
-void parse_error_impl( const char *file, unsigned line, const char *msg, ... )
+void ptx_recognizer::parse_error_impl( const char *file, unsigned line, const char *msg, ... )
 {
    va_list ap;
    char buf[1024];
@@ -259,13 +267,13 @@ void parse_error_impl( const char *file, unsigned line, const char *msg, ... )
    va_end(ap);
 
    g_error_detected = 1;
-   printf("%s:%u: Parse error: %s (%s:%u)\n\n", g_filename, ptx_lineno, buf, file, line);
-   ptx_error(NULL);
+   printf("%s:%u: Parse error: %s (%s:%u)\n\n", g_filename, ptx_get_lineno(scanner), buf, file, line);
+   ptx_error(scanner, NULL);
    abort();
    exit(1);
 }
 
-void parse_assert_impl( int test_value, const char *file, unsigned line, const char *msg, ... )
+void ptx_recognizer::parse_assert_impl( int test_value, const char *file, unsigned line, const char *msg, ... )
 {
    va_list ap;
    char buf[1024];
@@ -280,7 +288,7 @@ void parse_assert_impl( int test_value, const char *file, unsigned line, const c
 extern char linebuf[4096];
 
 
-void set_return()
+void ptx_recognizer::set_return()
 {
    parse_assert( (g_opcode == CALL_OP || g_opcode == CALLP_OP), "only call can have return value");
    g_operands.front().set_return();
@@ -300,7 +308,7 @@ const ptx_instruction *ptx_instruction_lookup( const char *filename, unsigned li
    return l->second; 
 }
 
-void add_instruction() 
+void ptx_recognizer::add_instruction()
 {
    PTX_PARSE_DPRINTF("add_instruction: %s", ((g_opcode>0)?g_opcode_string[g_opcode]:"<label>") );
    assert( g_shader_core_config != 0 );
@@ -316,15 +324,15 @@ void add_instruction()
                                              g_scalar_type,
                                              g_space_spec,
                                              g_filename,
-                                             ptx_lineno,
+                                             ptx_get_lineno(scanner),
                                              linebuf,
                                              g_shader_core_config );
    g_instructions.push_back(i);
-   g_inst_lookup[g_filename][ptx_lineno] = i;
+   g_inst_lookup[g_filename][ptx_get_lineno(scanner)] = i;
    init_instruction_state();
 }
 
-void add_variables() 
+void ptx_recognizer::add_variables()
 {
    PTX_PARSE_DPRINTF("add_variables");
    if ( !g_operands.empty() ) {
@@ -334,7 +342,7 @@ void add_variables()
    init_directive_state();
 }
 
-void set_variable_type()
+void ptx_recognizer::set_variable_type()
 {
    PTX_PARSE_DPRINTF("set_variable_type space_spec=%s scalar_type_spec=%s", 
            g_ptx_token_decode[g_space_spec.get_type()].c_str(), 
@@ -378,7 +386,7 @@ int pad_address (new_addr_type address, unsigned size, unsigned maxalign) {
     return alignto ? ((alignto - (address % alignto)) % alignto) : 0;
 }
 
-void add_identifier( const char *identifier, int array_dim, unsigned array_ident ) 
+void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsigned array_ident )
 {
    if(array_ident==ARRAY_IDENTIFIER){
        g_size *= array_dim;
@@ -426,7 +434,7 @@ void add_identifier( const char *identifier, int array_dim, unsigned array_ident
    default:
       break;
    }
-   g_last_symbol = g_current_symbol_table->add_variable(identifier,type,num_bits/8,g_filename,ptx_lineno);
+   g_last_symbol = g_current_symbol_table->add_variable(identifier,type,num_bits/8,g_filename,ptx_get_lineno(scanner));
    switch ( ti.get_memory_space().get_type() ) {
    case reg_space: {
       regnum = g_current_symbol_table->next_reg_num();
@@ -569,7 +577,7 @@ void add_identifier( const char *identifier, int array_dim, unsigned array_ident
    }
 }
 
-void add_constptr(const char* identifier1, const char* identifier2, int offset)
+void ptx_recognizer::add_constptr(const char* identifier1, const char* identifier2, int offset)
 {
    symbol *s1 = g_current_symbol_table->lookup(identifier1);
    const symbol *s2 = g_current_symbol_table->lookup(identifier2);
@@ -584,7 +592,7 @@ void add_constptr(const char* identifier1, const char* identifier2, int offset)
    s1->set_address( addr + offset );
 }
 
-void add_function_arg()
+void ptx_recognizer::add_function_arg()
 {
    assert(g_size>0);
    if( g_func_info ) {
@@ -597,20 +605,20 @@ void add_function_arg()
 
 }
 
-void add_extern_spec() 
+void ptx_recognizer::add_extern_spec()
 {
    PTX_PARSE_DPRINTF("add_extern_spec");
    g_extern_spec = 1;
 }
 
-void add_alignment_spec( int spec )
+void ptx_recognizer::add_alignment_spec( int spec )
 {
    PTX_PARSE_DPRINTF("add_alignment_spec");
    parse_assert( g_alignment_spec == -1, "multiple .align specifiers per variable declaration not allowed." );
    g_alignment_spec = spec;
 }
 
-void add_ptr_spec( enum _memory_space_t spec ) 
+void ptx_recognizer::add_ptr_spec( enum _memory_space_t spec )
 {
    PTX_PARSE_DPRINTF("add_ptr_spec \"%s\"", g_ptx_token_decode[spec].c_str() );
    parse_assert( g_ptr_spec == undefined_space, "multiple ptr space specifiers not allowed." );
@@ -618,7 +626,7 @@ void add_ptr_spec( enum _memory_space_t spec )
    g_ptr_spec = spec; 
 }
 
-void add_space_spec( enum _memory_space_t spec, int value ) 
+void ptx_recognizer::add_space_spec( enum _memory_space_t spec, int value )
 {
    PTX_PARSE_DPRINTF("add_space_spec \"%s\"", g_ptx_token_decode[spec].c_str() );
    parse_assert( g_space_spec == undefined_space, "multiple space specifiers not allowed." );
@@ -637,14 +645,14 @@ void add_space_spec( enum _memory_space_t spec, int value )
    }
 }
 
-void add_vector_spec(int spec ) 
+void ptx_recognizer::add_vector_spec(int spec )
 {
    PTX_PARSE_DPRINTF("add_vector_spec");
    parse_assert( g_vector_spec == -1, "multiple vector specifiers not allowed." );
    g_vector_spec = spec;
 }
 
-void add_scalar_type_spec( int type_spec ) 
+void ptx_recognizer::add_scalar_type_spec( int type_spec )
 {
    //save size of parameter
    switch ( type_spec ) {
@@ -682,23 +690,23 @@ void add_scalar_type_spec( int type_spec )
    g_scalar_type_spec = type_spec;
 }
 
-void add_label( const char *identifier ) 
+void ptx_recognizer::add_label( const char *identifier )
 {
    PTX_PARSE_DPRINTF("add_label");
    symbol *s = g_current_symbol_table->lookup(identifier);
    if ( s != NULL ) {
       g_label = s;
    } else {
-      g_label = g_current_symbol_table->add_variable(identifier,NULL,0,g_filename,ptx_lineno);
+      g_label = g_current_symbol_table->add_variable(identifier,NULL,0,g_filename,ptx_get_lineno(scanner));
    }
 }
 
-void add_opcode( int opcode ) 
+void ptx_recognizer::add_opcode( int opcode )
 {
    g_opcode = opcode;
 }
 
-void add_pred( const char *identifier, int neg, int predModifier ) 
+void ptx_recognizer::add_pred( const char *identifier, int neg, int predModifier )
 {
    PTX_PARSE_DPRINTF("add_pred");
    const symbol *s = g_current_symbol_table->lookup(identifier);
@@ -711,17 +719,17 @@ void add_pred( const char *identifier, int neg, int predModifier )
    g_pred_mod = predModifier;
 }
 
-void add_option( int option ) 
+void ptx_recognizer::add_option( int option )
 {
    PTX_PARSE_DPRINTF("add_option");
    g_options.push_back( option );
 }
-void add_wmma_option( int option ) 
+void ptx_recognizer::add_wmma_option( int option )
 {
    PTX_PARSE_DPRINTF("add_option");
    g_wmma_options.push_back( option );
 }
-void add_double_operand( const char *d1, const char *d2 )
+void ptx_recognizer::add_double_operand( const char *d1, const char *d2 )
 {
    //operands that access two variables.
    //eg. s[$ofs1+$r0], g[$ofs1+=$r0]
@@ -734,7 +742,7 @@ void add_double_operand( const char *d1, const char *d2 )
    g_operands.push_back( operand_info(s1,s2) );
 }
 
-void add_1vector_operand( const char *d1 ) 
+void ptx_recognizer::add_1vector_operand( const char *d1 )
 {
    // handles the single element vector operand ({%v1}) found in tex.1d instructions
    PTX_PARSE_DPRINTF("add_1vector_operand");
@@ -743,7 +751,7 @@ void add_1vector_operand( const char *d1 )
    g_operands.push_back( operand_info(s1,NULL,NULL,NULL) );
 }
 
-void add_2vector_operand( const char *d1, const char *d2 ) 
+void ptx_recognizer::add_2vector_operand( const char *d1, const char *d2 )
 {
    PTX_PARSE_DPRINTF("add_2vector_operand");
    const symbol *s1 = g_current_symbol_table->lookup(d1);
@@ -752,7 +760,7 @@ void add_2vector_operand( const char *d1, const char *d2 )
    g_operands.push_back( operand_info(s1,s2,NULL,NULL) );
 }
 
-void add_3vector_operand( const char *d1, const char *d2, const char *d3 ) 
+void ptx_recognizer::add_3vector_operand( const char *d1, const char *d2, const char *d3 )
 {
    PTX_PARSE_DPRINTF("add_3vector_operand");
    const symbol *s1 = g_current_symbol_table->lookup(d1);
@@ -762,7 +770,7 @@ void add_3vector_operand( const char *d1, const char *d2, const char *d3 )
    g_operands.push_back( operand_info(s1,s2,s3,NULL) );
 }
 
-void add_4vector_operand( const char *d1, const char *d2, const char *d3, const char *d4 ) 
+void ptx_recognizer::add_4vector_operand( const char *d1, const char *d2, const char *d3, const char *d4 )
 {
    PTX_PARSE_DPRINTF("add_4vector_operand");
    const symbol *s1 = g_current_symbol_table->lookup(d1);
@@ -776,7 +784,7 @@ void add_4vector_operand( const char *d1, const char *d2, const char *d3, const
    if ( s4 == null_op ) s4 = NULL;
    g_operands.push_back( operand_info(s1,s2,s3,s4) );
 }
-void add_8vector_operand( const char *d1, const char *d2, const char *d3, const char *d4,const char *d5,const char *d6,const char *d7,const char *d8 ) 
+void ptx_recognizer::add_8vector_operand( const char *d1, const char *d2, const char *d3, const char *d4,const char *d5,const char *d6,const char *d7,const char *d8 )
 {
    PTX_PARSE_DPRINTF("add_8vector_operand");
    const symbol *s1 = g_current_symbol_table->lookup(d1);
@@ -799,13 +807,13 @@ void add_8vector_operand( const char *d1, const char *d2, const char *d3, const
    g_operands.push_back( operand_info(s1,s2,s3,s4,s5,s6,s7,s8) );
 }
 
-void add_builtin_operand( int builtin, int dim_modifier ) 
+void ptx_recognizer::add_builtin_operand( int builtin, int dim_modifier )
 {
    PTX_PARSE_DPRINTF("add_builtin_operand");
    g_operands.push_back( operand_info(builtin,dim_modifier) );
 }
 
-void add_memory_operand() 
+void ptx_recognizer::add_memory_operand()
 {
    PTX_PARSE_DPRINTF("add_memory_operand");
    assert( !g_operands.empty() );
@@ -813,7 +821,7 @@ void add_memory_operand()
 }
 
 /*TODO: add other memory locations*/
-void change_memory_addr_space(const char *identifier) 
+void ptx_recognizer::change_memory_addr_space(const char *identifier)
 {
    /*0 = N/A, not reading from memory
     *1 = global memory
@@ -860,7 +868,7 @@ void change_memory_addr_space(const char *identifier)
    parse_assert(recognizedType, "Error: unrecognized memory type.");
 }
 
-void change_operand_lohi( int lohi )
+void ptx_recognizer::change_operand_lohi( int lohi )
 {
    /*0 = N/A, read entire operand
     *1 = lo, reading from lowest bits
@@ -874,14 +882,14 @@ void change_operand_lohi( int lohi )
 
 }
 
-void set_immediate_operand_type ()
+void ptx_recognizer::set_immediate_operand_type ()
 {
      PTX_PARSE_DPRINTF("set_immediate_operand_type");
      assert( !g_operands.empty() );
      g_operands.back().set_immediate_addr();
 }
 
-void change_double_operand_type( int operand_type )
+void ptx_recognizer::change_double_operand_type( int operand_type )
 {
    /*
     *-3 = reg / reg (set instruction, but both get same value)
@@ -913,7 +921,7 @@ void change_double_operand_type( int operand_type )
 
 }
 
-void change_operand_neg( )
+void ptx_recognizer::change_operand_neg( )
 {
    PTX_PARSE_DPRINTF("change_operand_neg");
    assert( !g_operands.empty() );
@@ -922,32 +930,32 @@ void change_operand_neg( )
 
 }
 
-void add_literal_int( int value ) 
+void ptx_recognizer::add_literal_int( int value )
 {
    PTX_PARSE_DPRINTF("add_literal_int");
    g_operands.push_back( operand_info(value) );
 }
 
-void add_literal_float( float value ) 
+void ptx_recognizer::add_literal_float( float value )
 {
    PTX_PARSE_DPRINTF("add_literal_float");
    g_operands.push_back( operand_info(value) );
 }
 
-void add_literal_double( double value ) 
+void ptx_recognizer::add_literal_double( double value )
 {
    PTX_PARSE_DPRINTF("add_literal_double");
    g_operands.push_back( operand_info(value) );
 }
 
-void add_scalar_operand( const char *identifier ) 
+void ptx_recognizer::add_scalar_operand( const char *identifier )
 {
    PTX_PARSE_DPRINTF("add_scalar_operand");
    const symbol *s = g_current_symbol_table->lookup(identifier);
    if ( s == NULL ) {
       if ( g_opcode == BRA_OP || g_opcode == CALLP_OP) {
          // forward branch target...
-         s = g_current_symbol_table->add_variable(identifier,NULL,0,g_filename,ptx_lineno);
+         s = g_current_symbol_table->add_variable(identifier,NULL,0,g_filename,ptx_get_lineno(scanner));
       } else {
          std::string msg = std::string("operand \"") + identifier + "\" has no declaration.";
          parse_error( msg.c_str() );
@@ -956,19 +964,19 @@ void add_scalar_operand( const char *identifier )
    g_operands.push_back( operand_info(s) );
 }
 
-void add_neg_pred_operand( const char *identifier ) 
+void ptx_recognizer::add_neg_pred_operand( const char *identifier )
 {
    PTX_PARSE_DPRINTF("add_neg_pred_operand");
    const symbol *s = g_current_symbol_table->lookup(identifier);
    if ( s == NULL ) {
-       s = g_current_symbol_table->add_variable(identifier,NULL,1,g_filename,ptx_lineno);
+       s = g_current_symbol_table->add_variable(identifier,NULL,1,g_filename,ptx_get_lineno(scanner));
    }
    operand_info op(s);
    op.set_neg_pred();
    g_operands.push_back( op );
 }
 
-void add_address_operand( const char *identifier, int offset ) 
+void ptx_recognizer::add_address_operand( const char *identifier, int offset )
 {
    PTX_PARSE_DPRINTF("add_address_operand");
    const symbol *s = g_current_symbol_table->lookup(identifier);
@@ -979,23 +987,23 @@ void add_address_operand( const char *identifier, int offset )
    g_operands.push_back( operand_info(s,offset) );
 }
 
-void add_address_operand2( int offset )
+void ptx_recognizer::add_address_operand2( int offset )
 {
    PTX_PARSE_DPRINTF("add_address_operand");
    g_operands.push_back( operand_info((unsigned)offset) );
 }
 
-void add_array_initializer()
+void ptx_recognizer::add_array_initializer()
 {
    g_last_symbol->add_initializer(g_operands);
 }
 
-void add_version_info( float ver, unsigned ext )
+void ptx_recognizer::add_version_info( float ver, unsigned ext )
 {
    g_global_symbol_table->set_ptx_version(ver,ext);
 }
 
-void add_file( unsigned num, const char *filename )
+void ptx_recognizer::add_file( unsigned num, const char *filename )
 {
    if( g_filename == NULL ) {
       char *b = strdup(filename);
@@ -1022,44 +1030,44 @@ void add_file( unsigned num, const char *filename )
    g_current_symbol_table = g_global_symbol_table;
 }
 
-void *reset_symtab()
+void * ptx_recognizer::reset_symtab()
 {
    void *result = g_current_symbol_table;
    g_current_symbol_table = g_global_symbol_table;
    return result;
 }
 
-void set_symtab(void*symtab)
+void ptx_recognizer::set_symtab(void*symtab)
 {
    g_current_symbol_table = (symbol_table*)symtab;
 }
 
-void add_pragma( const char *str )
+void ptx_recognizer::add_pragma( const char *str )
 {
    printf("GPGPU-Sim PTX: Warning -- ignoring pragma '%s'\n", str );
 }
 
-void version_header(double a) {}  //intentional dummy function
+void ptx_recognizer::version_header(double a) {}  //intentional dummy function
 
-void target_header(char* a) 
+void ptx_recognizer::target_header(char* a)
 {
    g_global_symbol_table->set_sm_target(a,NULL,NULL);
 }
 
-void target_header2(char* a, char* b) 
+void ptx_recognizer::target_header2(char* a, char* b)
 {
    g_global_symbol_table->set_sm_target(a,b,NULL);
 }
 
-void target_header3(char* a, char* b, char* c) 
+void ptx_recognizer::target_header3(char* a, char* b, char* c)
 {
    g_global_symbol_table->set_sm_target(a,b,c);
 }
 
-void maxnt_id(int x, int y, int z) {
+void ptx_recognizer::maxnt_id(int x, int y, int z) {
   g_func_info->set_maxnt_id(x * y * z);
 }
 
-void func_header(const char* a) {} //intentional dummy function
-void func_header_info(const char* a) {} //intentional dummy function
-void func_header_info_int(const char* a, int b) {} //intentional dummy function
+void ptx_recognizer::func_header(const char* a) {} //intentional dummy function
+void ptx_recognizer::func_header_info(const char* a) {} //intentional dummy function
+void ptx_recognizer::func_header_info_int(const char* a, int b) {} //intentional dummy function
-- 
cgit v1.3


From 6afdcb5a0c2f4295549c1b53187c08ad75004293 Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Fri, 17 May 2019 16:21:31 -0400
Subject: Fix linebuf and col

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 src/cuda-sim/ptx.l         | 14 ++++++--------
 src/cuda-sim/ptx.y         |  3 +++
 src/cuda-sim/ptx_parser.cc |  4 ++--
 3 files changed, 11 insertions(+), 10 deletions(-)

(limited to 'src/cuda-sim/ptx_parser.cc')

diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index d173451..cc83bd1 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -41,10 +41,8 @@ class ptx_recognizer;
 #include "ptx.tab.h"
 #include <string.h>
 
-#define LINEBUF_SIZE (64*1024)
-char linebuf[LINEBUF_SIZE];
-unsigned col = 0;
-#define TC col+=strlen(yytext);
+#define LINEBUF_SIZE (4*1024)
+#define TC yylval->col+=strlen(yytext);
 #define CHECK_UNSIGNED \
 	if( yytext[strlen(yytext)-1]=='U' ) { \
 		printf("GPGPU-Sim: ERROR ** U modifier not implemented\n"); \
@@ -416,7 +414,7 @@ breakaddr  TC; yylval->int_value = BREAKADDR_OP; return OPCODE;
 
 "//"[^\n]* TC;	// eat single
 
-\n.*  col=0; strncpy(linebuf, yytext + 1, LINEBUF_SIZE); yyless( 1 );
+\n.*  yylval->col=0; strncpy(yylval->linebuf, yytext + 1, LINEBUF_SIZE); yyless( 1 );
 
 " " TC;
 "\t" TC;
@@ -464,10 +462,10 @@ int ptx_error( yyscan_t yyscanner, ptx_recognizer* recognizer, const char *s )
 	fflush(stdout);
 	if( s != NULL )
 		printf("%s:%u Syntax error:\n\n", g_filename, yylineno );
-	printf("   %s\n", linebuf );
+	printf("   %s\n", yylval->linebuf );
 	printf("   ");
-	for( i=0; i < col-1; i++ ) {
-		if( linebuf[i] == '\t' ) printf("\t");
+	for( i=0; i < yylval->col-1; i++ ) {
+		if( yylval->linebuf[i] == '\t' ) printf("\t");
 		else printf(" ");
 	}
 			
diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y
index 366b0cc..38a627f 100644
--- a/src/cuda-sim/ptx.y
+++ b/src/cuda-sim/ptx.y
@@ -43,6 +43,9 @@ class ptx_recognizer;
   int    int_value;
   char * string_value;
   void * ptr_value;
+#define LINEBUF_SIZE (4*1024)
+  char linebuf[LINEBUF_SIZE];
+  unsigned col;
 }
 
 %token <string_value> STRING
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index a1f6e00..2ec8386 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -33,6 +33,7 @@ typedef void * yyscan_t;
 #include <stdarg.h>
 
 extern int ptx_get_lineno (yyscan_t yyscanner );
+extern YYSTYPE* ptx_get_lval (yyscan_t yyscanner );
 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 );
@@ -285,7 +286,6 @@ void ptx_recognizer::parse_assert_impl( int test_value, const char *file, unsign
       parse_error_impl(file,line, msg);
 }
 
-extern char linebuf[4096];
 
 
 void ptx_recognizer::set_return()
@@ -325,7 +325,7 @@ void ptx_recognizer::add_instruction()
                                              g_space_spec,
                                              g_filename,
                                              ptx_get_lineno(scanner),
-                                             linebuf,
+                                             ptx_get_lval(scanner)->linebuf,
                                              g_shader_core_config );
    g_instructions.push_back(i);
    g_inst_lookup[g_filename][ptx_get_lineno(scanner)] = i;
-- 
cgit v1.3


From 579047c1dd9fec1b69e7904aeb6a25553c8dbd69 Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Mon, 20 May 2019 02:17:58 -0400
Subject: Fix col and linebuf

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 src/cuda-sim/ptx.l         | 17 +++++++++--------
 src/cuda-sim/ptx.y         |  6 ++----
 src/cuda-sim/ptx_parser.cc |  4 +++-
 src/cuda-sim/ptx_parser.h  |  3 +++
 4 files changed, 17 insertions(+), 13 deletions(-)

(limited to 'src/cuda-sim/ptx_parser.cc')

diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index cc83bd1..2546a99 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -37,17 +37,19 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
 %{
 #include "opcodes.h"
-class ptx_recognizer;
+#include "ptx_parser.h"
 #include "ptx.tab.h"
 #include <string.h>
 
 #define LINEBUF_SIZE (4*1024)
-#define TC yylval->col+=strlen(yytext);
+#define TC recognizer->col+=strlen(yytext);
 #define CHECK_UNSIGNED \
 	if( yytext[strlen(yytext)-1]=='U' ) { \
 		printf("GPGPU-Sim: ERROR ** U modifier not implemented\n"); \
 		abort(); \
 	}
+#define YY_DECL int ptx_lex \
+	       (YYSTYPE * yylval_param , yyscan_t yyscanner, ptx_recognizer* recognizer)
 int ptx_error( yyscan_t yyscanner, ptx_recognizer* recognizer, const char *s );
 %}
 
@@ -414,7 +416,7 @@ breakaddr  TC; yylval->int_value = BREAKADDR_OP; return OPCODE;
 
 "//"[^\n]* TC;	// eat single
 
-\n.*  yylval->col=0; strncpy(yylval->linebuf, yytext + 1, LINEBUF_SIZE); yyless( 1 );
+\n.*  recognizer->col=0; strncpy(recognizer->linebuf, yytext + 1, LINEBUF_SIZE); yyless( 1 );
 
 " " TC;
 "\t" TC;
@@ -447,9 +449,8 @@ breakaddr  TC; yylval->int_value = BREAKADDR_OP; return OPCODE;
 
 <*>\t@@DWARF.*\n
 
-<INITIAL,NOT_OPCODE,IN_FUNC_DECL>.  TC;
+<INITIAL,NOT_OPCODE,IN_FUNC_DECL>.  TC; ptx_error(yyscanner, recognizer, (const char*)NULL);
 %%
-//ptx_error(yyscanner, (const char*)NULL);
 
 extern int g_error_detected;
 extern const char *g_filename;
@@ -462,10 +463,10 @@ int ptx_error( yyscan_t yyscanner, ptx_recognizer* recognizer, const char *s )
 	fflush(stdout);
 	if( s != NULL )
 		printf("%s:%u Syntax error:\n\n", g_filename, yylineno );
-	printf("   %s\n", yylval->linebuf );
+	printf("   %s\n", recognizer->linebuf );
 	printf("   ");
-	for( i=0; i < yylval->col-1; i++ ) {
-		if( yylval->linebuf[i] == '\t' ) printf("\t");
+	for( i=0; i < recognizer->col-1; i++ ) {
+		if( recognizer->linebuf[i] == '\t' ) printf("\t");
 		else printf(" ");
 	}
 			
diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y
index 38a627f..8447693 100644
--- a/src/cuda-sim/ptx.y
+++ b/src/cuda-sim/ptx.y
@@ -36,6 +36,7 @@ class ptx_recognizer;
 %parse-param {yyscan_t scanner}
 %parse-param {ptx_recognizer* recognizer}
 %lex-param {yyscan_t scanner}
+%lex-param {ptx_recognizer* recognizer}
 
 %union {
   double double_value;
@@ -43,9 +44,6 @@ class ptx_recognizer;
   int    int_value;
   char * string_value;
   void * ptr_value;
-#define LINEBUF_SIZE (4*1024)
-  char linebuf[LINEBUF_SIZE];
-  unsigned col;
 }
 
 %token <string_value> STRING
@@ -232,7 +230,7 @@ class ptx_recognizer;
 	#include <math.h>
 	void syntax_not_implemented(yyscan_t yyscanner, ptx_recognizer* recognizer);
 	extern int g_func_decl;
-	int ptx_lex(YYSTYPE * yylval_param, yyscan_t yyscanner);
+	int ptx_lex(YYSTYPE * yylval_param, yyscan_t yyscanner, ptx_recognizer* recognizer);
 	int ptx_error( yyscan_t yyscanner, ptx_recognizer* recognizer, const char *s );
 %}
 
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 2ec8386..c4cc911 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -37,6 +37,7 @@ extern YYSTYPE* ptx_get_lval (yyscan_t yyscanner );
 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 FILE *ptx_get_in (yyscan_t yyscanner );
 extern int ptx_parse(yyscan_t scanner, ptx_recognizer* recognizer);
 extern int ptx_lex_destroy(yyscan_t scanner);
 
@@ -182,6 +183,7 @@ symbol_table *init_parser( const char *ptx_filename )
    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);
    fclose(ptx_in);
    return g_global_symbol_table;
@@ -325,7 +327,7 @@ void ptx_recognizer::add_instruction()
                                              g_space_spec,
                                              g_filename,
                                              ptx_get_lineno(scanner),
-                                             ptx_get_lval(scanner)->linebuf,
+                                             linebuf,
                                              g_shader_core_config );
    g_instructions.push_back(i);
    g_inst_lookup[g_filename][ptx_get_lineno(scanner)] = i;
diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h
index 70c4dfb..fdff6ce 100644
--- a/src/cuda-sim/ptx_parser.h
+++ b/src/cuda-sim/ptx_parser.h
@@ -44,6 +44,9 @@ class ptx_recognizer {
 	    scanner = NULL;
 	}
 	yyscan_t scanner;
+#define LINEBUF_SIZE (4*1024)
+	char linebuf[LINEBUF_SIZE];
+	unsigned col;
 	void init_directive_state();
 	void init_instruction_state();
 	void start_function( int entry_point );
-- 
cgit v1.3


From 603e53ee948fecd623837b52c4e94ca2dbf0012b Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Thu, 30 May 2019 13:44:15 -0400
Subject: Move some variables in ptx_parser

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 src/cuda-sim/ptx_parser.cc | 5 -----
 src/cuda-sim/ptx_parser.h  | 7 +++++++
 2 files changed, 7 insertions(+), 5 deletions(-)

(limited to 'src/cuda-sim/ptx_parser.cc')

diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index c4cc911..a61dc16 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -67,7 +67,6 @@ memory_space_t g_ptr_spec = undefined_space;
 int g_scalar_type_spec = -1;
 int g_vector_spec = -1;
 int g_alignment_spec = -1;
-int g_size = -1;
 int g_extern_spec = 0;
 
 // variable declaration stuff:
@@ -201,10 +200,6 @@ void ptx_recognizer::start_function( int entry_point )
    g_entry_func_param_index=0;
 }
 
-char *g_add_identifier_cached__identifier = NULL;
-int g_add_identifier_cached__array_dim;
-int g_add_identifier_cached__array_ident;
-
 void ptx_recognizer::add_function_name( const char *name )
 {
    PTX_PARSE_DPRINTF("add_function_name %s %s", name,  ((g_entry_point==1)?"(entrypoint)":((g_entry_point==2)?"(extern)":"")));
diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h
index fdff6ce..fb1b7f2 100644
--- a/src/cuda-sim/ptx_parser.h
+++ b/src/cuda-sim/ptx_parser.h
@@ -42,11 +42,18 @@ class ptx_recognizer {
     public:
 	ptx_recognizer() {
 	    scanner = NULL;
+	    g_size = -1;
+	    g_add_identifier_cached__identifier = NULL;
 	}
 	yyscan_t scanner;
 #define LINEBUF_SIZE (4*1024)
 	char linebuf[LINEBUF_SIZE];
 	unsigned col;
+	int g_size;
+	char *g_add_identifier_cached__identifier;
+	int g_add_identifier_cached__array_dim;
+	int g_add_identifier_cached__array_ident;
+
 	void init_directive_state();
 	void init_instruction_state();
 	void start_function( int entry_point );
-- 
cgit v1.3


From 976ec40b2614fe4e935486ce8be932bf5e573447 Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Sun, 9 Jun 2019 21:36:23 -0400
Subject: Move ptx_parser vars

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 src/cuda-sim/Makefile      |  8 ++++----
 src/cuda-sim/ptx_parser.cc | 21 ---------------------
 src/cuda-sim/ptx_parser.h  | 30 ++++++++++++++++++++++++++++++
 3 files changed, 34 insertions(+), 25 deletions(-)

(limited to 'src/cuda-sim/ptx_parser.cc')

diff --git a/src/cuda-sim/Makefile b/src/cuda-sim/Makefile
index 999dad7..85d1c8c 100644
--- a/src/cuda-sim/Makefile
+++ b/src/cuda-sim/Makefile
@@ -79,16 +79,16 @@ libgpgpu_ptx_sim.a: $(OBJS)
 	ar rcs $(OUTPUT_DIR)/libgpgpu_ptx_sim.a $(OUTPUT_DIR)/ptx.tab.o $(OUTPUT_DIR)/lex.ptx_.o $(OUTPUT_DIR)/ptxinfo.tab.o $(OUTPUT_DIR)/lex.ptxinfo_.o $(OBJS)
 
 $(OUTPUT_DIR)/ptx.tab.o: $(OUTPUT_DIR)/ptx.tab.c
-	$(CPP) -c $(OPT) -DYYDEBUG $(OUTPUT_DIR)/ptx.tab.c -o $(OUTPUT_DIR)/ptx.tab.o
+	$(CPP) -c $(CXX_OPT) -DYYDEBUG $(OUTPUT_DIR)/ptx.tab.c -o $(OUTPUT_DIR)/ptx.tab.o
 
 $(OUTPUT_DIR)/lex.ptx_.o: $(OUTPUT_DIR)/lex.ptx_.c 
-	$(CPP) -c $(OPT) $(OUTPUT_DIR)/lex.ptx_.c -o $(OUTPUT_DIR)/lex.ptx_.o
+	$(CPP) -c $(CXX_OPT) $(OUTPUT_DIR)/lex.ptx_.c -o $(OUTPUT_DIR)/lex.ptx_.o
 
 $(OUTPUT_DIR)/ptxinfo.tab.o: $(OUTPUT_DIR)/ptxinfo.tab.c
-	$(CPP) -c $(OPT) -DYYDEBUG $(OUTPUT_DIR)/ptxinfo.tab.c -o $(OUTPUT_DIR)/ptxinfo.tab.o
+	$(CPP) -c $(CXX_OPT) -DYYDEBUG $(OUTPUT_DIR)/ptxinfo.tab.c -o $(OUTPUT_DIR)/ptxinfo.tab.o
 
 $(OUTPUT_DIR)/lex.ptxinfo_.o: $(OUTPUT_DIR)/lex.ptxinfo_.c $(OUTPUT_DIR)/ptxinfo.tab.c
-	$(CPP) -c $(OPT) $(OUTPUT_DIR)/lex.ptxinfo_.c -o $(OUTPUT_DIR)/lex.ptxinfo_.o
+	$(CPP) -c $(CXX_OPT) $(OUTPUT_DIR)/lex.ptxinfo_.c -o $(OUTPUT_DIR)/lex.ptxinfo_.o
 
 $(OUTPUT_DIR)/ptx.tab.c: ptx.y
 	bison --name-prefix=ptx_ -v -d ptx.y --file-prefix=$(OUTPUT_DIR)/ptx
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index a61dc16..84eec65 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -61,27 +61,6 @@ static symbol *g_last_symbol = NULL;
 
 int g_error_detected = 0;
 
-// type specifier stuff:
-memory_space_t g_space_spec = undefined_space;
-memory_space_t g_ptr_spec = undefined_space;
-int g_scalar_type_spec = -1;
-int g_vector_spec = -1;
-int g_alignment_spec = -1;
-int g_extern_spec = 0;
-
-// variable declaration stuff:
-type_info *g_var_type = NULL;
-
-// instruction definition stuff:
-const symbol *g_pred;
-int g_neg_pred;
-int g_pred_mod;
-symbol *g_label;
-int g_opcode = -1;
-std::list<operand_info> g_operands;
-std::list<int> g_options;
-std::list<int> g_wmma_options;
-std::list<int> g_scalar_type;
 
 #define PTX_PARSE_DPRINTF(...) \
    if( g_debug_ir_generation ) { \
diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h
index fb1b7f2..b2388fa 100644
--- a/src/cuda-sim/ptx_parser.h
+++ b/src/cuda-sim/ptx_parser.h
@@ -29,6 +29,7 @@
 #define ptx_parser_INCLUDED
 
 #include "../abstract_hardware_model.h"
+#include "ptx_ir.h"
 extern const char *g_filename;
 extern int g_error_detected;
 
@@ -44,7 +45,16 @@ class ptx_recognizer {
 	    scanner = NULL;
 	    g_size = -1;
 	    g_add_identifier_cached__identifier = NULL;
+	    g_alignment_spec = -1;
+	    g_var_type = NULL;
+	    g_opcode = -1;
+	    g_space_spec = undefined_space;
+	    g_ptr_spec = undefined_space;
+	    g_scalar_type_spec = -1;
+	    g_vector_spec = -1;
+	    g_extern_spec = 0;
 	}
+	// global list
 	yyscan_t scanner;
 #define LINEBUF_SIZE (4*1024)
 	char linebuf[LINEBUF_SIZE];
@@ -53,7 +63,27 @@ class ptx_recognizer {
 	char *g_add_identifier_cached__identifier;
 	int g_add_identifier_cached__array_dim;
 	int g_add_identifier_cached__array_ident;
+	int g_alignment_spec;
+	// variable declaration stuff:
+	type_info *g_var_type;
+	// instruction definition stuff:
+	const symbol *g_pred;
+	int g_neg_pred;
+	int g_pred_mod;
+	symbol *g_label;
+	int g_opcode;
+	std::list<operand_info> g_operands;
+	std::list<int> g_options;
+	std::list<int> g_wmma_options;
+	std::list<int> g_scalar_type;
+	// type specifier stuff:
+	memory_space_t g_space_spec;
+	memory_space_t g_ptr_spec;
+	int g_scalar_type_spec;
+	int g_vector_spec;
+	int g_extern_spec;
 
+	// member function list
 	void init_directive_state();
 	void init_instruction_state();
 	void start_function( int entry_point );
-- 
cgit v1.3


From dca532b090edcf3fb840a697a8a0919d301c3dd7 Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Sun, 9 Jun 2019 21:57:13 -0400
Subject: Add more vars for ptx_parser

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 src/cuda-sim/ptx_parser.cc | 5 -----
 src/cuda-sim/ptx_parser.h  | 8 ++++++++
 2 files changed, 8 insertions(+), 5 deletions(-)

(limited to 'src/cuda-sim/ptx_parser.cc')

diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 84eec65..897beaf 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -49,7 +49,6 @@ void set_ptx_warp_size(const struct core_config * warp_size)
 
 static bool g_debug_ir_generation=false;
 const char *g_filename;
-unsigned g_max_regs_per_thread = 0;
 
 // the program intermediate representation...
 static symbol_table *g_global_allfiles_symbol_table = NULL;
@@ -341,10 +340,6 @@ bool check_for_duplicates( const char *identifier )
 extern std::set<std::string>   g_globals;
 extern std::set<std::string>   g_constants;
 
-int g_func_decl = 0;
-int g_ident_add_uid = 0;
-unsigned g_const_alloc = 1;
-
 // Returns padding that needs to be inserted ahead of address to make it aligned to min(size, maxalign)
 /*
  * @param address the address in bytes
diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h
index b2388fa..c52fe20 100644
--- a/src/cuda-sim/ptx_parser.h
+++ b/src/cuda-sim/ptx_parser.h
@@ -53,6 +53,10 @@ class ptx_recognizer {
 	    g_scalar_type_spec = -1;
 	    g_vector_spec = -1;
 	    g_extern_spec = 0;
+	    g_func_decl = 0;
+	    g_ident_add_uid = 0;
+	    g_const_alloc = 1;
+	    g_max_regs_per_thread = 0;
 	}
 	// global list
 	yyscan_t scanner;
@@ -82,6 +86,10 @@ class ptx_recognizer {
 	int g_scalar_type_spec;
 	int g_vector_spec;
 	int g_extern_spec;
+	int g_func_decl;
+	int g_ident_add_uid;
+	unsigned g_const_alloc;
+	unsigned g_max_regs_per_thread;
 
 	// member function list
 	void init_directive_state();
-- 
cgit v1.3


From c146d3707216f235a1dde3bd978239fd55b111ed Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Wed, 12 Jun 2019 18:17:45 -0400
Subject: Move all ptx_parser inside gpgpu_context

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 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_parser.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 5530b51c7c572d25717ad7f5d81196011dfcb540 Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Wed, 12 Jun 2019 18:49:58 -0400
Subject: Move symbol table globals

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 libcuda/gpgpu_context.h    |  2 ++
 src/cuda-sim/ptx_parser.cc | 10 +++-------
 src/cuda-sim/ptx_parser.h  |  7 +++++++
 3 files changed, 12 insertions(+), 7 deletions(-)

(limited to 'src/cuda-sim/ptx_parser.cc')

diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h
index b586e47..fa5d02b 100644
--- a/libcuda/gpgpu_context.h
+++ b/libcuda/gpgpu_context.h
@@ -7,11 +7,13 @@
 class gpgpu_context {
     public:
 	gpgpu_context() {
+	    g_global_allfiles_symbol_table = NULL;
 	    api = new cuda_runtime_api();
 	    ptxinfo = new ptxinfo_data();
 	    ptx_parser = new ptx_recognizer();
 	}
 	// global list
+	symbol_table *g_global_allfiles_symbol_table;
 	// objects pointers for each file
 	cuda_runtime_api* api;
 	ptxinfo_data* ptxinfo;
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 7080ee9..32e12a1 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -52,12 +52,8 @@ static bool g_debug_ir_generation=false;
 const char *g_filename;
 
 // the program intermediate representation...
-static symbol_table *g_global_allfiles_symbol_table = NULL;
-static symbol_table *g_global_symbol_table = NULL;
 std::map<std::string,symbol_table*> g_sym_name_to_symbol_table;
-static symbol_table *g_current_symbol_table = NULL;
 static std::list<ptx_instruction*> g_instructions;
-static symbol *g_last_symbol = NULL;
 
 int g_error_detected = 0;
 
@@ -128,7 +124,7 @@ symbol_table * gpgpu_context::init_parser( const char *ptx_filename )
    g_filename = strdup(ptx_filename);
    if  (g_global_allfiles_symbol_table == NULL) {
        g_global_allfiles_symbol_table = new symbol_table("global_allfiles", 0, NULL);
-       g_global_symbol_table = g_current_symbol_table = g_global_allfiles_symbol_table;
+       ptx_parser->g_global_symbol_table = ptx_parser->g_current_symbol_table = g_global_allfiles_symbol_table;
    }
    /*else {
        g_global_symbol_table = g_current_symbol_table = new symbol_table("global",0,g_global_allfiles_symbol_table);
@@ -163,7 +159,7 @@ symbol_table * gpgpu_context::init_parser( const char *ptx_filename )
    ptx_in = ptx_get_in(ptx_parser->scanner);
    ptx_lex_destroy(ptx_parser->scanner);
    fclose(ptx_in);
-   return g_global_symbol_table;
+   return ptx_parser->g_global_symbol_table;
 }
 
 static int g_entry_point;
@@ -331,7 +327,7 @@ void ptx_recognizer::set_variable_type()
                                                   g_extern_spec );
 }
 
-bool check_for_duplicates( const char *identifier )
+bool ptx_recognizer::check_for_duplicates( const char *identifier )
 {
    const symbol *s = g_current_symbol_table->lookup(identifier);
    return ( s != NULL );
diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h
index 5293c2e..cc8fe18 100644
--- a/src/cuda-sim/ptx_parser.h
+++ b/src/cuda-sim/ptx_parser.h
@@ -56,6 +56,9 @@ class ptx_recognizer {
 	    g_ident_add_uid = 0;
 	    g_const_alloc = 1;
 	    g_max_regs_per_thread = 0;
+	    g_global_symbol_table = NULL;
+	    g_current_symbol_table = NULL;
+	    symbol *g_last_symbol = NULL;
 	}
 	// global list
 	yyscan_t scanner;
@@ -89,6 +92,9 @@ class ptx_recognizer {
 	int g_ident_add_uid;
 	unsigned g_const_alloc;
 	unsigned g_max_regs_per_thread;
+	symbol_table *g_global_symbol_table;
+	symbol_table *g_current_symbol_table;
+	symbol *g_last_symbol;
 
 	// member function list
 	void init_directive_state();
@@ -154,6 +160,7 @@ class ptx_recognizer {
 	//Jin: handle instructino group for cdp
 	void start_inst_group();
 	void end_inst_group();
+	bool check_for_duplicates( const char *identifier );
 
 };
 
-- 
cgit v1.3


From e5447646aaca133f2cc5cbafcaf7aeef4b384c62 Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Wed, 12 Jun 2019 19:13:55 -0400
Subject: Add some vars in ptx_parser

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 src/cuda-sim/ptx.l         | 3 +--
 src/cuda-sim/ptx_parser.cc | 5 -----
 src/cuda-sim/ptx_parser.h  | 5 +++++
 3 files changed, 6 insertions(+), 7 deletions(-)

(limited to 'src/cuda-sim/ptx_parser.cc')

diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index 2546a99..18952a9 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -452,14 +452,13 @@ breakaddr  TC; yylval->int_value = BREAKADDR_OP; return OPCODE;
 <INITIAL,NOT_OPCODE,IN_FUNC_DECL>.  TC; ptx_error(yyscanner, recognizer, (const char*)NULL);
 %%
 
-extern int g_error_detected;
 extern const char *g_filename;
 
 int ptx_error( yyscan_t yyscanner, ptx_recognizer* recognizer, const char *s )
 {
     struct yyguts_t * yyg = (struct yyguts_t*)yyscanner;
 	int i;
-	g_error_detected = 1;
+	recognizer->g_error_detected = 1;
 	fflush(stdout);
 	if( s != NULL )
 		printf("%s:%u Syntax error:\n\n", g_filename, yylineno );
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 32e12a1..a01d1e3 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -53,10 +53,6 @@ const char *g_filename;
 
 // the program intermediate representation...
 std::map<std::string,symbol_table*> g_sym_name_to_symbol_table;
-static std::list<ptx_instruction*> g_instructions;
-
-int g_error_detected = 0;
-
 
 #define PTX_PARSE_DPRINTF(...) \
    if( g_debug_ir_generation ) { \
@@ -67,7 +63,6 @@ int g_error_detected = 0;
       fflush(stdout); \
    }
 
-static unsigned g_entry_func_param_index=0;
 static function_info *g_func_info = NULL;
 static std::map<unsigned,std::string> g_ptx_token_decode;
 static operand_info g_return_var;
diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h
index cc8fe18..c4e0b89 100644
--- a/src/cuda-sim/ptx_parser.h
+++ b/src/cuda-sim/ptx_parser.h
@@ -59,6 +59,8 @@ class ptx_recognizer {
 	    g_global_symbol_table = NULL;
 	    g_current_symbol_table = NULL;
 	    symbol *g_last_symbol = NULL;
+	    g_error_detected = 0;
+	    g_entry_func_param_index=0;
 	}
 	// global list
 	yyscan_t scanner;
@@ -95,6 +97,9 @@ class ptx_recognizer {
 	symbol_table *g_global_symbol_table;
 	symbol_table *g_current_symbol_table;
 	symbol *g_last_symbol;
+	std::list<ptx_instruction*> g_instructions;
+	int g_error_detected;
+	unsigned g_entry_func_param_index;
 
 	// member function list
 	void init_directive_state();
-- 
cgit v1.3


From 632bdaab79b9c1fd51ef8152aa00209f76c03101 Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Wed, 12 Jun 2019 19:55:35 -0400
Subject: Fix some compiler warning

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 src/cuda-sim/ptx_loader.h  | 4 ++--
 src/cuda-sim/ptx_parser.cc | 2 --
 src/cuda-sim/ptx_parser.h  | 9 ++++++---
 3 files changed, 8 insertions(+), 7 deletions(-)

(limited to 'src/cuda-sim/ptx_parser.cc')

diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h
index 3637bff..ee09e16 100644
--- a/src/cuda-sim/ptx_loader.h
+++ b/src/cuda-sim/ptx_loader.h
@@ -29,12 +29,12 @@
 #define PTX_LOADER_H_INCLUDED
 #include <string>
 
-#define LINEBUF_SIZE 1024
+#define PTXINFO_LINEBUF_SIZE 1024
 typedef void * yyscan_t;
 class ptxinfo_data{
     public:
 	yyscan_t scanner;
-	char linebuf[LINEBUF_SIZE];
+	char linebuf[PTXINFO_LINEBUF_SIZE];
 	unsigned col;
 	const char *g_ptxinfo_filename;
 
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index a01d1e3..2872b84 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -63,9 +63,7 @@ std::map<std::string,symbol_table*> g_sym_name_to_symbol_table;
       fflush(stdout); \
    }
 
-static function_info *g_func_info = NULL;
 static std::map<unsigned,std::string> g_ptx_token_decode;
-static operand_info g_return_var;
 
 const char *decode_token( int type )
 {
diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h
index c4e0b89..25c01fe 100644
--- a/src/cuda-sim/ptx_parser.h
+++ b/src/cuda-sim/ptx_parser.h
@@ -58,14 +58,15 @@ class ptx_recognizer {
 	    g_max_regs_per_thread = 0;
 	    g_global_symbol_table = NULL;
 	    g_current_symbol_table = NULL;
-	    symbol *g_last_symbol = NULL;
+	    g_last_symbol = NULL;
 	    g_error_detected = 0;
 	    g_entry_func_param_index=0;
+	    g_func_info = NULL;
 	}
 	// global list
 	yyscan_t scanner;
-#define LINEBUF_SIZE (4*1024)
-	char linebuf[LINEBUF_SIZE];
+#define PTX_LINEBUF_SIZE (4*1024)
+	char linebuf[PTX_LINEBUF_SIZE];
 	unsigned col;
 	int g_size;
 	char *g_add_identifier_cached__identifier;
@@ -100,6 +101,8 @@ class ptx_recognizer {
 	std::list<ptx_instruction*> g_instructions;
 	int g_error_detected;
 	unsigned g_entry_func_param_index;
+	function_info *g_func_info;
+	operand_info g_return_var;
 
 	// member function list
 	void init_directive_state();
-- 
cgit v1.3


From cb678c3670de4a435a3260ed80dc476da3860082 Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Mon, 1 Jul 2019 16:23:33 -0400
Subject: Move g_debug_ir_generation and GPGPUSim_Init

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 libcuda/cuda_runtime_api.cc     | 674 ++++++++++++++++++++++------------------
 libcuda/gpgpu_context.h         |   2 +
 libopencl/opencl_runtime_api.cc |   4 +-
 src/cuda-sim/ptx_parser.cc      |   3 +-
 src/cuda-sim/ptx_parser.h       |   3 +
 src/gpgpusim_entrypoint.cc      |   5 +-
 src/gpgpusim_entrypoint.h       |   1 -
 7 files changed, 375 insertions(+), 317 deletions(-)

(limited to 'src/cuda-sim/ptx_parser.cc')

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 61f8415..a34727f 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -197,7 +197,7 @@ void register_ptx_function( const char *name, function_info *impl )
 # endif
 #endif
 
-struct _cuda_device_id *GPGPUSim_Init()
+struct _cuda_device_id *gpgpu_context::GPGPUSim_Init()
 {
 	//static _cuda_device_id *the_device = NULL;
 	_cuda_device_id *the_device = GPGPUsim_ctx_ptr()->the_cude_device;
@@ -255,9 +255,10 @@ struct _cuda_device_id *GPGPUSim_Init()
 static CUctx_st* GPGPUSim_Context()
 {
 	//static CUctx_st *the_context = NULL;
+	gpgpu_context *cur_ctx = GPGPU_Context();
 	CUctx_st *the_context = GPGPUsim_ctx_ptr()->the_context;
 	if( the_context == NULL ) {
-		_cuda_device_id *the_gpu = GPGPUSim_Init();
+		_cuda_device_id *the_gpu = cur_ctx->GPGPUSim_Init();
 		GPGPUsim_ctx_ptr()->the_context = new CUctx_st(the_gpu);
 		the_context = GPGPUsim_ctx_ptr()->the_context;
 	}
@@ -496,7 +497,7 @@ cudaError_t cudaSetDeviceInternal(int device, gpgpu_context* gpgpu_ctx = NULL)
 	    announce_call(__my_func__);
     }
 	//set the active device to run cuda
-	if ( device <= GPGPUSim_Init()->num_devices() ) {
+	if ( device <= ctx->GPGPUSim_Init()->num_devices() ) {
 		ctx->api->g_active_device = device;
 		return g_last_cudaError = cudaSuccess;
 	} else {
@@ -519,6 +520,55 @@ cudaError_t cudaGetDeviceInternal(int *device, gpgpu_context* gpgpu_ctx = NULL)
 	return g_last_cudaError = cudaSuccess;
 }
 
+__host__ cudaError_t CUDARTAPI cudaDeviceGetLimitInternal( size_t* pValue, cudaLimit limit, gpgpu_context* gpgpu_ctx = NULL )
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+    if(g_debug_execution >= 3){
+	announce_call(__my_func__);
+    }
+    _cuda_device_id *dev = ctx->GPGPUSim_Init();
+    const struct cudaDeviceProp *prop = dev->get_prop();
+    const gpgpu_sim_config& config=dev->get_gpgpu()->get_config();
+    switch(limit) {
+	case 0:  // cudaLimitStackSize
+	    *pValue=config.stack_limit();
+	    break;
+	case 2:  // cudaLimitMallocHeapSize
+	    *pValue=config.heap_limit();
+	    break;
+#if (CUDART_VERSION > 5050)
+	case 3: // cudaLimitDevRuntimeSyncDepth
+	    if(prop->major > 2){
+		*pValue=config.sync_depth_limit();
+		break;
+	    }
+	    else{
+		printf("ERROR:Limit %s is not supported on this architecture \n",limit);
+		abort();
+	    }
+	case 4: // cudaLimitDevRuntimePendingLaunchCount
+	    if(prop->major > 2){
+		*pValue=config.pending_launch_count_limit();
+		break;
+	    }
+	    else{
+		printf("ERROR:Limit %s is not supported on this architecture \n",limit);
+		abort();
+	    }
+#endif
+	default:
+	    printf("ERROR:Limit %s unimplemented \n",limit);
+	    abort();
+    }
+    return g_last_cudaError = cudaSuccess;
+
+}
+
 
 void** cudaRegisterFatBinaryInternal( void *fatCubin, gpgpu_context* gpgpu_ctx = NULL)
 {
@@ -739,6 +789,260 @@ cudaError_t cudaConfigureCallInternal(dim3 gridDim, dim3 blockDim, size_t shared
 	return g_last_cudaError = cudaSuccess;
 }
 
+__host__ cudaError_t CUDARTAPI cudaGetDeviceCountInternal(int *count, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+    if(g_debug_execution >= 3){
+	announce_call(__my_func__);
+    }
+    _cuda_device_id *dev = ctx->GPGPUSim_Init();
+    *count = dev->num_devices();
+    return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaGetDevicePropertiesInternal(struct cudaDeviceProp *prop, int device, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+    if(g_debug_execution >= 3){
+	announce_call(__my_func__);
+    }
+    _cuda_device_id *dev = ctx->GPGPUSim_Init();
+    if (device <= dev->num_devices() )  {
+	*prop= *dev->get_prop();
+	return g_last_cudaError = cudaSuccess;
+    } else {
+	return g_last_cudaError = cudaErrorInvalidDevice;
+    }
+}
+
+#if (CUDART_VERSION > 5000)
+__host__ cudaError_t CUDARTAPI cudaDeviceGetAttributeInternal(int *value, enum cudaDeviceAttr attr, int device, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+        const struct cudaDeviceProp *prop;
+        _cuda_device_id *dev = ctx->GPGPUSim_Init();
+        if (device <= dev->num_devices() )  {
+                prop = dev->get_prop();
+                switch (attr) {
+                case 1:
+                        *value= prop->maxThreadsDim[0] * prop->maxThreadsDim[1] * prop->maxThreadsDim[2] * prop->maxGridSize[0] * prop->maxGridSize[1] * prop->maxGridSize[2];
+                        break;
+                case 2:
+                        *value= prop->maxThreadsDim[0];
+                        break;
+                case 3:
+                        *value= prop->maxThreadsDim[1];
+                        break;
+                case 4:
+                        *value= prop->maxThreadsDim[2];
+                        break; 
+		case 5:
+                        *value= prop->maxGridSize[0];
+                        break;
+                case 6:
+                        *value= prop->maxGridSize[1];
+                        break;
+                case 7:
+                        *value= prop->maxGridSize[2];
+                        break;
+		case 8:
+                        *value= prop->sharedMemPerBlock;
+                        break;
+		case 9:
+                        *value= prop->totalConstMem;
+                        break;
+                case 10:
+                        *value= prop->warpSize;
+                        break;
+                case 11:
+                        *value= 16;//dummy value
+                        break;
+                case 12:
+                        *value= prop->regsPerBlock;
+                        break;
+                case 13:
+                        *value= 1480000;//for 1080ti
+                        break;
+                case 14:
+                        *value= prop->textureAlignment ;
+                        break;
+                case 15:
+                        *value = 0;
+                        break;
+                case 16:
+                        *value= prop->multiProcessorCount ;
+                        break;
+                case 17:
+                case 18:
+                case 19:
+                        *value = 0;
+                        break;
+                case 21:
+                case 22:
+                case 23:
+                case 24:
+                case 25:
+                case 26:
+                case 27:
+                case 28:
+                case 42:
+                case 45:
+                case 46:
+                case 47:
+                case 48:
+                case 49:
+                case 52:
+                case 53:
+                case 55:
+                case 56:
+                case 57:
+                case 58:
+                case 59:
+                case 60:
+                case 61:
+                case 62:
+                case 63:
+                case 64:
+                case 66:
+                case 67:
+                case 69:
+                case 70:
+                case 71:
+                case 73:
+                case 74:
+                case 77:
+                        *value = 1000;//dummy value
+                        break;
+                case 29:
+                case 43:
+                case 54:
+                case 65:
+                case 68:
+                case 72:
+                        *value = 10;//dummy value
+                        break;
+                case 30:
+                case 51:
+                        *value = 128;//dummy value
+                        break;
+                case 31:
+                        *value = 1;
+                        break;
+                case 32:
+                        *value = 0;
+                        break;
+                case 33:
+                case 50:
+                        *value = 0;//dummy value
+                        break;
+		        case 34:
+                        *value= 0;
+                        break;
+                case 35:
+                        *value = 0;
+                        break;
+                case 36:
+                        *value = 1250000;//CK value for 1080ti
+                        break;
+                case 37:
+                        *value = 352;//value for 1080ti
+                        break;
+                case 38:
+                        *value = 3000000;//value for 1080ti
+                        break;
+                case 39:
+                        *value= dev->get_gpgpu()->threads_per_core();
+                        break;
+                case 40:
+                        *value= 0;
+                        break;
+                case 41:
+                        *value= 0;
+                        break;
+                case 75://cudaDevAttrComputeCapabilityMajor
+                        *value= prop->major ;
+                        break;
+                case 76://cudaDevAttrComputeCapabilityMinor
+                        *value= prop->minor ;
+                        break;
+                case 78:
+                        *value= 0 ; //TODO: as of now, we dont support stream priorities.
+                        break;
+                case 79:
+                        *value= 0;
+                        break;
+                case 80:
+                        *value= 0;
+                        break;
+		#if (CUDART_VERSION > 5050)
+		case 81:
+                        *value= prop->sharedMemPerMultiprocessor;
+                        break;
+                case 82:
+                        *value= prop->regsPerMultiprocessor;
+                        break;
+		#endif
+                case 83:
+                case 84:
+                case 85:
+                case 86:
+                        *value= 0;
+                        break;
+                case 87:
+                        *value= 4;//dummy value
+                        break;
+                case 88:
+                case 89:
+                case 90:
+                case 91:
+                case 95:
+                        *value= 0;
+                        break;
+		default:
+			printf("ERROR: Attribute number %d unimplemented \n",attr);
+			abort();
+                }
+                return g_last_cudaError = cudaSuccess;
+        } else {
+                return g_last_cudaError = cudaErrorInvalidDevice;
+        }
+}
+#endif
+
+__host__ cudaError_t CUDARTAPI cudaChooseDeviceInternal(int *device, const struct cudaDeviceProp *prop, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	_cuda_device_id *dev = ctx->GPGPUSim_Init();
+	*device = dev->get_id();
+	return g_last_cudaError = cudaSuccess;
+}
+
 cudaError_t cudaSetupArgumentInternal(const void *arg, size_t size, size_t offset, gpgpu_context* gpgpu_ctx = NULL)
 {
     gpgpu_context *ctx;
@@ -1058,6 +1362,52 @@ cudaError_t cudaHostAllocInternal(void **pHost,  size_t bytes, unsigned int flag
 
 #endif
 
+size_t getMaxThreadsPerBlock(struct cudaFuncAttributes *attr, gpgpu_context *ctx) {
+  _cuda_device_id *dev = ctx->GPGPUSim_Init();
+  struct cudaDeviceProp prop;
+
+  prop = *dev->get_prop();
+
+  size_t max = prop.maxThreadsPerBlock;
+
+  if ((prop.regsPerBlock / attr->numRegs) < max)
+    max = prop.regsPerBlock / attr->numRegs;
+
+  return max;
+}
+
+cudaError_t CUDARTAPI cudaFuncGetAttributesInternal(struct cudaFuncAttributes *attr, const char *hostFun, gpgpu_context* gpgpu_ctx = NULL )
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	CUctx_st *context = GPGPUSim_Context();
+	function_info *entry = context->get_kernel(hostFun);
+	if( entry ) {
+		const struct gpgpu_ptx_sim_info *kinfo = entry->get_kernel_info();
+		attr->sharedSizeBytes = kinfo->smem;
+		attr->constSizeBytes  = kinfo->cmem;
+		attr->localSizeBytes  = kinfo->lmem;
+		attr->numRegs         = kinfo->regs;
+		if(kinfo->maxthreads > 0)
+		  attr->maxThreadsPerBlock = kinfo->maxthreads;
+		else
+		  attr->maxThreadsPerBlock = getMaxThreadsPerBlock(attr, ctx);
+#if CUDART_VERSION >= 3000
+		attr->ptxVersion      = kinfo->ptx_version;
+		attr->binaryVersion   = kinfo->sm_target;
+#endif
+	}
+	return g_last_cudaError = cudaSuccess;
+}
+
+
 /*******************************************************************************
  *                                                                              *
  *                                                                              *
@@ -1503,235 +1853,24 @@ __host__ cudaError_t CUDARTAPI cudaGetSymbolSize(size_t *size, const char *symbo
  *******************************************************************************/
 __host__ cudaError_t CUDARTAPI cudaGetDeviceCount(int *count)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	_cuda_device_id *dev = GPGPUSim_Init();
-	*count = dev->num_devices();
-	return g_last_cudaError = cudaSuccess;
+    return cudaGetDeviceCountInternal(count);
 }
-
-__host__ cudaError_t CUDARTAPI cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device)
-{
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	_cuda_device_id *dev = GPGPUSim_Init();
-	if (device <= dev->num_devices() )  {
-		*prop= *dev->get_prop();
-		return g_last_cudaError = cudaSuccess;
-	} else {
-		return g_last_cudaError = cudaErrorInvalidDevice;
-	}
+
+__host__ cudaError_t CUDARTAPI cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device)
+{
+    return cudaGetDevicePropertiesInternal(prop, device);
 }
 
 #if (CUDART_VERSION > 5000)
 __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDeviceAttr attr, int device)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-        const struct cudaDeviceProp *prop;
-        _cuda_device_id *dev = GPGPUSim_Init();
-        if (device <= dev->num_devices() )  {
-                prop = dev->get_prop();
-                switch (attr) {
-                case 1:
-                        *value= prop->maxThreadsDim[0] * prop->maxThreadsDim[1] * prop->maxThreadsDim[2] * prop->maxGridSize[0] * prop->maxGridSize[1] * prop->maxGridSize[2];
-                        break;
-                case 2:
-                        *value= prop->maxThreadsDim[0];
-                        break;
-                case 3:
-                        *value= prop->maxThreadsDim[1];
-                        break;
-                case 4:
-                        *value= prop->maxThreadsDim[2];
-                        break; 
-		case 5:
-                        *value= prop->maxGridSize[0];
-                        break;
-                case 6:
-                        *value= prop->maxGridSize[1];
-                        break;
-                case 7:
-                        *value= prop->maxGridSize[2];
-                        break;
-		case 8:
-                        *value= prop->sharedMemPerBlock;
-                        break;
-		case 9:
-                        *value= prop->totalConstMem;
-                        break;
-                case 10:
-                        *value= prop->warpSize;
-                        break;
-                case 11:
-                        *value= 16;//dummy value
-                        break;
-                case 12:
-                        *value= prop->regsPerBlock;
-                        break;
-                case 13:
-                        *value= 1480000;//for 1080ti
-                        break;
-                case 14:
-                        *value= prop->textureAlignment ;
-                        break;
-                case 15:
-                        *value = 0;
-                        break;
-                case 16:
-                        *value= prop->multiProcessorCount ;
-                        break;
-                case 17:
-                case 18:
-                case 19:
-                        *value = 0;
-                        break;
-                case 21:
-                case 22:
-                case 23:
-                case 24:
-                case 25:
-                case 26:
-                case 27:
-                case 28:
-                case 42:
-                case 45:
-                case 46:
-                case 47:
-                case 48:
-                case 49:
-                case 52:
-                case 53:
-                case 55:
-                case 56:
-                case 57:
-                case 58:
-                case 59:
-                case 60:
-                case 61:
-                case 62:
-                case 63:
-                case 64:
-                case 66:
-                case 67:
-                case 69:
-                case 70:
-                case 71:
-                case 73:
-                case 74:
-                case 77:
-                        *value = 1000;//dummy value
-                        break;
-                case 29:
-                case 43:
-                case 54:
-                case 65:
-                case 68:
-                case 72:
-                        *value = 10;//dummy value
-                        break;
-                case 30:
-                case 51:
-                        *value = 128;//dummy value
-                        break;
-                case 31:
-                        *value = 1;
-                        break;
-                case 32:
-                        *value = 0;
-                        break;
-                case 33:
-                case 50:
-                        *value = 0;//dummy value
-                        break;
-		        case 34:
-                        *value= 0;
-                        break;
-                case 35:
-                        *value = 0;
-                        break;
-                case 36:
-                        *value = 1250000;//CK value for 1080ti
-                        break;
-                case 37:
-                        *value = 352;//value for 1080ti
-                        break;
-                case 38:
-                        *value = 3000000;//value for 1080ti
-                        break;
-                case 39:
-                        *value= dev->get_gpgpu()->threads_per_core();
-                        break;
-                case 40:
-                        *value= 0;
-                        break;
-                case 41:
-                        *value= 0;
-                        break;
-                case 75://cudaDevAttrComputeCapabilityMajor
-                        *value= prop->major ;
-                        break;
-                case 76://cudaDevAttrComputeCapabilityMinor
-                        *value= prop->minor ;
-                        break;
-                case 78:
-                        *value= 0 ; //TODO: as of now, we dont support stream priorities.
-                        break;
-                case 79:
-                        *value= 0;
-                        break;
-                case 80:
-                        *value= 0;
-                        break;
-		#if (CUDART_VERSION > 5050)
-		case 81:
-                        *value= prop->sharedMemPerMultiprocessor;
-                        break;
-                case 82:
-                        *value= prop->regsPerMultiprocessor;
-                        break;
-		#endif
-                case 83:
-                case 84:
-                case 85:
-                case 86:
-                        *value= 0;
-                        break;
-                case 87:
-                        *value= 4;//dummy value
-                        break;
-                case 88:
-                case 89:
-                case 90:
-                case 91:
-                case 92:
-                case 93:
-                case 94:
-                case 95:
-                        *value= 0;
-                        break;
-		default:
-			printf("ERROR: Attribute number %d unimplemented \n",attr);
-			abort();
-                }
-                return g_last_cudaError = cudaSuccess;
-        } else {
-                return g_last_cudaError = cudaErrorInvalidDevice;
-        }
+    return cudaDeviceGetAttributeInternal(value, attr, device);
 }
 #endif
 
 __host__ cudaError_t CUDARTAPI cudaChooseDevice(int *device, const struct cudaDeviceProp *prop)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	_cuda_device_id *dev = GPGPUSim_Init();
-	*device = dev->get_id();
-	return g_last_cudaError = cudaSuccess;
+    return cudaChooseDeviceInternal(device, prop);
 }
 
 __host__ cudaError_t CUDARTAPI cudaSetDevice(int device)
@@ -1744,47 +1883,9 @@ __host__ cudaError_t CUDARTAPI cudaGetDevice(int *device)
     return cudaGetDeviceInternal(device);
 }
 
-__host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit limit )
+__host__ cudaError_t CUDARTAPI cudaDeviceGetLimit( size_t* pValue, cudaLimit limit )
 {
-	if(g_debug_execution >= 3){
-            announce_call(__my_func__);
-   	 }
-        _cuda_device_id *dev = GPGPUSim_Init();
-	const struct cudaDeviceProp *prop = dev->get_prop();
-	const gpgpu_sim_config& config=dev->get_gpgpu()->get_config();
-        switch(limit) {
-        case 0:  // cudaLimitStackSize
-        	*pValue=config.stack_limit();
-                break;
-        case 2:  // cudaLimitMallocHeapSize
-                *pValue=config.heap_limit();
-                break;
-#if (CUDART_VERSION > 5050)
-        case 3: // cudaLimitDevRuntimeSyncDepth
-		if(prop->major > 2){
-			*pValue=config.sync_depth_limit();
-	                break;
-		}
-		else{
-			printf("ERROR:Limit %s is not supported on this architecture \n",limit);
-			abort();
-		}
-        case 4: // cudaLimitDevRuntimePendingLaunchCount
-		if(prop->major > 2){
-     	        	*pValue=config.pending_launch_count_limit();
-	                break;
-		}
-		else{
-			printf("ERROR:Limit %s is not supported on this architecture \n",limit);
-			abort();
-		}
-#endif
-        default:
-                        printf("ERROR:Limit %s unimplemented \n",limit);
-                        abort();
-        }
-	return g_last_cudaError = cudaSuccess;
-
+    return cudaDeviceGetLimitInternal( pValue, limit );
 }
 
 __host__ cudaError_t CUDARTAPI cudaStreamGetPriority ( cudaStream_t hStream, int* priority )
@@ -3167,58 +3268,9 @@ cudaError_t CUDARTAPI cudaSetDeviceFlags( int flags )
     }
 }
 
-size_t getMaxThreadsPerBlock(struct cudaFuncAttributes *attr) {
-  _cuda_device_id *dev = GPGPUSim_Init();
-  struct cudaDeviceProp prop;
-
-  prop = *dev->get_prop();
-
-  size_t max = prop.maxThreadsPerBlock;
-
-  if ((prop.regsPerBlock / attr->numRegs) < max)
-    max = prop.regsPerBlock / attr->numRegs;
-
-  return max;
-}
-
 cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *attr, const char *hostFun )
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	CUctx_st *context = GPGPUSim_Context();
-	function_info *entry = context->get_kernel(hostFun);
-	if( entry ) {
-		const struct gpgpu_ptx_sim_info *kinfo = entry->get_kernel_info();
-		attr->sharedSizeBytes = kinfo->smem;
-		attr->constSizeBytes  = kinfo->cmem;
-		attr->localSizeBytes  = kinfo->lmem;
-		attr->numRegs         = kinfo->regs;
-		if(kinfo->maxthreads > 0)
-		  attr->maxThreadsPerBlock = kinfo->maxthreads;
-		else
-		  attr->maxThreadsPerBlock = getMaxThreadsPerBlock(attr);
-#if CUDART_VERSION >= 3000
-		attr->ptxVersion      = kinfo->ptx_version;
-		attr->binaryVersion   = kinfo->sm_target;
-#endif
-	}
-	return g_last_cudaError = cudaSuccess;
-}
-
-cudaError_t CUDARTAPI cudaEventCreateWithFlags(cudaEvent_t *event, int flags)
-{
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	CUevent_st *e = new CUevent_st(flags==cudaEventBlockingSync);
-	g_timer_events[e->get_uid()] = e;
-#if CUDART_VERSION >= 3000
-	*event = e;
-#else
-	*event = e->get_uid();
-#endif
-	return g_last_cudaError = cudaSuccess;
+    return cudaFuncGetAttributesInternal(attr, hostFun );
 }
 
 cudaError_t CUDARTAPI cudaDriverGetVersion(int *driverVersion)
diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h
index fa5d02b..dd8b51d 100644
--- a/libcuda/gpgpu_context.h
+++ b/libcuda/gpgpu_context.h
@@ -26,6 +26,8 @@ class gpgpu_context {
 	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*);
+	class gpgpu_sim *gpgpu_ptx_sim_init_perf();
+	struct _cuda_device_id *GPGPUSim_Init();
 };
 gpgpu_context* GPGPU_Context();
 
diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc
index 50a02fa..e91e2e0 100644
--- a/libopencl/opencl_runtime_api.cc
+++ b/libopencl/opencl_runtime_api.cc
@@ -648,7 +648,9 @@ class _cl_device_id *GPGPUSim_Init()
 {
    static _cl_device_id *the_device = NULL;
    if( !the_device ) { 
-      gpgpu_sim *the_gpu = gpgpu_ptx_sim_init_perf(); 
+       gpgpu_context *ctx;
+       ctx = GPGPU_Context();
+      gpgpu_sim *the_gpu = ctx->gpgpu_ptx_sim_init_perf(); 
       the_device = new _cl_device_id(the_gpu);
    } 
    start_sim_thread(2);
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 2872b84..9094ec3 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -48,7 +48,6 @@ void set_ptx_warp_size(const struct core_config * warp_size)
    g_shader_core_config=warp_size;
 }
 
-static bool g_debug_ir_generation=false;
 const char *g_filename;
 
 // the program intermediate representation...
@@ -70,7 +69,7 @@ const char *decode_token( int type )
    return g_ptx_token_decode[type].c_str();
 }
 
-void read_parser_environment_variables() 
+void ptx_recognizer::read_parser_environment_variables() 
 {
    g_filename = getenv("PTX_SIM_KERNELFILE"); 
    char *dbg_level = getenv("PTX_SIM_DEBUG");
diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h
index 25c01fe..bc9a872 100644
--- a/src/cuda-sim/ptx_parser.h
+++ b/src/cuda-sim/ptx_parser.h
@@ -62,6 +62,7 @@ class ptx_recognizer {
 	    g_error_detected = 0;
 	    g_entry_func_param_index=0;
 	    g_func_info = NULL;
+	    g_debug_ir_generation=false;
 	}
 	// global list
 	yyscan_t scanner;
@@ -103,6 +104,7 @@ class ptx_recognizer {
 	unsigned g_entry_func_param_index;
 	function_info *g_func_info;
 	operand_info g_return_var;
+	bool g_debug_ir_generation;
 
 	// member function list
 	void init_directive_state();
@@ -169,6 +171,7 @@ class ptx_recognizer {
 	void start_inst_group();
 	void end_inst_group();
 	bool check_for_duplicates( const char *identifier );
+	void read_parser_environment_variables(); 
 
 };
 
diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc
index b09674a..476a9d4 100644
--- a/src/gpgpusim_entrypoint.cc
+++ b/src/gpgpusim_entrypoint.cc
@@ -35,6 +35,7 @@
 #include "gpgpu-sim/gpu-sim.h"
 #include "gpgpu-sim/icnt_wrapper.h"
 #include "stream_manager.h"
+#include "../libcuda/gpgpu_context.h"
 
 #define MAX(a,b) (((a)>(b))?(a):(b))
 
@@ -208,12 +209,12 @@ void exit_simulation()
 
 extern bool g_cuda_launch_blocking;
 
-gpgpu_sim *gpgpu_ptx_sim_init_perf()
+gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf()
 {
    srand(1);
    print_splash();
    read_sim_environment_variables();
-   read_parser_environment_variables();
+   ptx_parser->read_parser_environment_variables();
    option_parser_t opp = option_parser_create();
 
    ptx_reg_options(opp);
diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h
index 9ce7fef..a443151 100644
--- a/src/gpgpusim_entrypoint.h
+++ b/src/gpgpusim_entrypoint.h
@@ -74,7 +74,6 @@ struct GPGPUsim_ctx {
 
 };
 
-class gpgpu_sim *gpgpu_ptx_sim_init_perf();
 void start_sim_thread(int api);
 
 class gpgpu_sim* g_the_gpu();
-- 
cgit v1.3


From cb60ae4893086175fee49dc6088d46665a7f088b Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Wed, 3 Jul 2019 12:04:22 -0400
Subject: move g_shader_core_config

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 src/cuda-sim/ptx_parser.cc | 4 +---
 src/cuda-sim/ptx_parser.h  | 3 +++
 src/gpgpu-sim/gpu-sim.cc   | 5 ++---
 3 files changed, 6 insertions(+), 6 deletions(-)

(limited to 'src/cuda-sim/ptx_parser.cc')

diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 9094ec3..09bc15d 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -42,8 +42,7 @@ extern FILE *ptx_get_in (yyscan_t yyscanner );
 extern int ptx_parse(yyscan_t scanner, ptx_recognizer* recognizer);
 extern int ptx_lex_destroy(yyscan_t scanner);
 
-static const struct core_config *g_shader_core_config;
-void set_ptx_warp_size(const struct core_config * warp_size)
+void ptx_recognizer::set_ptx_warp_size(const struct core_config * warp_size)
 {
    g_shader_core_config=warp_size;
 }
@@ -154,7 +153,6 @@ symbol_table * gpgpu_context::init_parser( const char *ptx_filename )
    return ptx_parser->g_global_symbol_table;
 }
 
-static int g_entry_point;
 
 void ptx_recognizer::start_function( int entry_point )
 {
diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h
index bc9a872..ec300a1 100644
--- a/src/cuda-sim/ptx_parser.h
+++ b/src/cuda-sim/ptx_parser.h
@@ -105,6 +105,8 @@ class ptx_recognizer {
 	function_info *g_func_info;
 	operand_info g_return_var;
 	bool g_debug_ir_generation;
+	int g_entry_point;
+	const struct core_config *g_shader_core_config;
 
 	// member function list
 	void init_directive_state();
@@ -172,6 +174,7 @@ class ptx_recognizer {
 	void end_inst_group();
 	bool check_for_duplicates( const char *identifier );
 	void read_parser_environment_variables(); 
+	void set_ptx_warp_size(const struct core_config * warp_size);
 
 };
 
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index c894278..2ff37d1 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -65,6 +65,7 @@
 #include "visualizer.h"
 #include "stats.h"
 #include "../cuda-sim/cuda_device_runtime.h"
+#include "../../libcuda/gpgpu_context.h"
 
 #ifdef GPGPUSIM_POWER_MODEL
 #include "power_interface.h"
@@ -697,15 +698,13 @@ void gpgpu_sim::stop_all_running_kernels(){
     }
 }
 
-void set_ptx_warp_size(const struct core_config * warp_size);
-
 gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config, gpgpu_context* ctx )
     : gpgpu_t(config), m_config(config)
 {
     gpgpu_ctx = ctx;
     m_shader_config = &m_config.m_shader_config;
     m_memory_config = &m_config.m_memory_config;
-    set_ptx_warp_size(m_shader_config);
+    ctx->ptx_parser->set_ptx_warp_size(m_shader_config);
     ptx_file_line_stats_create_exposed_latency_tracker(m_config.num_shader());
 
 #ifdef GPGPUSIM_POWER_MODEL
-- 
cgit v1.3


From 60c38b0f77378eb111e88f632702d19dc3746cc7 Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Wed, 3 Jul 2019 14:54:16 -0400
Subject: Remove g_filename

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 libcuda/gpgpu_context.h    |  5 +++--
 src/cuda-sim/ptx.l         |  5 ++---
 src/cuda-sim/ptx.y         |  5 ++---
 src/cuda-sim/ptx_loader.h  |  5 +++++
 src/cuda-sim/ptx_parser.cc | 26 ++++++++++++--------------
 src/cuda-sim/ptx_parser.h  |  7 +++++--
 src/cuda-sim/ptxinfo.l     |  4 ++--
 src/gpgpu-sim/gpu-sim.h    |  1 +
 8 files changed, 32 insertions(+), 26 deletions(-)

(limited to 'src/cuda-sim/ptx_parser.cc')

diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h
index 69341be..d819559 100644
--- a/libcuda/gpgpu_context.h
+++ b/libcuda/gpgpu_context.h
@@ -10,12 +10,13 @@ class gpgpu_context {
 	gpgpu_context() {
 	    g_global_allfiles_symbol_table = NULL;
 	    api = new cuda_runtime_api();
-	    ptxinfo = new ptxinfo_data();
-	    ptx_parser = new ptx_recognizer();
+	    ptxinfo = new ptxinfo_data(this);
+	    ptx_parser = new ptx_recognizer(this);
 	    the_gpgpusim = new GPGPUsim_ctx(this);
 	}
 	// global list
 	symbol_table *g_global_allfiles_symbol_table;
+	const char *g_filename;
 	// objects pointers for each file
 	cuda_runtime_api* api;
 	ptxinfo_data* ptxinfo;
diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index 18952a9..3a2a839 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -40,6 +40,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include "ptx_parser.h"
 #include "ptx.tab.h"
 #include <string.h>
+#include "../../libcuda/gpgpu_context.h"
 
 #define LINEBUF_SIZE (4*1024)
 #define TC recognizer->col+=strlen(yytext);
@@ -452,8 +453,6 @@ breakaddr  TC; yylval->int_value = BREAKADDR_OP; return OPCODE;
 <INITIAL,NOT_OPCODE,IN_FUNC_DECL>.  TC; ptx_error(yyscanner, recognizer, (const char*)NULL);
 %%
 
-extern const char *g_filename;
-
 int ptx_error( yyscan_t yyscanner, ptx_recognizer* recognizer, const char *s )
 {
     struct yyguts_t * yyg = (struct yyguts_t*)yyscanner;
@@ -461,7 +460,7 @@ int ptx_error( yyscan_t yyscanner, ptx_recognizer* recognizer, const char *s )
 	recognizer->g_error_detected = 1;
 	fflush(stdout);
 	if( s != NULL )
-		printf("%s:%u Syntax error:\n\n", g_filename, yylineno );
+		printf("%s:%u Syntax error:\n\n", recognizer->gpgpu_ctx->g_filename, yylineno );
 	printf("   %s\n", recognizer->linebuf );
 	printf("   ");
 	for( i=0; i < recognizer->col-1; i++ ) {
diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y
index 837fbe9..a01c3c6 100644
--- a/src/cuda-sim/ptx.y
+++ b/src/cuda-sim/ptx.y
@@ -30,6 +30,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 %{
 typedef void * yyscan_t;
 class ptx_recognizer;
+#include "../../libcuda/gpgpu_context.h"
 %}
 
 %define api.pure full
@@ -628,11 +629,9 @@ address_expression: IDENTIFIER { recognizer->add_address_operand($1,0); }
 
 %%
 
-extern const char *g_filename;
-
 void syntax_not_implemented(yyscan_t yyscanner, ptx_recognizer* recognizer)
 {
-	printf("Parse error (%s): this syntax is not (yet) implemented:\n",g_filename);
+	printf("Parse error (%s): this syntax is not (yet) implemented:\n", recognizer->gpgpu_ctx->g_filename);
 	ptx_error(yyscanner, recognizer, NULL);
 	abort();
 }
diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h
index 77f27c8..c214b95 100644
--- a/src/cuda-sim/ptx_loader.h
+++ b/src/cuda-sim/ptx_loader.h
@@ -30,13 +30,18 @@
 #include <string>
 
 #define PTXINFO_LINEBUF_SIZE 1024
+class gpgpu_context;
 typedef void * yyscan_t;
 class ptxinfo_data{
     public:
+	ptxinfo_data(gpgpu_context* ctx) {
+	    gpgpu_ctx = ctx;
+	}
 	yyscan_t scanner;
 	char linebuf[PTXINFO_LINEBUF_SIZE];
 	unsigned col;
 	const char *g_ptxinfo_filename;
+	class gpgpu_context* gpgpu_ctx;
 	void ptxinfo_addinfo();
 };
 
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 09bc15d..05fc618 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -47,14 +47,12 @@ void ptx_recognizer::set_ptx_warp_size(const struct core_config * warp_size)
    g_shader_core_config=warp_size;
 }
 
-const char *g_filename;
-
 // the program intermediate representation...
 std::map<std::string,symbol_table*> g_sym_name_to_symbol_table;
 
 #define PTX_PARSE_DPRINTF(...) \
    if( g_debug_ir_generation ) { \
-      printf(" %s:%u => ",g_filename,ptx_get_lineno(scanner)); \
+      printf(" %s:%u => ",gpgpu_ctx->g_filename,ptx_get_lineno(scanner)); \
       printf("   (%s:%u) ", __FILE__, __LINE__); \
       printf(__VA_ARGS__); \
       printf("\n"); \
@@ -70,7 +68,7 @@ const char *decode_token( int type )
 
 void ptx_recognizer::read_parser_environment_variables() 
 {
-   g_filename = getenv("PTX_SIM_KERNELFILE"); 
+   gpgpu_ctx->g_filename = getenv("PTX_SIM_KERNELFILE");
    char *dbg_level = getenv("PTX_SIM_DEBUG");
    if ( dbg_level && strlen(dbg_level) ) {
       int debug_execution=0;
@@ -180,7 +178,7 @@ void ptx_recognizer::add_function_name( const char *name )
    if( prior_decl ) {
       g_func_info->remove_args();
    }
-   g_global_symbol_table->add_function( g_func_info, g_filename, ptx_get_lineno(scanner) );
+   g_global_symbol_table->add_function( g_func_info, gpgpu_ctx->g_filename, ptx_get_lineno(scanner) );
 }
 
 //Jin: handle instruction group for cdp
@@ -229,7 +227,7 @@ void ptx_recognizer::parse_error_impl( const char *file, unsigned line, const ch
    va_end(ap);
 
    g_error_detected = 1;
-   printf("%s:%u: Parse error: %s (%s:%u)\n\n", g_filename, ptx_get_lineno(scanner), buf, file, line);
+   printf("%s:%u: Parse error: %s (%s:%u)\n\n", gpgpu_ctx->g_filename, ptx_get_lineno(scanner), buf, file, line);
    ptx_error(scanner, NULL);
    abort();
    exit(1);
@@ -284,12 +282,12 @@ void ptx_recognizer::add_instruction()
                                              g_wmma_options, 
                                              g_scalar_type,
                                              g_space_spec,
-                                             g_filename,
+                                             gpgpu_ctx->g_filename,
                                              ptx_get_lineno(scanner),
                                              linebuf,
                                              g_shader_core_config );
    g_instructions.push_back(i);
-   g_inst_lookup[g_filename][ptx_get_lineno(scanner)] = i;
+   g_inst_lookup[gpgpu_ctx->g_filename][ptx_get_lineno(scanner)] = i;
    init_instruction_state();
 }
 
@@ -391,7 +389,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi
    default:
       break;
    }
-   g_last_symbol = g_current_symbol_table->add_variable(identifier,type,num_bits/8,g_filename,ptx_get_lineno(scanner));
+   g_last_symbol = g_current_symbol_table->add_variable(identifier,type,num_bits/8,gpgpu_ctx->g_filename,ptx_get_lineno(scanner));
    switch ( ti.get_memory_space().get_type() ) {
    case reg_space: {
       regnum = g_current_symbol_table->next_reg_num();
@@ -654,7 +652,7 @@ void ptx_recognizer::add_label( const char *identifier )
    if ( s != NULL ) {
       g_label = s;
    } else {
-      g_label = g_current_symbol_table->add_variable(identifier,NULL,0,g_filename,ptx_get_lineno(scanner));
+      g_label = g_current_symbol_table->add_variable(identifier,NULL,0,gpgpu_ctx->g_filename,ptx_get_lineno(scanner));
    }
 }
 
@@ -912,7 +910,7 @@ void ptx_recognizer::add_scalar_operand( const char *identifier )
    if ( s == NULL ) {
       if ( g_opcode == BRA_OP || g_opcode == CALLP_OP) {
          // forward branch target...
-         s = g_current_symbol_table->add_variable(identifier,NULL,0,g_filename,ptx_get_lineno(scanner));
+         s = g_current_symbol_table->add_variable(identifier,NULL,0,gpgpu_ctx->g_filename,ptx_get_lineno(scanner));
       } else {
          std::string msg = std::string("operand \"") + identifier + "\" has no declaration.";
          parse_error( msg.c_str() );
@@ -926,7 +924,7 @@ void ptx_recognizer::add_neg_pred_operand( const char *identifier )
    PTX_PARSE_DPRINTF("add_neg_pred_operand");
    const symbol *s = g_current_symbol_table->lookup(identifier);
    if ( s == NULL ) {
-       s = g_current_symbol_table->add_variable(identifier,NULL,1,g_filename,ptx_get_lineno(scanner));
+       s = g_current_symbol_table->add_variable(identifier,NULL,1,gpgpu_ctx->g_filename,ptx_get_lineno(scanner));
    }
    operand_info op(s);
    op.set_neg_pred();
@@ -962,7 +960,7 @@ void ptx_recognizer::add_version_info( float ver, unsigned ext )
 
 void ptx_recognizer::add_file( unsigned num, const char *filename )
 {
-   if( g_filename == NULL ) {
+   if( gpgpu_ctx->g_filename == NULL ) {
       char *b = strdup(filename);
       char *l=b;
       char *n=b;
@@ -978,7 +976,7 @@ void ptx_recognizer::add_file( unsigned num, const char *filename )
 
       char *q = strtok(NULL,".");
       if( q && !strcmp(q,"cu") ) {
-          g_filename = strdup(buf);
+          gpgpu_ctx->g_filename = strdup(buf);
       }
 
       free( b );
diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h
index ec300a1..e09aab1 100644
--- a/src/cuda-sim/ptx_parser.h
+++ b/src/cuda-sim/ptx_parser.h
@@ -30,17 +30,17 @@
 
 #include "../abstract_hardware_model.h"
 #include "ptx_ir.h"
-extern const char *g_filename;
 extern int g_error_detected;
 
 #ifdef __cplusplus 
 const class ptx_instruction *ptx_instruction_lookup( const char *filename, unsigned linenumber );
 #endif
 
+class gpgpu_context;
 typedef void * yyscan_t;
 class ptx_recognizer {
     public:
-	ptx_recognizer() {
+	ptx_recognizer( gpgpu_context* ctx ) {
 	    scanner = NULL;
 	    g_size = -1;
 	    g_add_identifier_cached__identifier = NULL;
@@ -63,6 +63,7 @@ class ptx_recognizer {
 	    g_entry_func_param_index=0;
 	    g_func_info = NULL;
 	    g_debug_ir_generation=false;
+	    gpgpu_ctx = ctx;
 	}
 	// global list
 	yyscan_t scanner;
@@ -107,6 +108,8 @@ class ptx_recognizer {
 	bool g_debug_ir_generation;
 	int g_entry_point;
 	const struct core_config *g_shader_core_config;
+	// backward pointer
+	class gpgpu_context* gpgpu_ctx;
 
 	// member function list
 	void init_directive_state();
diff --git a/src/cuda-sim/ptxinfo.l b/src/cuda-sim/ptxinfo.l
index b0ada09..3a152b0 100644
--- a/src/cuda-sim/ptxinfo.l
+++ b/src/cuda-sim/ptxinfo.l
@@ -39,6 +39,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include "ptx_loader.h"
 #include "ptxinfo.tab.h"
 #include <string.h>
+#include "../../libcuda/gpgpu_context.h"
 
 #define LINEBUF_SIZE 1024
 #define TC if( (yylineno == 1) && (ptxinfo->col + strlen(yytext) < LINEBUF_SIZE) ) { \
@@ -88,7 +89,6 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 %%
 
 extern int g_ptxinfo_error_detected;
-extern const char *g_filename;
 
 int ptxinfo_error(yyscan_t yyscanner, ptxinfo_data* ptxinfo, const char* msg)
 {
@@ -98,7 +98,7 @@ int ptxinfo_error(yyscan_t yyscanner, ptxinfo_data* ptxinfo, const char* msg)
 	fflush(stdout);
 	printf("GPGPU-Sim: ERROR while parsing output of ptxas (used to capture resource usage information)\n");
 	if( msg != NULL )
-		printf("GPGPU-Sim:     %s (%s:%u) Syntax error:\n\n", g_filename, ptxinfo->g_ptxinfo_filename, yylineno );
+		printf("GPGPU-Sim:     %s (%s:%u) Syntax error:\n\n", ptxinfo->gpgpu_ctx->g_filename, ptxinfo->g_ptxinfo_filename, yylineno );
 	printf("   %s\n", ptxinfo->linebuf );
 	printf("   ");
 	for( i=0; i < ptxinfo->col-1; i++ ) {
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index b1d2e45..e2c913a 100644
--- a/src/gpgpu-sim/gpu-sim.h
+++ b/src/gpgpu-sim/gpu-sim.h
@@ -506,6 +506,7 @@ private:
    void gpgpu_debug();
 
 ///// data /////
+// backward pointer
    class gpgpu_context* gpgpu_ctx;
 
    class simt_core_cluster **m_cluster;
-- 
cgit v1.3


From b3e786e3d8d720217f36a214e9b5be9a19ab9dd2 Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Mon, 8 Jul 2019 12:37:52 -0400
Subject: Move opcode_latency_int thus pass gpgpu_context into many classes

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 libcuda/gpgpu_context.h           |  3 +++
 src/cuda-sim/cuda-sim.cc          |  7 ++++---
 src/cuda-sim/cuda-sim.h           | 11 +++++++++--
 src/cuda-sim/ptx_ir.cc            |  4 +++-
 src/cuda-sim/ptx_ir.h             |  7 ++++++-
 src/cuda-sim/ptx_parser.cc        |  3 ++-
 src/gpgpu-sim/gpu-sim.cc          |  2 +-
 src/gpgpu-sim/gpu-sim.h           | 11 ++++++++---
 src/gpgpu-sim/mem_latency_stat.cc |  2 +-
 src/gpgpu-sim/mem_latency_stat.h  |  4 ++--
 src/gpgpu-sim/power_interface.cc  |  2 +-
 src/gpgpu-sim/power_interface.h   |  2 +-
 src/gpgpu-sim/power_stat.cc       |  6 +++---
 src/gpgpu-sim/power_stat.h        |  6 +++---
 src/gpgpu-sim/shader.cc           |  7 ++++---
 src/gpgpu-sim/shader.h            | 18 ++++++++++++------
 src/gpgpusim_entrypoint.cc        |  4 ++--
 17 files changed, 65 insertions(+), 34 deletions(-)

(limited to 'src/cuda-sim/ptx_parser.cc')

diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h
index a8b60f4..2e21009 100644
--- a/libcuda/gpgpu_context.h
+++ b/libcuda/gpgpu_context.h
@@ -4,6 +4,7 @@
 #include "../src/cuda-sim/ptx_loader.h"
 #include "../src/cuda-sim/ptx_parser.h"
 #include "../src/gpgpusim_entrypoint.h"
+#include "../src/cuda-sim/cuda-sim.h"
 
 class gpgpu_context {
     public:
@@ -13,6 +14,7 @@ class gpgpu_context {
 	    ptxinfo = new ptxinfo_data(this);
 	    ptx_parser = new ptx_recognizer(this);
 	    the_gpgpusim = new GPGPUsim_ctx(this);
+	    func_sim = new cuda_sim();
 	}
 	// global list
 	symbol_table *g_global_allfiles_symbol_table;
@@ -22,6 +24,7 @@ class gpgpu_context {
 	ptxinfo_data* ptxinfo;
 	ptx_recognizer* ptx_parser;
 	GPGPUsim_ctx* the_gpgpusim;
+	cuda_sim* func_sim;
 	// member function list
 	void cuobjdumpParseBinary(unsigned int handle);
 	class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num );
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 261d605..df0bbd7 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -51,6 +51,7 @@ typedef void * yyscan_t;
 #include "decuda_pred_table/decuda_pred_table.h"
 #include "../stream_manager.h"
 #include "cuda_device_runtime.h"
+#include "../../libcuda/gpgpu_context.h"
 
 int gpgpu_ptx_instruction_classification;
 void ** g_inst_classification_stat = NULL;
@@ -66,12 +67,12 @@ int cp_cta_resume;
 unsigned g_ptx_sim_num_insn = 0;
 unsigned gpgpu_param_num_shaders = 0;
 
-char *opcode_latency_int, *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor;
+char *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor;
 char *opcode_initiation_int, *opcode_initiation_fp, *opcode_initiation_dp,*opcode_initiation_sfu,*opcode_initiation_tensor;
 char *cdp_latency_str;
 unsigned cdp_latency[5];
 
-void ptx_opcocde_latency_options (option_parser_t opp) {
+void cuda_sim::ptx_opcocde_latency_options (option_parser_t opp) {
 	option_parser_register(opp, "-ptx_opcode_latency_int", OPT_CSTR, &opcode_latency_int,
 			"Opcode latencies for integers <ADD,MAX,MUL,MAD,DIV>"
 			"Default 1,1,19,25,145",
@@ -667,7 +668,7 @@ void ptx_instruction::set_opcode_and_latency()
 	 * [3] MAD
 	 * [4] DIV
 	 */
-	sscanf(opcode_latency_int, "%u,%u,%u,%u,%u",
+	sscanf(gpgpu_ctx->func_sim->opcode_latency_int, "%u,%u,%u,%u,%u",
 			&int_latency[0],&int_latency[1],&int_latency[2],
 			&int_latency[3],&int_latency[4]);
 	sscanf(opcode_latency_fp, "%u,%u,%u,%u,%u",
diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h
index e690356..96d34f5 100644
--- a/src/cuda-sim/cuda-sim.h
+++ b/src/cuda-sim/cuda-sim.h
@@ -47,10 +47,9 @@ extern int g_debug_thread_uid;
 extern void ** g_inst_classification_stat;
 extern void ** g_inst_op_classification_stat;
 extern int g_ptx_kernel_count; // used for classification stat collection purposes 
-extern char *opcode_latency_int, *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor;
+extern char *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor;
 
 
-void ptx_opcocde_latency_options (option_parser_t opp);
 extern class kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,
                                             gpgpu_ptx_sim_arg_list_t args, 
                                             struct dim3 gridDim, 
@@ -134,4 +133,12 @@ void print_ptxinfo();
 void clear_ptxinfo();
 struct gpgpu_ptx_sim_info get_ptxinfo();
 
+class cuda_sim {
+    public:
+	//global variables
+	char *opcode_latency_int;
+	//global functions
+	void ptx_opcocde_latency_options (option_parser_t opp);
+};
+
 #endif
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index 8cedf79..1bd409e 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -1095,8 +1095,10 @@ ptx_instruction::ptx_instruction( int opcode,
                                   const char *file, 
                                   unsigned line,
                                   const char *source,
-                                  const core_config *config ) : warp_inst_t(config)
+                                  const core_config *config,
+				  gpgpu_context* ctx ) : warp_inst_t(config)
 {
+   gpgpu_ctx = ctx;
    m_uid = ++g_num_ptx_inst_uid;
    m_PC = 0;
    m_opcode = opcode;
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index fd869c6..1604551 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -43,6 +43,8 @@
 
 #include "memory.h"
 
+class gpgpu_context;
+
 class type_info_key {
 public:
    type_info_key()
@@ -931,7 +933,8 @@ public:
                     const char *file, 
                     unsigned line,
                     const char *source,
-                    const core_config *config );
+                    const core_config *config,
+		    gpgpu_context* ctx);
 
    void print_insn() const;
    virtual void print_insn( FILE *fp ) const;
@@ -1187,6 +1190,8 @@ private:
    virtual void pre_decode();
    friend class function_info;
    static unsigned g_num_ptx_inst_uid;
+   // backward pointer
+   class gpgpu_context* gpgpu_ctx;
 };
 
 class param_info {
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 05fc618..0139534 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -285,7 +285,8 @@ void ptx_recognizer::add_instruction()
                                              gpgpu_ctx->g_filename,
                                              ptx_get_lineno(scanner),
                                              linebuf,
-                                             g_shader_core_config );
+                                             g_shader_core_config,
+					     gpgpu_ctx );
    g_instructions.push_back(i);
    g_inst_lookup[gpgpu_ctx->g_filename][ptx_get_lineno(scanner)] = i;
    init_instruction_state();
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index 2ff37d1..39acdd9 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -1826,7 +1826,7 @@ void gpgpu_sim::dump_pipeline( int mask, int s, int m ) const
    fflush(stdout);
 }
 
-const struct shader_core_config * gpgpu_sim::getShaderCoreConfig()
+const shader_core_config * gpgpu_sim::getShaderCoreConfig()
 {
    return m_shader_config;
 }
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index e2c913a..7eeb7dd 100644
--- a/src/gpgpu-sim/gpu-sim.h
+++ b/src/gpgpu-sim/gpu-sim.h
@@ -295,7 +295,10 @@ extern bool g_interactive_debugger_enabled;
 
 class gpgpu_sim_config : public power_config, public gpgpu_functional_sim_config {
 public:
-    gpgpu_sim_config() { m_valid = false; }
+    gpgpu_sim_config(gpgpu_context* ctx): m_shader_config(ctx) {
+	m_valid = false;
+	gpgpu_ctx = ctx;
+    }
     void reg_options(class OptionParser * opp);
     void init() 
     {
@@ -341,6 +344,8 @@ private:
     void init_clock_domains(void ); 
 
 
+    // backward pointer
+    class gpgpu_context* gpgpu_ctx;
     bool m_valid;
     shader_core_config m_shader_config;
     memory_config m_memory_config;
@@ -473,7 +478,7 @@ public:
    /*!
     * Returning the configuration of the shader core, used by the functional simulation only so far
     */
-   const struct shader_core_config * getShaderCoreConfig();
+   const shader_core_config * getShaderCoreConfig();
    
    
    //! Get shader core Memory Configuration
@@ -537,7 +542,7 @@ private:
    const gpgpu_sim_config &m_config;
   
    const struct cudaDeviceProp     *m_cuda_properties;
-   const struct shader_core_config *m_shader_config;
+   const shader_core_config *m_shader_config;
    const struct memory_config      *m_memory_config;
 
    // stats
diff --git a/src/gpgpu-sim/mem_latency_stat.cc b/src/gpgpu-sim/mem_latency_stat.cc
index 04dc75b..d08ba39 100644
--- a/src/gpgpu-sim/mem_latency_stat.cc
+++ b/src/gpgpu-sim/mem_latency_stat.cc
@@ -42,7 +42,7 @@
 #include <stdlib.h>
 #include <stdio.h>
 
-memory_stats_t::memory_stats_t( unsigned n_shader, const struct shader_core_config *shader_config, const struct memory_config *mem_config, const class gpgpu_sim* gpu )
+memory_stats_t::memory_stats_t( unsigned n_shader, const shader_core_config *shader_config, const struct memory_config *mem_config, const class gpgpu_sim* gpu )
 {
    assert( mem_config->m_valid );
    assert( shader_config->m_valid );
diff --git a/src/gpgpu-sim/mem_latency_stat.h b/src/gpgpu-sim/mem_latency_stat.h
index b9285c1..6ce568d 100644
--- a/src/gpgpu-sim/mem_latency_stat.h
+++ b/src/gpgpu-sim/mem_latency_stat.h
@@ -35,7 +35,7 @@
 class memory_stats_t {
 public:
    memory_stats_t( unsigned n_shader, 
-                   const struct shader_core_config *shader_config, 
+                   const class shader_core_config *shader_config, 
                    const struct memory_config *mem_config,
 				   const class gpgpu_sim* gpu);
 
@@ -53,7 +53,7 @@ public:
 
    unsigned m_n_shader;
 
-   const struct shader_core_config *m_shader_config;
+   const shader_core_config *m_shader_config;
    const struct memory_config *m_memory_config;
    const class gpgpu_sim* m_gpu;
 
diff --git a/src/gpgpu-sim/power_interface.cc b/src/gpgpu-sim/power_interface.cc
index 3861b6a..0272aa6 100644
--- a/src/gpgpu-sim/power_interface.cc
+++ b/src/gpgpu-sim/power_interface.cc
@@ -38,7 +38,7 @@ void init_mcpat(const gpgpu_sim_config &config, class gpgpu_sim_wrapper *wrapper
 
 }
 
-void mcpat_cycle(const gpgpu_sim_config &config, const struct shader_core_config *shdr_config, class gpgpu_sim_wrapper *wrapper, class power_stat_t *power_stats, unsigned stat_sample_freq, unsigned tot_cycle, unsigned cycle, unsigned tot_inst, unsigned inst){
+void mcpat_cycle(const gpgpu_sim_config &config, const shader_core_config *shdr_config, class gpgpu_sim_wrapper *wrapper, class power_stat_t *power_stats, unsigned stat_sample_freq, unsigned tot_cycle, unsigned cycle, unsigned tot_inst, unsigned inst){
 
 	static bool mcpat_init=true;
 
diff --git a/src/gpgpu-sim/power_interface.h b/src/gpgpu-sim/power_interface.h
index afac22b..a388c23 100644
--- a/src/gpgpu-sim/power_interface.h
+++ b/src/gpgpu-sim/power_interface.h
@@ -36,7 +36,7 @@
 #include "gpgpu_sim_wrapper.h"
 
 void init_mcpat(const gpgpu_sim_config &config, class gpgpu_sim_wrapper *wrapper, unsigned stat_sample_freq, unsigned tot_inst, unsigned inst);
-void mcpat_cycle(const gpgpu_sim_config &config, const struct shader_core_config *shdr_config, class gpgpu_sim_wrapper *wrapper, class power_stat_t *power_stats,
+void mcpat_cycle(const gpgpu_sim_config &config, const shader_core_config *shdr_config, class gpgpu_sim_wrapper *wrapper, class power_stat_t *power_stats,
         unsigned stat_sample_freq, unsigned tot_cycle, unsigned cycle, unsigned tot_inst, unsigned inst);
 void mcpat_reset_perf_count(class gpgpu_sim_wrapper *wrapper);
 
diff --git a/src/gpgpu-sim/power_stat.cc b/src/gpgpu-sim/power_stat.cc
index 4c995e9..007b4c6 100644
--- a/src/gpgpu-sim/power_stat.cc
+++ b/src/gpgpu-sim/power_stat.cc
@@ -42,7 +42,7 @@
 
 
 
-power_mem_stat_t::power_mem_stat_t(const struct memory_config *mem_config, const struct shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats){
+power_mem_stat_t::power_mem_stat_t(const struct memory_config *mem_config, const shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats){
 	   assert( mem_config->m_valid );
 	   m_mem_stats = mem_stats;
 	   m_config = mem_config;
@@ -125,7 +125,7 @@ void power_mem_stat_t::print (FILE *fout) const {
 }
 
 
-power_core_stat_t::power_core_stat_t( const struct shader_core_config *shader_config, shader_core_stats *core_stats )
+power_core_stat_t::power_core_stat_t( const shader_core_config *shader_config, shader_core_stats *core_stats )
 {
      	assert( shader_config->m_valid );
         m_config = shader_config;
@@ -266,7 +266,7 @@ for(unsigned i=0; i<m_config->num_shader(); ++i){
     }
 }
 
-power_stat_t::power_stat_t( const struct shader_core_config *shader_config,float * average_pipeline_duty_cycle,float *active_sms,shader_core_stats * shader_stats, const struct memory_config *mem_config,memory_stats_t * memory_stats)
+power_stat_t::power_stat_t( const shader_core_config *shader_config,float * average_pipeline_duty_cycle,float *active_sms,shader_core_stats * shader_stats, const struct memory_config *mem_config,memory_stats_t * memory_stats)
 {
 	assert( shader_config->m_valid );
 	assert( mem_config->m_valid );
diff --git a/src/gpgpu-sim/power_stat.h b/src/gpgpu-sim/power_stat.h
index 20af2e5..91fade9 100644
--- a/src/gpgpu-sim/power_stat.h
+++ b/src/gpgpu-sim/power_stat.h
@@ -73,7 +73,7 @@ struct shader_core_power_stats_pod {
 
 class power_core_stat_t : public shader_core_power_stats_pod {
 public:
-   power_core_stat_t(const struct shader_core_config *shader_config, shader_core_stats *core_stats);
+   power_core_stat_t(const shader_core_config *shader_config, shader_core_stats *core_stats);
    void visualizer_print( gzFile visualizer_file );
    void print (FILE *fout);
    void init();
@@ -113,7 +113,7 @@ struct mem_power_stats_pod{
 
 class power_mem_stat_t : public mem_power_stats_pod{
 public:
-   power_mem_stat_t(const struct memory_config *mem_config, const struct shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats);
+   power_mem_stat_t(const struct memory_config *mem_config, const shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats);
    void visualizer_print( gzFile visualizer_file );
    void print (FILE *fout) const;
    void init();
@@ -128,7 +128,7 @@ private:
 
 class power_stat_t {
 public:
-   power_stat_t( const struct shader_core_config *shader_config,float * average_pipeline_duty_cycle,float * active_sms,shader_core_stats * shader_stats, const struct memory_config *mem_config,memory_stats_t * memory_stats);
+   power_stat_t( const shader_core_config *shader_config,float * average_pipeline_duty_cycle,float * active_sms,shader_core_stats * shader_stats, const struct memory_config *mem_config,memory_stats_t * memory_stats);
    void visualizer_print( gzFile visualizer_file );
    void print (FILE *fout) const;
    void save_stats(){
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index 69b619a..4d12068 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -46,6 +46,7 @@
 #include <limits.h>
 #include "traffic_breakdown.h"
 #include "shader_trace.h"
+#include "../../libcuda/gpgpu_context.h"
 
 #define PRIORITIZE_MSHR_OVER_WB 1
 #define MAX(a,b) (((a)>(b))?(a):(b))
@@ -69,7 +70,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu,
                                   class simt_core_cluster *cluster,
                                   unsigned shader_id,
                                   unsigned tpc_id,
-                                  const struct shader_core_config *config,
+                                  const shader_core_config *config,
                                   const struct memory_config *mem_config,
                                   shader_core_stats *stats )
    : core_t( gpu, NULL, config->warp_size, config->n_thread_per_shader ),
@@ -3018,7 +3019,7 @@ void shader_core_config::set_pipeline_latency() {
 			 * [3] MAD
 			 * [4] DIV
 			 */
-			sscanf(opcode_latency_int, "%u,%u,%u,%u,%u",
+			sscanf(gpgpu_ctx->func_sim->opcode_latency_int, "%u,%u,%u,%u,%u",
 					&int_latency[0],&int_latency[1],&int_latency[2],
 					&int_latency[3],&int_latency[4]);
 			sscanf(opcode_latency_fp, "%u,%u,%u,%u,%u",
@@ -3786,7 +3787,7 @@ void opndcoll_rfu_t::collector_unit_t::dispatch()
 
 simt_core_cluster::simt_core_cluster( class gpgpu_sim *gpu, 
                                       unsigned cluster_id, 
-                                      const struct shader_core_config *config, 
+                                      const shader_core_config *config, 
                                       const struct memory_config *mem_config,
                                       shader_core_stats *stats, 
                                       class memory_stats_t *mstats )
diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h
index 25b9607..e0cefac 100644
--- a/src/gpgpu-sim/shader.h
+++ b/src/gpgpu-sim/shader.h
@@ -69,6 +69,8 @@
 
 #define WRITE_MASK_SIZE 8
 
+class gpgpu_context;
+
 enum exec_unit_type_t
 {
   NONE = 0,
@@ -294,7 +296,7 @@ typedef std::bitset<WARP_PER_CTA_MAX> warp_set_t;
 int register_bank(int regnum, int wid, unsigned num_banks, unsigned bank_warp_shift, bool sub_core_model, unsigned banks_per_sched, unsigned sched_id );
 
 class shader_core_ctx;
-struct shader_core_config;
+class shader_core_config;
 class shader_core_stats;
 
 enum scheduler_prioritization_type
@@ -1032,7 +1034,7 @@ struct ifetch_buffer_t {
     unsigned m_warp_id;
 };
 
-struct shader_core_config;
+class shader_core_config;
 
 class simd_function_unit {
 public:
@@ -1362,10 +1364,12 @@ const char* const pipeline_stage_name_decode[] = {
     "N_PIPELINE_STAGES" 
 };
 
-struct shader_core_config : public core_config
+class shader_core_config : public core_config
 {
-    shader_core_config(){
+    public:
+    shader_core_config(gpgpu_context* ctx){
 	pipeline_widths_string = NULL;
+	gpgpu_ctx = ctx;
     }
 
     void init()
@@ -1425,6 +1429,8 @@ struct shader_core_config : public core_config
     unsigned cid_to_sid( unsigned cid, unsigned cluster_id ) const { return cluster_id*n_simt_cores_per_cluster + cid; }
     void set_pipeline_latency();
 
+    // backward pointer
+    class gpgpu_context* gpgpu_ctx;
 // data
     char *gpgpu_shader_core_pipeline_opt;
     bool gpgpu_perfect_mem;
@@ -1770,7 +1776,7 @@ public:
                      class simt_core_cluster *cluster,
                      unsigned shader_id,
                      unsigned tpc_id,
-                     const struct shader_core_config *config,
+                     const shader_core_config *config,
                      const struct memory_config *mem_config,
                      shader_core_stats *stats );
 
@@ -2065,7 +2071,7 @@ class simt_core_cluster {
 public:
     simt_core_cluster( class gpgpu_sim *gpu, 
                        unsigned cluster_id, 
-                       const struct shader_core_config *config, 
+                       const shader_core_config *config, 
                        const struct memory_config *mem_config,
                        shader_core_stats *stats,
                        memory_stats_t *mstats );
diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc
index b54f20c..5018305 100644
--- a/src/gpgpusim_entrypoint.cc
+++ b/src/gpgpusim_entrypoint.cc
@@ -218,10 +218,10 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf()
    option_parser_t opp = option_parser_create();
 
    ptx_reg_options(opp);
-   ptx_opcocde_latency_options(opp);
+   func_sim->ptx_opcocde_latency_options(opp);
 
    icnt_reg_options(opp);
-   GPGPUsim_ctx_ptr()->g_the_gpu_config = new gpgpu_sim_config();
+   GPGPUsim_ctx_ptr()->g_the_gpu_config = new gpgpu_sim_config(this);
    GPGPUsim_ctx_ptr()->g_the_gpu_config->reg_options(opp); // register GPU microrachitecture options
 
    option_parser_cmdline(opp, sg_argc, sg_argv); // parse configuration options
-- 
cgit v1.3


From c58d52ccde1eea6c45d0f8d2d5411bcf1ebab4f9 Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Wed, 10 Jul 2019 00:45:39 -0400
Subject: Move g_globals and g_constants

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 libopencl/opencl_runtime_api.cc | 10 +++++-----
 src/cuda-sim/cuda-sim.cc        |  5 +----
 src/cuda-sim/cuda-sim.h         |  4 +++-
 src/cuda-sim/ptx_parser.cc      |  7 ++-----
 src/stream_manager.cc           |  5 +++--
 5 files changed, 14 insertions(+), 17 deletions(-)

(limited to 'src/cuda-sim/ptx_parser.cc')

diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc
index aaaec4f..7f029c7 100644
--- a/libopencl/opencl_runtime_api.cc
+++ b/libopencl/opencl_runtime_api.cc
@@ -950,11 +950,11 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
 
    gpgpu_t *gpu = command_queue->get_device()->the_device();
    if (kernel->get_implementation()->get_ptx_version().ver() <3.0){
-	   gpgpu_ptx_sim_memcpy_symbol( "%_global_size", _global_size, 3 * sizeof(int), 0, 1, gpu );
-	   gpgpu_ptx_sim_memcpy_symbol( "%_work_dim", &work_dim, 1 * sizeof(int), 0, 1, gpu  );
-	   gpgpu_ptx_sim_memcpy_symbol( "%_global_num_groups", &GridDim, 3 * sizeof(int), 0, 1, gpu );
-	   gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
-	   gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
+	   ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_size", _global_size, 3 * sizeof(int), 0, 1, gpu );
+	   ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_work_dim", &work_dim, 1 * sizeof(int), 0, 1, gpu  );
+	   ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_num_groups", &GridDim, 3 * sizeof(int), 0, 1, gpu );
+	   ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
+	   ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
    }
    kernel_info_t *grid = ctx->func_sim->gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu);
    if ( ctx->func_sim->g_ptx_sim_mode )
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 7e8aab9..3d1da62 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -1942,9 +1942,6 @@ void print_splash()
    }
 }
 
-std::set<std::string>   g_globals;
-std::set<std::string>   g_constants;
-
 void cuda_sim::gpgpu_ptx_sim_register_const_variable(void *hostVar, const char *deviceName, size_t size )
 {
    printf("GPGPU-Sim PTX registering constant %s (%zu bytes) to name mapping\n", deviceName, size );
@@ -1957,7 +1954,7 @@ void cuda_sim::gpgpu_ptx_sim_register_global_variable(void *hostVar, const char
    g_global_name_lookup[hostVar] = deviceName;
 }
 
-void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu )
+void cuda_sim::gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu )
 {
    printf("GPGPU-Sim PTX: starting gpgpu_ptx_sim_memcpy_symbol with hostVar 0x%p\n", hostVar);
    bool found_sym = false;
diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h
index 628e434..5c95100 100644
--- a/src/cuda-sim/cuda-sim.h
+++ b/src/cuda-sim/cuda-sim.h
@@ -46,7 +46,6 @@ extern int g_debug_execution;
 extern int g_debug_thread_uid;
 
 extern void   print_splash();
-extern void   gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu );
 
 extern void ptxinfo_opencl_addinfo( std::map<std::string,function_info*> &kernels );
 unsigned ptx_sim_init_thread( kernel_info_t &kernel,
@@ -161,6 +160,8 @@ class cuda_sim {
 	bool g_cuda_launch_blocking;
 	void ** g_inst_classification_stat;
 	void ** g_inst_op_classification_stat;
+	std::set<std::string>   g_globals;
+	std::set<std::string>   g_constants;
 	// backward pointer
 	class gpgpu_context* gpgpu_ctx;
 	//global functions
@@ -179,6 +180,7 @@ class cuda_sim {
 	void set_param_gpgpu_num_shaders(int num_shaders);
 	struct rec_pts find_reconvergence_points( function_info *finfo );
 	address_type get_converge_point( address_type pc );
+	void   gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu );
 };
 
 #endif
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 0139534..269ec4d 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -322,9 +322,6 @@ bool ptx_recognizer::check_for_duplicates( const char *identifier )
    return ( s != NULL );
 }
 
-extern std::set<std::string>   g_globals;
-extern std::set<std::string>   g_constants;
-
 // Returns padding that needs to be inserted ahead of address to make it aligned to min(size, maxalign)
 /*
  * @param address the address in bytes
@@ -453,7 +450,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi
          g_current_symbol_table->alloc_global( num_bits/8 + addr_pad ); 
       }
       if( g_current_symbol_table == g_global_symbol_table ) { 
-         g_constants.insert( identifier ); 
+         gpgpu_ctx->func_sim->g_constants.insert( identifier );
       }
       assert( g_current_symbol_table != NULL );
       g_sym_name_to_symbol_table[ identifier ] = g_current_symbol_table;
@@ -471,7 +468,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi
       fflush(stdout);
       g_last_symbol->set_address( addr+addr_pad );
       g_current_symbol_table->alloc_global( num_bits/8 + addr_pad );
-      g_globals.insert( identifier );
+      gpgpu_ctx->func_sim->g_globals.insert( identifier );
       assert( g_current_symbol_table != NULL );
       g_sym_name_to_symbol_table[ identifier ] = g_current_symbol_table;
       break;
diff --git a/src/stream_manager.cc b/src/stream_manager.cc
index e07f4e4..be3dd71 100644
--- a/src/stream_manager.cc
+++ b/src/stream_manager.cc
@@ -29,6 +29,7 @@
 #include "gpgpusim_entrypoint.h"
 #include "cuda-sim/cuda-sim.h"
 #include "gpgpu-sim/gpu-sim.h"
+#include "../libcuda/gpgpu_context.h"
 
 unsigned CUstream_st::sm_next_stream_uid = 0;
 
@@ -150,13 +151,13 @@ bool stream_operation::do_operation( gpgpu_sim *gpu )
     case stream_memcpy_to_symbol:
         if(g_debug_execution >= 3)
             printf("memcpy to symbol\n");
-        gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_src,m_cnt,m_offset,1,gpu);
+        gpu->gpgpu_ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_src,m_cnt,m_offset,1,gpu);
         m_stream->record_next_done();
         break;
     case stream_memcpy_from_symbol:
         if(g_debug_execution >= 3)
             printf("memcpy from symbol\n");
-        gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_dst,m_cnt,m_offset,0,gpu);
+        gpu->gpgpu_ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_dst,m_cnt,m_offset,0,gpu);
         m_stream->record_next_done();
         break;
     case stream_kernel_launch:
-- 
cgit v1.3


From 87726d32ada00fcd93f2cf24ccae4ba593c4f9ec Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Wed, 10 Jul 2019 15:03:16 -0400
Subject: Move g_pc_to_finfo

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 src/abstract_hardware_model.cc |  3 ++-
 src/abstract_hardware_model.h  | 13 ++++++++-----
 src/cuda-sim/cuda-sim.cc       |  7 +++----
 src/cuda-sim/cuda-sim.h        |  5 +++--
 src/cuda-sim/ptx_ir.cc         | 12 +++++++-----
 src/cuda-sim/ptx_ir.h          |  9 +++++++--
 src/cuda-sim/ptx_parser.cc     |  2 +-
 src/gpgpu-sim/shader.cc        |  4 ++--
 src/gpgpu-sim/shader.h         |  2 +-
 9 files changed, 34 insertions(+), 23 deletions(-)

(limited to 'src/cuda-sim/ptx_parser.cc')

diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc
index 248e7a5..ef09051 100644
--- a/src/abstract_hardware_model.cc
+++ b/src/abstract_hardware_model.cc
@@ -39,6 +39,7 @@
 #include <sys/stat.h>
 #include <sstream>
 #include <iostream>
+#include "../libcuda/gpgpu_context.h"
 
 unsigned mem_access_t::sm_next_access_uid = 0;   
 unsigned warp_inst_t::sm_next_uid = 0;
@@ -945,7 +946,7 @@ void simt_stack::print (FILE *fout) const
         } else {
             fprintf(fout," " );
         }
-        ptx_print_insn( stack_entry.m_pc, fout );
+        m_gpu->gpgpu_ctx->func_sim->ptx_print_insn( stack_entry.m_pc, fout );
         fprintf(fout,"\n");
     }
 
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index da29a11..8ef8376 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -347,9 +347,11 @@ public:
    mutable bool volta_cache_config_set;
 };
 
-struct core_config {
-    core_config() 
-    { 
+class core_config {
+    public:
+    core_config(gpgpu_context* ctx)
+    {
+	gpgpu_ctx = ctx;
         m_valid = false; 
         num_shmem_bank=16; 
         shmem_limited_broadcast = false; 
@@ -361,6 +363,8 @@ struct core_config {
 
     bool m_valid;
     unsigned warp_size;
+    // backward pointer
+    class gpgpu_context* gpgpu_ctx;
 
     // off-chip memory request architecture parameters
     int gpgpu_coalesce_arch;
@@ -934,7 +938,7 @@ public:
         m_empty=true; 
         m_config=NULL; 
     }
-    warp_inst_t( const core_config *config ) 
+    warp_inst_t( const core_config *config )
     { 
         m_uid=0;
         assert(config->warp_size<=MAX_WARP_SIZE); 
@@ -1105,7 +1109,6 @@ public:
     unsigned get_uid() const { return m_uid; }
     unsigned get_schd_id() const { return m_scheduler_id; }
 
-
 protected:
 
     unsigned m_uid;
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 3d1da62..fb9bc9e 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -220,7 +220,6 @@ void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref)
 }
 
 unsigned g_assemble_code_next_pc=0; 
-std::map<unsigned,function_info*> g_pc_to_finfo;
 std::vector<ptx_instruction*> function_info::s_g_pc_to_insn;
 
 #define MAX_INST_SIZE 8 /*bytes*/
@@ -257,7 +256,7 @@ void function_info::ptx_assemble()
          const symbol *l = pI->get_label();
          labels[l->name()] = n;
       } else {
-         g_pc_to_finfo[PC] = this;
+         gpgpu_ctx->func_sim->g_pc_to_finfo[PC] = this;
          m_instr_mem[n] = pI;
          s_g_pc_to_insn.push_back(pI);
          assert(pI == s_g_pc_to_insn[PC]);
@@ -497,7 +496,7 @@ void gpgpu_t::gpu_memset( size_t dst_start_addr, int c, size_t count )
    }
 }
 
-void ptx_print_insn( address_type pc, FILE *fp )
+void cuda_sim::ptx_print_insn( address_type pc, FILE *fp )
 {
    std::map<unsigned,function_info*>::iterator f = g_pc_to_finfo.find(pc);
    if( f == g_pc_to_finfo.end() ) {
@@ -509,7 +508,7 @@ void ptx_print_insn( address_type pc, FILE *fp )
    finfo->print_insn(pc,fp);
 }
 
-std::string ptx_get_insn_str( address_type pc )
+std::string cuda_sim::ptx_get_insn_str( address_type pc )
 {
    std::map<unsigned,function_info*>::iterator f = g_pc_to_finfo.find(pc);
    if( f == g_pc_to_finfo.end() ) {
diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h
index 5c95100..c578524 100644
--- a/src/cuda-sim/cuda-sim.h
+++ b/src/cuda-sim/cuda-sim.h
@@ -61,8 +61,6 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel,
                               bool functionalSimulationMode = false);
 const warp_inst_t *ptx_fetch_inst( address_type pc );
 const struct gpgpu_ptx_sim_info* ptx_sim_kernel_info(const class function_info *kernel);
-void ptx_print_insn( address_type pc, FILE *fp );
-std::string ptx_get_insn_str( address_type pc );
 
 
 /*!
@@ -162,6 +160,7 @@ class cuda_sim {
 	void ** g_inst_op_classification_stat;
 	std::set<std::string>   g_globals;
 	std::set<std::string>   g_constants;
+	std::map<unsigned,function_info*> g_pc_to_finfo;
 	// backward pointer
 	class gpgpu_context* gpgpu_ctx;
 	//global functions
@@ -181,6 +180,8 @@ class cuda_sim {
 	struct rec_pts find_reconvergence_points( function_info *finfo );
 	address_type get_converge_point( address_type pc );
 	void   gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu );
+	void ptx_print_insn( address_type pc, FILE *fp );
+	std::string ptx_get_insn_str( address_type pc );
 };
 
 #endif
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index 1bd409e..c537091 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -83,8 +83,9 @@ symbol_table::symbol_table()
    assert(0); 
 }
 
-symbol_table::symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent )
+symbol_table::symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent, gpgpu_context* ctx )
 {
+   gpgpu_ctx = ctx;
    m_scope_name = std::string(scope_name);
    m_reg_allocator=0;
    m_shared_next = 0;
@@ -183,7 +184,7 @@ symbol_table* symbol_table::start_inst_group() {
 
    //previous added
    assert(m_inst_group_symtab.find(std::string(inst_group_name)) == m_inst_group_symtab.end());
-   symbol_table *sym_table = new symbol_table(inst_group_name, 3/*inst group*/, this );
+   symbol_table *sym_table = new symbol_table(inst_group_name, 3/*inst group*/, this, gpgpu_ctx );
  
    sym_table->m_global_next = m_global_next;
    sym_table->m_shared_next = m_shared_next;
@@ -221,7 +222,7 @@ bool symbol_table::add_function_decl( const char *name, int entry_point, functio
       *func_info = m_function_info_lookup[key];
       prior_decl = true;
    } else {
-      *func_info = new function_info(entry_point);
+      *func_info = new function_info(entry_point, gpgpu_ctx);
       (*func_info)->set_name(name);
       (*func_info)->set_maxnt_id(0);
       m_function_info_lookup[key] = *func_info;
@@ -232,7 +233,7 @@ bool symbol_table::add_function_decl( const char *name, int entry_point, functio
       *sym_table = m_function_symtab_lookup[key];
    } else {
       assert( !prior_decl );
-      *sym_table = new symbol_table( "", entry_point, this );
+      *sym_table = new symbol_table( "", entry_point, this, gpgpu_ctx );
       
       // Initial setup code to support a register represented as "_".
       // This register is used when an instruction operand is
@@ -1373,8 +1374,9 @@ std::string ptx_instruction::to_string() const
 
 unsigned function_info::sm_next_uid = 1;
 
-function_info::function_info(int entry_point ) 
+function_info::function_info(int entry_point, gpgpu_context* ctx )
 {
+   gpgpu_ctx = ctx;
    m_uid = sm_next_uid++;
    m_entry_point = (entry_point==1)?true:false;
    m_extern = (entry_point==2)?true:false;
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index 1604551..babd54b 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -306,7 +306,7 @@ private:
 class symbol_table {
 public:
    symbol_table();
-   symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent );
+   symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent, gpgpu_context* ctx);
    void set_name( const char *name );
    const ptx_version &get_ptx_version() const;
    unsigned get_sm_target() const;
@@ -348,6 +348,9 @@ public:
    symbol_table* start_inst_group();
    symbol_table* end_inst_group();
 
+   // backward pointer
+   class gpgpu_context* gpgpu_ctx;
+
 private:
    unsigned m_reg_allocator;
    unsigned m_shared_next;
@@ -1233,7 +1236,7 @@ private:
 
 class function_info {
 public:
-   function_info(int entry_point );
+   function_info(int entry_point, gpgpu_context* ctx );
    const ptx_version &get_ptx_version() const { return m_symtab->get_ptx_version(); }
    unsigned get_sm_target() const { return m_symtab->get_sm_target(); }
    bool is_extern() const { return m_extern; }
@@ -1403,6 +1406,8 @@ public:
 
    void set_maxnt_id(unsigned maxthreads) { maxnt_id = maxthreads;}
    unsigned get_maxnt_id() { return maxnt_id;}
+   // backward pointer
+   class gpgpu_context* gpgpu_ctx;
 
 private:
    unsigned maxnt_id;
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 269ec4d..5a94679 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -112,7 +112,7 @@ symbol_table * gpgpu_context::init_parser( const char *ptx_filename )
 {
    g_filename = strdup(ptx_filename);
    if  (g_global_allfiles_symbol_table == NULL) {
-       g_global_allfiles_symbol_table = new symbol_table("global_allfiles", 0, NULL);
+       g_global_allfiles_symbol_table = new symbol_table("global_allfiles", 0, NULL, this);
        ptx_parser->g_global_symbol_table = ptx_parser->g_current_symbol_table = g_global_allfiles_symbol_table;
    }
    /*else {
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index 69790fc..6cd6d8f 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -1025,7 +1025,7 @@ void scheduler_unit::cycle()
             m_simt_stack[warp_id]->get_pdom_stack_top_info(&pc,&rpc);
             SCHED_DPRINTF( "Warp (warp_id %u, dynamic_warp_id %u) has valid instruction (%s)\n",
                            (*iter)->get_warp_id(), (*iter)->get_dynamic_warp_id(),
-                           ptx_get_insn_str( pc).c_str() );
+                           m_shader->m_config->gpgpu_ctx->func_sim->ptx_get_insn_str( pc).c_str() );
             if( pI ) {
                 assert(valid);
                 if( pc != pI->pc ) {
@@ -2690,7 +2690,7 @@ void warp_inst_t::print( FILE *fout ) const
     for (unsigned j=0; j<m_config->warp_size; j++)
         fprintf(fout, "%c", (active(j)?'1':'0') );
     fprintf(fout, "]: ");
-    ptx_print_insn( pc, fout );
+    m_config->gpgpu_ctx->func_sim->ptx_print_insn( pc, fout );
     fprintf(fout, "\n");
 }
 void shader_core_ctx::incexecstat(warp_inst_t *&inst)
diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h
index e0cefac..2837f1b 100644
--- a/src/gpgpu-sim/shader.h
+++ b/src/gpgpu-sim/shader.h
@@ -1367,7 +1367,7 @@ const char* const pipeline_stage_name_decode[] = {
 class shader_core_config : public core_config
 {
     public:
-    shader_core_config(gpgpu_context* ctx){
+    shader_core_config(gpgpu_context* ctx):core_config(ctx){
 	pipeline_widths_string = NULL;
 	gpgpu_ctx = ctx;
     }
-- 
cgit v1.3


From e6f475eaa9c27cd601033c37321466010299dabf Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Thu, 11 Jul 2019 01:47:04 -0400
Subject: Move g_inst_lookup

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 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_parser.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<std::string,std::map<unsigned,const ptx_instruction*> > 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<std::string,std::map<unsigned,const ptx_instruction*> >::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<std::string,std::map<unsigned,const ptx_instruction*> > 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 33d064f5a38b60fb0ed4b4ad78212c6e5d9bfc0f Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Thu, 11 Jul 2019 13:57:08 -0400
Subject: g_sym_name_to_symbol_table

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 src/cuda-sim/cuda-sim.cc   | 4 ++--
 src/cuda-sim/ptx_parser.cc | 2 --
 src/cuda-sim/ptx_parser.h  | 2 ++
 3 files changed, 4 insertions(+), 4 deletions(-)

(limited to 'src/cuda-sim/ptx_parser.cc')

diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index d098f2b..e4ff0f9 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -1995,8 +1995,8 @@ void cuda_sim::gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src,
    const char *mem_name = NULL;
    memory_space *mem = NULL;
 
-   std::map<std::string,symbol_table*>::iterator st = g_sym_name_to_symbol_table.find(sym_name.c_str());
-   assert( st != g_sym_name_to_symbol_table.end() );
+   std::map<std::string,symbol_table*>::iterator st = gpgpu_ctx->ptx_parser->g_sym_name_to_symbol_table.find(sym_name.c_str());
+   assert( st != gpgpu_ctx->ptx_parser->g_sym_name_to_symbol_table.end() );
    symbol_table *symtab = st->second;
 
    symbol *sym = symtab->lookup(sym_name.c_str());
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 8b2b72e..1d38edf 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -47,8 +47,6 @@ void ptx_recognizer::set_ptx_warp_size(const struct core_config * warp_size)
    g_shader_core_config=warp_size;
 }
 
-// the program intermediate representation...
-std::map<std::string,symbol_table*> g_sym_name_to_symbol_table;
 
 #define PTX_PARSE_DPRINTF(...) \
    if( g_debug_ir_generation ) { \
diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h
index 7e87ec8..1e559a2 100644
--- a/src/cuda-sim/ptx_parser.h
+++ b/src/cuda-sim/ptx_parser.h
@@ -104,6 +104,8 @@ class ptx_recognizer {
 	int g_entry_point;
 	const struct core_config *g_shader_core_config;
 	std::map<std::string,std::map<unsigned,const ptx_instruction*> > g_inst_lookup;
+	// the program intermediate representation...
+	std::map<std::string,symbol_table*> g_sym_name_to_symbol_table;
 	// backward pointer
 	class gpgpu_context* gpgpu_ctx;
 
-- 
cgit v1.3


From 67950eb12eec495b1976934acc763db481d955f8 Mon Sep 17 00:00:00 2001
From: Mengchi Zhang <zhan2308@purdue.edu>
Date: Mon, 15 Jul 2019 11:58:10 -0400
Subject: Move operand_info::sm_next_uid

Signed-off-by: Mengchi Zhang <zhan2308@purdue.edu>
---
 libcuda/gpgpu_context.h    |  2 ++
 src/cuda-sim/ptx_ir.cc     | 20 ++++++++++-------
 src/cuda-sim/ptx_ir.h      | 54 ++++++++++++++++++++++++----------------------
 src/cuda-sim/ptx_parser.cc | 30 +++++++++++++-------------
 src/cuda-sim/ptx_parser.h  |  2 +-
 src/cuda-sim/ptx_sim.cc    |  8 +++----
 6 files changed, 62 insertions(+), 54 deletions(-)

(limited to 'src/cuda-sim/ptx_parser.cc')

diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h
index 0dd6311..9c5ae7b 100644
--- a/libcuda/gpgpu_context.h
+++ b/libcuda/gpgpu_context.h
@@ -13,6 +13,7 @@ class gpgpu_context {
 	    g_global_allfiles_symbol_table = NULL;
 	    sm_next_access_uid=0;
 	    warp_inst_sm_next_uid=0;
+	    operand_info_sm_next_uid = 1;
 	    api = new cuda_runtime_api(this);
 	    ptxinfo = new ptxinfo_data(this);
 	    ptx_parser = new ptx_recognizer(this);
@@ -25,6 +26,7 @@ class gpgpu_context {
 	const char *g_filename;
 	unsigned sm_next_access_uid;
 	unsigned warp_inst_sm_next_uid;
+	unsigned operand_info_sm_next_uid;//uid for operand_info
 	// objects pointers for each file
 	cuda_runtime_api* api;
 	ptxinfo_data* ptxinfo;
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index c537091..0fa29eb 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -39,6 +39,7 @@ typedef void * yyscan_t;
 #include "assert.h"
 
 #include "cuda-sim.h"
+#include "../../libcuda/gpgpu_context.h"
 
 #define STR_SIZE 1024
 
@@ -322,11 +323,9 @@ void symbol_table::dump()
    printf("\n");
 }
 
-unsigned operand_info::sm_next_uid=1;
-
 unsigned operand_info::get_uid()
 {
-   unsigned result = sm_next_uid++;
+   unsigned result = (gpgpu_ctx->operand_info_sm_next_uid)++;
    return result;
 }
 
@@ -1015,7 +1014,7 @@ void copy_buffer_to_frame(ptx_thread_info * thread, const arg_buffer_t &a)
 {
    if( a.is_reg() ) {
       ptx_reg_t value = a.get_reg();
-      operand_info dst_reg = operand_info(a.get_dst()); 
+      operand_info dst_reg = operand_info(a.get_dst(), thread->get_gpu()->gpgpu_ctx); 
       thread->set_reg(dst_reg.get_symbol(),value);
    } else {  
       const void *buffer = a.get_param_buffer();
@@ -1039,7 +1038,8 @@ void copy_buffer_list_into_frame(ptx_thread_info * thread, arg_buffer_list_t &ar
 
 static std::list<operand_info> check_operands( int opcode,
                                         const std::list<int> &scalar_type,
-                                        const std::list<operand_info> &operands )
+                                        const std::list<operand_info> &operands,
+					gpgpu_context* ctx)
 {
    static int g_warn_literal_operands_two_type_inst;
     if( (opcode == CVT_OP) || (opcode == SET_OP) || (opcode == SLCT_OP) || (opcode == TEX_OP) || (opcode==MMA_OP) || (opcode == DP4A_OP)) {
@@ -1066,7 +1066,7 @@ static std::list<operand_info> check_operands( int opcode,
                     if( (op.get_type() == double_op_t) && (inst_type == F32_TYPE) ) {
                         ptx_reg_t v = op.get_literal_value();
                         float u = (float)v.f64;
-                        operand_info n(u);
+                        operand_info n(u, ctx);
                         result.push_back(n);
                     } else {
                         result.push_back(op);
@@ -1097,7 +1097,7 @@ ptx_instruction::ptx_instruction( int opcode,
                                   unsigned line,
                                   const char *source,
                                   const core_config *config,
-				  gpgpu_context* ctx ) : warp_inst_t(config)
+				  gpgpu_context* ctx ) : warp_inst_t(config), m_return_var(ctx)
 {
    gpgpu_ctx = ctx;
    m_uid = ++g_num_ptx_inst_uid;
@@ -1107,7 +1107,7 @@ ptx_instruction::ptx_instruction( int opcode,
    m_neg_pred = neg_pred;
    m_pred_mod = pred_mod;
    m_label = label;
-   const std::list<operand_info> checked_operands = check_operands(opcode,scalar_type,operands);
+   const std::list<operand_info> checked_operands = check_operands(opcode,scalar_type,operands, ctx);
    m_operands.insert(m_operands.begin(), checked_operands.begin(), checked_operands.end() );
    m_return_var = return_var;
    m_options = options;
@@ -1371,6 +1371,10 @@ std::string ptx_instruction::to_string() const
                            m_source.c_str() );
    return std::string( buf );
 }
+operand_info ptx_instruction::get_pred() const
+{
+    return operand_info( m_pred, gpgpu_ctx);
+}
 
 unsigned function_info::sm_next_uid = 1;
 
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index babd54b..8ed94d9 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -377,9 +377,9 @@ private:
 
 class operand_info {
 public:
-   operand_info()
+   operand_info(gpgpu_context* ctx)
    {
-      init();
+      init(ctx);
       m_is_non_arch_reg = false;
       m_addr_space = undefined_space;
       m_operand_lohi = 0;
@@ -392,9 +392,9 @@ public:
       m_addr_offset = 0;
       m_value.m_symbolic=NULL;
    }
-   operand_info( const symbol *addr )
+   operand_info( const symbol *addr, gpgpu_context* ctx )
    {
-      init();
+      init(ctx);
       m_is_non_arch_reg = false;
       m_addr_space = undefined_space;
       m_operand_lohi = 0;
@@ -435,9 +435,9 @@ public:
       m_is_return_var = false;
       m_immediate_address=false;
    }
-   operand_info( const symbol *addr1, const symbol *addr2 )
+   operand_info( const symbol *addr1, const symbol *addr2, gpgpu_context* ctx )
    {
-      init();
+      init(ctx);
       m_is_non_arch_reg = false;
       m_addr_space = undefined_space;
       m_operand_lohi = 0;
@@ -462,9 +462,9 @@ public:
       m_is_return_var = false;
       m_immediate_address=false;
    }
-   operand_info( int builtin_id, int dim_mod )
+   operand_info( int builtin_id, int dim_mod, gpgpu_context* ctx )
    {
-      init();
+      init(ctx);
       m_is_non_arch_reg = false;
       m_addr_space = undefined_space;
       m_operand_lohi = 0;
@@ -481,9 +481,9 @@ public:
       m_is_return_var = false;
       m_immediate_address=false;
    }
-   operand_info( const symbol *addr, int offset )
+   operand_info( const symbol *addr, int offset, gpgpu_context* ctx )
    {
-      init();
+      init(ctx);
       m_is_non_arch_reg = false;
       m_addr_space = undefined_space;
       m_operand_lohi = 0;
@@ -500,9 +500,9 @@ public:
       m_is_return_var = false;
       m_immediate_address=false;
    }
-   operand_info( unsigned x )
+   operand_info( unsigned x, gpgpu_context* ctx )
    {
-      init();
+      init(ctx);
       m_is_non_arch_reg = false;
       m_addr_space = undefined_space;
       m_operand_lohi = 0;
@@ -519,9 +519,9 @@ public:
       m_is_return_var = false;
       m_immediate_address=true;
    }
-   operand_info( int x )
+   operand_info( int x, gpgpu_context* ctx )
    {
-      init();
+      init(ctx);
       m_is_non_arch_reg = false;
       m_addr_space = undefined_space;
       m_operand_lohi = 0;
@@ -538,9 +538,9 @@ public:
       m_is_return_var = false;
       m_immediate_address=false;
    }
-   operand_info( float x )
+   operand_info( float x, gpgpu_context* ctx )
    {
-      init();
+      init(ctx);
       m_is_non_arch_reg = false;
       m_addr_space = undefined_space;
       m_operand_lohi = 0;
@@ -557,9 +557,9 @@ public:
       m_is_return_var = false;
       m_immediate_address=false;
    }
-   operand_info( double x )
+   operand_info( double x, gpgpu_context* ctx )
    {
-      init();
+      init(ctx);
       m_is_non_arch_reg = false;
       m_addr_space = undefined_space;
       m_operand_lohi = 0;
@@ -576,9 +576,9 @@ public:
       m_is_return_var = false;
       m_immediate_address=false;
    }
-   operand_info( const symbol *s1, const symbol *s2, const symbol *s3, const symbol *s4 )
+   operand_info( const symbol *s1, const symbol *s2, const symbol *s3, const symbol *s4, gpgpu_context* ctx )
    {
-      init();
+      init(ctx);
       m_is_non_arch_reg = false;
       m_addr_space = undefined_space;
       m_operand_lohi = 0;
@@ -603,9 +603,9 @@ public:
       m_is_return_var = false;
       m_immediate_address=false;
    }
-   operand_info( const symbol *s1, const symbol *s2, const symbol *s3, const symbol *s4 ,const symbol *s5,const symbol *s6,const symbol *s7, const symbol *s8)
+   operand_info( const symbol *s1, const symbol *s2, const symbol *s3, const symbol *s4 ,const symbol *s5,const symbol *s6,const symbol *s7, const symbol *s8, gpgpu_context* ctx)
    {
-      init();
+      init(ctx);
       m_is_non_arch_reg = false;
       m_addr_space = undefined_space;
       m_operand_lohi = 0;
@@ -631,8 +631,9 @@ public:
       m_immediate_address=false;
    }
 
-   void init()
+   void init(gpgpu_context* ctx)
    {
+       gpgpu_ctx = ctx;
        m_uid=(unsigned)-1;
        m_valid=false;
        m_vector=false;
@@ -842,6 +843,7 @@ public:
    bool is_non_arch_reg() const { return m_is_non_arch_reg; }
 
 private:
+   gpgpu_context* gpgpu_ctx;
    unsigned m_uid;
    bool m_valid;
    bool m_vector;
@@ -957,7 +959,7 @@ public:
    unsigned source_line() const { return m_source_line;}
    unsigned get_num_operands() const { return m_operands.size();}
    bool has_pred() const { return m_pred != NULL;}
-   operand_info get_pred() const { return operand_info( m_pred );}
+   operand_info get_pred() const;
    bool get_pred_neg() const { return m_neg_pred;}
    int get_pred_mod() const { return m_pred_mod;}
    const char *get_source() const { return m_source.c_str();}
@@ -1448,14 +1450,14 @@ private:
 
 class arg_buffer_t {
 public:
-   arg_buffer_t()
+   arg_buffer_t(gpgpu_context* ctx) : m_src_op(ctx)
    {
       m_is_reg=false;
       m_is_param=false;
       m_param_value=NULL;
       m_reg_value=ptx_reg_t();
    }
-   arg_buffer_t( const arg_buffer_t &another )
+   arg_buffer_t( const arg_buffer_t &another, gpgpu_context* ctx ) : m_src_op(ctx)
    {
       make_copy(another);
    }
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 1d38edf..81b70af 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -102,7 +102,7 @@ void ptx_recognizer::init_instruction_state()
    g_opcode = -1;
    g_options.clear();
    g_wmma_options.clear();
-   g_return_var = operand_info();
+   g_return_var = operand_info(gpgpu_ctx);
    init_directive_state();
 }
 
@@ -689,7 +689,7 @@ void ptx_recognizer::add_double_operand( const char *d1, const char *d2 )
    const symbol *s1 = g_current_symbol_table->lookup(d1);
    const symbol *s2 = g_current_symbol_table->lookup(d2);
    parse_assert( s1 != NULL && s2 != NULL, "component(s) missing declarations.");
-   g_operands.push_back( operand_info(s1,s2) );
+   g_operands.push_back( operand_info(s1,s2,gpgpu_ctx) );
 }
 
 void ptx_recognizer::add_1vector_operand( const char *d1 )
@@ -698,7 +698,7 @@ void ptx_recognizer::add_1vector_operand( const char *d1 )
    PTX_PARSE_DPRINTF("add_1vector_operand");
    const symbol *s1 = g_current_symbol_table->lookup(d1);
    parse_assert( s1 != NULL, "component(s) missing declarations.");
-   g_operands.push_back( operand_info(s1,NULL,NULL,NULL) );
+   g_operands.push_back( operand_info(s1,NULL,NULL,NULL,gpgpu_ctx) );
 }
 
 void ptx_recognizer::add_2vector_operand( const char *d1, const char *d2 )
@@ -707,7 +707,7 @@ void ptx_recognizer::add_2vector_operand( const char *d1, const char *d2 )
    const symbol *s1 = g_current_symbol_table->lookup(d1);
    const symbol *s2 = g_current_symbol_table->lookup(d2);
    parse_assert( s1 != NULL && s2 != NULL, "v2 component(s) missing declarations.");
-   g_operands.push_back( operand_info(s1,s2,NULL,NULL) );
+   g_operands.push_back( operand_info(s1,s2,NULL,NULL,gpgpu_ctx) );
 }
 
 void ptx_recognizer::add_3vector_operand( const char *d1, const char *d2, const char *d3 )
@@ -717,7 +717,7 @@ void ptx_recognizer::add_3vector_operand( const char *d1, const char *d2, const
    const symbol *s2 = g_current_symbol_table->lookup(d2);
    const symbol *s3 = g_current_symbol_table->lookup(d3);
    parse_assert( s1 != NULL && s2 != NULL && s3 != NULL, "v3 component(s) missing declarations.");
-   g_operands.push_back( operand_info(s1,s2,s3,NULL) );
+   g_operands.push_back( operand_info(s1,s2,s3,NULL,gpgpu_ctx) );
 }
 
 void ptx_recognizer::add_4vector_operand( const char *d1, const char *d2, const char *d3, const char *d4 )
@@ -732,7 +732,7 @@ void ptx_recognizer::add_4vector_operand( const char *d1, const char *d2, const
    if ( s2 == null_op ) s2 = NULL;
    if ( s3 == null_op ) s3 = NULL;
    if ( s4 == null_op ) s4 = NULL;
-   g_operands.push_back( operand_info(s1,s2,s3,s4) );
+   g_operands.push_back( operand_info(s1,s2,s3,s4,gpgpu_ctx) );
 }
 void ptx_recognizer::add_8vector_operand( const char *d1, const char *d2, const char *d3, const char *d4,const char *d5,const char *d6,const char *d7,const char *d8 )
 {
@@ -754,13 +754,13 @@ void ptx_recognizer::add_8vector_operand( const char *d1, const char *d2, const
    if ( s6 == null_op ) s6 = NULL;
    if ( s7 == null_op ) s7 = NULL;
    if ( s8 == null_op ) s8 = NULL;
-   g_operands.push_back( operand_info(s1,s2,s3,s4,s5,s6,s7,s8) );
+   g_operands.push_back( operand_info(s1,s2,s3,s4,s5,s6,s7,s8,gpgpu_ctx) );
 }
 
 void ptx_recognizer::add_builtin_operand( int builtin, int dim_modifier )
 {
    PTX_PARSE_DPRINTF("add_builtin_operand");
-   g_operands.push_back( operand_info(builtin,dim_modifier) );
+   g_operands.push_back( operand_info(builtin,dim_modifier,gpgpu_ctx) );
 }
 
 void ptx_recognizer::add_memory_operand()
@@ -883,19 +883,19 @@ void ptx_recognizer::change_operand_neg( )
 void ptx_recognizer::add_literal_int( int value )
 {
    PTX_PARSE_DPRINTF("add_literal_int");
-   g_operands.push_back( operand_info(value) );
+   g_operands.push_back( operand_info(value,gpgpu_ctx) );
 }
 
 void ptx_recognizer::add_literal_float( float value )
 {
    PTX_PARSE_DPRINTF("add_literal_float");
-   g_operands.push_back( operand_info(value) );
+   g_operands.push_back( operand_info(value,gpgpu_ctx) );
 }
 
 void ptx_recognizer::add_literal_double( double value )
 {
    PTX_PARSE_DPRINTF("add_literal_double");
-   g_operands.push_back( operand_info(value) );
+   g_operands.push_back( operand_info(value,gpgpu_ctx) );
 }
 
 void ptx_recognizer::add_scalar_operand( const char *identifier )
@@ -911,7 +911,7 @@ void ptx_recognizer::add_scalar_operand( const char *identifier )
          parse_error( msg.c_str() );
       }
    }
-   g_operands.push_back( operand_info(s) );
+   g_operands.push_back( operand_info(s,gpgpu_ctx) );
 }
 
 void ptx_recognizer::add_neg_pred_operand( const char *identifier )
@@ -921,7 +921,7 @@ void ptx_recognizer::add_neg_pred_operand( const char *identifier )
    if ( s == NULL ) {
        s = g_current_symbol_table->add_variable(identifier,NULL,1,gpgpu_ctx->g_filename,ptx_get_lineno(scanner));
    }
-   operand_info op(s);
+   operand_info op(s, gpgpu_ctx);
    op.set_neg_pred();
    g_operands.push_back( op );
 }
@@ -934,13 +934,13 @@ void ptx_recognizer::add_address_operand( const char *identifier, int offset )
       std::string msg = std::string("operand \"") + identifier + "\" has no declaration.";
       parse_error( msg.c_str() );
    }
-   g_operands.push_back( operand_info(s,offset) );
+   g_operands.push_back( operand_info(s,offset,gpgpu_ctx) );
 }
 
 void ptx_recognizer::add_address_operand2( int offset )
 {
    PTX_PARSE_DPRINTF("add_address_operand");
-   g_operands.push_back( operand_info((unsigned)offset) );
+   g_operands.push_back( operand_info((unsigned)offset,gpgpu_ctx) );
 }
 
 void ptx_recognizer::add_array_initializer()
diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h
index 1e559a2..11a3d20 100644
--- a/src/cuda-sim/ptx_parser.h
+++ b/src/cuda-sim/ptx_parser.h
@@ -35,7 +35,7 @@ class gpgpu_context;
 typedef void * yyscan_t;
 class ptx_recognizer {
     public:
-	ptx_recognizer( gpgpu_context* ctx ) {
+	ptx_recognizer( gpgpu_context* ctx ) : g_return_var(ctx) {
 	    scanner = NULL;
 	    g_size = -1;
 	    g_add_identifier_cached__identifier = NULL;
diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc
index b6cf4bd..62bb529 100644
--- a/src/cuda-sim/ptx_sim.cc
+++ b/src/cuda-sim/ptx_sim.cc
@@ -422,9 +422,9 @@ bool ptx_thread_info::callstack_pop()
    assert( !((rv_src != NULL) ^ (rv_dst != NULL)) ); // ensure caller and callee agree on whether there is a return value
 
    // read return value from callee frame
-   arg_buffer_t buffer;
+   arg_buffer_t buffer(m_gpu->gpgpu_ctx);
    if( rv_src != NULL ) 
-      buffer = copy_arg_to_buffer(this, operand_info(rv_src), rv_dst );
+      buffer = copy_arg_to_buffer(this, operand_info(rv_src, m_gpu->gpgpu_ctx), rv_dst );
 
    m_symbol_table = m_callstack.back().m_symbol_table;
    m_NPC = m_callstack.back().m_PC;
@@ -456,9 +456,9 @@ bool ptx_thread_info::callstack_pop_plus()
    assert( !((rv_src != NULL) ^ (rv_dst != NULL)) ); // ensure caller and callee agree on whether there is a return value
 
    // read return value from callee frame
-   arg_buffer_t buffer;
+   arg_buffer_t buffer(m_gpu->gpgpu_ctx);
    if( rv_src != NULL )
-      buffer = copy_arg_to_buffer(this, operand_info(rv_src), rv_dst );
+      buffer = copy_arg_to_buffer(this, operand_info(rv_src, m_gpu->gpgpu_ctx), rv_dst );
 
    m_symbol_table = m_callstack.back().m_symbol_table;
    m_NPC = m_callstack.back().m_PC;
-- 
cgit v1.3