remove temporarily generated files

1 files changed, 0 insertions, 2515 deletions

diff --git a/libcuda/cuda_runtime_api.cc~ b/libcuda/cuda_runtime_api.cc~
deleted file mode 100644
index de7f5e9..0000000
--- a/libcuda/cuda_runtime_api.cc~
+++ /dev/null
@@ -1,2515 +0,0 @@
-// This file created from cuda_runtime_api.h distributed with CUDA 1.1
-// Changes Copyright 2009,  Tor M. Aamodt, Ali Bakhoda and George L. Yuan
-// University of British Columbia
-
-/* 
- * cuda_runtime_api.cc
- *
- * Copyright © 2009 by Tor M. Aamodt, Wilson W. L. Fung, Ali Bakhoda, 
- * George L. Yuan and the University of British Columbia, Vancouver, 
- * BC V6T 1Z4, All Rights Reserved.
- * 
- * THIS IS A LEGAL DOCUMENT BY DOWNLOADING GPGPU-SIM, YOU ARE AGREEING TO THESE
- * TERMS AND CONDITIONS.
- * 
- * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
- * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
- * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
- * ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNERS OR CONTRIBUTORS BE
- * LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
- * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
- * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
- * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
- * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
- * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
- * POSSIBILITY OF SUCH DAMAGE.
- * 
- * NOTE: The files libcuda/cuda_runtime_api.c and src/cuda-sim/cuda-math.h
- * are derived from the CUDA Toolset available from http://www.nvidia.com/cuda
- * (property of NVIDIA).  The files benchmarks/BlackScholes/ and 
- * benchmarks/template/ are derived from the CUDA SDK available from 
- * http://www.nvidia.com/cuda (also property of NVIDIA).  The files from 
- * src/intersim/ are derived from Booksim (a simulator provided with the 
- * textbook "Principles and Practices of Interconnection Networks" available 
- * from http://cva.stanford.edu/books/ppin/). As such, those files are bound by 
- * the corresponding legal terms and conditions set forth separately (original 
- * copyright notices are left in files from these sources and where we have 
- * modified a file our copyright notice appears before the original copyright 
- * notice).  
- * 
- * Using this version of GPGPU-Sim requires a complete installation of CUDA 
- * which is distributed seperately by NVIDIA under separate terms and 
- * conditions.  To use this version of GPGPU-Sim with OpenCL requires a
- * recent version of NVIDIA's drivers which support OpenCL.
- * 
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions are met:
- * 
- * 1. Redistributions of source code must retain the above copyright notice,
- * this list of conditions and the following disclaimer.
- * 
- * 2. Redistributions in binary form must reproduce the above copyright notice,
- * this list of conditions and the following disclaimer in the documentation
- * and/or other materials provided with the distribution.
- * 
- * 3. Neither the name of the University of British Columbia nor the names of
- * its contributors may be used to endorse or promote products derived from
- * this software without specific prior written permission.
- * 
- * 4. This version of GPGPU-SIM is distributed freely for non-commercial use only.  
- *  
- * 5. No nonprofit user may place any restrictions on the use of this software,
- * including as modified by the user, by any other authorized user.
- * 
- * 6. GPGPU-SIM was developed primarily by Tor M. Aamodt, Wilson W. L. Fung, 
- * Ali Bakhoda, George L. Yuan, at the University of British Columbia, 
- * Vancouver, BC V6T 1Z4
- */
-
-/*
- * Copyright 1993-2007 NVIDIA Corporation.  All rights reserved.
- *
- * NOTICE TO USER:   
- *
- * This source code is subject to NVIDIA ownership rights under U.S. and 
- * international Copyright laws.  Users and possessors of this source code 
- * are hereby granted a nonexclusive, royalty-free license to use this code 
- * in individual and commercial software.
- *
- * NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE 
- * CODE FOR ANY PURPOSE.  IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR 
- * IMPLIED WARRANTY OF ANY KIND.  NVIDIA DISCLAIMS ALL WARRANTIES WITH 
- * REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF 
- * MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
- * IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL, 
- * OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS 
- * OF USE, DATA OR PROFITS,  WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE 
- * OR OTHER TORTIOUS ACTION,  ARISING OUT OF OR IN CONNECTION WITH THE USE 
- * OR PERFORMANCE OF THIS SOURCE CODE.  
- *
- * U.S. Government End Users.   This source code is a "commercial item" as 
- * that term is defined at  48 C.F.R. 2.101 (OCT 1995), consisting  of 
- * "commercial computer  software"  and "commercial computer software 
- * documentation" as such terms are  used in 48 C.F.R. 12.212 (SEPT 1995) 
- * and is provided to the U.S. Government only as a commercial end item.  
- * Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through 
- * 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the 
- * source code with only those rights set forth herein. 
- *
- * Any use of this source code in individual and commercial software must 
- * include, in the user documentation and internal comments to the code,
- * the above Disclaimer and U.S. Government End Users Notice.
- */
-
-#include <stdlib.h>
-#include <stdio.h>
-#include <string.h>
-#include <assert.h>
-#include <time.h>
-#include <stdarg.h>
-#include <iostream>
-#include <string>
-#include <regex>
-#include <sstream>
-#include <fstream>
-#ifdef OPENGL_SUPPORT
-#define GL_GLEXT_PROTOTYPES
-#ifdef __APPLE__
-#include <GLUT/glut.h> // Apple's version of GLUT is here
-#else
-#include <GL/gl.h>
-#endif
-#endif
-
-#define __CUDA_RUNTIME_API_H__
-
-#include "host_defines.h"
-#include "builtin_types.h"
-#include "driver_types.h"
-#if (CUDART_VERSION < 8000)
-#include "__cudaFatFormat.h"
-#endif
-#include "../src/gpgpu-sim/gpu-sim.h"
-#include "../src/cuda-sim/ptx_loader.h"
-#include "../src/cuda-sim/cuda-sim.h"
-#include "../src/cuda-sim/ptx_ir.h"
-#include "../src/cuda-sim/ptx_parser.h"
-#include "../src/gpgpusim_entrypoint.h"
-#include "../src/stream_manager.h"
-#include "../src/abstract_hardware_model.h"
-
-#include <pthread.h>
-#include <semaphore.h>
-
-#ifdef __APPLE__
-#include <mach-o/dyld.h>
-#endif
-
-std::map<void *,void **> pinned_memory; //support for pinned memories added
-std::map<void *, size_t> pinned_memory_size;
-int no_of_ptx=0;
-
-extern void synchronize();
-extern void exit_simulation();
-
-static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr, gpgpu_t *gpu );
-static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu );
-
-static kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, 
-		gpgpu_ptx_sim_arg_list_t args,
-		struct dim3 gridDim,
-		struct dim3 blockDim,
-		struct CUctx_st* context );
-
-/*DEVICE_BUILTIN*/
-struct cudaArray
-{
-	void *devPtr;
-	int devPtr32;
-	struct cudaChannelFormatDesc desc;
-	int width;
-	int height;
-	int size; //in bytes
-	unsigned dimensions;
-};
-
-#if !defined(__dv)
-#if defined(__cplusplus)
-#define __dv(v) \
-		= v
-#else /* __cplusplus */
-#define __dv(v)
-#endif /* __cplusplus */
-#endif /* !__dv */
-
-cudaError_t g_last_cudaError = cudaSuccess;
-
-extern stream_manager *g_stream_manager;
-
-void register_ptx_function( const char *name, function_info *impl )
-{
-	// no longer need this
-}
-
-#if defined __APPLE__
-#   define __my_func__    __PRETTY_FUNCTION__
-#else
-# if defined __cplusplus ? __GNUC_PREREQ (2, 6) : __GNUC_PREREQ (2, 4)
-#   define __my_func__    __PRETTY_FUNCTION__
-# else
-#  if defined __STDC_VERSION__ && __STDC_VERSION__ >= 199901L
-#   define __my_func__    __func__
-#  else
-#   define __my_func__    ((__const char *) 0)
-#  endif
-# endif
-#endif
-
-struct _cuda_device_id {
-	_cuda_device_id(gpgpu_sim* gpu) {m_id = 0; m_next = NULL; m_gpgpu=gpu;}
-	struct _cuda_device_id *next() { return m_next; }
-	unsigned num_shader() const { return m_gpgpu->get_config().num_shader(); }
-	int num_devices() const {
-		if( m_next == NULL ) return 1;
-		else return 1 + m_next->num_devices();
-	}
-	struct _cuda_device_id *get_device( unsigned n )
-	{
-		assert( n < (unsigned)num_devices() );
-		struct _cuda_device_id *p=this;
-		for(unsigned i=0; i<n; i++)
-			p = p->m_next;
-		return p;
-	}
-	const struct cudaDeviceProp *get_prop() const
-	{
-		return m_gpgpu->get_prop();
-	}
-	unsigned get_id() const { return m_id; }
-
-	gpgpu_sim *get_gpgpu() { return m_gpgpu; }
-private:
-	unsigned m_id;
-	class gpgpu_sim *m_gpgpu;
-	struct _cuda_device_id *m_next;
-};
-
-struct CUctx_st {
-	CUctx_st( _cuda_device_id *gpu )
-	{
-		m_gpu = gpu;
-		m_binary_info.cmem = 0;
-		m_binary_info.gmem = 0;
-	}
-
-	_cuda_device_id *get_device() { return m_gpu; }
-
-	void add_binary( symbol_table *symtab, unsigned fat_cubin_handle )
-	{
-		m_code[fat_cubin_handle] = symtab;
-		m_last_fat_cubin_handle = fat_cubin_handle;
-	}
-
-	void add_ptxinfo( const char *deviceFun, const struct gpgpu_ptx_sim_info &info )
-	{
-		symbol *s = m_code[m_last_fat_cubin_handle]->lookup(deviceFun);
-		assert( s != NULL );
-		function_info *f = s->get_pc();
-		assert( f != NULL );
-		f->set_kernel_info(info);
-	}
-
-	void add_ptxinfo( const struct gpgpu_ptx_sim_info &info )
-	{
-		m_binary_info = info;
-	}
-
-	void register_function( unsigned fat_cubin_handle, const char *hostFun, const char *deviceFun )
-	{
-		if( m_code.find(fat_cubin_handle) != m_code.end() ) {
-			symbol *s = m_code[fat_cubin_handle]->lookup(deviceFun);
-			if(s != NULL) {
-				function_info *f = s->get_pc();
-				assert( f != NULL );
-				m_kernel_lookup[hostFun] = f;
-			}
-			else {
-				printf("Warning: cannot find deviceFun %s\n", deviceFun);
-				m_kernel_lookup[hostFun] = NULL;
-			}
-	//		assert( s != NULL );
-	//		function_info *f = s->get_pc();
-	//		assert( f != NULL );
-	//		m_kernel_lookup[hostFun] = f;
-		} else {
-			m_kernel_lookup[hostFun] = NULL;
-		}
-	}
-
-	function_info *get_kernel(const char *hostFun)
-	{
-		std::map<const void*,function_info*>::iterator i=m_kernel_lookup.find(hostFun);
-		assert( i != m_kernel_lookup.end() );
-		return i->second;
-	}
-
-private:
-	_cuda_device_id *m_gpu; // selected gpu
-	std::map<unsigned,symbol_table*> m_code; // fat binary handle => global symbol table
-	unsigned m_last_fat_cubin_handle;
-	std::map<const void*,function_info*> m_kernel_lookup; // unique id (CUDA app function address) => kernel entry point
-	struct gpgpu_ptx_sim_info m_binary_info;
-
-};
-
-class kernel_config {
-public:
-	kernel_config( dim3 GridDim, dim3 BlockDim, size_t sharedMem, struct CUstream_st *stream )
-	{
-		m_GridDim=GridDim;
-		m_BlockDim=BlockDim;
-		m_sharedMem=sharedMem;
-		m_stream = stream;
-	}
-	void set_arg( const void *arg, size_t size, size_t offset )
-	{
-		m_args.push_front( gpgpu_ptx_sim_arg(arg,size,offset) );
-	}
-	dim3 grid_dim() const { return m_GridDim; }
-	dim3 block_dim() const { return m_BlockDim; }
-	gpgpu_ptx_sim_arg_list_t get_args() { return m_args; }
-	struct CUstream_st *get_stream() { return m_stream; }
-
-private:
-	dim3 m_GridDim;
-	dim3 m_BlockDim;
-	size_t m_sharedMem;
-	struct CUstream_st *m_stream;
-	gpgpu_ptx_sim_arg_list_t m_args;
-};
-
-class _cuda_device_id *GPGPUSim_Init()
-{
-	static _cuda_device_id *the_device = NULL;
-	if( !the_device ) {
-		gpgpu_sim *the_gpu = gpgpu_ptx_sim_init_perf();
-
-		cudaDeviceProp *prop = (cudaDeviceProp *) calloc(sizeof(cudaDeviceProp),1);
-		snprintf(prop->name,256,"GPGPU-Sim_v%s", g_gpgpusim_version_string );
-		prop->major = 5;
-		prop->minor = 2;
-		prop->totalGlobalMem = 0x80000000 /* 2 GB */;
-		prop->memPitch = 0;
-		prop->maxThreadsPerBlock = 512;
-		prop->maxThreadsDim[0] = 512;
-		prop->maxThreadsDim[1] = 512;
-		prop->maxThreadsDim[2] = 512;
-		prop->maxGridSize[0] = 0x40000000;
-		prop->maxGridSize[1] = 0x40000000;
-		prop->maxGridSize[2] = 0x40000000;
-		prop->totalConstMem = 0x40000000;
-		prop->textureAlignment = 0;
-		prop->sharedMemPerBlock = the_gpu->shared_mem_size();
-		prop->regsPerBlock = the_gpu->num_registers_per_core();
-		prop->warpSize = the_gpu->wrp_size();
-		prop->clockRate = the_gpu->shader_clock();
-#if (CUDART_VERSION >= 2010)
-		prop->multiProcessorCount = the_gpu->get_config().num_shader();
-#endif
-		the_gpu->set_prop(prop);
-		the_device = new _cuda_device_id(the_gpu);
-	}
-	start_sim_thread(1);
-	return the_device;
-}
-
-static CUctx_st* GPGPUSim_Context()
-{
-	static CUctx_st *the_context = NULL;
-	if( the_context == NULL ) {
-		_cuda_device_id *the_gpu = GPGPUSim_Init();
-		the_context = new CUctx_st(the_gpu);
-	}
-	return the_context;
-}
-
- void ptxinfo_addinfo()
-{
-	 if(!get_ptxinfo_kname()){
-		 /* This info is not per kernel (since CUDA 5.0 some info (e.g. gmem, and cmem) is added at the beginning for the whole binary ) */
-		CUctx_st *context = GPGPUSim_Context();
-		print_ptxinfo();
-		context->add_ptxinfo(get_ptxinfo());
-		clear_ptxinfo();
-		return;
-	 }
-	 if( !strcmp("__cuda_dummy_entry__",get_ptxinfo_kname()) ) {
-		// this string produced by ptxas for empty ptx files (e.g., bandwidth test)
-		clear_ptxinfo();
-		return;
-	}
-	CUctx_st *context = GPGPUSim_Context();
-	print_ptxinfo();
-	context->add_ptxinfo( get_ptxinfo_kname(), get_ptxinfo() );
-	clear_ptxinfo();
-}
-
-void cuda_not_implemented( const char* func, unsigned line )
-{
-	fflush(stdout);
-	fflush(stderr);
-	printf("\n\nGPGPU-Sim PTX: Execution error: CUDA API function \"%s()\" has not been implemented yet.\n"
-			"                 [$GPGPUSIM_ROOT/libcuda/%s around line %u]\n\n\n",
-			func,__FILE__, line );
-	fflush(stdout);
-	abort();
-}
-
-
-#define gpgpusim_ptx_error(msg, ...) gpgpusim_ptx_error_impl(__func__, __FILE__,__LINE__, msg, ##__VA_ARGS__)
-#define gpgpusim_ptx_assert(cond,msg, ...) gpgpusim_ptx_assert_impl((cond),__func__, __FILE__,__LINE__, msg, ##__VA_ARGS__)
-
-void gpgpusim_ptx_error_impl( const char *func, const char *file, unsigned line, const char *msg, ... )
-{
-	va_list ap;
-	char buf[1024];
-	va_start(ap,msg);
-	vsnprintf(buf,1024,msg,ap);
-	va_end(ap);
-
-	printf("GPGPU-Sim CUDA API: %s\n", buf);
-	printf("                    [%s:%u : %s]\n", file, line, func );
-	abort();
-}
-
-void gpgpusim_ptx_assert_impl( int test_value, const char *func, const char *file, unsigned line, const char *msg, ... )
-{
-	va_list ap;
-	char buf[1024];
-	va_start(ap,msg);
-	vsnprintf(buf,1024,msg,ap);
-	va_end(ap);
-
-	if ( test_value == 0 )
-		gpgpusim_ptx_error_impl(func, file, line, msg);
-}
-
-
-typedef std::map<unsigned,CUevent_st*> event_tracker_t;
-
-int CUevent_st::m_next_event_uid;
-event_tracker_t g_timer_events;
-int g_active_device = 0; //active gpu that runs the code
-std::list<kernel_config> g_cuda_launch_stack;
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-
-extern "C" {
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-cudaError_t cudaPeekAtLastError(void)
-{
-	return g_last_cudaError;
-}
-
-__host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) 
-{
-	CUctx_st* context = GPGPUSim_Context();
-	*devPtr = context->get_device()->get_gpgpu()->gpu_malloc(size);
-	if(g_debug_execution >= 3)
-		printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *devPtr);
-	if ( *devPtr  ) {
-		return g_last_cudaError = cudaSuccess;
-	} else {
-		return g_last_cudaError = cudaErrorMemoryAllocation;
-	}
-}
-
-__host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size)
-{
-	GPGPUSim_Context();
-	*ptr = malloc(size);
-	if ( *ptr  ) {
-		//track pinned memory size allocated in the host so that same amount of memory is also allocated in GPU.
-		pinned_memory_size[*ptr]=size;
-		return g_last_cudaError = cudaSuccess;
-	} else {
-		return g_last_cudaError = cudaErrorMemoryAllocation;
-	}
-}
-__host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr, size_t *pitch, size_t width, size_t height)
-{
-	unsigned malloc_width_inbytes = width;
-	printf("GPGPU-Sim PTX: cudaMallocPitch (width = %d)\n", malloc_width_inbytes);
-	CUctx_st* ctx = GPGPUSim_Context();
-	*devPtr = ctx->get_device()->get_gpgpu()->gpu_malloc(malloc_width_inbytes*height);
-	pitch[0] = malloc_width_inbytes;
-	if ( *devPtr  ) {
-		return g_last_cudaError = cudaSuccess;
-	} else {
-		return g_last_cudaError = cudaErrorMemoryAllocation;
-	}
-}
-
-__host__ cudaError_t CUDARTAPI cudaMallocArray(struct cudaArray **array, const struct cudaChannelFormatDesc *desc, size_t width, size_t height __dv(1))
-{
-	unsigned size = width * height * ((desc->x + desc->y + desc->z + desc->w)/8);
-	CUctx_st* context = GPGPUSim_Context();
-	(*array) = (struct cudaArray*) malloc(sizeof(struct cudaArray));
-	(*array)->desc = *desc;
-	(*array)->width = width;
-	(*array)->height = height;
-	(*array)->size = size;
-	(*array)->dimensions = 2;
-	((*array)->devPtr32)= (int) (long long)context->get_device()->get_gpgpu()->gpu_mallocarray(size);
-	printf("GPGPU-Sim PTX: cudaMallocArray: devPtr32 = %d\n", ((*array)->devPtr32));
-	((*array)->devPtr) = (void*) (long long) ((*array)->devPtr32);
-	if ( ((*array)->devPtr) ) {
-		return g_last_cudaError = cudaSuccess;
-	} else {
-		return g_last_cudaError = cudaErrorMemoryAllocation;
-	}
-}
-
-__host__ cudaError_t CUDARTAPI cudaFree(void *devPtr)
-{
-	// TODO...  manage g_global_mem space?
-	return g_last_cudaError = cudaSuccess;
-}
-__host__ cudaError_t CUDARTAPI cudaFreeHost(void *ptr)
-{
-	free (ptr);  // this will crash the system if called twice
-	return g_last_cudaError = cudaSuccess;
-}
-
-__host__ cudaError_t CUDARTAPI cudaFreeArray(struct cudaArray *array)
-{
-	// TODO...  manage g_global_mem space?
-	return g_last_cudaError = cudaSuccess;
-};
-
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-
-__host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind)
-{
-	//CUctx_st *context = GPGPUSim_Context();
-	//gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	if(g_debug_execution >= 3)
-		printf("GPGPU-Sim PTX: cudaMemcpy(): devPtr = %p\n", dst);
-	if( kind == cudaMemcpyHostToDevice )
-		g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) );
-	else if( kind == cudaMemcpyDeviceToHost )
-		g_stream_manager->push( stream_operation((size_t)src,dst,count,0) );
-	else if( kind == cudaMemcpyDeviceToDevice )
-		g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) );
-	else if ( kind == cudaMemcpyDefault ) {
-		if ((size_t)src >= GLOBAL_HEAP_START) {
-			if ((size_t)dst >= GLOBAL_HEAP_START)
-				g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) ); // device to device
-			else
-				g_stream_manager->push( stream_operation((size_t)src,dst,count,0) ); // device to host
-		}
-		else {
-			if ((size_t)dst >= GLOBAL_HEAP_START)
-				g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) );
-			else {
-				printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported transfer: host to host\n");
-				abort();
-			}
-		}
-	}
-	else {
-		printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported cudaMemcpyKind\n");
-		abort();
-	}
-	return g_last_cudaError = cudaSuccess;
-}
-
-__host__ cudaError_t CUDARTAPI cudaMemcpyToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind)
-{ 
-	CUctx_st *context = GPGPUSim_Context();
-	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	size_t size = count;
-	printf("GPGPU-Sim PTX: cudaMemcpyToArray\n");
-	if( kind == cudaMemcpyHostToDevice )
-		gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size);
-	else if( kind == cudaMemcpyDeviceToHost )
-		gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size);
-	else if( kind == cudaMemcpyDeviceToDevice )
-		gpu->memcpy_gpu_to_gpu( (size_t)(dst->devPtr), (size_t)src, size);
-	else {
-		printf("GPGPU-Sim PTX: cudaMemcpyToArray - ERROR : unsupported cudaMemcpyKind\n");
-		abort();
-	}
-	dst->devPtr32 = (unsigned) (size_t)(dst->devPtr);
-	return g_last_cudaError = cudaSuccess;
-}
-
-
-__host__ cudaError_t CUDARTAPI cudaMemcpyFromArray(void *dst, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t count, enum cudaMemcpyKind kind)
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-
-__host__ cudaError_t CUDARTAPI cudaMemcpyArrayToArray(struct cudaArray *dst, size_t wOffsetDst, size_t hOffsetDst, const struct cudaArray *src, size_t wOffsetSrc, size_t hOffsetSrc, size_t count, enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToDevice))
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-
-__host__ cudaError_t CUDARTAPI cudaMemcpy2D(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind)
-{
-	CUctx_st *context = GPGPUSim_Context();
-	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	size_t size = spitch*height;
-	gpgpusim_ptx_assert( (dpitch==spitch), "different src and dst pitch not supported yet" );
-	if( kind == cudaMemcpyHostToDevice )
-		gpu->memcpy_to_gpu( (size_t)dst, src, size );
-	else if( kind == cudaMemcpyDeviceToHost )
-		gpu->memcpy_from_gpu( dst, (size_t)src, size );
-	else if( kind == cudaMemcpyDeviceToDevice )
-		gpu->memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, size);
-	else {
-		printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n");
-		abort();
-	}
-	return g_last_cudaError = cudaSuccess;
-}
-
-
-__host__ cudaError_t CUDARTAPI cudaMemcpy2DToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind)
-{
-	CUctx_st *context = GPGPUSim_Context();
-	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	size_t size = spitch*height;
-	size_t channel_size = dst->desc.w+dst->desc.x+dst->desc.y+dst->desc.z;
-	gpgpusim_ptx_assert( ((channel_size%8) == 0), "none byte multiple destination channel size not supported (sz=%u)", channel_size );
-	unsigned elem_size = channel_size/8;
-	gpgpusim_ptx_assert( (dst->dimensions==2), "copy to none 2D array not supported" );
-	gpgpusim_ptx_assert( (wOffset==0), "non-zero wOffset not yet supported" );
-	gpgpusim_ptx_assert( (hOffset==0), "non-zero hOffset not yet supported" );
-	gpgpusim_ptx_assert( (dst->height == (int)height), "partial copy not supported" );
-	gpgpusim_ptx_assert( (elem_size*dst->width == width), "partial copy not supported" );
-	gpgpusim_ptx_assert( (spitch == width), "spitch != width not supported" );
-	if( kind == cudaMemcpyHostToDevice )
-		gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size);
-	else if( kind == cudaMemcpyDeviceToHost )
-		gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size);
-	else if( kind == cudaMemcpyDeviceToDevice )
-		gpu->memcpy_gpu_to_gpu( (size_t)dst->devPtr, (size_t)src, size);
-	else {
-		printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n");
-		abort();
-	}
-	dst->devPtr32 = (unsigned) (size_t)(dst->devPtr);
-	return g_last_cudaError = cudaSuccess;
-}
-
-
-__host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArray(void *dst, size_t dpitch, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t width, size_t height, enum cudaMemcpyKind kind)
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-
-__host__ cudaError_t CUDARTAPI cudaMemcpy2DArrayToArray(struct cudaArray *dst, size_t wOffsetDst, size_t hOffsetDst, const struct cudaArray *src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToDevice))
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-
-__host__ cudaError_t CUDARTAPI cudaMemcpyToSymbol(const char *symbol, const void *src, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyHostToDevice))
-{
-	//CUctx_st *context = GPGPUSim_Context();
-	assert(kind == cudaMemcpyHostToDevice);
-	printf("GPGPU-Sim PTX: cudaMemcpyToSymbol: symbol = %p\n", symbol);
-	//stream_operation( const char *symbol, const void *src, size_t count, size_t offset )
-	g_stream_manager->push( stream_operation(src,symbol,count,offset,0) );
-	//gpgpu_ptx_sim_memcpy_symbol(symbol,src,count,offset,1,context->get_device()->get_gpgpu());
-	return g_last_cudaError = cudaSuccess;
-}
-
-
-__host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbol(void *dst, const char *symbol, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost))
-{
-	//CUctx_st *context = GPGPUSim_Context();
-	assert(kind == cudaMemcpyDeviceToHost);
-	printf("GPGPU-Sim PTX: cudaMemcpyFromSymbol: symbol = %p\n", symbol);
-	g_stream_manager->push( stream_operation(symbol,dst,count,offset,0) );
-	//gpgpu_ptx_sim_memcpy_symbol(symbol,dst,count,offset,0,context->get_device()->get_gpgpu());
-	return g_last_cudaError = cudaSuccess;
-}
-
-
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-
-__host__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream)
-{
-	struct CUstream_st *s = (struct CUstream_st *)stream;
-	switch( kind ) {
-	case cudaMemcpyHostToDevice: g_stream_manager->push( stream_operation(src,(size_t)dst,count,s) ); break;
-	case cudaMemcpyDeviceToHost: g_stream_manager->push( stream_operation((size_t)src,dst,count,s) ); break;
-	case cudaMemcpyDeviceToDevice: g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,s) ); break;
-	default:
-		abort();
-	}
-	return g_last_cudaError = cudaSuccess;
-}
-
-
-__host__ cudaError_t CUDARTAPI cudaMemcpyToArrayAsync(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream)
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-
-__host__ cudaError_t CUDARTAPI cudaMemcpyFromArrayAsync(void *dst, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream)
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-
-__host__ cudaError_t CUDARTAPI cudaMemcpy2DAsync(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream)
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-
-__host__ cudaError_t CUDARTAPI cudaMemcpy2DToArrayAsync(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream)
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-
-__host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArrayAsync(void *dst, size_t dpitch, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream)
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-
-__host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count)
-{
-	CUctx_st *context = GPGPUSim_Context();
-	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	gpu->gpu_memset((size_t)mem, c, count);
-	return g_last_cudaError = cudaSuccess;
-}
-
-//memset operation is done but i think its not async?
-__host__ cudaError_t CUDARTAPI cudaMemsetAsync(void *mem, int c, size_t count, 	cudaStream_t stream=0)
-{
-	printf("GPGPU-Sim PTX: WARNING: Asynchronous memset not supported (%s)\n", __my_func__);
-	CUctx_st *context = GPGPUSim_Context();
-	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	gpu->gpu_memset((size_t)mem, c, count);
-	return g_last_cudaError = cudaSuccess;
-}
-
-__host__ cudaError_t CUDARTAPI cudaMemset2D(void *mem, size_t pitch, int c, size_t width, size_t height)
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-
-__host__ cudaError_t CUDARTAPI cudaGetSymbolAddress(void **devPtr, const char *symbol)
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-
-__host__ cudaError_t CUDARTAPI cudaGetSymbolSize(size_t *size, const char *symbol)
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-__host__ cudaError_t CUDARTAPI cudaGetDeviceCount(int *count)
-{
-	_cuda_device_id *dev = GPGPUSim_Init();
-	*count = dev->num_devices();
-	return g_last_cudaError = cudaSuccess;
-}
-
-__host__ cudaError_t CUDARTAPI cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device)
-{
-	_cuda_device_id *dev = GPGPUSim_Init();
-	if (device <= dev->num_devices() )  {
-		*prop= *dev->get_prop();
-		return g_last_cudaError = cudaSuccess;
-	} else {
-		return g_last_cudaError = cudaErrorInvalidDevice;
-	}
-}
-
-#if (CUDART_VERSION > 5000)
-__host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDeviceAttr attr, int device)
-{
-        const struct cudaDeviceProp *prop;
-        _cuda_device_id *dev = GPGPUSim_Init();
-        if (device <= dev->num_devices() )  {
-                prop = dev->get_prop();
-                switch (attr) {
-                case 5:
-                        *value= prop->maxGridSize[0];
-                        break;
-                case 6:
-                        *value= prop->maxGridSize[1];
-                        break;
-                case 7:
-                        *value= prop->maxGridSize[2];
-                        break;
-                case 10:
-                        *value= prop->warpSize;
-                        break;
-                case 12:
-                        *value= prop->regsPerBlock;
-                        break;
-                case 14:
-                        *value= prop->textureAlignment ;
-                        break;
-                case 16:
-                        *value= prop->multiProcessorCount ;
-                        break;
-                case 39:
-                        *value= dev->get_gpgpu()->threads_per_core();
-                        break;
-                case 75:
-                        *value= 8 ;
-                        break;
-                case 76:
-                        *value= 3 ;
-                        break;
-                case 78:
-                        *value= 0 ; //TODO: as of now, we dont support stream priorities.
-                        break;
-		default:
-			printf("ERROR: implement the attribute numbered %d \n",attr);
-			abort();
-                }
-                return g_last_cudaError = cudaSuccess;
-        } else {
-                return g_last_cudaError = cudaErrorInvalidDevice;
-        }
-}
-#endif
-
-__host__ cudaError_t CUDARTAPI cudaChooseDevice(int *device, const struct cudaDeviceProp *prop)
-{
-	_cuda_device_id *dev = GPGPUSim_Init();
-	*device = dev->get_id();
-	return g_last_cudaError = cudaSuccess;
-}
-
-__host__ cudaError_t CUDARTAPI cudaSetDevice(int device)
-{
-	//set the active device to run cuda
-	if ( device <= GPGPUSim_Init()->num_devices() ) {
-		g_active_device = device;
-		return g_last_cudaError = cudaSuccess;
-	} else {
-		return g_last_cudaError = cudaErrorInvalidDevice;
-	}
-}
-
-__host__ cudaError_t CUDARTAPI cudaGetDevice(int *device)
-{
-	*device = g_active_device;
-	return g_last_cudaError = cudaSuccess;
-}
-
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-
-__host__ cudaError_t CUDARTAPI cudaBindTexture(size_t *offset,
-		const struct textureReference *texref,
-		const void *devPtr,
-		const struct cudaChannelFormatDesc *desc,
-		size_t size __dv(UINT_MAX))
-{
-	CUctx_st *context = GPGPUSim_Context();
-	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	printf("GPGPU-Sim PTX: in cudaBindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference));
-	struct cudaArray *array;
-	array = (struct cudaArray*) malloc(sizeof(struct cudaArray));
-	array->desc = *desc;
-	array->size = size;
-	array->width = size;
-	array->height = 1;
-	array->dimensions = 1;
-	array->devPtr = (void*)devPtr;
-	array->devPtr32 = (int)(long long)devPtr;
-	offset = 0;
-	printf("GPGPU-Sim PTX:   size = %zu\n", size);
-	printf("GPGPU-Sim PTX:   texref = %p, array = %p\n", texref, array);
-	printf("GPGPU-Sim PTX:   devPtr32 = %x\n", array->devPtr32);
-	printf("GPGPU-Sim PTX:   Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref));
-	printf("GPGPU-Sim PTX:   ChannelFormatDesc: x=%d, y=%d, z=%d, w=%d\n", desc->x, desc->y, desc->z, desc->w);
-	printf("GPGPU-Sim PTX:   Texture Normalized? = %d\n", texref->normalized);
-	gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array);
-	devPtr = (void*)(long long)array->devPtr32;
-	printf("GPGPU-Sim PTX: devPtr = %p\n", devPtr);
-	return g_last_cudaError = cudaSuccess;
-}
-
-
-__host__ cudaError_t CUDARTAPI cudaBindTextureToArray(const struct textureReference *texref, const struct cudaArray *array, const struct cudaChannelFormatDesc *desc)
-{
-	CUctx_st *context = GPGPUSim_Context();
-	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	printf("GPGPU-Sim PTX: in cudaBindTextureToArray: %p %p\n", texref, array);
-	printf("GPGPU-Sim PTX:   devPtr32 = %x\n", array->devPtr32);
-	printf("GPGPU-Sim PTX:   Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref));
-	printf("GPGPU-Sim PTX:   Texture Normalized? = %d\n", texref->normalized);
-	gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array);
-	return g_last_cudaError = cudaSuccess;
-}
-
-__host__ cudaError_t CUDARTAPI cudaUnbindTexture(const struct textureReference *texref)
-{
-	return g_last_cudaError = cudaSuccess;
-}
-
-__host__ cudaError_t CUDARTAPI cudaGetTextureAlignmentOffset(size_t *offset, const struct textureReference *texref)
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-__host__ cudaError_t CUDARTAPI cudaGetTextureReference(const struct textureReference **texref, const char *symbol)
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-__host__ cudaError_t CUDARTAPI cudaGetChannelDesc(struct cudaChannelFormatDesc *desc, const struct cudaArray *array)
-{
-	*desc = array->desc;
-	return g_last_cudaError = cudaSuccess;
-}
-
-
-__host__ struct cudaChannelFormatDesc CUDARTAPI cudaCreateChannelDesc(int x, int y, int z, int w, enum cudaChannelFormatKind f)
-{
-	struct cudaChannelFormatDesc dummy;
-	dummy.x = x;
-	dummy.y = y;
-	dummy.z = z;
-	dummy.w = w;
-	dummy.f = f;
-	return dummy;
-}
-
-__host__ cudaError_t CUDARTAPI cudaGetLastError(void)
-{
-	return g_last_cudaError;
-}
-
-__host__ const char* CUDARTAPI cudaGetErrorString(cudaError_t error)
-{
-	if( g_last_cudaError == cudaSuccess )
-		return "no error";
-	char buf[1024];
-	snprintf(buf,1024,"<<GPGPU-Sim PTX: there was an error (code = %d)>>", g_last_cudaError);
-	return strdup(buf);
-}
-
-__host__ cudaError_t CUDARTAPI cudaConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, cudaStream_t stream)
-{
-	struct CUstream_st *s = (struct CUstream_st *)stream;
-	g_cuda_launch_stack.push_back( kernel_config(gridDim,blockDim,sharedMem,s) );
-	return g_last_cudaError = cudaSuccess;
-}
-
-__host__ cudaError_t CUDARTAPI cudaSetupArgument(const void *arg, size_t size, size_t offset)
-{
-	gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
-	kernel_config &config = g_cuda_launch_stack.back();
-	config.set_arg(arg,size,offset);
-
-	return g_last_cudaError = cudaSuccess;
-}
-
-
-__host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun )
-{
-	CUctx_st* context = GPGPUSim_Context();
-	char *mode = getenv("PTX_SIM_MODE_FUNC");
-	if( mode )
-		sscanf(mode,"%u", &g_ptx_sim_mode);
-	gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
-	kernel_config config = g_cuda_launch_stack.back();
-	struct CUstream_st *stream = config.get_stream();
-	printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun,
-			g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 );
-	kernel_info_t *grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context);
-        //do dynamic PDOM analysis for performance simulation scenario
-	std::string kname = grid->name();
-	function_info *kernel_func_info = grid->entry();
-	if (kernel_func_info->is_pdom_set()) {
-    		printf("GPGPU-Sim PTX: PDOM analysis already done for %s \n", kname.c_str() );
-    	} else {
-    		printf("GPGPU-Sim PTX: finding reconvergence points for \'%s\'...\n", kname.c_str() );
-		kernel_func_info->do_pdom();
-		kernel_func_info->set_pdom();
-    	} 
-	dim3 gridDim = config.grid_dim();
-	dim3 blockDim = config.block_dim();
-	printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n",
-			kname.c_str(), stream?stream->get_uid():0, gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z );
-	stream_operation op(grid,g_ptx_sim_mode,stream);
-	g_stream_manager->push(op);
-	g_cuda_launch_stack.pop_back();
-	return g_last_cudaError = cudaSuccess;
-}
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-
-__host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream)
-{
-	printf("GPGPU-Sim PTX: cudaStreamCreate\n");
-#if (CUDART_VERSION >= 3000)
-	*stream = new struct CUstream_st();
-	g_stream_manager->add_stream(*stream);
-#else
-	*stream = 0;
-	printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__);
-#endif
-	return g_last_cudaError = cudaSuccess;
-}
-
-//TODO: introduce priorities
-__host__ cudaError_t CUDARTAPI cudaStreamCreateWithPriority(cudaStream_t *stream, unsigned int flags, int  priority) {
-        return cudaStreamCreate(stream);
-}
-
-__host__ cudaError_t CUDARTAPI cudaDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPriority) {
-       	return cudaSuccess;	
-}
-
-__host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t *pStream, unsigned int flags) {
-	return cudaStreamCreate(pStream);
-}
-
-__host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream)
-{
-#if (CUDART_VERSION >= 3000)
-	g_stream_manager->destroy_stream(stream);
-#endif
-	return g_last_cudaError = cudaSuccess;
-}
-
-__host__ cudaError_t CUDARTAPI cudaStreamSynchronize(cudaStream_t stream)
-{
-#if (CUDART_VERSION >= 3000)
-	if( stream == NULL )
-		synchronize();
-		return g_last_cudaError = cudaSuccess;
-	stream->synchronize();
-#else
-	printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__);
-#endif
-	return g_last_cudaError = cudaSuccess;
-}
-
-__host__ cudaError_t CUDARTAPI cudaStreamQuery(cudaStream_t stream)
-{
-#if (CUDART_VERSION >= 3000)
-	if( stream == NULL )
-		return g_last_cudaError = cudaErrorInvalidResourceHandle;
-	return g_last_cudaError = stream->empty()?cudaSuccess:cudaErrorNotReady;
-#else
-	printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__);
-	return g_last_cudaError = cudaSuccess; // it is always success because all cuda calls are synchronous
-#endif
-}
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-
-__host__ cudaError_t CUDARTAPI cudaEventCreate(cudaEvent_t *event)
-{
-	CUevent_st *e = new CUevent_st(false);
-	g_timer_events[e->get_uid()] = e;
-#if CUDART_VERSION >= 3000
-	*event = e;
-#else
-	*event = e->get_uid();
-#endif
-	return g_last_cudaError = cudaSuccess;
-}
-
-CUevent_st *get_event(cudaEvent_t event)
-{
-	unsigned event_uid;
-#if CUDART_VERSION >= 3000
-	event_uid = event->get_uid();
-#else
-	event_uid = event;
-#endif
-	event_tracker_t::iterator e = g_timer_events.find(event_uid);
-	if( e == g_timer_events.end() )
-		return NULL;
-	return e->second;
-}
-
-__host__ cudaError_t CUDARTAPI cudaEventRecord(cudaEvent_t event, cudaStream_t stream)
-{
-	CUevent_st *e = get_event(event);
-	if( !e ) return g_last_cudaError = cudaErrorUnknown;
-	struct CUstream_st *s = (struct CUstream_st *)stream;
-	stream_operation op(e,s);
-	g_stream_manager->push(op);
-	return g_last_cudaError = cudaSuccess;
-}
-
-__host__ cudaError_t CUDARTAPI cudaEventQuery(cudaEvent_t event)
-{
-	CUevent_st *e = get_event(event);
-	if( e == NULL ) {
-		return g_last_cudaError = cudaErrorInvalidValue;
-	} else if( e->done() ) {
-		return g_last_cudaError = cudaSuccess;
-	} else {
-		return g_last_cudaError = cudaErrorNotReady;
-	}
-}
-
-__host__ cudaError_t CUDARTAPI cudaEventSynchronize(cudaEvent_t event)
-{
-	printf("GPGPU-Sim API: cudaEventSynchronize ** waiting for event\n");
-	fflush(stdout);
-	CUevent_st *e = (CUevent_st*) event;
-	while( !e->done() )
-		;
-	printf("GPGPU-Sim API: cudaEventSynchronize ** event detected\n");
-	fflush(stdout);
-	return g_last_cudaError = cudaSuccess;
-}
-
-__host__ cudaError_t CUDARTAPI cudaEventDestroy(cudaEvent_t event)
-{
-	CUevent_st *e = get_event(event);
-	unsigned event_uid = e->get_uid();
-	event_tracker_t::iterator pe = g_timer_events.find(event_uid);
-	if( pe == g_timer_events.end() )
-		return g_last_cudaError = cudaErrorInvalidValue;
-	g_timer_events.erase(pe);
-	return g_last_cudaError = cudaSuccess;
-}
-
-
-__host__ cudaError_t CUDARTAPI cudaEventElapsedTime(float *ms, cudaEvent_t start, cudaEvent_t end)
-{
-	time_t elapsed_time;
-	CUevent_st *s = get_event(start);
-	CUevent_st *e = get_event(end);
-	if( s==NULL || e==NULL )
-		return g_last_cudaError = cudaErrorUnknown;
-	elapsed_time = e->clock() - s->clock();
-	*ms = 1000*elapsed_time;
-	return g_last_cudaError = cudaSuccess;
-}
-
-
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-
-__host__ cudaError_t CUDARTAPI cudaThreadExit(void)
-{
-	exit_simulation();
-	return g_last_cudaError = cudaSuccess;
-}
-
-__host__ cudaError_t CUDARTAPI cudaThreadSynchronize(void)
-{
-	//Called on host side
-	synchronize();
-	return g_last_cudaError = cudaSuccess;
-};
-
-int CUDARTAPI __cudaSynchronizeThreads(void**, void*)
-{
-	return cudaThreadExit();
-}
-
-
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-
-#if (CUDART_VERSION >= 3010)
-
-typedef struct CUuuid_st {                                /**< CUDA definition of UUID */
-    char bytes[16];
-} CUuuid;
-
-/**
- * CUDA UUID types
- */
-// typedef __device_builtin__ struct CUuuid_st cudaUUID_t;
-
-__host__ cudaError_t CUDARTAPI cudaGetExportTable(const void **ppExportTable, const cudaUUID_t *pExportTableId)
-{
-	printf("cudaGetExportTable: UUID = "); 
-	for (int s = 0; s < 16; s++) {
-		printf("%#2x ", (unsigned char) (pExportTableId->bytes[s])); 
-	}
-	printf("\n"); 
-	return g_last_cudaError = cudaSuccess;
-}
-
-#endif
-
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-
-//#include "../../cuobjdump_to_ptxplus/cuobjdump_parser.h"
-
-enum cuobjdumpSectionType {
-	PTXSECTION=0,
-	ELFSECTION
-};
-
-
-class cuobjdumpSection {
-public:
-	//Constructor
-	cuobjdumpSection() {
-		arch = 0;
-		identifier = "";
-	}
-	virtual ~cuobjdumpSection() {}
-	unsigned getArch() {return arch;}
-	void setArch(unsigned a) {arch = a;}
-	std::string getIdentifier() {return identifier;}
-	void setIdentifier(std::string i) {identifier = i;}
-	virtual void print(){std::cout << "cuobjdump Section: unknown type" << std::endl;}
-private:
-	unsigned arch;
-	std::string identifier;
-};
-
-class cuobjdumpELFSection : public cuobjdumpSection
-{
-public:
-	cuobjdumpELFSection() {}
-	virtual ~cuobjdumpELFSection() {
-		elffilename = "";
-		sassfilename = "";
-	}
-	std::string getELFfilename() {return elffilename;}
-	void setELFfilename(std::string f) {elffilename = f;}
-	std::string getSASSfilename() {return sassfilename;}
-	void setSASSfilename(std::string f) {sassfilename = f;}
-	virtual void print() {
-		std::cout << "ELF Section:" << std::endl;
-		std::cout << "arch: sm_" << getArch() << std::endl;
-		std::cout << "identifier: " << getIdentifier() << std::endl;
-		std::cout << "elf filename: " << getELFfilename() << std::endl;
-		std::cout << "sass filename: " << getSASSfilename() << std::endl;
-		std::cout << std::endl;
-	}
-private:
-	std::string elffilename;
-	std::string sassfilename;
-};
-
-class cuobjdumpPTXSection : public cuobjdumpSection
-{
-public:
-	cuobjdumpPTXSection(){
-		ptxfilename = "";
-	}
-	std::string getPTXfilename() {return ptxfilename;}
-	void setPTXfilename(std::string f) {ptxfilename = f;}
-	virtual void print() {
-		std::cout << "PTX Section:" << std::endl;
-		std::cout << "arch: sm_" << getArch() << std::endl;
-		std::cout << "identifier: " << getIdentifier() << std::endl;
-		std::cout << "ptx filename: " << getPTXfilename() << std::endl;
-		std::cout << std::endl;
-	}
-private:
-	std::string ptxfilename;
-};
-
-std::list<cuobjdumpSection*> cuobjdumpSectionList;
-std::list<cuobjdumpSection*> libSectionList;
-
-// sectiontype: 0 for ptx, 1 for elf
-void addCuobjdumpSection(int sectiontype){
-	if (sectiontype)
-		cuobjdumpSectionList.push_front(new cuobjdumpELFSection());
-	else
-		cuobjdumpSectionList.push_front(new cuobjdumpPTXSection());
-	printf("## Adding new section %s\n", sectiontype?"ELF":"PTX");
-}
-
-void setCuobjdumparch(const char* arch){
-	unsigned archnum;
-	sscanf(arch, "sm_%u", &archnum);
-	assert (archnum && "cannot have sm_0");
-	printf("Adding arch: %s\n", arch);
-	cuobjdumpSectionList.front()->setArch(archnum);
-}
-
-void setCuobjdumpidentifier(const char* identifier){
-	printf("Adding identifier: %s\n", identifier);
-	cuobjdumpSectionList.front()->setIdentifier(identifier);
-}
-
-void setCuobjdumpptxfilename(const char* filename){
-	printf("Adding ptx filename: %s\n", filename);
-	cuobjdumpSection* x = cuobjdumpSectionList.front();
-	if (dynamic_cast<cuobjdumpPTXSection*>(x) == NULL){
-		assert (0 && "You shouldn't be trying to add a ptxfilename to an elf section");
-	}
-	(dynamic_cast<cuobjdumpPTXSection*>(x))->setPTXfilename(filename);
-}
-
-void setCuobjdumpelffilename(const char* filename){
-	if (dynamic_cast<cuobjdumpELFSection*>(cuobjdumpSectionList.front()) == NULL){
-		assert (0 && "You shouldn't be trying to add a elffilename to an ptx section");
-	}
-	(dynamic_cast<cuobjdumpELFSection*>(cuobjdumpSectionList.front()))->setELFfilename(filename);
-}
-
-void setCuobjdumpsassfilename(const char* filename){
-	if (dynamic_cast<cuobjdumpELFSection*>(cuobjdumpSectionList.front()) == NULL){
-		assert (0 && "You shouldn't be trying to add a sassfilename to an ptx section");
-	}
-	(dynamic_cast<cuobjdumpELFSection*>(cuobjdumpSectionList.front()))->setSASSfilename(filename);
-}
-extern int cuobjdump_parse();
-extern FILE *cuobjdump_in;
-
-//! Return the executable file of the process containing the PTX/SASS code 
-//!
-//! This Function returns the executable file ran by the process.  This
-//! executable is supposed to contain the PTX/SASS code.  It provides workaround
-//! for processes running on valgrind by dereferencing /proc/<pid>/exe within the
-//! GPGPU-Sim process before calling cuobjdump to extract PTX/SASS.  This is
-//! needed because valgrind uses x86 emulation to detect memory leak.  Other
-//! processes (e.g. cuobjdump) reading /proc/<pid>/exe will see the emulator
-//! executable instead of the application binary.  
-//! 
-std::string get_app_binary(){
-   char self_exe_path[1025];
-#ifdef __APPLE__
-   uint32_t size = sizeof(self_exe_path);
-   if( _NSGetExecutablePath(self_exe_path,&size) != 0 ) {
-	   printf("GPGPU-Sim ** ERROR: _NSGetExecutablePath input buffer too small\n");
-	   exit(1);
-   }
-#else
-   std::stringstream exec_link;
-   exec_link << "/proc/self/exe";
-
-   ssize_t path_length = readlink(exec_link.str().c_str(), self_exe_path, 1024); 
-   assert(path_length != -1); 
-   self_exe_path[path_length] = '\0'; 
-#endif
-
-   printf("self exe links to: %s\n", self_exe_path); 
-   return self_exe_path; 
-}
-
-//above func gives abs path whereas this give just the name of application.
-char* get_app_binary_name(std::string abs_path){
-   char *self_exe_path;
-#ifdef __APPLE__
-   //TODO: get apple device and check the result.
-   printf("WARNING: not tested for Apple-mac devices \n");
-   abort();
-#else
-   char* buf = strdup(abs_path.c_str());
-   char *token = strtok(buf, "/");
-   while(token !=NULL){
-	self_exe_path = token;
-	token = strtok(NULL,"/");
-   }
-#endif
-   self_exe_path = strtok(self_exe_path, ".");
-   printf("self exe links to: %s\n", self_exe_path);
-   return self_exe_path;
-}
-
-//! Call cuobjdump to extract everything (-elf -sass -ptx)
-/*!
- *	This Function extract the whole PTX (for all the files) using cuobjdump
- *	to _cuobjdump_complete_output_XXXXXX then runs a parser to chop it up with each binary in
- *	its own file
- *	It is also responsible for extracting the libraries linked to the binary if the option is
- *	enabled
- * */
-void extract_code_using_cuobjdump(){
-    CUctx_st *context = GPGPUSim_Context();
-    unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
-
-    //prevent the dumping by cuobjdump everytime we execute the code!
-    const char *override_cuobjdump = getenv("CUOBJDUMP_SIM_FILE"); 
-    char command[1000], ptx_file[1000];
-    std::string app_binary = get_app_binary(); 
-    //Running cuobjdump using dynamic link to current process
-    snprintf(command,1000,"md5sum %s ", app_binary.c_str());
-    printf("Running md5sum using \"%s\"\n", command);
-    system(command);
-    // Running cuobjdump using dynamic link to current process
-    // Needs the option '-all' to extract PTX from CDP-enabled binary 
-    extern bool g_cdp_enabled;
-
-    //dump ptx for all individial ptx files into sepearte files which is later used by ptxas.
-    int result=0;
-#if (CUDART_VERSION >= 6000)
-    char fname2[1024];
-    snprintf(fname2,1024,"_cuobjdump_list_ptx_XXXXXX");
-    int fd2=mkstemp(fname2);
-    close(fd2);
-    snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -lptx -arch=sm_%u %s > %s", forced_max_capability, app_binary.c_str(), fname2);
-    result = system(command);
-    if( result != 0 ) {
-        printf("WARNING: Failed to execute cuobjdump to get list of ptx files \n");
-        exit(0);
-    } else {
-	/*
-          as we got list of ptx files, we need to extract one by one into seperate files so that ptxas can understand it. 
-          In this way, the duplicate definitions in a single embedded file can be prevented. 
-          No of lines in the file is equal to no of ptx fileis available.
-        */
-        FILE *fp = fopen(fname2,"r");
-        if (fp==NULL) {
-            printf("WARNING: cuobjdump file error! Could not open file %s \n", fname2);
-            exit(0);
-        } else {
-            for (char c = getc(fp); c != EOF; c = getc(fp))
-                if (c == '\n')
-                    no_of_ptx = no_of_ptx + 1;
-                    fclose(fp);
-        }
-	if(no_of_ptx==0){
-		printf("WARNING: Number of ptx in the executable file are 0. One of the reasons might be\n");
-		printf("\t1. CDP is enabled\n");
-		printf("\t2. cuobjdump -lptx doesnt recognize sm_%u\n",forced_max_capability);
-		printf("\t3. the application was not compiled iwth nvcc flag sm_%u\n",forced_max_capability);
-	}
-    }
-    if(!g_cdp_enabled) {
-        //based on the list above, dump ptx files individually. Format of dumped ptx file is prog_name.unique_no.sm_<>.ptx
-        for (int index=1; index<= no_of_ptx; index++){
-            snprintf(ptx_file, 1000, "%s.%d.sm_%u.ptx", get_app_binary_name(app_binary), index, forced_max_capability);
-            printf("Extracting specific PTX file named %s \n",ptx_file);
-            snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -arch=sm_%u -xptx %s %s", forced_max_capability, ptx_file, app_binary.c_str());
-            if (system(command)!=0) {
-                printf("ERROR: command: %s failed \n",command);
-                exit(0);
-            }
-        }
-    }
-#endif
-    //TODO: redundant to dump twice. how can it be prevented?
-    //dump only for specific arch
-    char fname[1024];
-    if ((override_cuobjdump == NULL) || (strlen(override_cuobjdump)==0)) {
-	snprintf(fname,1024,"_cuobjdump_complete_output_XXXXXX");
-	int fd=mkstemp(fname);
-	close(fd);
-	if(!g_cdp_enabled)
-            snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass %s > %s", app_binary.c_str(), fname);
-	else
-            snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass -all %s > %s", app_binary.c_str(), fname);
-	bool parse_output = true; 
-	result = system(command);
-	if(result) {
-            if (context->get_device()->get_gpgpu()->get_config().experimental_lib_support() && (result == 65280)) {  
-                // Some CUDA application may exclusively use kernels provided by CUDA
-                // libraries (e.g. CUBLAS).  Skipping cuobjdump extraction from the
-                // executable for this case. 
-                // 65280 is the return code from cuobjdump denoting the specific error (tested on CUDA 4.0/4.1/4.2)
-                printf("WARNING: Failed to execute: %s\n", command); 
-                printf("         Executable binary does not contain any GPU kernel.\n"); 
-                parse_output = false; 
-            } else {
-                printf("ERROR: Failed to execute: %s\n", command); 
-                exit(1);
-            }
-        }
-
-        if (parse_output) {
-            printf("Parsing file %s\n", fname);
-            cuobjdump_in = fopen(fname, "r");
-
-            cuobjdump_parse();
-            fclose(cuobjdump_in);
-            printf("Done parsing!!!\n");
-        } else {
-            printf("Parsing skipped for %s\n", fname); 
-        }
-
-        if (context->get_device()->get_gpgpu()->get_config().experimental_lib_support()){
-            //Experimental library support
-            //Currently only for cufft
-
-            std::stringstream cmd;
-            cmd << "ldd " << app_binary << " | grep $CUDA_INSTALL_PATH | awk \'{print $3}\' > _tempfile_.txt";
-            int result = system(cmd.str().c_str());
-            if(result){
-                std::cout << "Failed to execute: " << cmd.str() << std::endl;
-                exit(1);
-            }
-            std::ifstream libsf;
-            libsf.open("_tempfile_.txt");
-            if(!libsf.is_open()) {
-                std::cout << "Failed to open: _tempfile_.txt" << std::endl;
-                exit(1);
-            }
-
-            //Save the original section list
-            std::list<cuobjdumpSection*> tmpsl = cuobjdumpSectionList;
-            cuobjdumpSectionList.clear();
-
-            std::string line;
-            std::getline(libsf, line);
-            std::cout << "DOING: " << line << std::endl;
-            int cnt=1;
-            while(libsf.good()){
-                std::stringstream libcodfn;
-                libcodfn << "_cuobjdump_complete_lib_" << cnt << "_";
-                cmd.str(""); //resetting
-                cmd << "$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass ";
-                cmd << line;
-                cmd << " > ";
-                cmd << libcodfn.str();
-                std::cout << "Running cuobjdump on " << line << std::endl;
-                std::cout << "Using command: " << cmd.str() << std::endl;
-                result = system(cmd.str().c_str());
-                if(result) {printf("ERROR: Failed to execute: %s\n", command); exit(1);}
-                    std::cout << "Done" << std::endl;
-
-                    std::cout << "Trying to parse " << libcodfn.str() << std::endl;
-                    cuobjdump_in = fopen(libcodfn.str().c_str(), "r");
-                    cuobjdump_parse();
-                    fclose(cuobjdump_in);
-                    std::getline(libsf, line);
-            }
-            libSectionList = cuobjdumpSectionList;
-
-            //Restore the original section list
-            cuobjdumpSectionList = tmpsl;
-        }
-    } else {
-        printf("GPGPU-Sim PTX: overriding cuobjdump with '%s' (CUOBJDUMP_SIM_FILE is set)\n", override_cuobjdump);
-        snprintf(fname,1024, "%s",override_cuobjdump);
-    }
-}
-
-//! Read file into char*
-//TODO: convert this to C++ streams, will be way cleaner
-char* readfile (const std::string filename){
-	assert (filename != "");
-	FILE* fp = fopen(filename.c_str(),"r");
-	if (!fp) {
-		std::cout << "ERROR: Could not open file %s for reading\n" << filename << std::endl;
-		assert (0);
-	}
-	// finding size of the file
-	int filesize= 0;
-	fseek (fp , 0 , SEEK_END);
-
-	filesize = ftell (fp);
-	fseek (fp, 0, SEEK_SET);
-	// allocate and copy the entire ptx
-	char* ret = (char*)malloc((filesize +1)* sizeof(char));
-	fread(ret,1,filesize,fp);
-	ret[filesize]='\0';
-	fclose(fp);
-	return ret;
-}
-
-//! Function that helps debugging
-void printSectionList(std::list<cuobjdumpSection*> sl) {
-	std::list<cuobjdumpSection*>::iterator iter;
-	for (	iter = sl.begin();
-			iter != sl.end();
-			iter++
-	){
-		(*iter)->print();
-	}
-}
-
-//! Remove unecessary sm versions from the section list
-std::list<cuobjdumpSection*> pruneSectionList(std::list<cuobjdumpSection*> cuobjdumpSectionList, CUctx_st *context) {
-	unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
-
-	//For ptxplus, force the max capability to 19 if it's higher or unspecified(0)
-	if (context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus()){
-		if (	(forced_max_capability == 0) ||
-				(forced_max_capability >= 20)){
-			printf("GPGPU-Sim: WARNING: Capability >= 20 are not supported in PTXPlus\n\tSetting forced_max_capability to 19\n");
-			forced_max_capability = 19;
-		}
-	}
-
-	std::list<cuobjdumpSection*> prunedList;
-
-	//Find the highest capability (that is lower than the forced maximum) for each cubin file
-	//and set it in cuobjdumpSectionMap. Do this only for ptx sections
-	std::map<std::string, unsigned> cuobjdumpSectionMap;
-	int min_ptx_capability_found=0;
-	for (	std::list<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin();
-			iter != cuobjdumpSectionList.end();
-			iter++){
-		unsigned capability = (*iter)->getArch();
-		if(dynamic_cast<cuobjdumpPTXSection*>(*iter) != NULL){
-			if(capability<min_ptx_capability_found || min_ptx_capability_found==0)
-				min_ptx_capability_found=capability;
-			if (capability <= forced_max_capability ||	forced_max_capability==0) {
-				if((cuobjdumpSectionMap.find((*iter)->getIdentifier())==cuobjdumpSectionMap.end())
-						|| (cuobjdumpSectionMap[(*iter)->getIdentifier()] < capability))
-						cuobjdumpSectionMap[(*iter)->getIdentifier()] = capability;
-			}
-		}
-	}
-
-	//Throw away the sections with the lower capabilites and push those with the highest in
-	//the pruned list
-	for (	std::list<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin();
-			iter != cuobjdumpSectionList.end();
-			iter++){
-		unsigned capability = (*iter)->getArch();
-		if(capability == cuobjdumpSectionMap[(*iter)->getIdentifier()]){
-			prunedList.push_back(*iter);
-		} else {
-			delete *iter;
-		}
-	}
-	if(prunedList.empty()){
-		printf("Error: No PTX sections found with sm capability that is lower than current forced maximum capability \n minimum ptx capability found = %u, maximum forced ptx capability = %u \n User might want to change either the forced maximum capability from gpgpusim configuration or update the compilation to generate the required PTX version\n",min_ptx_capability_found,forced_max_capability);
-		abort();
-	}
-	return prunedList;
-}
-
-//! Merge all PTX sections that have a specific identifier into one file
-std::list<cuobjdumpSection*> mergeMatchingSections(std::list<cuobjdumpSection*> cuobjdumpSectionList, std::string identifier){
-	const char *ptxcode = "";
-	std::list<cuobjdumpSection*>::iterator old_iter;
-	cuobjdumpPTXSection* old_ptxsection = NULL;
-	cuobjdumpPTXSection* ptxsection;
-	std::list<cuobjdumpSection*> mergedList;
-
-	for (	std::list<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin();
-			iter != cuobjdumpSectionList.end();
-			iter++){
-		if((ptxsection=dynamic_cast<cuobjdumpPTXSection*>(*iter)) != NULL &&
-			strcmp(ptxsection->getIdentifier().c_str(), identifier.c_str()) == 0){
-			// Read and remove the last PTX section
-			if (old_ptxsection != NULL) {
-				ptxcode = readfile(old_ptxsection->getPTXfilename());
-				// remove ptx file?
-				delete *old_iter;
-			}
-
-			// Append all the PTX from the last PTX section into the current PTX section
-			// Add 50 to ptxcode to ignore the information regarding version/target/address_size
-			if (strlen(ptxcode) >= 50) {
-				FILE *ptxfile = fopen((ptxsection->getPTXfilename()).c_str(), "a");
-				fprintf(ptxfile, "%s", ptxcode + 50);
-				fclose(ptxfile);
-			}
-
-			old_iter = iter;
-			old_ptxsection = ptxsection;
-		}
-		// Store all non-PTX sections and PTX sections with non-matching identifiers
-		else {
-			mergedList.push_back(*iter);
-		}
-	}
-
-	// Store the final PTX section
-	mergedList.push_back(*old_iter);
-
-	return mergedList;
-}
-
-//! Merge any PTX sections with matching identifiers
-std::list<cuobjdumpSection*> mergeSections(std::list<cuobjdumpSection*> cuobjdumpSectionList){
-	std::vector<std::string> identifier;
-	cuobjdumpPTXSection* ptxsection;
-
-	// Add all identifiers present in PTX sections to a vector
-	for (	std::list<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin();
-			iter != cuobjdumpSectionList.end();
-			iter++){
-		if((ptxsection=dynamic_cast<cuobjdumpPTXSection*>(*iter)) != NULL){
-			std::string current_id = ptxsection->getIdentifier();
-
-			// If we haven't yet seen a given identifier, add it to the vector
-			if (std::find(identifier.begin(), identifier.end(), current_id) == identifier.end()) {
-				identifier.push_back(current_id);
-			}
-		}
-	}
-
-	// Call mergeMatchingSections on all identifiers in the vector
-	for (	std::vector<std::string>::iterator iter = identifier.begin();
-			iter != identifier.end();
-			iter++) {
-		cuobjdumpSectionList = mergeMatchingSections(cuobjdumpSectionList, *iter);
-	}
-
-	return cuobjdumpSectionList;
-}
-
-
-//! Within the section list, find the ELF section corresponding to a given identifier
-cuobjdumpELFSection* findELFSectionInList(std::list<cuobjdumpSection*> sectionlist, const std::string identifier){
-
-	std::list<cuobjdumpSection*>::iterator iter;
-	for (	iter = sectionlist.begin();
-			iter != sectionlist.end();
-			iter++
-	){
-		cuobjdumpELFSection* elfsection;
-		if((elfsection=dynamic_cast<cuobjdumpELFSection*>(*iter)) != NULL){
-			if(elfsection->getIdentifier() == identifier)
-				return elfsection;
-		}
-	}
-	return NULL;
-}
-
-//! Find an ELF section in all the known lists
-cuobjdumpELFSection* findELFSection(const std::string identifier){
-	cuobjdumpELFSection* sec = findELFSectionInList(cuobjdumpSectionList, identifier);
-	if (sec!=NULL)return sec;
-	sec = findELFSectionInList(libSectionList, identifier);
-	if (sec!=NULL)return sec;
-	std::cout << "Could not find " << identifier << std::endl;
-	assert(0 && "Could not find the required ELF section");
-	return NULL;
-}
-
-//! Within the section list, find the PTX section corresponding to a given identifier
-cuobjdumpPTXSection* findPTXSectionInList(std::list<cuobjdumpSection*> sectionlist, const std::string identifier){
-	std::list<cuobjdumpSection*>::iterator iter;
-	for (	iter = sectionlist.begin();
-			iter != sectionlist.end();
-			iter++
-	){
-		cuobjdumpPTXSection* ptxsection;
-		if((ptxsection=dynamic_cast<cuobjdumpPTXSection*>(*iter)) != NULL){
-			if(ptxsection->getIdentifier() == identifier)
-				return ptxsection;
-			else {
-				extern bool g_cdp_enabled;
-				if(g_cdp_enabled) {
-					printf("Warning: __cudaRegisterFatBinary needs %s, but find PTX section with %s\n",
-						identifier.c_str(), ptxsection->getIdentifier().c_str());
-					return ptxsection;
-				}
-			}
-		}
-	}
-	return NULL;
-}
-
-//! Find an PTX section in all the known lists
-cuobjdumpPTXSection* findPTXSection(const std::string identifier){
-	cuobjdumpPTXSection* sec = findPTXSectionInList(cuobjdumpSectionList, identifier);
-	if (sec!=NULL)return sec;
-	sec = findPTXSectionInList(libSectionList, identifier);
-	if (sec!=NULL)return sec;
-	std::cout << "Could not find " << identifier << std::endl;
-	assert(0 && "Could not find the required PTX section");
-	return NULL;
-}
-
-
-
-//! Extract the code using cuobjdump and remove unnecessary sections
-void cuobjdumpInit(){
-	CUctx_st *context = GPGPUSim_Context();
-	extract_code_using_cuobjdump(); //extract all the output of cuobjdump to _cuobjdump_*.*
-	const char* pre_load = getenv("CUOBJDUMP_SIM_FILE");
-	if (pre_load ==NULL || strlen(pre_load)==0){
-		cuobjdumpSectionList = pruneSectionList(cuobjdumpSectionList, context);
-		cuobjdumpSectionList = mergeSections(cuobjdumpSectionList);
-	}
-}
-
-std::map<int, std::string> fatbinmap;
-std::map<int, bool>fatbin_registered;
-std::map<std::string, symbol_table*> name_symtab;
-
-//! Keep track of the association between filename and cubin handle
-void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename){
-	fatbinmap[handle] = filename;
-}
-
-//! Either submit PTX for simulation or convert SASS to PTXPlus and submit it
-void cuobjdumpParseBinary(unsigned int handle){
-
-	if(fatbin_registered[handle]) return;
-	fatbin_registered[handle] = true;
-	CUctx_st *context = GPGPUSim_Context();
-	std::string fname = fatbinmap[handle];
-
-	if (name_symtab.find(fname) != name_symtab.end()) {
-		symbol_table *symtab = name_symtab[fname];
-		context->add_binary(symtab, handle);
-		return;
-	}
-
-	unsigned max_capability = 0;
-	for (	std::list<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin();
-			iter != cuobjdumpSectionList.end();
-			iter++){
-		unsigned capability = (*iter)->getArch();
-		if (capability > max_capability) max_capability = capability;
-	}
-	if (max_capability > 20) printf("WARNING: No guarantee that PTX will be parsed for SM version %u\n", max_capability);
-
-	cuobjdumpPTXSection* ptx = NULL;
-	const char* pre_load = getenv("CUOBJDUMP_SIM_FILE");
-	if(pre_load==NULL || strlen(pre_load)==0)
-		ptx = findPTXSection(fname);
-	symbol_table *symtab;
-	char *ptxcode;
-	const char *override_ptx_name = getenv("PTX_SIM_KERNELFILE"); 
-	if (override_ptx_name == NULL or getenv("PTX_SIM_USE_PTX_FILE") == NULL or strlen(getenv("PTX_SIM_USE_PTX_FILE"))==0) {
-		ptxcode = readfile(ptx->getPTXfilename());
-	} else {
-		printf("GPGPU-Sim PTX: overriding embedded ptx with '%s' (PTX_SIM_USE_PTX_FILE is set)\n", override_ptx_name);
-		ptxcode = readfile(override_ptx_name);
-	}
-	if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) {
-		cuobjdumpELFSection* elfsection = findELFSection(ptx->getIdentifier());
-		assert (elfsection!= NULL);
-		char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(
-				ptx->getPTXfilename(),
-				elfsection->getELFfilename(),
-				elfsection->getSASSfilename());
-		symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxplus_str, handle);
-		printf("Adding %s with cubin handle %u\n", ptx->getPTXfilename().c_str(), handle);
-		context->add_binary(symtab, handle);
-		gpgpu_ptxinfo_load_from_string( ptxcode, handle, max_capability );
-		delete[] ptxplus_str;
-	} else {
-		symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxcode, handle);
-		//if CUOBJDUMP_SIM_FILE is not set, ptx is NULL. So comment below.
-		//printf("Adding %s with cubin handle %u\n", ptx->getPTXfilename().c_str(), handle);
-		context->add_binary(symtab, handle);
-		gpgpu_ptxinfo_load_from_string( ptxcode, handle, max_capability );
-	}
-	load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
-	load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
-	name_symtab[fname] = symtab;
-
-	//TODO: Remove temporarily files as per configurations
-}
-
-void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
-{
-#if (CUDART_VERSION < 2010)
-	printf("GPGPU-Sim PTX: ERROR ** this version of GPGPU-Sim requires CUDA 2.1 or higher\n");
-	exit(1);
-#endif
-	CUctx_st *context = GPGPUSim_Context();
-	static unsigned next_fat_bin_handle = 1;
-	if(context->get_device()->get_gpgpu()->get_config().use_cuobjdump()) {
-		// The following workaround has only been verified on 64-bit systems. 
-		if (sizeof(void*) == 4) 
-			printf("GPGPU-Sim PTX: FatBin file name extraction has not been tested on 32-bit system.\n"); 
-
-		#if (CUDART_VERSION <= 6000)
-		// FatBin handle from the .fatbin.c file (one of the intermediate files generated by NVCC)
-		typedef struct {int m; int v; const unsigned long long* d; char* f;} __fatDeviceText __attribute__ ((aligned (8))); 
-		__fatDeviceText * fatDeviceText = (__fatDeviceText *) fatCubin;
-
-		// Extract the source code file name that generate the given FatBin. 
-		// - Obtains the pointer to the actual fatbin structure from the FatBin handle (fatCubin).
-		// - An integer inside the fatbin structure contains the relative offset to the source code file name.
-		// - This offset differs among different CUDA and GCC versions. 
-		char * pfatbin = (char*) fatDeviceText->d; 
-		int offset = *((int*)(pfatbin+48)); 
-		char * filename = (pfatbin+16+offset); 
-		#else
-		const char * filename = "default";
-		#endif
-		// The extracted file name is associated with a fat_cubin_handle passed
-		// into cudaLaunch().  Inside cudaLaunch(), the associated file name is
-		// used to find the PTX/SASS section from cuobjdump, which contains the
-		// PTX/SASS code for the launched kernel function.  
-		// This allows us to work around the fact that cuobjdump only outputs the
-		// file name associated with each section. 
-		unsigned long long fat_cubin_handle = next_fat_bin_handle;
-		next_fat_bin_handle++;
-		printf("GPGPU-Sim PTX: __cudaRegisterFatBinary, fat_cubin_handle = %llu, filename=%s\n", fat_cubin_handle, filename);
-		/*!
-		 * This function extracts all data from all files in first call
-		 * then for next calls, only returns the appropriate number
-		 */
-		assert(fat_cubin_handle >= 1);
-		if (fat_cubin_handle==1) cuobjdumpInit();
-		cuobjdumpRegisterFatBinary(fat_cubin_handle, filename);
-
-		return (void**)fat_cubin_handle;
-	} 
-#if (CUDART_VERSION < 8000)
-	else {
-		static unsigned source_num=1;
-		unsigned long long fat_cubin_handle = next_fat_bin_handle++;
-		__cudaFatCudaBinary *info =   (__cudaFatCudaBinary *)fatCubin;
-		assert( info->version >= 3 );
-		unsigned num_ptx_versions=0;
-		unsigned max_capability=0;
-		unsigned selected_capability=0;
-		bool found=false;
-		unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
-		if (!info->ptx){
-			printf("ERROR: Cannot find ptx code in cubin file\n"
-					"\tIf you are using CUDA 4.0 or higher, please enable -gpgpu_ptx_use_cuobjdump or downgrade to CUDA 3.1\n");
-			exit(1);
-		}
-		while( info->ptx[num_ptx_versions].gpuProfileName != NULL ) {
-			unsigned capability=0;
-			sscanf(info->ptx[num_ptx_versions].gpuProfileName,"compute_%u",&capability);
-			printf("GPGPU-Sim PTX: __cudaRegisterFatBinary found PTX versions for '%s', ", info->ident);
-			printf("capability = %s\n", info->ptx[num_ptx_versions].gpuProfileName );
-			if( forced_max_capability ) {
-				if( capability > max_capability && capability <= forced_max_capability ) {
-					found = true;
-					max_capability=capability;
-					selected_capability = num_ptx_versions;
-				}
-			} else {
-				if( capability > max_capability ) {
-					found = true;
-					max_capability=capability;
-					selected_capability = num_ptx_versions;
-				}
-			}
-			num_ptx_versions++;
-		}
-		if( found  ) {
-			printf("GPGPU-Sim PTX: Loading PTX for %s, capability = %s\n",
-					info->ident, info->ptx[selected_capability].gpuProfileName );
-			symbol_table *symtab;
-			const char *ptx = info->ptx[selected_capability].ptx;
-			if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) {
-				printf("GPGPU-Sim PTX: ERROR ** PTXPlus is only supported through cuobjdump\n"
-						"\tEither enable cuobjdump or disable PTXPlus in your configuration file\n");
-				exit(1);
-			} else {
-				symtab=gpgpu_ptx_sim_load_ptx_from_string(ptx,source_num);
-				context->add_binary(symtab,fat_cubin_handle);
-				gpgpu_ptxinfo_load_from_string( ptx, source_num, max_capability );
-			}
-			source_num++;
-			load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
-			load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
-		} else {
-			printf("GPGPU-Sim PTX: warning -- did not find an appropriate PTX in cubin\n");
-		}
-		return (void**)fat_cubin_handle;
-	}
-#else
-        else {
-		printf("ERROR **  __cudaRegisterFatBinary() needs to be updated\n");
-		abort();
-        }
-#endif
-}
-
-void __cudaUnregisterFatBinary(void **fatCubinHandle)
-{
-	;
-}
-
-cudaError_t cudaDeviceReset ( void ) {
-	// Should reset the simulated GPU
-	return g_last_cudaError = cudaSuccess;
-}
-cudaError_t CUDARTAPI cudaDeviceSynchronize(void){
-	// I don't know what this should do
-	return g_last_cudaError = cudaSuccess;
-}
-
-
-void CUDARTAPI __cudaRegisterFunction(
-		void   **fatCubinHandle,
-		const char    *hostFun,
-		char    *deviceFun,
-		const char    *deviceName,
-		int      thread_limit,
-		uint3   *tid,
-		uint3   *bid,
-		dim3    *bDim,
-		dim3    *gDim
-)
-{
-	CUctx_st *context = GPGPUSim_Context();
-	unsigned fat_cubin_handle = (unsigned)(unsigned long long)fatCubinHandle;
-	printf("GPGPU-Sim PTX: __cudaRegisterFunction %s : hostFun 0x%p, fat_cubin_handle = %u\n",
-			deviceFun, hostFun, fat_cubin_handle);
-	if(context->get_device()->get_gpgpu()->get_config().use_cuobjdump())
-		cuobjdumpParseBinary(fat_cubin_handle);
-	context->register_function( fat_cubin_handle, hostFun, deviceFun );
-}
-
-extern void __cudaRegisterVar(
-		void **fatCubinHandle,
-		char *hostVar, //pointer to...something
-		char *deviceAddress, //name of variable
-		const char *deviceName, //name of variable (same as above)
-		int ext,
-		int size,
-		int constant,
-		int global )
-{
-	printf("GPGPU-Sim PTX: __cudaRegisterVar: hostVar = %p; deviceAddress = %s; deviceName = %s\n", hostVar, deviceAddress, deviceName);
-	printf("GPGPU-Sim PTX: __cudaRegisterVar: Registering const memory space of %d bytes\n", size);
-	if(GPGPUSim_Context()->get_device()->get_gpgpu()->get_config().use_cuobjdump())
-		cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle);
-	fflush(stdout);
-	if ( constant && !global && !ext ) {
-		gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size);
-	} else if ( !constant && !global && !ext ) {
-		gpgpu_ptx_sim_register_global_variable(hostVar,deviceName,size);
-	} else cuda_not_implemented(__my_func__,__LINE__);
-}
-
-
-void __cudaRegisterShared(
-		void **fatCubinHandle,
-		void **devicePtr
-)
-{
-	// we don't do anything here
-	printf("GPGPU-Sim PTX: __cudaRegisterShared\n" );
-}
-
-void CUDARTAPI __cudaRegisterSharedVar(
-		void   **fatCubinHandle,
-		void   **devicePtr,
-		size_t   size,
-		size_t   alignment,
-		int      storage
-)
-{
-	// we don't do anything here
-	printf("GPGPU-Sim PTX: __cudaRegisterSharedVar\n" );
-}
-
-void __cudaRegisterTexture(
-		void **fatCubinHandle,
-		const struct textureReference *hostVar,
-		const void **deviceAddress,
-		const char *deviceName,
-		int dim,
-		int norm,
-		int ext
-) //passes in a newly created textureReference
-{
-	std::string devStr (deviceName);
-	#if (CUDART_VERSION > 4020)
-	if (devStr.size() > 2 && devStr.data()[0] == ':' && devStr.data()[1] == ':')
-		devStr = devStr.replace(0, 2, "");
-	#endif
-	CUctx_st *context = GPGPUSim_Context();
-	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	printf("GPGPU-Sim PTX: in __cudaRegisterTexture:\n");
-	gpu->gpgpu_ptx_sim_bindNameToTexture(devStr.data(), hostVar, dim, norm, ext);
-	printf("GPGPU-Sim PTX:   int dim = %d\n", dim);
-	printf("GPGPU-Sim PTX:   int norm = %d\n", norm);
-	printf("GPGPU-Sim PTX:   int ext = %d\n", ext);
-	printf("GPGPU-Sim PTX:   Execution warning: Not finished implementing \"%s\"\n", __my_func__ );
-}
-
-#ifndef OPENGL_SUPPORT
-typedef unsigned long GLuint;
-#endif
-
-cudaError_t cudaGLRegisterBufferObject(GLuint bufferObj)
-{
-	printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ );
-	return g_last_cudaError = cudaSuccess;
-}
-
-struct glbmap_entry {
-	GLuint m_bufferObj;
-	void *m_devPtr;
-	size_t m_size;
-	struct glbmap_entry *m_next;
-};
-typedef struct glbmap_entry glbmap_entry_t;
-
-glbmap_entry_t* g_glbmap = NULL;
-
-cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj) 
-{
-#ifdef OPENGL_SUPPORT
-	GLint buffer_size=0;
-	CUctx_st* ctx = GPGPUSim_Context();
-
-	glbmap_entry_t *p = g_glbmap;
-	while ( p && p->m_bufferObj != bufferObj )
-		p = p->m_next;
-	if ( p == NULL ) {
-		glBindBuffer(GL_ARRAY_BUFFER,bufferObj);
-		glGetBufferParameteriv(GL_ARRAY_BUFFER,GL_BUFFER_SIZE,&buffer_size);
-		assert( buffer_size != 0 );
-		*devPtr = ctx->get_device()->get_gpgpu()->gpu_malloc(buffer_size);
-
-		// create entry and insert to front of list
-		glbmap_entry_t *n = (glbmap_entry_t *) calloc(1,sizeof(glbmap_entry_t));
-		n->m_next = g_glbmap;
-		g_glbmap = n;
-
-		// initialize entry
-		n->m_bufferObj = bufferObj;
-		n->m_devPtr = *devPtr;
-		n->m_size = buffer_size;
-
-		p = n;
-	} else {
-		buffer_size = p->m_size;
-		*devPtr = p->m_devPtr;
-	}
-
-	if ( *devPtr  ) {
-		char *data = (char *) calloc(p->m_size,1);
-		glGetBufferSubData(GL_ARRAY_BUFFER,0,buffer_size,data);
-		memcpy_to_gpu( (size_t) *devPtr, data, buffer_size );
-		free(data);
-		printf("GPGPU-Sim PTX: cudaGLMapBufferObject %zu bytes starting at 0x%llx..\n", (size_t)buffer_size,
-				(unsigned long long) *devPtr);
-		return g_last_cudaError = cudaSuccess;
-	} else {
-		return g_last_cudaError = cudaErrorMemoryAllocation;
-	}
-
-	return g_last_cudaError = cudaSuccess;
-#else
-	fflush(stdout);
-	fflush(stderr);
-	printf("GPGPU-Sim PTX: GPGPU-Sim support for OpenGL integration disabled -- exiting\n");
-	fflush(stdout);
-	exit(50);
-#endif
-}
-
-cudaError_t cudaGLUnmapBufferObject(GLuint bufferObj)
-{
-#ifdef OPENGL_SUPPORT
-	glbmap_entry_t *p = g_glbmap;
-	while ( p && p->m_bufferObj != bufferObj )
-		p = p->m_next;
-	if ( p == NULL )
-		return g_last_cudaError = cudaErrorUnknown;
-
-	char *data = (char *) calloc(p->m_size,1);
-	memcpy_from_gpu( data,(size_t)p->m_devPtr,p->m_size );
-	glBufferSubData(GL_ARRAY_BUFFER,0,p->m_size,data);
-	free(data);
-
-	return g_last_cudaError = cudaSuccess;
-#else
-	fflush(stdout);
-	fflush(stderr);
-	printf("GPGPU-Sim PTX: support for OpenGL integration disabled -- exiting\n");
-	fflush(stdout);
-	exit(50);
-#endif
-}
-
-cudaError_t cudaGLUnregisterBufferObject(GLuint bufferObj) 
-{
-	printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ );
-	return g_last_cudaError = cudaSuccess;
-}
-
-#if (CUDART_VERSION >= 2010)
-
-cudaError_t CUDARTAPI cudaHostAlloc(void **pHost,  size_t bytes, unsigned int flags)
-{
-	*pHost = malloc(bytes);
-	//need to track the size allocated so that cudaHostGetDevicePointer() can function properly.
-	//TODO: vary this function behavior based on flags value (following nvidia documentation)
-	pinned_memory_size[*pHost]=bytes;
-	if( *pHost )
-		return g_last_cudaError = cudaSuccess;
-	else
-		return g_last_cudaError = cudaErrorMemoryAllocation;
-}
-
-cudaError_t CUDARTAPI cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags)
-{
-	//only cpu memory allocation happens in cudaHostAlloc. Linking with device pointer to pinned memory happens here.
-	//TODO: once kernel is executed, the contents in global pointer of GPU must be copied back to CPU host pointer!
-	flags=0;
-	CUctx_st* context = GPGPUSim_Context();
-	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	std::map<void *, size_t>::const_iterator i = pinned_memory_size.find(pHost);
-	assert(i != pinned_memory_size.end());
-	size_t size = i->second;
-	*pDevice = gpu->gpu_malloc(size);
-	if(g_debug_execution >= 3)
-		printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *pDevice);
-	if ( *pDevice  ) {
-		pinned_memory[pHost]=pDevice;
-		//Copy contents in cpu to gpu
-		gpu->memcpy_to_gpu((size_t)*pDevice,pHost,size);
-		return g_last_cudaError = cudaSuccess;
-	} else {
-		return g_last_cudaError = cudaErrorMemoryAllocation;
-	}
-}
-
-cudaError_t CUDARTAPI cudaSetValidDevices(int *device_arr, int len)
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-cudaError_t CUDARTAPI cudaSetDeviceFlags( int flags )
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *attr, const char *hostFun )
-{
-	CUctx_st *context = GPGPUSim_Context();
-	function_info *entry = context->get_kernel(hostFun);
-	if( entry ) {
-		const struct gpgpu_ptx_sim_info *kinfo = entry->get_kernel_info();
-		attr->sharedSizeBytes = kinfo->smem;
-		attr->constSizeBytes  = kinfo->cmem;
-		attr->localSizeBytes  = kinfo->lmem;
-		attr->numRegs         = kinfo->regs;
-		attr->maxThreadsPerBlock = 0; // from pragmas?
-#if CUDART_VERSION >= 3000
-		attr->ptxVersion      = kinfo->ptx_version;
-		attr->binaryVersion   = kinfo->sm_target;
-#endif
-	}
-	return g_last_cudaError = cudaSuccess;
-}
-
-cudaError_t CUDARTAPI cudaEventCreateWithFlags(cudaEvent_t *event, int flags)
-{
-	CUevent_st *e = new CUevent_st(flags==cudaEventBlockingSync);
-	g_timer_events[e->get_uid()] = e;
-#if CUDART_VERSION >= 3000
-	*event = e;
-#else
-	*event = e->get_uid();
-#endif
-	return g_last_cudaError = cudaSuccess;
-}
-
-cudaError_t CUDARTAPI cudaDriverGetVersion(int *driverVersion)
-{
-	*driverVersion = CUDART_VERSION;
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-cudaError_t CUDARTAPI cudaRuntimeGetVersion(int *runtimeVersion)
-{
-	*runtimeVersion = CUDART_VERSION;
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-#if CUDART_VERSION >= 3000
-__host__ cudaError_t CUDARTAPI cudaFuncSetCacheConfig(const char *func, enum cudaFuncCache  cacheConfig )
-{
-	CUctx_st *context = GPGPUSim_Context();
-	context->get_device()->get_gpgpu()->set_cache_config(context->get_kernel(func)->get_name(), (FuncCache)cacheConfig);
-	return g_last_cudaError = cudaSuccess;
-}
-
-//Jin: hack for cdp
-__host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t value) {
-    return g_last_cudaError = cudaSuccess;
-}
-#endif
-
-#endif
-
-cudaError_t CUDARTAPI cudaGLSetGLDevice(int device)
-{
-	printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s\"\n", __my_func__ );
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-typedef void* HGPUNV;
-
-cudaError_t CUDARTAPI cudaWGLGetDevice(int *device, HGPUNV hGpu)
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-void CUDARTAPI __cudaMutexOperation(int lock)
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-}
-
-void  CUDARTAPI __cudaTextureFetch(const void *tex, void *index, int integer, void *val) 
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-}
-
-}
-
-namespace cuda_math {
-
-void CUDARTAPI __cudaMutexOperation(int lock)
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-}
-
-void  CUDARTAPI __cudaTextureFetch(const void *tex, void *index, int integer, void *val) 
-{
-	cuda_not_implemented(__my_func__,__LINE__);
-}
-
-int CUDARTAPI __cudaSynchronizeThreads(void**, void*)
-{
-	//TODO This function should syncronize if we support Asyn kernel calls
-	return g_last_cudaError = cudaSuccess;
-}
-
-}
-
-////////
-
-extern int ptx_parse();
-extern int ptx__scan_string(const char*);
-extern FILE *ptx_in;
-
-extern int ptxinfo_parse();
-extern int ptxinfo_debug;
-extern FILE *ptxinfo_in;
-
-/// static functions
-
-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" );
-	fflush(stdout);
-	int ng_bytes=0;
-	symbol_table::iterator g=symtab->global_iterator_begin();
-
-	for ( ; g!=symtab->global_iterator_end(); g++) {
-		symbol *global = *g;
-		if ( global->has_initializer() ) {
-			printf( "GPGPU-Sim PTX:     initializing '%s' ... ", global->name().c_str() );
-			unsigned addr=global->get_address();
-			const type_info *type = global->type();
-			type_info_key ti=type->get_key();
-			size_t size;
-			int t;
-			ti.type_decode(size,t);
-			int nbytes = size/8;
-			int offset=0;
-			std::list<operand_info> init_list = global->get_initializer();
-			for ( std::list<operand_info>::iterator i=init_list.begin(); i!=init_list.end(); i++ ) {
-				operand_info op = *i;
-				ptx_reg_t value = op.get_literal_value();
-				assert( (addr+offset+nbytes) < min_gaddr ); // min_gaddr is start of "heap" for cudaMalloc
-				gpu->get_global_memory()->write(addr+offset,nbytes,&value,NULL,NULL); // assuming little endian here
-				offset+=nbytes;
-				ng_bytes+=nbytes;
-			}
-			printf(" wrote %u bytes\n", offset );
-		}
-	}
-	printf( "GPGPU-Sim PTX: finished loading globals (%u bytes total).\n", ng_bytes );
-	fflush(stdout);
-	return ng_bytes;
-}
-
-static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu ) 
-{
-	printf( "GPGPU-Sim PTX: loading constants with explicit initializers... " );
-	fflush(stdout);
-	int nc_bytes = 0;
-	symbol_table::iterator g=symtab->const_iterator_begin();
-
-	for ( ; g!=symtab->const_iterator_end(); g++) {
-		symbol *constant = *g;
-		if ( constant->is_const() && constant->has_initializer() ) {
-
-			// get the constant element data size
-			int basic_type;
-			size_t num_bits;
-			constant->type()->get_key().type_decode(num_bits,basic_type);
-
-			std::list<operand_info> init_list = constant->get_initializer();
-			int nbytes_written = 0;
-			for ( std::list<operand_info>::iterator i=init_list.begin(); i!=init_list.end(); i++ ) {
-				operand_info op = *i;
-				ptx_reg_t value = op.get_literal_value();
-				int nbytes = num_bits/8;
-				switch ( op.get_type() ) {
-				case int_t: assert(nbytes >= 1); break;
-				case float_op_t: assert(nbytes == 4); break;
-				case double_op_t: assert(nbytes >= 4); break; // account for double DEMOTING
-				default:
-					abort();
-				}
-				unsigned addr=constant->get_address() + nbytes_written;
-				assert( addr+nbytes < min_gaddr );
-
-				gpu->get_global_memory()->write(addr,nbytes,&value,NULL,NULL); // assume little endian (so u8 is the first byte in u32)
-				nc_bytes+=nbytes;
-				nbytes_written += nbytes;
-			}
-		}
-	}
-	printf( " done.\n");
-	fflush(stdout);
-	return nc_bytes;
-}
-
-kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, 
-		gpgpu_ptx_sim_arg_list_t args,
-		struct dim3 gridDim,
-		struct dim3 blockDim,
-		CUctx_st* context )
-{
-	function_info *entry = context->get_kernel(hostFun);
-	kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry);
-	if( entry == NULL ) {
-		printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found for %p\n", hostFun);
-		abort();
-	}
-	unsigned argcount=args.size();
-	unsigned argn=1;
-	for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) {
-		entry->add_param_data(argcount-argn,&(*a));
-		argn++;
-	}
-
-	entry->finalize(result->get_param_memory());
-	g_ptx_kernel_count++;
-	fflush(stdout);
-
-	return result;
-}