From b64d38562079a7d4720c15c9f6309912f4090795 Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Fri, 8 Oct 2010 13:05:25 -0800 Subject: 1. modify shader_core_ctx::execute_pipe() to model instruction throughput correctly (off-by-one error) 2. adding code to dump_pipeline to display reason for memory stage stalling 3. removing dead code resulting from prior changes correlation vs. GT200 is 0.95 (need to add back modeling of memory writebacks, shared memory latency) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7831] --- libcuda/cuda_runtime_api.cc | 35 ----------------------------------- 1 file changed, 35 deletions(-) (limited to 'libcuda/cuda_runtime_api.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index caf0ad2..df8ef4f 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1479,8 +1479,6 @@ int CUDARTAPI __cudaSynchronizeThreads(void**, void*) //////// - - extern "C" int ptx_parse(); extern "C" int ptx__scan_string(const char*); extern "C" FILE *ptx_in; @@ -1492,39 +1490,6 @@ extern "C" FILE *ptxinfo_in; /// static functions - - -std::map *g_kernel_name_to_function_lookup=NULL; -std::map *g_host_to_kernel_entrypoint_name_lookup=NULL; - -function_info *get_kernel(const char *kernel_key, std::string &kernel_func_name_mangled ) -{ - if ( g_host_to_kernel_entrypoint_name_lookup->find(kernel_key) == - g_host_to_kernel_entrypoint_name_lookup->end() ) { - printf("GPGPU-Sim PTX: ERROR ** cannot locate __global__ function from hostPtr\n" ); - printf("GPGPU-Sim PTX: registered PTX kernels: \n"); - std::map::iterator i_eptr = g_host_to_kernel_entrypoint_name_lookup->begin(); - for (; i_eptr != g_host_to_kernel_entrypoint_name_lookup->end(); ++i_eptr) { - printf("GPGPU-Sim PTX: (%p,%s)\n", i_eptr->first, i_eptr->second.c_str()); - } - printf("\n"); - abort(); - } - kernel_func_name_mangled = (*g_host_to_kernel_entrypoint_name_lookup)[kernel_key]; - if ( g_kernel_name_to_function_lookup->find(kernel_func_name_mangled) == - g_kernel_name_to_function_lookup->end() ) { - printf("GPGPU-Sim PTX: ERROR ** function \'%s\' not found in ptx file\n", - kernel_func_name_mangled.c_str() ); - abort(); - } - return (*g_kernel_name_to_function_lookup)[kernel_func_name_mangled]; -} - - -void ptxinfo_cuda_addinfo( CUctx_st *context ) -{ -} - static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr, gpgpu_t *gpu ) { printf( "GPGPU-Sim PTX: loading globals with explicit initializers... \n" ); -- cgit v1.3