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] --- libcuda/cuda_runtime_api.cc | 15 +++++++++++++-- 1 file changed, 13 insertions(+), 2 deletions(-) (limited to 'libcuda/cuda_runtime_api.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index df59506..27b4bd3 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -167,9 +167,20 @@ struct gpgpu_ptx_sim_arg { struct gpgpu_ptx_sim_arg *g_ptx_sim_params; cudaError_t g_last_cudaError; +#include "../src/cuda-sim/ptx_loader.h" #include "../src/cuda-sim/cuda-sim.h" #include "../src/gpgpusim_entrypoint.h" +void register_ptx_function( const char *name, function_info *impl ) +{ + register_function_implementation( name, impl ); +} + +extern "C" void ptxinfo_addinfo() +{ + ptxinfo_cuda_addinfo(); +} + #if defined __APPLE__ # define __my_func__ __PRETTY_FUNCTION__ #else @@ -822,9 +833,9 @@ __host__ cudaError_t CUDARTAPI cudaLaunch(const char *symbol ) printf("GPGPU-Sim PTX: cudaLaunch for %p (mode=%s)\n", symbol, g_ptx_sim_mode?"functional simulation":"performance simulation"); if( g_ptx_sim_mode ) - gpgpu_ptx_sim_main_func( symbol, g_cudaGridDim, g_cudaBlockDim, g_ptx_sim_params ); + gpgpu_cuda_ptx_sim_main_func( symbol, g_cudaGridDim, g_cudaBlockDim, g_ptx_sim_params ); else - gpgpu_ptx_sim_main_perf( symbol, g_cudaGridDim, g_cudaBlockDim, g_ptx_sim_params ); + gpgpu_cuda_ptx_sim_main_perf( symbol, g_cudaGridDim, g_cudaBlockDim, g_ptx_sim_params ); g_ptx_sim_params=NULL; return g_last_cudaError = cudaSuccess; } -- cgit v1.3