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/cuda-sim.h | 11 +++++++---- 1 file changed, 7 insertions(+), 4 deletions(-) (limited to 'src/cuda-sim/cuda-sim.h') diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 6662b4c..2c5e0a7 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -16,10 +16,11 @@ extern memory_space *g_global_mem; extern int g_debug_execution; extern int g_debug_thread_uid; extern std::map *g_kernel_name_to_function_lookup; -extern std::map g_kernel_name_to_symtab_lookup; -extern void gpgpu_ptx_sim_add_ptxstring( const char *ptx, const char *source_fname ); -extern void gpgpu_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, struct gpgpu_ptx_sim_arg *); +extern void gpgpu_cuda_ptx_sim_init_grid(const char *kernel_key,struct gpgpu_ptx_sim_arg *args, struct dim3 gridDim, struct dim3 blockDim ); +extern void gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,struct gpgpu_ptx_sim_arg *args, struct dim3 gridDim, struct dim3 blockDim ); +extern void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, struct gpgpu_ptx_sim_arg *); +extern void print_splash(); extern void* gpgpu_ptx_sim_malloc( size_t count ); extern void* gpgpu_ptx_sim_mallocarray( size_t count ); extern void gpgpu_ptx_sim_memcpy_to_gpu( size_t dst_start_addr, const void *src, size_t count ); @@ -39,6 +40,8 @@ extern int gpgpu_ptx_sim_sizeofTexture(const char* name); extern const char* gpgpu_ptx_sim_findNamefromTexture(const struct textureReference* texref); extern const struct textureReference* gpgpu_ptx_sim_accessTextureofName(const char* name); extern void read_sim_environment_variables(); -extern void register_ptx_function( const char *name, function_info *impl, symbol_table *symtab ); +extern void register_function_implementation( const char *name, function_info *impl ); +extern void ptxinfo_cuda_addinfo(); +extern void ptxinfo_opencl_addinfo( std::map &kernels ); #endif -- cgit v1.3