diff options
| author | Tor Aamodt <[email protected]> | 2010-10-08 13:05:25 -0800 |
|---|---|---|
| committer | Tor Aamodt <[email protected]> | 2010-10-08 13:05:25 -0800 |
| commit | b64d38562079a7d4720c15c9f6309912f4090795 (patch) | |
| tree | c5f2f444610a5d5bc9605edfdca227733088f219 /libcuda/cuda_runtime_api.cc | |
| parent | f07ae1260b5950fc658d19b8c9920daae76fcaec (diff) | |
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]
Diffstat (limited to 'libcuda/cuda_runtime_api.cc')
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 35 |
1 files changed, 0 insertions, 35 deletions
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<std::string,function_info*> *g_kernel_name_to_function_lookup=NULL; -std::map<const void*,std::string> *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<const void*,std::string>::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" ); |
