- bug fix: parsing hex int constants didn't have a-fA-F!?!

- modify template to use printf from kernel (this works on our Fermi card--want it working on simulator for debugging)
- add support for extern function declarations (so don't try to assemble them)... printf is extern

[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6862]

6 files changed, 30 insertions, 14 deletions

diff --git a/benchmarks/CUDA/template/template_kernel.cu b/benchmarks/CUDA/template/template_kernel.cu
index f65507b..4ced538 100644
--- a/benchmarks/CUDA/template/template_kernel.cu
+++ b/benchmarks/CUDA/template/template_kernel.cu
@@ -47,6 +47,7 @@ testKernel( float* g_idata, float* g_odata)
   // emulation
   SDATA(tid) = g_idata[tid];
   __syncthreads();
+  printf("thread tid=%u reads %f from g_idata[]\n", tid, sdata[tid]);
 
   // perform some computations
   SDATA(tid) = (float) num_threads * SDATA( tid);
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 4480221..e2326f2 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -436,6 +436,23 @@ unsigned g_assemble_code_next_pc=1;
 std::map<unsigned,function_info*> g_pc_to_finfo;
 std::vector<ptx_instruction*> function_info::s_g_pc_to_insn;
 
+unsigned function_info::sm_next_uid = 1;
+
+function_info::function_info(int entry_point ) 
+{
+   m_uid = sm_next_uid++;
+   m_entry_point = (entry_point==1)?true:false;
+   m_extern = (entry_point==2)?true:false;
+   num_reconvergence_pairs = 0;
+   m_symtab = NULL;
+   m_assembled = false;
+   m_return_var_sym = NULL; 
+   m_kernel_info.cmem = 0;
+   m_kernel_info.lmem = 0;
+   m_kernel_info.regs = 0;
+   m_kernel_info.smem = 0;
+}
+
 void function_info::print_insn( unsigned pc, FILE * fp ) const
 {
    unsigned index = pc - m_start_PC;
@@ -1765,6 +1782,10 @@ void gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num )
 void gpgpu_ptx_assemble( std::string kname, void *kinfo )
 {
     function_info *func_info = (function_info *)kinfo;
+    if( func_info->is_extern() ) {
+       printf("GPGPU-Sim: skipping assembly for extern declared function \'%s\'", func_info->get_name().c_str() );
+       return;
+    }
     g_func_info = func_info;
     g_current_symbol_table = g_kernel_name_to_symtab_lookup[ kname ];
     if( g_current_symbol_table == NULL ) {
diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l
index 0d66bfd..9e45bc3 100644
--- a/src/cuda-sim/ptx.l
+++ b/src/cuda-sim/ptx.l
@@ -225,7 +225,7 @@ xor     TC; ptx_lval.int_value = XOR_OP; return OPCODE;
 
 [0-9]+\.[0-9]+ 	 TC; sscanf(yytext,"%lf", &ptx_lval.double_value); return DOUBLE_OPERAND;
 	
-0[xX][0-9]+U?   TC; CHECK_UNSIGNED; sscanf(yytext,"%x", &ptx_lval.int_value); return INT_OPERAND;
+0[xX][0-9a-fA-F]+U? TC; CHECK_UNSIGNED; sscanf(yytext,"%x", &ptx_lval.int_value); return INT_OPERAND;
 0[0-7]+U?   	TC; printf("GPGPU-Sim: ERROR ** parsing octal not (yet) implemented\n"); abort(); return INT_OPERAND;
 0[bB][01]+U?  	TC; printf("GPGPU-Sim: ERROR ** parsing binary not (yet) implemented\n"); abort(); return INT_OPERAND;
 [-]?[0-9]+U?    TC; CHECK_UNSIGNED; ptx_lval.int_value =  atoi(yytext); return INT_OPERAND;
diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y
index 9ddb214..4f1502a 100644
--- a/src/cuda-sim/ptx.y
+++ b/src/cuda-sim/ptx.y
@@ -235,6 +235,7 @@ function_ident_param: IDENTIFIER { add_function_name($1); } LEFT_PAREN param_lis
 
 function_decl_header: ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; }
 	| FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; }
+	| EXTERN_DIRECTIVE FUNC_DIRECTIVE { $$ = 2; g_func_decl=1; }
 	;
 
 param_list: param_entry { add_directive(); }
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index c981802..79dc2b0 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -195,7 +195,7 @@ int g_add_identifier_cached__array_ident;
 
 void add_function_name( const char *name ) 
 {
-   DPRINTF("add_function_name %s %s", name,  ((g_entry_point)?"(entrypoint)":""));
+   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 );
    if( g_entry_point ) {
       g_entrypoint_func_info = g_func_info;
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index aca3556..67b274c 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -879,18 +879,8 @@ private:
 
 class function_info {
 public:
-   function_info(int entry_point ) 
-   {
-      m_entry_point = entry_point?true:false;
-      num_reconvergence_pairs = 0;
-      m_symtab = NULL;
-      m_assembled = false;
-      m_return_var_sym = NULL; 
-      m_kernel_info.cmem = 0;
-      m_kernel_info.lmem = 0;
-      m_kernel_info.regs = 0;
-      m_kernel_info.smem = 0;
-   }
+   function_info(int entry_point );
+   bool is_extern() const { return m_extern; }
    void set_name(const char *name)
    {
       m_name = name;
@@ -1079,7 +1069,9 @@ public:
    }
 
 private:
+   unsigned m_uid;
    bool m_entry_point;
+   bool m_extern;
    bool m_assembled;
    std::string m_name;
    ptx_instruction **m_instr_mem;
@@ -1101,6 +1093,7 @@ private:
    symbol_table *m_symtab;
 
    static std::vector<ptx_instruction*> s_g_pc_to_insn; // a direct mapping from PC to instruction
+   static unsigned sm_next_uid;
 };
 
 /*******************************/