summaryrefslogtreecommitdiff
path: root/libcuda
diff options
context:
space:
mode:
authorTor Aamodt <[email protected]>2010-10-08 13:05:25 -0800
committerTor Aamodt <[email protected]>2010-10-08 13:05:25 -0800
commitb64d38562079a7d4720c15c9f6309912f4090795 (patch)
treec5f2f444610a5d5bc9605edfdca227733088f219 /libcuda
parentf07ae1260b5950fc658d19b8c9920daae76fcaec (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')
-rw-r--r--libcuda/cuda_runtime_api.cc35
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" );