From b400820185db543ff98fd823a15958444e12fee6 Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Sun, 8 Aug 2010 21:15:00 -0800 Subject: ** fully decouple kernel compilation from kernel launch for opencl changes required to do this (this enables oclReduction SDK 3.1 to pass and is a prerequisite to Fermi style concurrent kernel launching) - separate kernel launch for cuda and opencl (and init_grid functions) - add pgm_info structure for holding kernel function_info pointer in _cl_program object after ptx compilation ** added support for clEnqueueNDRangeKernel with local_work_size==NULL which is part of OpenCL spec and used in oclSortingNetworks (however, this is a braindead implementation that only handles the case where global_work_size[0] is smaller than the max number of threads per shader; moreover oclSortingNetworks is still not working but for what looks like another reason) ** refactoring / cleanup - g_global_symbol_table made static to ptx_parser.cc - remove g_kernel_name_to_symtab_lookup (not really being used) - moving various function prototypes into headers (e.g., if foo() defined in bar.cc, then place prototype in bar.h) - adding icnt_reg_options() [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7168] --- src/cuda-sim/ptx_parser.cc | 7 ++++--- 1 file changed, 4 insertions(+), 3 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 17e4ca2..05ce397 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -73,7 +73,7 @@ const char *g_filename; unsigned g_max_regs_per_thread = 0; // the program intermediate representation... -symbol_table *g_global_symbol_table = NULL; +static symbol_table *g_global_symbol_table = NULL; std::map g_sym_name_to_symbol_table; static symbol_table *g_current_symbol_table = NULL; static std::list g_instructions; @@ -133,7 +133,7 @@ void read_parser_environment_variables() } } -void init_parser( const char *ptx_filename ) +symbol_table *init_parser( const char *ptx_filename ) { g_filename = strdup(ptx_filename); g_global_symbol_table = g_current_symbol_table = new symbol_table("global",0,NULL); @@ -142,6 +142,8 @@ void init_parser( const char *ptx_filename ) #define DEF(X,Y) g_ptx_token_decode[X] = Y; #include "ptx_parser_decode.def" #undef DEF + + return g_global_symbol_table; } void init_directive_state() @@ -192,7 +194,6 @@ void add_function_name( const char *name ) 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; - g_entrypoint_symbol_table = g_current_symbol_table; } if( g_add_identifier_cached__identifier ) { add_identifier( g_add_identifier_cached__identifier, -- cgit v1.3