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_ir.h | 10 +++------- 1 file changed, 3 insertions(+), 7 deletions(-) (limited to 'src/cuda-sim/ptx_ir.h') diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index fc7f192..1b1d864 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1015,8 +1015,8 @@ public: return &m_kernel_info; } - const void set_kernel_info (const struct gpgpu_ptx_sim_kernel_info *info) { - m_kernel_info = *info; + const void set_kernel_info (const struct gpgpu_ptx_sim_kernel_info &info) { + m_kernel_info = info; } symbol_table *get_symtab() { @@ -1254,14 +1254,11 @@ struct textureInfo { extern function_info *g_func_info; -extern int g_error_detected; -extern symbol_table *g_entrypoint_symbol_table; extern function_info *g_entrypoint_func_info; -extern symbol_table *g_global_symbol_table; extern std::map g_sym_name_to_symbol_table; + #define GLOBAL_HEAP_START 0x10000000 // start allocating from this address (lower values used for allocating globals in .ptx file) - #define SHARED_MEM_SIZE_MAX (64*1024) #define LOCAL_MEM_SIZE_MAX 1024 #define MAX_STREAMING_MULTIPROCESSORS 64 @@ -1274,7 +1271,6 @@ extern std::map g_sym_name_to_symbol_table; #define STATIC_ALLOC_LIMIT (GLOBAL_HEAP_START - (TOTAL_LOCAL_MEM+TOTAL_SHARED_MEM)) -void gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); void gpgpu_ptx_assemble( std::string kname, void *kinfo ); #include "../option_parser.h" void ptx_reg_options(option_parser_t opp); -- cgit v1.3