Merge branch 'dev' of https://github.com/purdue-aalp/gpgpu-sim_distribution into dev-private

28 files changed, 1199 insertions, 908 deletions

diff --git a/Jenkinsfile b/Jenkinsfile
index a345bc0..8a3ad19 100644
--- a/Jenkinsfile
+++ b/Jenkinsfile
@@ -13,11 +13,11 @@ pipeline {
                 parallel "4.2": {
                     sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/4.2_env_setup.sh &&\
                     source `pwd`/setup_environment &&\
-                    make -j'
+                    make -j 10'
                 }, "10.1" : {
                     sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/10.1_env_setup.sh &&\
                     source `pwd`/setup_environment &&\
-                    make -j'
+                    make -j 10'
                 }
             }
         }
@@ -32,13 +32,13 @@ pipeline {
                     source `pwd`/setup_environment &&\
                     cd gpgpu-sim_simulations && \
                     source ./benchmarks/src/setup_environment && \
-                    make -j -C ./benchmarks/src rodinia_2.0-ft sdk-4.2 && \
+                    make -j 10 -C ./benchmarks/src rodinia_2.0-ft sdk-4.2 && \
                     make -C ./benchmarks/src data'
                 sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/10.1_env_setup.sh &&\
                     source `pwd`/setup_environment &&\
                     cd gpgpu-sim_simulations && \
                     source ./benchmarks/src/setup_environment && \
-                    make -j -C ./benchmarks/src/ rodinia_2.0-ft sdk-4.2 && \
+                    make -j 10 -C ./benchmarks/src/ rodinia_2.0-ft sdk-4.2 && \
                     make -C ./benchmarks/src data'
             }
         }
@@ -77,8 +77,8 @@ pipeline {
                         ./gpgpu-sim_simulations/util/plotting/plot-get-stats.py -c per-app-merge-10.1.csv -P cuda-10.1 &&\
                         ./gpgpu-sim_simulations/util/plotting/merge-stats.py -c ./gpgpu-sim-results-repo/${JOB_NAME}/stats-per-kernel-4.2-ptxplus.csv,./stats-per-kernel-4.2-ptxplus.csv -R > per-kernel-merge-4.2-ptxplus.csv &&\
                         ./gpgpu-sim_simulations/util/plotting/merge-stats.py -c ./gpgpu-sim-results-repo/${JOB_NAME}/stats-per-kernel-10.1.csv,./stats-per-kernel-10.1.csv -R > per-kernel-merge-10.1.csv &&\
-                        ./gpgpu-sim_simulations/util/plotting/plot-correlation.py -c per-kernel-merge-4.2-ptxplus.csv -p cuda-4.2 | grep "Correl=" | tee correl.4.2.txt &&\
-                        ./gpgpu-sim_simulations/util/plotting/plot-correlation.py -c per-kernel-merge-10.1.csv -p cuda-10.1 | grep "Correl=" | tee correl.10.1.txt &&\
+                        ./gpgpu-sim_simulations/util/plotting/plot-correlation.py -c per-kernel-merge-4.2-ptxplus.csv -p cuda-4.2 | grep -B 1 "Correl=" | tee correl.4.2.txt &&\
+                        ./gpgpu-sim_simulations/util/plotting/plot-correlation.py -c per-kernel-merge-10.1.csv -p cuda-10.1 | grep -B 1 "Correl=" | tee correl.10.1.txt &&\
                         mkdir -p ./gpgpu-sim-results-repo/${JOB_NAME}/ && cp stats-per-*.csv ./gpgpu-sim-results-repo/${JOB_NAME}/ &&\
                         cd ./gpgpu-sim-results-repo &&\
                         git diff --quiet && git diff --staged --quiet || git commit -am "Jenkins automated checkin ${JOB_NAME} Build:${BUILD_NUMBER}" &&\
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index e71db4c..5cefe60 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -150,9 +150,6 @@
 #endif
 
 
-extern void synchronize();
-extern void exit_simulation();
-
 /*DEVICE_BUILTIN*/
 struct cudaArray
 {
@@ -176,8 +173,6 @@ struct cudaArray
 
 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
@@ -199,8 +194,7 @@ void register_ptx_function( const char *name, function_info *impl )
 
 struct _cuda_device_id *gpgpu_context::GPGPUSim_Init()
 {
-	//static _cuda_device_id *the_device = NULL;
-	_cuda_device_id *the_device = GPGPUsim_ctx_ptr()->the_cude_device;
+	_cuda_device_id *the_device = the_gpgpusim->the_cude_device;
 	if( !the_device ) {
 		gpgpu_sim *the_gpu = gpgpu_ptx_sim_init_perf();
 
@@ -245,22 +239,21 @@ struct _cuda_device_id *gpgpu_context::GPGPUSim_Init()
 		prop->maxThreadsPerMultiProcessor = the_gpu->threads_per_core();
 #endif
 		the_gpu->set_prop(prop);
-		GPGPUsim_ctx_ptr()->the_cude_device = new _cuda_device_id(the_gpu);
-		the_device = GPGPUsim_ctx_ptr()->the_cude_device;
+		the_gpgpusim->the_cude_device = new _cuda_device_id(the_gpu);
+		the_device = the_gpgpusim->the_cude_device;
 	}
 	start_sim_thread(1);
 	return the_device;
 }
 
-static CUctx_st* GPGPUSim_Context()
+CUctx_st* GPGPUSim_Context(gpgpu_context * ctx)
 {
 	//static CUctx_st *the_context = NULL;
-	gpgpu_context *cur_ctx = GPGPU_Context();
-	CUctx_st *the_context = GPGPUsim_ctx_ptr()->the_context;
+	CUctx_st *the_context = ctx->the_gpgpusim->the_context;
 	if( the_context == NULL ) {
-		_cuda_device_id *the_gpu = cur_ctx->GPGPUSim_Init();
-		GPGPUsim_ctx_ptr()->the_context = new CUctx_st(the_gpu);
-		the_context = GPGPUsim_ctx_ptr()->the_context;
+		_cuda_device_id *the_gpu = ctx->GPGPUSim_Init();
+		ctx->the_gpgpusim->the_context = new CUctx_st(the_gpu);
+		the_context = ctx->the_gpgpusim->the_context;
 	}
 	return the_context;
 }
@@ -276,9 +269,9 @@ gpgpu_context* GPGPU_Context()
 
  void ptxinfo_data::ptxinfo_addinfo()
 {
+	CUctx_st *context = GPGPUSim_Context(gpgpu_ctx);
 	 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();
@@ -289,7 +282,6 @@ gpgpu_context* GPGPU_Context()
 		clear_ptxinfo();
 		return;
 	}
-	CUctx_st *context = GPGPUSim_Context();
 	print_ptxinfo();
 	context->add_ptxinfo( get_ptxinfo_kname(), get_ptxinfo() );
 	clear_ptxinfo();
@@ -548,7 +540,7 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimitInternal( size_t* pValue, cudaL
 		break;
 	    }
 	    else{
-		printf("ERROR:Limit %s is not supported on this architecture \n",limit);
+		printf("ERROR:Limit %d is not supported on this architecture \n", limit);
 		abort();
 	    }
 	case 4: // cudaLimitDevRuntimePendingLaunchCount
@@ -557,12 +549,12 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimitInternal( size_t* pValue, cudaL
 		break;
 	    }
 	    else{
-		printf("ERROR:Limit %s is not supported on this architecture \n",limit);
+		printf("ERROR:Limit %d is not supported on this architecture \n",limit);
 		abort();
 	    }
 #endif
 	default:
-	    printf("ERROR:Limit %s unimplemented \n",limit);
+	    printf("ERROR:Limit %d unimplemented \n",limit);
 	    abort();
     }
     return g_last_cudaError = cudaSuccess;
@@ -585,7 +577,7 @@ void** cudaRegisterFatBinaryInternal( void *fatCubin, gpgpu_context* gpgpu_ctx =
 	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();
+	CUctx_st *context = GPGPUSim_Context(ctx);
 	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. 
@@ -732,7 +724,7 @@ void cudaRegisterFunctionInternal(
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	CUctx_st *context = GPGPUSim_Context();
+	CUctx_st *context = GPGPUSim_Context(ctx);
 	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);
@@ -763,7 +755,7 @@ void cudaRegisterVarInternal(
     }
 	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())
+	if(GPGPUSim_Context(ctx)->get_device()->get_gpgpu()->get_config().use_cuobjdump())
 		ctx->cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle);
 	fflush(stdout);
 	if ( constant && !global && !ext ) {
@@ -872,7 +864,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx =
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	CUctx_st* context = GPGPUSim_Context();
+	CUctx_st* context = GPGPUSim_Context(ctx);
 	char *mode = getenv("PTX_SIM_MODE_FUNC");
 	if( mode )
 		sscanf(mode,"%u", &(ctx->func_sim->g_ptx_sim_mode));
@@ -944,7 +936,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx =
 	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,ctx->func_sim->g_ptx_sim_mode,stream);
-	g_stream_manager()->push(op);
+	ctx->the_gpgpusim->g_stream_manager->push(op);
 	ctx->api->g_cuda_launch_stack.pop_back();
 	return g_last_cudaError = cudaSuccess;
 }
@@ -960,7 +952,7 @@ cudaError_t cudaMallocInternal(void **devPtr, size_t size, gpgpu_context* gpgpu_
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	CUctx_st* context = GPGPUSim_Context();
+	CUctx_st* context = GPGPUSim_Context(ctx);
 	*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);
@@ -994,6 +986,29 @@ cudaError_t cudaMallocHostInternal(void **ptr, size_t size, gpgpu_context* gpgpu
 	}
 }
 
+__host__ cudaError_t CUDARTAPI cudaMallocPitchInternal(void **devPtr, size_t *pitch, size_t width, size_t height, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	unsigned malloc_width_inbytes = width;
+	printf("GPGPU-Sim PTX: cudaMallocPitch (width = %d)\n", malloc_width_inbytes);
+	CUctx_st* context = GPGPUSim_Context(ctx);
+	*devPtr = context->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;
+	}
+}
+
 cudaError_t cudaHostGetDevicePointerInternal(void **pDevice, void *pHost, unsigned int flags, gpgpu_context* gpgpu_ctx = NULL)
 {
     gpgpu_context *ctx;
@@ -1011,7 +1026,7 @@ cudaError_t cudaHostGetDevicePointerInternal(void **pDevice, void *pHost, unsign
 	//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();
+	CUctx_st* context = GPGPUSim_Context(ctx);
 	gpgpu_t *gpu = context->get_device()->get_gpgpu();
 	std::map<void *, size_t>::const_iterator i = ctx->api->pinned_memory_size.find(pHost);
 	assert(i != ctx->api->pinned_memory_size.end());
@@ -1031,7 +1046,7 @@ cudaError_t cudaHostGetDevicePointerInternal(void **pDevice, void *pHost, unsign
 	}
 }
 
-cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, gpgpu_context* gpgpu_ctx = NULL)
+__host__ cudaError_t CUDARTAPI cudaMallocArrayInternal(struct cudaArray **array, const struct cudaChannelFormatDesc *desc, size_t width, size_t height __dv(1), gpgpu_context* gpgpu_ctx = NULL)
 {
     gpgpu_context *ctx;
     if (gpgpu_ctx){
@@ -1042,226 +1057,8 @@ cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, gpgpu
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-#ifdef OPENGL_SUPPORT
-	GLint buffer_size=0;
-	CUctx_st* context = GPGPUSim_Context();
-
-	glbmap_entry_t *p = ctx->api->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 = context->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 = ctx->api->g_glbmap;
-		ctx->api->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
-}
-
-#if CUDART_VERSION >= 6050
-CUresult
-cuLinkAddFileInternal(CUlinkState state, CUjitInputType type, const char *path,
-    unsigned int numOptions, CUjit_option *options, void **optionValues, gpgpu_context* gpgpu_ctx = NULL)
-{
-    gpgpu_context *ctx;
-    if (gpgpu_ctx){
-	ctx = gpgpu_ctx;
-    } else {
-	ctx = GPGPU_Context();
-    }
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-    static bool addedFile = false;
-    if (addedFile){
-        printf("GPGPU-Sim PTX: ERROR: cuLinkAddFile does not support multiple files\n");
-        abort();
-    }
-
-    //blocking
-    assert(type==CU_JIT_INPUT_PTX);
-	CUctx_st *context = GPGPUSim_Context();
-    char *file = getenv("PTX_JIT_PATH");
-    if(file==NULL){
-        printf("GPGPU-Sim PTX: ERROR: PTX_JIT_PATH has not been set\n");
-        abort();
-    }
-    strcat(file,"/");
-    strcat(file,path);
-	symbol_table *symtab = ctx->gpgpu_ptx_sim_load_ptx_from_filename( file );
-    std::string fname(path);
-    ctx->api->name_symtab[fname] = symtab;
-    context->add_binary(symtab, 1);
-    ctx->api->load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
-    ctx->api->load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
-    addedFile = true;
-	return CUDA_SUCCESS;
-}
-#endif
-
-#if (CUDART_VERSION >= 2010)
-
-cudaError_t cudaHostAllocInternal(void **pHost,  size_t bytes, unsigned int flags, gpgpu_context* gpgpu_ctx = NULL)
-{
-    gpgpu_context *ctx;
-    if (gpgpu_ctx){
-	ctx = gpgpu_ctx;
-    } else {
-	ctx = GPGPU_Context();
-    }
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	*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)
-	ctx->api->pinned_memory_size[*pHost]=bytes;
-	if( *pHost )
-		return g_last_cudaError = cudaSuccess;
-	else
-		return g_last_cudaError = cudaErrorMemoryAllocation;
-}
-
-#endif
-
-size_t getMaxThreadsPerBlock(struct cudaFuncAttributes *attr, gpgpu_context *ctx) {
-  _cuda_device_id *dev = ctx->GPGPUSim_Init();
-  struct cudaDeviceProp prop;
-
-  prop = *dev->get_prop();
-
-  size_t max = prop.maxThreadsPerBlock;
-
-  if (attr->numRegs && (prop.regsPerBlock / attr->numRegs) < max)
-    max = prop.regsPerBlock / attr->numRegs;
-
-  if (attr->sharedSizeBytes && (prop.sharedMemPerBlock / attr->sharedSizeBytes) < max)
-      max = prop.sharedMemPerBlock / attr->sharedSizeBytes;
-
-  return max;
-}
-
-cudaError_t CUDARTAPI cudaFuncGetAttributesInternal(struct cudaFuncAttributes *attr, const char *hostFun, gpgpu_context* gpgpu_ctx = NULL )
-{
-    gpgpu_context *ctx;
-    if (gpgpu_ctx){
-	ctx = gpgpu_ctx;
-    } else {
-	ctx = GPGPU_Context();
-    }
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	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;
-		if(kinfo->maxthreads > 0)
-		  attr->maxThreadsPerBlock = kinfo->maxthreads;
-		else
-		  attr->maxThreadsPerBlock = getMaxThreadsPerBlock(attr, ctx);
-#if CUDART_VERSION >= 3000
-		attr->ptxVersion      = kinfo->ptx_version;
-		attr->binaryVersion   = kinfo->sm_target;
-#endif
-	}
-	return g_last_cudaError = cudaSuccess;
-}
-
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-
-extern "C" {
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-cudaError_t cudaPeekAtLastError(void)
-{
-	return g_last_cudaError;
-}
-
-__host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) 
-{
-    return cudaMallocInternal(devPtr, size);
-}
-
-__host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size)
-{
-    return cudaMallocHostInternal(ptr, size);
-}
-__host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr, size_t *pitch, size_t width, size_t height)
-{
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	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))
-{
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
 	unsigned size = width * height * ((desc->x + desc->y + desc->z + desc->w)/8);
-	CUctx_st* context = GPGPUSim_Context();
+	CUctx_st* context = GPGPUSim_Context(ctx);
 	(*array) = (struct cudaArray*) malloc(sizeof(struct cudaArray));
 	(*array)->desc = *desc;
 	(*array)->width = width;
@@ -1278,41 +1075,14 @@ __host__ cudaError_t CUDARTAPI cudaMallocArray(struct cudaArray **array, const s
 	}
 }
 
-__host__ cudaError_t CUDARTAPI cudaFree(void *devPtr)
-{
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	// TODO...  manage g_global_mem space?
-	return g_last_cudaError = cudaSuccess;
-}
-__host__ cudaError_t CUDARTAPI cudaFreeHost(void *ptr)
+__host__ cudaError_t CUDARTAPI cudaMemcpyInternal(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, gpgpu_context* gpgpu_ctx = NULL)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
     }
-	free (ptr);  // this will crash the system if called twice
-	return g_last_cudaError = cudaSuccess;
-}
-
-__host__ cudaError_t CUDARTAPI cudaFreeArray(struct cudaArray *array)
-{
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	// 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)
-{
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
@@ -1321,21 +1091,21 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t cou
 	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) );
+		ctx->the_gpgpusim->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) );
+		ctx->the_gpgpusim->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) );
+		ctx->the_gpgpusim->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
+				ctx->the_gpgpusim->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
+				ctx->the_gpgpusim->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) );
+				ctx->the_gpgpusim->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();
@@ -1349,12 +1119,18 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t cou
 	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)
+__host__ cudaError_t CUDARTAPI cudaMemcpyToArrayInternal(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind, gpgpu_context* gpgpu_ctx = NULL)
 { 
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	CUctx_st *context = GPGPUSim_Context();
+	CUctx_st *context = GPGPUSim_Context(ctx);
 	gpgpu_t *gpu = context->get_device()->get_gpgpu();
 	size_t size = count;
 	printf("GPGPU-Sim PTX: cudaMemcpyToArray\n");
@@ -1372,33 +1148,18 @@ __host__ cudaError_t CUDARTAPI cudaMemcpyToArray(struct cudaArray *dst, size_t w
 	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)
-{
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	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))
+__host__ cudaError_t CUDARTAPI cudaMemcpy2DInternal(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, gpgpu_context* gpgpu_ctx = NULL)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
     }
-	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)
-{
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	CUctx_st *context = GPGPUSim_Context();
+	CUctx_st *context = GPGPUSim_Context(ctx);
 	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" );
@@ -1415,13 +1176,18 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2D(void *dst, size_t dpitch, const void
 	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)
+__host__ cudaError_t CUDARTAPI cudaMemcpy2DToArrayInternal(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, gpgpu_context* gpgpu_ctx = NULL)
 {
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	CUctx_st *context = GPGPUSim_Context();
+	CUctx_st *context = GPGPUSim_Context(ctx);
 	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;
@@ -1447,29 +1213,14 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2DToArray(struct cudaArray *dst, size_t
 	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)
-{
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	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))
+__host__ cudaError_t CUDARTAPI cudaMemcpyToSymbolInternal(const char *symbol, const void *src, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyHostToDevice), gpgpu_context* gpgpu_ctx = NULL)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
     }
-	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))
-{
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
@@ -1477,52 +1228,47 @@ __host__ cudaError_t CUDARTAPI cudaMemcpyToSymbol(const char *symbol, const void
 	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) );
+	ctx->the_gpgpusim->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))
+__host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbolInternal(void *dst, const char *symbol, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost), gpgpu_context* gpgpu_ctx = NULL)
 {
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
 	//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) );
+	ctx->the_gpgpusim->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 cudaMemGetInfo (size_t *free, size_t *total){
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	//placeholder; should interact with cudaMalloc and cudaFree?
-	*free = 10000000000;
-	*total = 10000000000;
-
-	return g_last_cudaError = cudaSuccess;
-}
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-
-__host__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream)
+__host__ cudaError_t CUDARTAPI cudaMemcpyAsyncInternal(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream, gpgpu_context* gpgpu_ctx = NULL)
 {
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
 	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;
+	case cudaMemcpyHostToDevice: ctx->the_gpgpusim->g_stream_manager->push( stream_operation(src,(size_t)dst,count,s) ); break;
+	case cudaMemcpyDeviceToHost: ctx->the_gpgpusim->g_stream_manager->push( stream_operation((size_t)src,dst,count,s) ); break;
+	case cudaMemcpyDeviceToDevice: ctx->the_gpgpusim->g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,s) ); break;
 	default:
 		abort();
 	}
@@ -1530,98 +1276,249 @@ __host__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_
 }
 
 
-__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)
+#if (CUDART_VERSION >= 8000)
+cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlagsInternal(int* numBlocks, const char *hostFunc, int blockSize, size_t dynamicSMemSize, unsigned int flags, gpgpu_context* gpgpu_ctx = NULL)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
     }
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
+	printf("GPGPU-Sim PTX: cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags %p\n", hostFunc);
+	CUctx_st *context = GPGPUSim_Context(ctx);
+	function_info *entry = context->get_kernel(hostFunc);
+	printf("Calculate Maxium Active Block with function ptr=%p, blockSize=%d, SMemSize=%d\n", hostFunc, blockSize, dynamicSMemSize);
+	if (flags == cudaOccupancyDefault) {
+		//create kernel_info based on entry
+		dim3 gridDim(context->get_device()->get_gpgpu()->max_cta_per_core()
+                        * context->get_device()->get_gpgpu()->get_config().num_shader());
+		dim3 blockDim(blockSize);
+		kernel_info_t result(gridDim, blockDim, entry);
+		//if(entry == NULL){
+		//	*numBlocks = 1;
+		//	return g_last_cudaError = cudaErrorUnknown;
+		//}
+		*numBlocks = context->get_device()->get_gpgpu()->get_max_cta(result);
+		printf("Maximum size is %d with gridDim %d and blockDim %d\n", *numBlocks, gridDim.x, blockDim.x);
+		return g_last_cudaError = cudaSuccess;
+	} else {
+		cuda_not_implemented(__my_func__,__LINE__);
+		return g_last_cudaError = cudaErrorUnknown;
+	}
 }
 
+#endif
 
-__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)
+__host__ cudaError_t CUDARTAPI cudaMemsetInternal(void *mem, int c, size_t count, gpgpu_context* gpgpu_ctx = NULL)
 {
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
+	CUctx_st *context = GPGPUSim_Context(ctx);
+	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 cudaMemcpy2DAsync(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream)
+//memset operation is done but i think its not async?
+__host__ cudaError_t CUDARTAPI cudaMemsetAsyncInternal(void *mem, int c, size_t count, 	cudaStream_t stream=0, gpgpu_context* gpgpu_ctx = NULL)
 {
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
+	printf("GPGPU-Sim PTX: WARNING: Asynchronous memset not supported (%s)\n", __my_func__);
+	CUctx_st *context = GPGPUSim_Context(ctx);
+	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 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)
+cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, gpgpu_context* gpgpu_ctx = NULL)
 {
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
+#ifdef OPENGL_SUPPORT
+	gpgpu_context *ctx;
+	if (gpgpu_ctx){
+	    ctx = gpgpu_ctx;
+	} else {
+	    ctx = GPGPU_Context();
+	}
+	GLint buffer_size=0;
+	CUctx_st* context = GPGPUSim_Context(ctx);
+
+	glbmap_entry_t *p = ctx->api->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 = context->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 = ctx->api->g_glbmap;
+		ctx->api->g_glbmap = n;
 
-__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)
+		// 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
+}
+
+#if CUDART_VERSION >= 6050
+CUresult
+cuLinkAddFileInternal(CUlinkState state, CUjitInputType type, const char *path,
+    unsigned int numOptions, CUjit_option *options, void **optionValues, gpgpu_context* gpgpu_ctx = NULL)
 {
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
+    static bool addedFile = false;
+    if (addedFile){
+        printf("GPGPU-Sim PTX: ERROR: cuLinkAddFile does not support multiple files\n");
+        abort();
+    }
+
+    //blocking
+    assert(type==CU_JIT_INPUT_PTX);
+	CUctx_st *context = GPGPUSim_Context(ctx);
+    char *file = getenv("PTX_JIT_PATH");
+    if(file==NULL){
+        printf("GPGPU-Sim PTX: ERROR: PTX_JIT_PATH has not been set\n");
+        abort();
+    }
+    strcat(file,"/");
+    strcat(file,path);
+	symbol_table *symtab = ctx->gpgpu_ptx_sim_load_ptx_from_filename( file );
+    std::string fname(path);
+    ctx->api->name_symtab[fname] = symtab;
+    context->add_binary(symtab, 1);
+    ctx->api->load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
+    ctx->api->load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
+    addedFile = true;
+	return CUDA_SUCCESS;
 }
+#endif
 
-#if (CUDART_VERSION >= 8000)
-cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks, const char *hostFunc, int blockSize, size_t dynamicSMemSize, unsigned int flags)
+#if (CUDART_VERSION >= 2010)
+
+cudaError_t cudaHostAllocInternal(void **pHost,  size_t bytes, unsigned int flags, gpgpu_context* gpgpu_ctx = NULL)
 {
-	printf("GPGPU-Sim PTX: cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags %p\n", hostFunc);
-	CUctx_st *context = GPGPUSim_Context();
-	function_info *entry = context->get_kernel(hostFunc);
-	printf("Calculate Maxium Active Block with function ptr=%p, blockSize=%d, SMemSize=%d\n", hostFunc, blockSize, dynamicSMemSize);
-	if (flags == cudaOccupancyDefault) {
-		//create kernel_info based on entry
-		dim3 gridDim(context->get_device()->get_gpgpu()->max_cta_per_core()
-                        * context->get_device()->get_gpgpu()->get_config().num_shader());
-		dim3 blockDim(blockSize);
-		kernel_info_t result(gridDim, blockDim, entry);
-		//if(entry == NULL){
-		//	*numBlocks = 1;
-		//	return g_last_cudaError = cudaErrorUnknown;
-		//}
-		*numBlocks = context->get_device()->get_gpgpu()->get_max_cta(result);
-		printf("Maximum size is %d with gridDim %d and blockDim %d\n", *numBlocks, gridDim.x, blockDim.x);
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	*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)
+	ctx->api->pinned_memory_size[*pHost]=bytes;
+	if( *pHost )
 		return g_last_cudaError = cudaSuccess;
-	} else {
-		cuda_not_implemented(__my_func__,__LINE__);
-		return g_last_cudaError = cudaErrorUnknown;
-	}
+	else
+		return g_last_cudaError = cudaErrorMemoryAllocation;
 }
 
 #endif
 
+size_t getMaxThreadsPerBlock(struct cudaFuncAttributes *attr, gpgpu_context *ctx) {
+  _cuda_device_id *dev = ctx->GPGPUSim_Init();
+  struct cudaDeviceProp prop;
 
+  prop = *dev->get_prop();
 
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-__host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count)
+  size_t max = prop.maxThreadsPerBlock;
+
+  if (attr->numRegs && (prop.regsPerBlock / attr->numRegs) < max)
+    max = prop.regsPerBlock / attr->numRegs;
+
+  if (attr->sharedSizeBytes && (prop.sharedMemPerBlock / attr->sharedSizeBytes) < max)
+      max = prop.sharedMemPerBlock / attr->sharedSizeBytes;
+
+  return max;
+}
+
+cudaError_t CUDARTAPI cudaFuncGetAttributesInternal(struct cudaFuncAttributes *attr, const char *hostFun, gpgpu_context* gpgpu_ctx = NULL )
 {
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	CUctx_st *context = GPGPUSim_Context();
-	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	gpu->gpu_memset((size_t)mem, c, count);
+	CUctx_st *context = GPGPUSim_Context(ctx);
+	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;
+		if(kinfo->maxthreads > 0)
+		  attr->maxThreadsPerBlock = kinfo->maxthreads;
+		else
+		  attr->maxThreadsPerBlock = getMaxThreadsPerBlock(attr, ctx);
+#if CUDART_VERSION >= 3000
+		attr->ptxVersion      = kinfo->ptx_version;
+		attr->binaryVersion   = kinfo->sm_target;
+#endif
+	}
 	return g_last_cudaError = cudaSuccess;
 }
 
@@ -1826,19 +1723,652 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttributeInternal(int *value, enum c
 }
 #endif
 
-//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)
+__host__ cudaError_t CUDARTAPI cudaBindTextureInternal(size_t *offset,
+		const struct textureReference *texref,
+		const void *devPtr,
+		const struct cudaChannelFormatDesc *desc,
+		size_t size __dv(UINT_MAX),
+		gpgpu_context* gpgpu_ctx = NULL)
 {
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	printf("GPGPU-Sim PTX: WARNING: Asynchronous memset not supported (%s)\n", __my_func__);
-	CUctx_st *context = GPGPUSim_Context();
+	CUctx_st *context = GPGPUSim_Context(ctx);
 	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	gpu->gpu_memset((size_t)mem, c, count);
+	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 cudaBindTextureToArrayInternal(const struct textureReference *texref, const struct cudaArray *array, const struct cudaChannelFormatDesc *desc, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	CUctx_st *context = GPGPUSim_Context(ctx);
+	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 cudaUnbindTextureInternal(const struct textureReference *texref, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+   CUctx_st *context = GPGPUSim_Context(ctx);
+   gpgpu_t *gpu = context->get_device()->get_gpgpu();
+   printf("GPGPU-Sim PTX: in cudaUnbindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference));
+   printf("GPGPU-Sim PTX:   Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref));
+
+   gpu->gpgpu_ptx_sim_unbindTexture(texref);
+   return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaLaunchKernelInternal( const char* hostFun, dim3 gridDim, dim3 blockDim, const void** args, size_t sharedMem, cudaStream_t stream, gpgpu_context* gpgpu_ctx = NULL )
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+
+    if(g_debug_execution >= 3){
+        announce_call(__my_func__);
+    }
+    CUctx_st *context = GPGPUSim_Context(ctx);
+    function_info *entry = context->get_kernel(hostFun);
+#if CUDART_VERSION < 10000
+    cudaConfigureCallInternal(gridDim, blockDim, sharedMem, stream, ctx);
+#endif
+    for(unsigned i = 0; i < entry->num_args(); i++){
+        std::pair<size_t, unsigned> p = entry->get_param_config(i);
+        cudaSetupArgumentInternal(args[i], p.first, p.second);
+    }
+
+    cudaLaunchInternal(hostFun);
+    return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaStreamCreateInternal(cudaStream_t *stream, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	printf("GPGPU-Sim PTX: cudaStreamCreate\n");
+#if (CUDART_VERSION >= 3000)
+	*stream = new struct CUstream_st();
+	ctx->the_gpgpusim->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;
+}
+
+__host__ cudaError_t CUDARTAPI cudaStreamDestroyInternal(cudaStream_t stream, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+#if (CUDART_VERSION >= 3000)
+	//per-stream synchronization required for application using external libraries without explicit synchronization in the code to 
+	//avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. 
+	stream->synchronize();
+	ctx->the_gpgpusim->g_stream_manager->destroy_stream(stream);
+#endif
+	return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaStreamSynchronizeInternal(cudaStream_t stream, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+#if (CUDART_VERSION >= 3000)
+	if( stream == NULL )
+		ctx->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;
+}
+
+void __cudaRegisterTextureInternal(
+		void **fatCubinHandle,
+		const struct textureReference *hostVar,
+		const void **deviceAddress,
+		const char *deviceName,
+		int dim,
+		int norm,
+		int ext,
+		gpgpu_context* gpgpu_ctx = NULL
+) //passes in a newly created textureReference
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	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(ctx);
+	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__ );
+}
+
+cudaError_t cudaGLUnmapBufferObjectInternal(GLuint bufferObj, gpgpu_context* gpgpu_ctx = NULL)
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+#ifdef OPENGL_SUPPORT
+	gpgpu_context *ctx;
+	if (gpgpu_ctx){
+	    ctx = gpgpu_ctx;
+	} else {
+	    ctx = GPGPU_Context();
+	}
+	CUctx_st* ctx = GPGPUSim_Context(ctx);
+	glbmap_entry_t *p = ctx->api->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
+}
+
+#if CUDART_VERSION >= 3000
+
+__host__ cudaError_t CUDARTAPI cudaFuncSetCacheConfigInternal(const char *func, enum cudaFuncCache  cacheConfig, gpgpu_context* gpgpu_ctx = NULL )
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	CUctx_st *context = GPGPUSim_Context(ctx);
+	context->get_device()->get_gpgpu()->set_cache_config(context->get_kernel(func)->get_name(), (FuncCache)cacheConfig);
+	return g_last_cudaError = cudaSuccess;
+}
+
+#endif
+
+#if CUDART_VERSION >= 4000
+CUresult CUDAAPI cuLaunchKernelInternal(CUfunction f,
+                                unsigned int gridDimX,
+                                unsigned int gridDimY,
+                                unsigned int gridDimZ,
+                                unsigned int blockDimX,
+                                unsigned int blockDimY,
+                                unsigned int blockDimZ,
+                                unsigned int sharedMemBytes,
+                                CUstream hStream,
+                                void **kernelParams,
+                                void **extra,
+				gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+    if (extra!=NULL){
+	    printf("GPGPU-Sim CUDA DRIVER API: ERROR: Currently do not support void** extra.\n");
+        abort();
+    }
+    const char *hostFun = (const char*) f;
+	CUctx_st *context = GPGPUSim_Context(ctx);
+	function_info *entry = context->get_kernel(hostFun);
+    cudaConfigureCallInternal(dim3(gridDimX, gridDimY, gridDimZ), dim3(blockDimX, blockDimY, blockDimZ), sharedMemBytes, (cudaStream_t) hStream, ctx);
+    for(unsigned i = 0; i < entry->num_args(); i++){
+        std::pair<size_t, unsigned> p = entry->get_param_config(i);
+        cudaSetupArgumentInternal(kernelParams[i], p.first, p.second, ctx);
+    }
+    cudaLaunchInternal(hostFun, ctx);
+	return CUDA_SUCCESS;
+}
+#endif /* CUDART_VERSION >= 4000 */
+
+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 cudaEventRecordInternal(cudaEvent_t event, cudaStream_t stream, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	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);
+	ctx->the_gpgpusim->g_stream_manager->push(op);
+	return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaStreamWaitEventInternal(cudaStream_t stream, cudaEvent_t event, unsigned int flags, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+   //reference: https://www.cs.cmu.edu/afs/cs/academic/class/15668-s11/www/cuda-doc/html/group__CUDART__STREAM_gfe68d207dc965685d92d3f03d77b0876.html 
+	CUevent_st *e = get_event(event);
+	if( !e ){
+	   printf("GPGPU-Sim API: Warning: cudaEventRecord has not been called on event before calling cudaStreamWaitEvent.\nNothing to be done.\n");
+      return g_last_cudaError = cudaSuccess;
+   }
+   if (!stream){
+     ctx->the_gpgpusim->g_stream_manager->pushCudaStreamWaitEventToAllStreams(e, flags);
+   } else {
+	   struct CUstream_st *s = (struct CUstream_st *)stream;
+	   stream_operation op(s,e,flags);
+	   ctx->the_gpgpusim->g_stream_manager->push(op);
+   }
+	return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaThreadExitInternal(gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	ctx->exit_simulation();
+	return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaThreadSynchronizeInternal(gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	//Called on host side
+	ctx->synchronize();
+	return g_last_cudaError = cudaSuccess;
+}
+
+cudaError_t CUDARTAPI cudaDeviceSynchronizeInternal(gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	//Blocks until the device has completed all preceding requested tasks
+	ctx->synchronize();
+	return g_last_cudaError = cudaSuccess;
+}
+
+/*******************************************************************************
+ *                                                                              *
+ *                                                                              *
+ *                                                                              *
+ *******************************************************************************/
+
+extern "C" {
+
+/*******************************************************************************
+ *                                                                              *
+ *                                                                              *
+ *                                                                              *
+ *******************************************************************************/
+cudaError_t cudaPeekAtLastError(void)
+{
+	return g_last_cudaError;
+}
+
+__host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) 
+{
+    return cudaMallocInternal(devPtr, size);
+}
+
+__host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size)
+{
+    return cudaMallocHostInternal(ptr, size);
+}
+__host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr, size_t *pitch, size_t width, size_t height)
+{
+    return cudaMallocPitchInternal(devPtr, pitch, width, height);
+}
+
+__host__ cudaError_t CUDARTAPI cudaMallocArray(struct cudaArray **array, const struct cudaChannelFormatDesc *desc, size_t width, size_t height __dv(1))
+{
+    return cudaMallocArrayInternal(array, desc, width, height);
+}
+
+__host__ cudaError_t CUDARTAPI cudaFree(void *devPtr)
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	// TODO...  manage g_global_mem space?
+	return g_last_cudaError = cudaSuccess;
+}
+__host__ cudaError_t CUDARTAPI cudaFreeHost(void *ptr)
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	free (ptr);  // this will crash the system if called twice
+	return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaFreeArray(struct cudaArray *array)
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	// 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)
+{
+    return cudaMemcpyInternal(dst, src, count, kind);
+}
+
+__host__ cudaError_t CUDARTAPI cudaMemcpyToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind)
+{ 
+    return cudaMemcpyToArrayInternal(dst, wOffset, hOffset, src, count, kind);
+}
+
+
+__host__ cudaError_t CUDARTAPI cudaMemcpyFromArray(void *dst, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t count, enum cudaMemcpyKind kind)
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	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))
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	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)
+{
+    return cudaMemcpy2DInternal(dst, dpitch, src, spitch, width, height, kind);
+}
+
+__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)
+{
+    return cudaMemcpy2DToArrayInternal(dst, wOffset, hOffset, src, spitch, width, height, kind);
+}
+
+__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)
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	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))
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	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))
+{
+    return cudaMemcpyToSymbolInternal(symbol, src, count, offset, kind);
+}
+
+
+__host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbol(void *dst, const char *symbol, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost))
+{
+    return cudaMemcpyFromSymbolInternal(dst, symbol, count, offset, kind);
+}
+
+__host__ cudaError_t CUDARTAPI cudaMemGetInfo (size_t *free, size_t *total){
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	//placeholder; should interact with cudaMalloc and cudaFree?
+	*free = 10000000000;
+	*total = 10000000000;
+
+	return g_last_cudaError = cudaSuccess;
+}
+
+/*******************************************************************************
+ *                                                                              *
+ *                                                                              *
+ *                                                                              *
+ *******************************************************************************/
+
+__host__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream)
+{
+    return cudaMemcpyAsyncInternal(dst, src, count, kind, stream);
+}
+
+
+__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)
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	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)
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	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)
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	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)
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	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)
+{
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	cuda_not_implemented(__my_func__,__LINE__);
+	return g_last_cudaError = cudaErrorUnknown;
+}
+
+#if (CUDART_VERSION >= 8000)
+cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks, const char *hostFunc, int blockSize, size_t dynamicSMemSize, unsigned int flags)
+{
+    return cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlagsInternal(numBlocks, hostFunc, blockSize, dynamicSMemSize, flags);
+}
+
+#endif
+
+
+
+/*******************************************************************************
+ *                                                                              *
+ *                                                                              *
+ *                                                                              *
+ *******************************************************************************/
+__host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count)
+{
+    return cudaMemsetInternal(mem, c, count);
+}
+
+//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)
+{
+    return cudaMemsetAsyncInternal(mem, c, count, stream=0);
+}
+
 __host__ cudaError_t CUDARTAPI cudaMemset2D(void *mem, size_t pitch, int c, size_t width, size_t height)
 {
 	if(g_debug_execution >= 3){
@@ -1986,61 +2516,18 @@ __host__ cudaError_t CUDARTAPI cudaBindTexture(size_t *offset,
 		const struct cudaChannelFormatDesc *desc,
 		size_t size __dv(UINT_MAX))
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	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;
+    return cudaBindTextureInternal(offset, texref, devPtr, desc, size __dv(UINT_MAX));
 }
 
 
 __host__ cudaError_t CUDARTAPI cudaBindTextureToArray(const struct textureReference *texref, const struct cudaArray *array, const struct cudaChannelFormatDesc *desc)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	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;
+    return cudaBindTextureToArrayInternal(texref, array, desc);
 }
 
-__host__ cudaError_t CUDARTAPI cudaUnbindTexture(const struct textureReference *texref){
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-   CUctx_st *context = GPGPUSim_Context();
-   gpgpu_t *gpu = context->get_device()->get_gpgpu();
-   printf("GPGPU-Sim PTX: in cudaUnbindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference));
-   printf("GPGPU-Sim PTX:   Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref));
-
-   gpu->gpgpu_ptx_sim_unbindTexture(texref);
-   return g_last_cudaError = cudaSuccess;
+__host__ cudaError_t CUDARTAPI cudaUnbindTexture(const struct textureReference *texref)
+{
+    return cudaUnbindTextureInternal(texref);
 }
 
 __host__ cudaError_t CUDARTAPI cudaGetTextureAlignmentOffset(size_t *offset, const struct textureReference *texref)
@@ -2125,24 +2612,9 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun )
     return cudaLaunchInternal( hostFun );
 }
 
-__host__ cudaError_t CUDARTAPI cudaLaunchKernel ( const char* hostFun, dim3 gridDim, dim3 blockDim, const void** args, size_t sharedMem, cudaStream_t stream )
+__host__ cudaError_t CUDARTAPI cudaLaunchKernel( const char* hostFun, dim3 gridDim, dim3 blockDim, const void** args, size_t sharedMem, cudaStream_t stream )
 {
-
-    if(g_debug_execution >= 3){
-        announce_call(__my_func__);
-    }
-    CUctx_st *context = GPGPUSim_Context();
-    function_info *entry = context->get_kernel(hostFun);
-#if CUDART_VERSION < 10000
-    cudaConfigureCallInternal(gridDim, blockDim, sharedMem, stream);
-#endif
-    for(unsigned i = 0; i < entry->num_args(); i++){
-        std::pair<size_t, unsigned> p = entry->get_param_config(i);
-        cudaSetupArgumentInternal(args[i], p.first, p.second);
-    }
-
-    cudaLaunchInternal(hostFun);
-    return g_last_cudaError = cudaSuccess;
+    return cudaLaunchKernelInternal(hostFun, gridDim, blockDim, args, sharedMem, stream);
 }
 
 
@@ -2154,18 +2626,7 @@ __host__ cudaError_t CUDARTAPI cudaLaunchKernel ( const char* hostFun, dim3 grid
 
 __host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	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;
+    return cudaStreamCreateInternal(stream);
 }
 
 //TODO: introduce priorities
@@ -2192,32 +2653,12 @@ __host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t
 
 __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-#if (CUDART_VERSION >= 3000)
-	//per-stream synchronization required for application using external libraries without explicit synchronization in the code to 
-	//avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. 
-	stream->synchronize();
-	g_stream_manager()->destroy_stream(stream);
-#endif
-	return g_last_cudaError = cudaSuccess;
+    return cudaStreamDestroyInternal(stream);
 }
 
 __host__ cudaError_t CUDARTAPI cudaStreamSynchronize(cudaStream_t stream)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-#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;
+    return cudaStreamSynchronizeInternal(stream);
 }
 
 __host__ cudaError_t CUDARTAPI cudaStreamQuery(cudaStream_t stream)
@@ -2256,52 +2697,14 @@ __host__ cudaError_t CUDARTAPI cudaEventCreate(cudaEvent_t *event)
 	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)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	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;
+    return cudaEventRecordInternal(event, stream);
 }
 
 __host__ cudaError_t CUDARTAPI cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event, unsigned int flags)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-   //reference: https://www.cs.cmu.edu/afs/cs/academic/class/15668-s11/www/cuda-doc/html/group__CUDART__STREAM_gfe68d207dc965685d92d3f03d77b0876.html 
-	CUevent_st *e = get_event(event);
-	if( !e ){
-	   printf("GPGPU-Sim API: Warning: cudaEventRecord has not been called on event before calling cudaStreamWaitEvent.\nNothing to be done.\n");
-      return g_last_cudaError = cudaSuccess;
-   }
-   if (!stream){
-      g_stream_manager()->pushCudaStreamWaitEventToAllStreams(e, flags);
-   } else {
-	   struct CUstream_st *s = (struct CUstream_st *)stream;
-	   stream_operation op(s,e,flags);
-	   g_stream_manager()->push(op);
-   }
-	return g_last_cudaError = cudaSuccess;
+    return cudaStreamWaitEventInternal(stream, event, flags);
 }
 
 __host__ cudaError_t CUDARTAPI cudaEventQuery(cudaEvent_t event)
@@ -2374,22 +2777,13 @@ __host__ cudaError_t CUDARTAPI cudaEventElapsedTime(float *ms, cudaEvent_t start
 
 __host__ cudaError_t CUDARTAPI cudaThreadExit(void)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	exit_simulation();
-	return g_last_cudaError = cudaSuccess;
+    return cudaThreadExitInternal();
 }
 
 __host__ cudaError_t CUDARTAPI cudaThreadSynchronize(void)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	//Called on host side
-	synchronize();
-	return g_last_cudaError = cudaSuccess;
-};
+    return cudaThreadSynchronizeInternal();
+}
 
 int CUDARTAPI __cudaSynchronizeThreads(void**, void*)
 {
@@ -2502,7 +2896,6 @@ void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context){
     while (std::getline(infile, line))
     {
          //int pos = line.find(std::string(get_app_binary_name(app_binary)));
-         const char *ptx_file = line.c_str();
          int pos1 = line.find("sm_");
          int pos2 = line.find_last_of(".");
          if (pos1==std::string::npos&&pos2==std::string::npos){
@@ -2529,12 +2922,11 @@ void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context){
  *	enabled
  * */
 void cuda_runtime_api::extract_code_using_cuobjdump(){
-    CUctx_st *context = GPGPUSim_Context();
-    unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
+    CUctx_st *context = GPGPUSim_Context(gpgpu_ctx);
 
     //prevent the dumping by cuobjdump everytime we execute the code!
     const char *override_cuobjdump = getenv("CUOBJDUMP_SIM_FILE"); 
-    char command[1000], ptx_file[1000];
+    char command[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());
@@ -2888,7 +3280,7 @@ cuobjdumpPTXSection* cuda_runtime_api::findPTXSection(const std::string identifi
 
 //! Extract the code using cuobjdump and remove unnecessary sections
 void cuda_runtime_api::cuobjdumpInit(){
-	CUctx_st *context = GPGPUSim_Context();
+	CUctx_st *context = GPGPUSim_Context(gpgpu_ctx);
 	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){
@@ -2901,7 +3293,7 @@ void cuda_runtime_api::cuobjdumpInit(){
 //! Either submit PTX for simulation or convert SASS to PTXPlus and submit it
 void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){
 
-	CUctx_st *context = GPGPUSim_Context();
+	CUctx_st *context = GPGPUSim_Context(this);
 	if(api->fatbin_registered[handle]) return;
 	api->fatbin_registered[handle] = true;
 	std::string fname = api->fatbinmap[handle];
@@ -3095,15 +3487,11 @@ cudaError_t cudaDeviceReset ( void ) {
     }
 	return g_last_cudaError = cudaSuccess;
 }
-cudaError_t CUDARTAPI cudaDeviceSynchronize(void){
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	//Blocks until the device has completed all preceding requested tasks
-	synchronize();
-	return g_last_cudaError = cudaSuccess;
-}
 
+cudaError_t CUDARTAPI cudaDeviceSynchronize(void)
+{
+    return cudaDeviceSynchronizeInternal();
+}
 
 void __cudaRegisterShared(
 		void **fatCubinHandle,
@@ -3142,22 +3530,7 @@ void __cudaRegisterTexture(
 		int ext
 ) //passes in a newly created textureReference
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	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__ );
+    __cudaRegisterTextureInternal(fatCubinHandle, hostVar, deviceAddress, deviceName, dim, norm, ext);
 }
 
 
@@ -3189,30 +3562,7 @@ cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj)
 
 cudaError_t cudaGLUnmapBufferObject(GLuint bufferObj)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-#ifdef OPENGL_SUPPORT
-	CUctx_st* ctx = GPGPUSim_Context();
-	glbmap_entry_t *p = ctx->api->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
+    return cudaGLUnmapBufferObjectInternal(bufferObj);
 }
 
 cudaError_t cudaGLUnregisterBufferObject(GLuint bufferObj) 
@@ -3324,12 +3674,7 @@ cudaError_t CUDARTAPI cudaRuntimeGetVersion(int *runtimeVersion)
 #if CUDART_VERSION >= 3000
 __host__ cudaError_t CUDARTAPI cudaFuncSetCacheConfig(const char *func, enum cudaFuncCache  cacheConfig )
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	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;
+    return cudaFuncSetCacheConfigInternal(func, cacheConfig);
 }
 
 //Jin: hack for cdp
@@ -5035,23 +5380,7 @@ CUresult CUDAAPI cuLaunchKernel(CUfunction f,
                                 void **kernelParams,
                                 void **extra)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-    if (extra!=NULL){
-	    printf("GPGPU-Sim CUDA DRIVER API: ERROR: Currently do not support void** extra.\n");
-        abort();
-    }
-    const char *hostFun = (const char*) f;
-	CUctx_st *context = GPGPUSim_Context();
-	function_info *entry = context->get_kernel(hostFun);
-    cudaConfigureCallInternal(dim3(gridDimX, gridDimY, gridDimZ), dim3(blockDimX, blockDimY, blockDimZ), sharedMemBytes, (cudaStream_t) hStream);
-    for(unsigned i = 0; i < entry->num_args(); i++){
-        std::pair<size_t, unsigned> p = entry->get_param_config(i);
-        cudaSetupArgument(kernelParams[i], p.first, p.second);
-    }
-    cudaLaunchInternal(hostFun);
-	return CUDA_SUCCESS;
+    return cuLaunchKernelInternal(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra);
 }
 #endif /* CUDART_VERSION >= 4000 */
 
diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h
index 45c5cdd..61d7507 100644
--- a/libcuda/gpgpu_context.h
+++ b/libcuda/gpgpu_context.h
@@ -52,6 +52,10 @@ class gpgpu_context {
 	cuda_device_runtime* device_runtime;
 	ptx_stats* stats;
 	// member function list
+	void synchronize();
+	void exit_simulation();
+	void print_simulation_time();
+	int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid );
 	void cuobjdumpParseBinary(unsigned int handle);
 	class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num );
 	class symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename );
@@ -60,6 +64,7 @@ class gpgpu_context {
 	void print_ptx_file( const char *p, unsigned source_num, const char *filename );
 	class symbol_table* init_parser(const char*);
 	class gpgpu_sim *gpgpu_ptx_sim_init_perf();
+	void start_sim_thread(int api);
 	struct _cuda_device_id *GPGPUSim_Init();
 	void ptx_reg_options(option_parser_t opp);
 	const ptx_instruction* pc_to_instruction(unsigned pc); 
diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc
index 7f029c7..b032c05 100644
--- a/libopencl/opencl_runtime_api.cc
+++ b/libopencl/opencl_runtime_api.cc
@@ -647,13 +647,13 @@ unsigned _cl_kernel::sm_context_uid = 0;
 class _cl_device_id *GPGPUSim_Init()
 {
    static _cl_device_id *the_device = NULL;
+   gpgpu_context *ctx;
+   ctx = GPGPU_Context();
    if( !the_device ) { 
-       gpgpu_context *ctx;
-       ctx = GPGPU_Context();
       gpgpu_sim *the_gpu = ctx->gpgpu_ptx_sim_init_perf(); 
       the_device = new _cl_device_id(the_gpu);
    } 
-   start_sim_thread(2);
+   ctx->start_sim_thread(2);
    return the_device;
 }
 
@@ -960,7 +960,7 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
    if ( ctx->func_sim->g_ptx_sim_mode )
       ctx->func_sim->gpgpu_opencl_ptx_sim_main_func( grid );
    else
-      gpgpu_opencl_ptx_sim_main_perf( grid );
+      ctx->gpgpu_opencl_ptx_sim_main_perf( grid );
    return CL_SUCCESS;
 }
 
diff --git a/nightly.jenkinsfile b/nightly.jenkinsfile
index e5ffa57..c841fb0 100644
--- a/nightly.jenkinsfile
+++ b/nightly.jenkinsfile
@@ -17,7 +17,7 @@ pipeline {
             steps {
                 sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\
                 source `pwd`/setup_environment &&\
-                make -j'
+                make -j 10'
             }
         }
         stage('nightly-simulations-build'){
@@ -34,7 +34,7 @@ pipeline {
                     source `pwd`/setup_environment &&\
                     cd gpgpu-sim_simulations && \
                     source ./benchmarks/src/setup_environment && \
-                    make -i -j -C ./benchmarks/src/ all && \
+                    make -i -j 10 -C ./benchmarks/src/ all && \
                     make -C ./benchmarks/src data'
             }
         }
@@ -46,7 +46,7 @@ pipeline {
                     sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\
                         source `pwd`/setup_environment &&\
                         ./gpgpu-sim_simulations/util/job_launching/run_simulations.py \
-                            -B `cat ./gpgpu-sim_simulations/util/job_launching/apps/all-apps.list` -C TITANV-2B -N nightly-$$ && \
+                            -B `cat ./gpgpu-sim_simulations/util/job_launching/apps/correlation-apps.list` -C QV100-2B -N nightly-$$ && \
                         ./gpgpu-sim_simulations/util/job_launching/monitor_func_test.py -I -S 1800 -v \
                             -s stats-per-app-9.1.csv -T 12 -K -N nightly-$$ && \
                         ./gpgpu-sim_simulations/util/plotting/plot-get-stats.py -c stats-per-app-9.1.csv -P cuda-9.1.nightly'
@@ -65,8 +65,8 @@ pipeline {
 
                 sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\
                     ./gpgpu-sim_simulations/util/job_launching/get_stats.py -R -K -k \
-                        -B `cat ./gpgpu-sim_simulations/util/job_launching/apps/all-apps.list` \
-                        -C TITANV-2B > stats-per-kernel-9.1.csv &&\
+                        -B `cat ./gpgpu-sim_simulations/util/job_launching/apps/correlation-apps.list` \
+                        -C QV100-2B > stats-per-kernel-9.1.csv &&\
                     ./gpgpu-sim_simulations/util/plotting/merge-stats.py \
                         -c ./gpgpu-sim-results-repo/${JOB_NAME}/stats-per-app-9.1.csv,./stats-per-app-9.1.csv -R \
                         > per-app-merge-9.1.csv &&\
@@ -76,7 +76,7 @@ pipeline {
                     source `pwd`/setup_environment &&\
                     ./gpgpu-sim_simulations/util/plotting/plot-correlation.py -c per-kernel-merge-9.1.csv \
                         -p cuda-9.1.nightly -b ./gpgpu-sim_simulations/util/plotting/known.correlation.outliers.list \
-                        | grep "Correl=" | tee correl.9.1.txt &&\
+                        | grep "Correl=" -B 1 | tee correl.9.1.txt &&\
                     cp stats-per-*.csv ./gpgpu-sim-results-repo/${JOB_NAME}/ &&\
                     cd ./gpgpu-sim-results-repo &&\
                     git pull &&\
diff --git a/setup_environment b/setup_environment
index 58f4713..ca60d6b 100644
--- a/setup_environment
+++ b/setup_environment
@@ -7,13 +7,10 @@ export GPGPUSIM_ROOT="$( cd "$( dirname "$BASH_SOURCE" )" && pwd )"
 
 GPGPUSIM_VERSION_STRING=`cat $GPGPUSIM_ROOT/version | awk '/Version/ {print $8}'`
 #Detect Git branch and commit #
-CURRENT_PWD=`pwd`
-cd $GPGPUSIM_ROOT
-GIT_COMMIT=`git log -n 1 | head -1 | sed -re 's/commit (.*)/\1/'`
-GIT_FILES_CHANGED=`git diff --numstat | wc | sed -re 's/^\s+([0-9]+).*/\1./'`
-GIT_FILES_CHANGED+=`git diff --numstat --cached | wc | sed -re 's/^\s+([0-9]+).*/\1/'`
+GIT_COMMIT=`git --git-dir=$GPGPUSIM_ROOT/.git log -n 1 | head -1 | sed -re 's/commit (.*)/\1/'`
+GIT_FILES_CHANGED=`git --git-dir=$GPGPUSIM_ROOT/.git diff --numstat | wc | sed -re 's/^\s+([0-9]+).*/\1./'`
+GIT_FILES_CHANGED+=`git --git-dir=$GPGPUSIM_ROOT/.git diff --numstat --cached | wc | sed -re 's/^\s+([0-9]+).*/\1/'`
 GPGPUSIM_BUILD_STRING="gpgpu-sim_git-commit-$GIT_COMMIT-modified_$GIT_FILES_CHANGED"
-cd $CURRENT_PWD
 
 echo -n "GPGPU-Sim version $GPGPUSIM_VERSION_STRING (build $GPGPUSIM_BUILD_STRING) ";
 
diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc
index 9aa1d73..e8e4b51 100644
--- a/src/abstract_hardware_model.cc
+++ b/src/abstract_hardware_model.cc
@@ -808,14 +808,14 @@ void kernel_info_t::notify_parent_finished() {
    if(m_parent_kernel) {
        m_kernel_entry->gpgpu_ctx->device_runtime->g_total_param_size -= ((m_kernel_entry->get_args_aligned_size() + 255)/256*256);
        m_parent_kernel->remove_child(this);
-       g_stream_manager()->register_finished_kernel(m_parent_kernel->get_uid());
+       m_kernel_entry->gpgpu_ctx->the_gpgpusim->g_stream_manager->register_finished_kernel(m_parent_kernel->get_uid());
    }
 }
 
 CUstream_st * kernel_info_t::create_stream_cta(dim3 ctaid) {
     assert(get_default_stream_cta(ctaid));
     CUstream_st * stream = new CUstream_st();
-    g_stream_manager()->add_stream(stream);
+    m_kernel_entry->gpgpu_ctx->the_gpgpusim->g_stream_manager->add_stream(stream);
     assert(m_cta_streams.find(ctaid) != m_cta_streams.end());
     assert(m_cta_streams[ctaid].size() >= 1); //must have default stream
     m_cta_streams[ctaid].push_back(stream);
@@ -831,7 +831,7 @@ CUstream_st * kernel_info_t::get_default_stream_cta(dim3 ctaid) {
     else {
       m_cta_streams[ctaid] = std::list<CUstream_st *>();
       CUstream_st * stream = new CUstream_st();
-      g_stream_manager()->add_stream(stream);
+      m_kernel_entry->gpgpu_ctx->the_gpgpusim->g_stream_manager->add_stream(stream);
       m_cta_streams[ctaid].push_back(stream);
       return stream;
     }
@@ -863,7 +863,7 @@ void kernel_info_t::destroy_cta_streams() {
      for(auto s = m_cta_streams.begin(); s != m_cta_streams.end(); s++) {
         stream_size += s->second.size();
         for(auto ss = s->second.begin(); ss != s->second.end(); ss++)
-        g_stream_manager()->destroy_stream(*ss);
+        m_kernel_entry->gpgpu_ctx->the_gpgpusim->g_stream_manager->destroy_stream(*ss);
         s->second.clear();
      }
      printf("size %lu\n", stream_size);
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index 60f7f68..341e44c 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -447,19 +447,25 @@ protected:
     class gpgpu_sim * m_gpu;
 };
 
-#define GLOBAL_HEAP_START 0xC0000000
-   // start allocating from this address (lower values used for allocating globals in .ptx file)
-#define SHARED_MEM_SIZE_MAX (96*1024)
-#define LOCAL_MEM_SIZE_MAX (16*1024)
-#define MAX_STREAMING_MULTIPROCESSORS 80   //scale it to Volta
-#define MAX_THREAD_PER_SM 2048
-#define MAX_WARP_PER_SM 64
-#define TOTAL_LOCAL_MEM_PER_SM (MAX_THREAD_PER_SM*LOCAL_MEM_SIZE_MAX)
-#define TOTAL_SHARED_MEM (MAX_STREAMING_MULTIPROCESSORS*SHARED_MEM_SIZE_MAX)
-#define TOTAL_LOCAL_MEM (MAX_STREAMING_MULTIPROCESSORS*MAX_THREAD_PER_SM*LOCAL_MEM_SIZE_MAX)
-#define SHARED_GENERIC_START (GLOBAL_HEAP_START-TOTAL_SHARED_MEM)
-#define LOCAL_GENERIC_START (SHARED_GENERIC_START-TOTAL_LOCAL_MEM)
-#define STATIC_ALLOC_LIMIT (GLOBAL_HEAP_START - (TOTAL_LOCAL_MEM+TOTAL_SHARED_MEM))
+// Let's just upgrade to C++11 so we can use constexpr here...
+// start allocating from this address (lower values used for allocating globals in .ptx file)
+const unsigned long long GLOBAL_HEAP_START = 0xC0000000;
+// Volta max shmem size is 96kB
+const unsigned long long SHARED_MEM_SIZE_MAX = 96 * (1 << 10);
+// Volta max local mem is 16kB
+const unsigned long long LOCAL_MEM_SIZE_MAX = 1 << 14;
+// Volta Titan V has 80 SMs
+const unsigned MAX_STREAMING_MULTIPROCESSORS = 80;
+// Max 2048 threads / SM
+const unsigned MAX_THREAD_PER_SM = 1 << 11;
+// MAX 64 warps / SM
+const unsigned MAX_WARP_PER_SM = 1 << 6;
+const unsigned long long TOTAL_LOCAL_MEM_PER_SM = MAX_THREAD_PER_SM * LOCAL_MEM_SIZE_MAX;
+const unsigned long long TOTAL_SHARED_MEM = MAX_STREAMING_MULTIPROCESSORS * SHARED_MEM_SIZE_MAX;
+const unsigned long long TOTAL_LOCAL_MEM = MAX_STREAMING_MULTIPROCESSORS * MAX_THREAD_PER_SM * LOCAL_MEM_SIZE_MAX;
+const unsigned long long SHARED_GENERIC_START = GLOBAL_HEAP_START - TOTAL_SHARED_MEM;
+const unsigned long long LOCAL_GENERIC_START = SHARED_GENERIC_START - TOTAL_LOCAL_MEM;
+const unsigned long long STATIC_ALLOC_LIMIT = GLOBAL_HEAP_START - (TOTAL_LOCAL_MEM + TOTAL_SHARED_MEM);
 
 #if !defined(__CUDA_RUNTIME_API_H__)
 
@@ -525,10 +531,10 @@ private:
     int checkpoint_option;
     int checkpoint_kernel;
     int checkpoint_CTA;
-    int resume_option;
-    int resume_kernel;
-    int resume_CTA;
-    int checkpoint_CTA_t;
+    unsigned resume_option;
+    unsigned resume_kernel;
+    unsigned resume_CTA;
+    unsigned checkpoint_CTA_t;
     int checkpoint_insn_Y;
     int   g_ptx_inst_debug_to_file;
     char* g_ptx_inst_debug_file;
@@ -546,10 +552,10 @@ public:
     int checkpoint_option;
     int checkpoint_kernel;
     int checkpoint_CTA;
-    int resume_option;
-    int resume_kernel;
-    int resume_CTA;
-    int checkpoint_CTA_t;
+    unsigned resume_option;
+    unsigned resume_kernel;
+    unsigned resume_CTA;
+    unsigned checkpoint_CTA_t;
     int checkpoint_insn_Y;
 
     //Move some cycle core stats here instead of being global
@@ -998,7 +1004,7 @@ public:
 			printf("Printing mem access generated\n");
 			std::list<mem_access_t>::iterator it;	
 			for (it = m_accessq.begin(); it != m_accessq.end(); ++it){
-   				 printf("MEM_TXN_GEN:%s:%x, Size:%d \n",mem_access_type_str(it->get_type()), it->get_addr(),it->get_size());
+   				 printf("MEM_TXN_GEN:%s:%llx, Size:%d \n",mem_access_type_str(it->get_type()), it->get_addr(),it->get_size());
 			}	
 		}
     }   
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 1b0e841..7a130ea 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -429,7 +429,7 @@ void gpgpu_t::memcpy_to_gpu( size_t dst_start_addr, const void *src, size_t coun
 
    // Copy into the performance model.
    //extern gpgpu_sim* g_the_gpu;
-   g_the_gpu()->perf_memcpy_to_gpu(dst_start_addr, count);
+   gpgpu_ctx->the_gpgpusim->g_the_gpu->perf_memcpy_to_gpu(dst_start_addr, count);
    if(g_debug_execution >= 3) {
       printf( " done.\n");
       fflush(stdout);
@@ -448,7 +448,7 @@ void gpgpu_t::memcpy_from_gpu( void *dst, size_t src_start_addr, size_t count )
 
    // Copy into the performance model.
    //extern gpgpu_sim* g_the_gpu;
-   g_the_gpu()->perf_memcpy_to_gpu(src_start_addr, count);
+   gpgpu_ctx->the_gpgpusim->g_the_gpu->perf_memcpy_to_gpu(src_start_addr, count);
    if(g_debug_execution >= 3) {
       printf( " done.\n");
       fflush(stdout);
@@ -1254,7 +1254,7 @@ void function_info::param_to_shared( memory_space *shared_mem, symbol_table *sym
 {
    // TODO: call this only for PTXPlus with GT200 models 
    //extern gpgpu_sim* g_the_gpu;
-   if (not g_the_gpu()->get_config().convert_to_ptxplus()) return;
+   if (not gpgpu_ctx->the_gpgpusim->g_the_gpu->get_config().convert_to_ptxplus()) return;
 
    // copies parameters into simulated shared memory
    for( std::map<unsigned,param_info>::iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) {
@@ -1309,7 +1309,7 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt
     char buff[1024];
     std::string filename_c(filename+"_c");
     snprintf(buff,1024,"c++filt %s > %s", get_name().c_str(), filename_c.c_str());
-    system(buff);
+    assert(system(buff) != NULL);
     FILE *fp = fopen(filename_c.c_str(), "r");
     fgets(buff, 1024, fp);
     fclose(fp);
@@ -1432,13 +1432,13 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt
     fout = fopen(ptx_config_fn.c_str(), "a");
     assert(fout!=NULL);
     for (unsigned i = 0; i<line_number; i++){
-        fgets(buff, 1024, fin);
+        assert(fgets(buff, 1024, fin) != NULL);
         assert(!feof(fin));
     }
     fprintf(fout, "\n\n");
     do{
         fprintf(fout, "%s", buff);
-        fgets(buff, 1024, fin);
+        assert(fgets(buff, 1024, fin) != NULL);
         if(feof(fin)){
             break;
         }
@@ -1491,7 +1491,6 @@ static unsigned get_tex_datasize( const ptx_instruction *pI, ptx_thread_info *th
    const operand_info &src1 = pI->src1(); //the name of the texture
    std::string texname = src1.name();
 
-   gpgpu_t *gpu = thread->get_gpu();
    /*
      For programs with many streams, textures can be bound and unbound
      asynchronously.  This means we need to use the kernel's "snapshot" of
@@ -1577,7 +1576,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
       }
       
       //Tensorcore is warp synchronous operation. So these instructions needs to be executed only once. To make the simulation faster removing the redundant tensorcore operation
-      if(!tensorcore_op(inst_opcode)||(tensorcore_op(inst_opcode))&&(lane_id==0)){
+      if(!tensorcore_op(inst_opcode)||((tensorcore_op(inst_opcode))&&(lane_id==0))){
 	      switch ( inst_opcode ) {
 	#define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,this); op_classification = CLASSIFICATION; break;
 	#define OP_W_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,get_core(),inst); op_classification = CLASSIFICATION; break;
@@ -2138,19 +2137,13 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL
 	 kernel_func_info->set_pdom();
     }
 
-    unsigned max_cta_tot = max_cta(kernel_info,kernel.threads_per_cta(), g_the_gpu()->getShaderCoreConfig()->warp_size, g_the_gpu()->getShaderCoreConfig()->n_thread_per_shader, g_the_gpu()->getShaderCoreConfig()->gpgpu_shmem_size, g_the_gpu()->getShaderCoreConfig()->gpgpu_shader_registers, g_the_gpu()->getShaderCoreConfig()->max_cta_per_core);
+    unsigned max_cta_tot = max_cta(kernel_info,kernel.threads_per_cta(), gpgpu_ctx->the_gpgpusim->g_the_gpu->getShaderCoreConfig()->warp_size, gpgpu_ctx->the_gpgpusim->g_the_gpu->getShaderCoreConfig()->n_thread_per_shader, gpgpu_ctx->the_gpgpusim->g_the_gpu->getShaderCoreConfig()->gpgpu_shmem_size, gpgpu_ctx->the_gpgpusim->g_the_gpu->getShaderCoreConfig()->gpgpu_shader_registers, gpgpu_ctx->the_gpgpusim->g_the_gpu->getShaderCoreConfig()->max_cta_per_core);
     printf("Max CTA : %d\n",max_cta_tot);
 
-    
-
-
-      
-    int inst_count=50;
-    int cp_op= g_the_gpu()->checkpoint_option;
-    int cp_CTA = g_the_gpu()->checkpoint_CTA;
-    int cp_kernel= g_the_gpu()->checkpoint_kernel;
-    cp_count= g_the_gpu()->checkpoint_insn_Y;
-    cp_cta_resume= g_the_gpu()->checkpoint_CTA_t;
+    int cp_op= gpgpu_ctx->the_gpgpusim->g_the_gpu->checkpoint_option;
+    int cp_kernel= gpgpu_ctx->the_gpgpusim->g_the_gpu->checkpoint_kernel;
+    cp_count= gpgpu_ctx->the_gpgpusim->g_the_gpu->checkpoint_insn_Y;
+    cp_cta_resume= gpgpu_ctx->the_gpgpusim->g_the_gpu->checkpoint_CTA_t;
     int cta_launched =0;
 
     //we excute the kernel one CTA (Block) at the time, as synchronization functions work block wise
@@ -2162,8 +2155,8 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL
         {
            functionalCoreSim cta(
                &kernel,
-               g_the_gpu(),
-               g_the_gpu()->getShaderCoreConfig()->warp_size
+               gpgpu_ctx->the_gpgpusim->g_the_gpu,
+               gpgpu_ctx->the_gpgpusim->g_the_gpu->getShaderCoreConfig()->warp_size
            );
            cta.execute(cp_count,temp);
 
@@ -2184,7 +2177,7 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL
 	{
       char f1name[2048];
       snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", kernel.get_uid() );
-      g_checkpoint->store_global_mem(g_the_gpu()->get_global_memory(), f1name , "%08x");
+      g_checkpoint->store_global_mem(gpgpu_ctx->the_gpgpusim->g_the_gpu->get_global_memory(), f1name , (char *)"%08x");
 	}
 
 
@@ -2195,7 +2188,7 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL
    //openCL kernel simulation calls don't register the kernel so we don't register its exit
    if(!openCL) {
       //extern stream_manager *g_stream_manager;
-      g_stream_manager()->register_finished_kernel(kernel.get_uid());
+      gpgpu_ctx->the_gpgpusim->g_stream_manager->register_finished_kernel(kernel.get_uid());
    }
 
    //******PRINTING*******
@@ -2210,7 +2203,7 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL
    //g_simulation_starttime is initilized by gpgpu_ptx_sim_init_perf() in gpgpusim_entrypoint.cc upon starting gpgpu-sim
    time_t end_time, elapsed_time, days, hrs, minutes, sec;
    end_time = time((time_t *)NULL);
-   elapsed_time = MAX(end_time - GPGPUsim_ctx_ptr()->g_simulation_starttime, 1);
+   elapsed_time = MAX(end_time - gpgpu_ctx->the_gpgpusim->g_simulation_starttime, 1);
 	
 
    //calculating and printing simulation time in terms of days, hours, minutes and seconds
@@ -2312,18 +2305,15 @@ void functionalCoreSim::execute(int inst_count, unsigned ctaid_cp)
     checkpoint *g_checkpoint;
     g_checkpoint = new checkpoint();
     
-    symbol * sym;
     ptx_reg_t regval;
     regval.u64= 123;
-    symbol_table * symtab= m_kernel->entry()->get_symtab();
-
 
     unsigned ctaid =m_kernel->get_next_cta_id_single();
     if(m_gpu->checkpoint_option==1 && (m_kernel->get_uid()==m_gpu->checkpoint_kernel) && (ctaid_cp>=m_gpu->checkpoint_CTA) && (ctaid_cp<m_gpu->checkpoint_CTA_t))
    {
        char fname[2048];
        snprintf(fname,2048,"checkpoint_files/shared_mem_%d.txt",ctaid-1 );
-       g_checkpoint->store_global_mem(m_thread[0]->m_shared_mem, fname , "%08x");
+       g_checkpoint->store_global_mem(m_thread[0]->m_shared_mem, fname , (char *)"%08x");
       for(int i=0; i<32*m_warp_count;i++)
       {
          char fname[2048];
@@ -2331,7 +2321,7 @@ void functionalCoreSim::execute(int inst_count, unsigned ctaid_cp)
           m_thread[i]->print_reg_thread(fname);
           char f1name[2048];
          snprintf(f1name,2048,"checkpoint_files/local_mem_thread_%d_%d_reg.txt",i,ctaid-1 );
-         g_checkpoint->store_global_mem(m_thread[i]->m_local_mem, f1name , "%08x");
+         g_checkpoint->store_global_mem(m_thread[i]->m_local_mem, f1name , (char *)"%08x");
          m_thread[i]->set_done();
          m_thread[i]->exitCore();
          m_thread[i]->registerExit();
diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc
index dc3adc3..4baced5 100644
--- a/src/cuda-sim/cuda_device_runtime.cc
+++ b/src/cuda-sim/cuda_device_runtime.cc
@@ -27,7 +27,6 @@
    }
 
 
-//extern stream_manager *g_stream_manager();
 
 //Handling device runtime api:
 //void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize)
@@ -285,7 +284,7 @@ void cuda_device_runtime::launch_one_device_kernel() {
         device_launch_operation_t &op = g_cuda_device_launch_op.front();
 
         stream_operation stream_op = stream_operation(op.grid, gpgpu_ctx->func_sim->g_ptx_sim_mode, op.stream);
-        g_stream_manager()->push(stream_op);
+        gpgpu_ctx->the_gpgpusim->g_stream_manager->push(stream_op);
         g_cuda_device_launch_op.pop_front();
     }
 }
diff --git a/src/cuda-sim/half.h b/src/cuda-sim/half.h
index 8f1a8eb..9f74bb7 100644
--- a/src/cuda-sim/half.h
+++ b/src/cuda-sim/half.h
@@ -642,10 +642,10 @@ namespace half_float
 			if(exp > 16)

 			{

 				if(R == std::round_toward_infinity)

-					return hbits | 0x7C00 - (hbits>>15);

+					return hbits | (0x7C00 - (hbits>>15));

 				else if(R == std::round_toward_neg_infinity)

-					return hbits | 0x7BFF + (hbits>>15);

-				return hbits | 0x7BFF + (R!=std::round_toward_zero);

+					return hbits | (0x7BFF + (hbits>>15));

+				return hbits | (0x7BFF + (R!=std::round_toward_zero));

 			}

 			if(exp < -13)

 				value = std::ldexp(value, 24);

diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index 58a077e..014e588 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -203,7 +203,7 @@ void ptx_thread_info::print_reg_thread(char * fname)
           const std::string &name = it->first->name();
           const std::string &dec= it->first->decl_location();
           unsigned size = it->first->get_size_in_bytes();
-          fprintf(fp,"%s %llu %s %d\n",name.c_str(),it->second, dec.c_str(),size );
+          fprintf(fp,"%s %llu %s %d\n", name.c_str(), it->second, dec.c_str(), size);
           
         }
    //m_regs.pop_back();     
@@ -224,7 +224,6 @@ void ptx_thread_info::resume_reg_thread(char * fname, symbol_table * symtab)
       {
           symbol *reg;
           char * pch;
-          unsigned size;
           pch = strtok (line," ");
           char * name =pch;
           reg= symtab->lookup(name);
@@ -232,11 +231,7 @@ void ptx_thread_info::resume_reg_thread(char * fname, symbol_table * symtab)
           pch = strtok (NULL," "); 
           data = atoi(pch);
           pch = strtok (NULL," ");
-          char * decl= pch;
           pch = strtok (NULL," ");
-          size = atoi(pch);
-
-
           m_regs.back()[reg] = data;
       }
       fclose ( fp2 );
@@ -1819,9 +1814,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
    	ptx_reg_t matrix_d[16][16];
    	ptx_reg_t src_data;
 	ptx_thread_info *thread;
-	int stride;
 
-	unsigned wmma_type = pI->get_wmma_type();
 	unsigned a_layout = pI->get_wmma_layout(0);
 	unsigned b_layout = pI->get_wmma_layout(1);
 	unsigned type = pI->get_type();
@@ -1833,7 +1826,6 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
          	tid= inst.warp_id_func()*core->get_warp_size();
         else
          	tid= inst.warp_id()*core->get_warp_size();
-	unsigned thread_group_index;
 	float temp;
 	half temp2; 	
 	
@@ -1847,9 +1839,9 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
          		ptx_reg_t v[8];
          		thread->get_vector_operand_values( src_a, v, nelem );
 			if(core->get_gpu()->gpgpu_ctx->debug_tensorcore){
-				printf("Thread%d_Iteration=%d\n:",thrd,operand_num);
-				for(k=0;k<nelem;k++){
-					printf("%x ",v[k].u64);
+				printf("Thread%d_Iteration=%d\n:", thrd, operand_num);
+				for(k = 0; k < nelem; k++){
+					printf("%llx ",v[k].u64);
 				}
 				printf("\n");
 			}
@@ -2027,7 +2019,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
 
 				printf("thread%d:",thrd);
 				for(k=0;k<8;k++){
-					printf("%x ",matrix_d[row_t[k]][col_t[k]].f16);
+					printf("%x ", (unsigned int)matrix_d[row_t[k]][col_t[k]].f16);
 				}
 				printf("\n");
 			}
@@ -2038,7 +2030,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst )
 			nw_data4.s64=((matrix_d[row_t[6]][col_t[6]].s64   & 0xffff))|((matrix_d[row_t[7]][col_t[7]].s64&0xffff)<<16);
    			thread->set_vector_operand_values(dst,nw_data1,nw_data2,nw_data3,nw_data4);
 			if(core->get_gpu()->gpgpu_ctx->debug_tensorcore)
-		 		printf("thread%d=%x,%x,%x,%x",thrd,nw_data1.s64,nw_data2.s64,nw_data3.s64,nw_data4.s64);
+		 		printf("thread%d=%llx,%llx,%llx,%llx", thrd, nw_data1.s64, nw_data2.s64, nw_data3.s64, nw_data4.s64);
 		
 		}
 		else{
@@ -2298,9 +2290,8 @@ unsigned int saturatei(unsigned int a, unsigned int max)
 
 ptx_reg_t f2x( ptx_reg_t x, unsigned from_width, unsigned to_width, int to_sign, int rounding_mode, int saturation_mode )
 {
- half mytemp;
- float myfloat;
-  half_float::half tmp_h;
+   half mytemp;
+   half_float::half tmp_h;
    //assert( from_width == 32); 
 
    enum cudaRoundMode mode = cudaRoundZero;
@@ -3085,7 +3076,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst )
    size_t size;
    unsigned smid;
    int t;
-   int thrd,odd,inx,k;
+   int thrd, k;
    ptx_thread_info *thread;
    
    const operand_info &src  = pI->operand_lookup(1);
@@ -3105,15 +3096,13 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst )
    _memory_op_t insn_memory_op = pI->has_memory_read() ? memory_load : memory_store;
    for (thrd=0; thrd < core->get_warp_size(); thrd++) {
 	thread = core->get_thread_info()[tid+thrd];
-	odd=thrd%2;
-	inx=thrd/2;
-    	ptx_reg_t addr_reg = thread->get_operand_value(src1, src, type, thread, 1);
+    ptx_reg_t addr_reg = thread->get_operand_value(src1, src, type, thread, 1);
    	ptx_reg_t src2_data = thread->get_operand_value(src2, src, type, thread, 1);
 	const operand_info &src_a=  pI->operand_lookup(1);
         unsigned nelem = src_a.get_vect_nelem();
         ptx_reg_t* v= new ptx_reg_t[8]; 
        	thread->get_vector_operand_values( src_a, v, nelem );
-  	stride=src2_data.u32;
+  	stride = src2_data.u32;
  
    	memory_space_t space = pI->get_space();
 
@@ -3130,9 +3119,9 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst )
 	}
    	decode_space(space,thread,src1,mem,addr);
 
-   	type_info_key::type_decode(type,size,t);
+   	type_info_key::type_decode(type, size, t);
    	if(core->get_gpu()->gpgpu_ctx->debug_tensorcore)
-		printf("mma_st: thrd=%d,addr=%x, fp(size=%d), stride=%d\n",thrd,addr_reg.u32,size,src2_data.u32);
+		printf("mma_st: thrd=%d, addr=%x, fp(size=%zu), stride=%d\n", thrd, addr_reg.u32, size, src2_data.u32);
 	addr_t new_addr = addr+thread_group_offset(thrd,wmma_type,wmma_layout,type,stride)*size/8;  
 	addr_t push_addr;
 
@@ -3152,7 +3141,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst )
 			mem_txn_addr[num_mem_txn++]=push_addr;
 	
 			if(core->get_gpu()->gpgpu_ctx->debug_tensorcore){
-				printf("wmma:store:thread%d=%x,%x,%x,%x,%x,%x,%x,%x\n",thrd,v[0].s64,v[1].s64,v[2].s64,v[3].s64,v[4].s64,v[5].s64,v[6].s64,v[7].s64);   
+				printf("wmma:store:thread%d=%llx,%llx,%llx,%llx,%llx,%llx,%llx,%llx\n",thrd,v[0].s64,v[1].s64,v[2].s64,v[3].s64,v[4].s64,v[5].s64,v[6].s64,v[7].s64);   
 				float temp;
 				int l;
 				printf("thread=%d:",thrd);
@@ -3179,7 +3168,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst )
 			}
 	
 			if(core->get_gpu()->gpgpu_ctx->debug_tensorcore)
-				printf("wmma:store:thread%d=%x,%x,%x,%x,%x,%x,%x,%x\n",thrd,nw_v[0].s64,nw_v[1].s64,nw_v[2].s64,nw_v[3].s64,nw_v[4].s64,nw_v[5].s64,nw_v[6].s64,nw_v[7].s64);   
+				printf("wmma:store:thread%d=%llx,%llx,%llx,%llx,%llx,%llx,%llx,%llx\n",thrd,nw_v[0].s64,nw_v[1].s64,nw_v[2].s64,nw_v[3].s64,nw_v[4].s64,nw_v[5].s64,nw_v[6].s64,nw_v[7].s64);   
 		}
 	}
    	
@@ -3238,11 +3227,11 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst )
 	}
 
 	decode_space(space,thread,src1,mem,addr);
-   	type_info_key::type_decode(type,size,t);
+   	type_info_key::type_decode(type, size, t);
 	
 	ptx_reg_t data[16];
    	if(core->get_gpu()->gpgpu_ctx->debug_tensorcore)	
-		printf("mma_ld: thrd=%d,addr=%x, fpsize=%d, stride=%d\n",thrd,src1_data.u32,size,src2_data.u32);
+		printf("mma_ld: thrd=%d,addr=%x, fpsize=%zu, stride=%d\n", thrd, src1_data.u32, size, src2_data.u32);
 	
 	addr_t new_addr = addr+thread_group_offset(thrd,wmma_type,wmma_layout,type,stride)*size/8;  
 	addr_t fetch_addr;
@@ -3341,7 +3330,7 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst )
 		if(type==F16_TYPE){
 			printf("\nmma_ld:thread%d= ",thrd);
 			for(i=0;i<16;i++){
-				printf("%x ",data[i].u64);
+				printf("%llx ",data[i].u64);
 			}
 			printf("\n");
 			
@@ -3361,7 +3350,7 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst )
 			printf("\n");
 			printf("\nmma_ld:thread%d= ",thrd);
 			for(i=0;i<8;i++){
-				printf("%x ",data[i].u64);
+				printf("%llx ",data[i].u64);
 			}
 			printf("\n");
 		}
@@ -3388,15 +3377,15 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst )
 		else
    			thread->set_wmma_vector_operand_values(dst,nw_data[0],nw_data[1],nw_data[2],nw_data[3],nw_data[4],nw_data[5],nw_data[6],nw_data[7]);
 		if(core->get_gpu()->gpgpu_ctx->debug_tensorcore){	
-			printf("mma_ld:data[0].s64=%x,data[1].s64=%x,new_data[0].s64=%x\n",data[0].u64,data[1].u64,nw_data[0].u64);	
-			printf("mma_ld:data[2].s64=%x,data[3].s64=%x,new_data[1].s64=%x\n",data[2].u64,data[3].u64,nw_data[1].u64);	
-			printf("mma_ld:data[4].s64=%x,data[5].s64=%x,new_data[2].s64=%x\n",data[4].u64,data[5].u64,nw_data[2].u64);	
-			printf("mma_ld:data[6].s64=%x,data[7].s64=%x,new_data[3].s64=%x\n",data[6].u64,data[7].u64,nw_data[3].u64);	
+			printf("mma_ld:data[0].s64=%llx,data[1].s64=%llx,new_data[0].s64=%llx\n",data[0].u64,data[1].u64,nw_data[0].u64);	
+			printf("mma_ld:data[2].s64=%llx,data[3].s64=%llx,new_data[1].s64=%llx\n",data[2].u64,data[3].u64,nw_data[1].u64);	
+			printf("mma_ld:data[4].s64=%llx,data[5].s64=%llx,new_data[2].s64=%llx\n",data[4].u64,data[5].u64,nw_data[2].u64);	
+			printf("mma_ld:data[6].s64=%llx,data[7].s64=%llx,new_data[3].s64=%llx\n",data[6].u64,data[7].u64,nw_data[3].u64);	
 			if(wmma_type!=LOAD_C){
-			printf("mma_ld:data[8].s64=%x,data[9].s64=%x,new_data[4].s64=%x\n",data[8].u64,data[9].u64,nw_data[4].s64);	
-			printf("mma_ld:data[10].s64=%x,data[11].s64=%x,new_data[5].s64=%x\n",data[10].u64,data[11].u64,nw_data[5].u64);	
-			printf("mma_ld:data[12].s64=%x,data[13].s64=%x,new_data[6].s64=%x\n",data[12].u64,data[13].u64,nw_data[6].u64);	
-			printf("mma_ld:data[14].s64=%x,data[15].s64=%x,new_data[7].s64=%x\n",data[14].u64,data[15].u64,nw_data[3].u64);	
+			printf("mma_ld:data[8].s64=%llx,data[9].s64=%llx,new_data[4].s64=%llx\n",data[8].u64,data[9].u64,nw_data[4].s64);	
+			printf("mma_ld:data[10].s64=%llx,data[11].s64=%llx,new_data[5].s64=%llx\n",data[10].u64,data[11].u64,nw_data[5].u64);	
+			printf("mma_ld:data[12].s64=%llx,data[13].s64=%llx,new_data[6].s64=%llx\n",data[12].u64,data[13].u64,nw_data[6].u64);	
+			printf("mma_ld:data[14].s64=%llx,data[15].s64=%llx,new_data[7].s64=%llx\n",data[14].u64,data[15].u64,nw_data[3].u64);	
 			}
 		}
 	}
@@ -4132,9 +4121,9 @@ int prmt_mode_present(int mode)
 	}
 	return returnval;
 }
-int read_byte(int mode,int control,int d_sel_index,signed long long value){
+int read_byte(int mode, int control, int d_sel_index, signed long long value){
 
-	int returnval;
+	int returnval = 0;
 	int prmt_f4e_mode[4][4]={{0,1,2,3},{1,2,3,4},{2,3,4,5},{3,4,5,6}};
 	int prmt_b4e_mode[4][4]={{0,7,6,5},{1,0,7,6},{2,1,0,7},{3,2,1,0}};
 	int prmt_rc8_mode[4][4]={{0,0,0,0},{1,1,1,1},{2,2,2,2},{3,3,3,3}};
@@ -4157,11 +4146,12 @@ int read_byte(int mode,int control,int d_sel_index,signed long long value){
 			case PRMT_RC8_MODE:	returnval=prmt_rc8_mode[control][d_sel_index];break;
 			case PRMT_ECL_MODE:	returnval=prmt_ecl_mode[control][d_sel_index];break;
 			case PRMT_ECR_MODE:	returnval=prmt_ecr_mode[control][d_sel_index];break;
-			case PRMT_RC16_MODE:	returnval=prmt_rc16_mode[control][d_sel_index];break;
-			default: printf("ERROR\n");break;
+			case PRMT_RC16_MODE: returnval=prmt_rc16_mode[control][d_sel_index];break;
+            // Change the default from printing "ERROR" to just asserting
+			default: assert(false);
 		}
 	}	
-	return (returnval<<8*d_sel_index);
+	return (returnval << 8 * d_sel_index);
 }
 
 void prmt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { 
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index 6978cc1..d8943d2 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -1415,7 +1415,7 @@ unsigned function_info::print_insn( unsigned pc, FILE * fp ) const
    snprintf(command,1024,"c++filt -p %s",m_name.c_str());
    FILE *p = popen(command,"r");
    buffer[0]=0;
-   fgets(buffer, 1023, p);
+   assert(fgets(buffer, 1023, p) != NULL);
    // Remove trailing "\n" in buffer
    char *c;
    if ((c=strchr(buffer, '\n')) != NULL) *c = '\0';
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index dca3cec..33bcf45 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -214,7 +214,8 @@ void fix_duplicate_errors(char fname2[1024]) {
 	long filesize = ftell(ptxsource);
 	rewind(ptxsource);
 	char *ptxdata = (char*)malloc((filesize+1)*sizeof(char));
-	fread(ptxdata, filesize, 1, ptxsource);
+    // Fail if we do not read the file
+	assert(fread(ptxdata, filesize, 1, ptxsource) == 1);
 	fclose(ptxsource);
 
 	FILE *ptxdest = fopen(fname2,"w");
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 81b70af..a4f4a0c 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -421,7 +421,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi
          assert( (num_bits%8) == 0  );
          addr = g_current_symbol_table->get_sstarr_next();
          addr_pad = pad_address(addr, num_bits/8, 128);
-         printf("from 0x%x to 0x%lx (sstarr memory space)\n",
+         printf("from 0x%llx to 0x%llx (sstarr memory space)\n",
                  addr+addr_pad,
                  addr+addr_pad + num_bits/8);
             fflush(stdout);
diff --git a/src/gpgpu-sim/addrdec.cc b/src/gpgpu-sim/addrdec.cc
index 3262456..f225933 100644
--- a/src/gpgpu-sim/addrdec.cc
+++ b/src/gpgpu-sim/addrdec.cc
@@ -184,7 +184,7 @@ void linear_to_raw_address_translation::addrdec_tlx(new_addr_type addr, addrdec_
 			  }
 
 			  assert(tlx->chip < m_n_channel);
-			  assert(tlx->sub_partition < m_n_channel*m_n_sub_partition_in_channel);
+			  assert(tlx->sub_partition < m_n_channel * m_n_sub_partition_in_channel);
 			  return;
 			break;
 		}
diff --git a/src/gpgpu-sim/addrdec.h b/src/gpgpu-sim/addrdec.h
index a5333fb..c9a1420 100644
--- a/src/gpgpu-sim/addrdec.h
+++ b/src/gpgpu-sim/addrdec.h
@@ -92,7 +92,7 @@ private:
    new_addr_type sub_partition_id_mask; 
 
    unsigned int gap;
-   int m_n_channel;
+   unsigned m_n_channel;
    int m_n_sub_partition_in_channel; 
 
 };
diff --git a/src/gpgpu-sim/dram.cc b/src/gpgpu-sim/dram.cc
index d443d79..9c33822 100644
--- a/src/gpgpu-sim/dram.cc
+++ b/src/gpgpu-sim/dram.cc
@@ -482,7 +482,6 @@ void dram_t::cycle()
 
         bool memory_pending_rw_found=false;
         for (unsigned j=0;j<m_config->nbk;j++) {
-     	   unsigned grp = get_bankgrp_number(j);
      	   if (bk[j]->mrq && (((bk[j]->curr_row == bk[j]->mrq->row) &&
      		  (bk[j]->mrq->rw == READ)  &&
      		  (bk[j]->state == BANK_ACTIVE))
@@ -817,10 +816,10 @@ void dram_t::visualize() const
 
 void dram_t::print_stat( FILE* simFile ) 
 {
-   fprintf(simFile,"DRAM (%llu): n_cmd=%llu n_nop=%llu n_act=%llu n_pre=%llu n_ref=%llu n_req=%llu n_rd=%llu n_write=%llu bw_util=%.4g ",
+   fprintf(simFile,"DRAM (%u): n_cmd=%llu n_nop=%llu n_act=%llu n_pre=%llu n_ref=%llu n_req=%llu n_rd=%llu n_write=%llu bw_util=%.4g ",
            id, n_cmd, n_nop, n_act, n_pre, n_ref, n_req, n_rd, n_wr,
            (float)bwutil/n_cmd);
-   fprintf(simFile, "mrqq: %d %.4g mrqsmax=%d ", max_mrqs, (float)ave_mrqs/n_cmd, max_mrqs_temp);
+   fprintf(simFile, "mrqq: %d %.4g mrqsmax=%llu ", max_mrqs, (float)ave_mrqs/n_cmd, max_mrqs_temp);
    fprintf(simFile, "\n");
    fprintf(simFile, "dram_util_bins:");
    for (unsigned i=0;i<10;i++) fprintf(simFile, " %d", dram_util_bins[i]);
@@ -899,10 +898,10 @@ void dram_t::set_dram_power_stats(	unsigned &cmd,
 unsigned dram_t::get_bankgrp_number(unsigned i)
 {
 	if(m_config->dram_bnkgrp_indexing_policy == HIGHER_BITS) { //higher bits
-		return i>>m_config->bk_tag_length;
+		return i >> m_config->bk_tag_length;
 	}
 	else if (m_config->dram_bnkgrp_indexing_policy == LOWER_BITS) { //lower bits
-		return i&((m_config->nbkgrp-1));
+		return i & ((m_config->nbkgrp - 1));
 	}
 	else {
 		assert(1);
diff --git a/src/gpgpu-sim/gpu-cache.cc b/src/gpgpu-sim/gpu-cache.cc
index b7d0ac3..65f2f6d 100644
--- a/src/gpgpu-sim/gpu-cache.cc
+++ b/src/gpgpu-sim/gpu-cache.cc
@@ -790,7 +790,7 @@ void cache_stats::print_stats(FILE *fout, const char *cache_name) const{
     }
     for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) {
     	 if(total_access[type] > 0)
-    	  fprintf(fout, "\t%s[%s][%s] = %llu\n",
+    	  fprintf(fout, "\t%s[%s][%s] = %u\n",
 				m_cache_name.c_str(),
 				mem_access_type_str((enum mem_access_type)type),
 				"TOTAL_ACCESS",
@@ -803,7 +803,7 @@ void cache_stats::print_fail_stats(FILE *fout, const char *cache_name) const{
 	    for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) {
 	        for (unsigned fail = 0; fail < NUM_CACHE_RESERVATION_FAIL_STATUS; ++fail) {
 	            if(m_fail_stats[type][fail] > 0){
-	                fprintf(fout, "\t%s[%s][%s] = %u\n",
+	                fprintf(fout, "\t%s[%s][%s] = %llu\n",
 	                    m_cache_name.c_str(),
 	                    mem_access_type_str((enum mem_access_type)type),
 						cache_fail_status_str((enum cache_reservation_fail_reason)fail),
@@ -1430,8 +1430,6 @@ data_cache::wr_miss_wa_lazy_fetch_on_read( new_addr_type addr,
 {
 
 	    new_addr_type block_addr = m_config.block_addr(addr);
-	    new_addr_type mshr_addr = m_config.mshr_addr(mf->get_addr());
-
 
 		//if the request writes to the whole cache line/sector, then, write and set cache line Modified.
 		//and no need to send read request to memory or reserve mshr
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index a151379..ce5c4cb 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -475,10 +475,10 @@ void gpgpu_sim_config::reg_options(option_parser_t opp)
     m_shader_config.reg_options(opp);
     m_memory_config.reg_options(opp);
     power_config::reg_options(opp);
-   option_parser_register(opp, "-gpgpu_max_cycle", OPT_INT32, &gpu_max_cycle_opt, 
+   option_parser_register(opp, "-gpgpu_max_cycle", OPT_INT64, &gpu_max_cycle_opt, 
                "terminates gpu simulation early (0 = no limit)",
                "0");
-   option_parser_register(opp, "-gpgpu_max_insn", OPT_INT32, &gpu_max_insn_opt, 
+   option_parser_register(opp, "-gpgpu_max_insn", OPT_INT64, &gpu_max_insn_opt, 
                "terminates gpu simulation early (0 = no limit)",
                "0");
    option_parser_register(opp, "-gpgpu_max_cta", OPT_INT32, &gpu_max_cta_opt, 
@@ -1117,8 +1117,8 @@ void gpgpu_sim::gpu_print_stat()
    printf("gpu_tot_sim_insn = %lld\n", gpu_tot_sim_insn+gpu_sim_insn);
    printf("gpu_tot_ipc = %12.4f\n", (float)(gpu_tot_sim_insn+gpu_sim_insn) / (gpu_tot_sim_cycle+gpu_sim_cycle));
    printf("gpu_tot_issued_cta = %lld\n", gpu_tot_issued_cta + m_total_cta_launched);
-   printf("gpu_occupancy = %.4f\% \n", gpu_occupancy.get_occ_fraction() * 100);
-   printf("gpu_tot_occupancy = %.4f\% \n", (gpu_occupancy + gpu_tot_occupancy).get_occ_fraction() * 100);
+   printf("gpu_occupancy = %.4f%% \n", gpu_occupancy.get_occ_fraction() * 100);
+   printf("gpu_tot_occupancy = %.4f%% \n", (gpu_occupancy + gpu_tot_occupancy).get_occ_fraction() * 100);
 
 
    fprintf(statfout, "max_total_param_size = %llu\n", gpgpu_ctx->device_runtime->g_max_total_param_size);
@@ -1144,7 +1144,7 @@ void gpgpu_sim::gpu_print_stat()
 
    time_t curr_time;
    time(&curr_time);
-   unsigned long long elapsed_time = MAX( curr_time - GPGPUsim_ctx_ptr()->g_simulation_starttime, 1 );
+   unsigned long long elapsed_time = MAX( curr_time - gpgpu_ctx->the_gpgpusim->g_simulation_starttime, 1 );
    printf( "gpu_total_sim_rate=%u\n", (unsigned)( ( gpu_tot_sim_insn + gpu_sim_insn ) / elapsed_time ) );
 
    //shader_print_l1_miss_stat( stdout );
@@ -1359,7 +1359,7 @@ bool shader_core_ctx::occupy_shader_resource_1block(kernel_info_t & k, bool occu
        m_occupied_regs += (padded_cta_size * ((kernel_info->regs+3)&~3));
        m_occupied_ctas++;
 
-      SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Occupied %d threads, %d shared mem, %d registers, %d ctas\n",
+      SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Occupied %u threads, %u shared mem, %u registers, %u ctas\n",
             m_occupied_n_threads, m_occupied_shmem, m_occupied_regs, m_occupied_ctas);  
    }
 
@@ -1476,7 +1476,7 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel )
         nthreads_in_block += ptx_sim_init_thread(kernel,&m_thread[i],m_sid,i,cta_size-(i-start_thread),m_config->n_thread_per_shader,this,free_cta_hw_id,warp_id,m_cluster->get_gpu());
         m_threadState[i].m_active = true; 
         // load thread local memory and register file
-        if(m_gpu->resume_option==1 && kernel.get_uid()==m_gpu->resume_kernel && ctaid>=m_gpu->resume_CTA && ctaid<m_gpu->checkpoint_CTA_t )
+        if(m_gpu->resume_option == 1 && kernel.get_uid() == m_gpu->resume_kernel && ctaid >= m_gpu->resume_CTA && ctaid < m_gpu->checkpoint_CTA_t )
         {
             char fname[2048];
             snprintf(fname,2048,"checkpoint_files/thread_%d_%d_reg.txt",i%cta_size,ctaid );
@@ -1491,7 +1491,7 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel )
     assert( nthreads_in_block > 0 && nthreads_in_block <= m_config->n_thread_per_shader); // should be at least one, but less than max
     m_cta_status[free_cta_hw_id]=nthreads_in_block;
 
-    if(m_gpu->resume_option==1 && kernel.get_uid()==m_gpu->resume_kernel && ctaid>=m_gpu->resume_CTA && ctaid<m_gpu->checkpoint_CTA_t )
+    if(m_gpu->resume_option == 1 && kernel.get_uid() == m_gpu->resume_kernel && ctaid >= m_gpu->resume_CTA && ctaid < m_gpu->checkpoint_CTA_t )
     {
         char f1name[2048];
         snprintf(f1name,2048,"checkpoint_files/shared_mem_%d.txt", ctaid);
@@ -1717,7 +1717,7 @@ void gpgpu_sim::cycle()
          time_t days, hrs, minutes, sec;
          time_t curr_time;
          time(&curr_time);
-         unsigned long long  elapsed_time = MAX(curr_time - GPGPUsim_ctx_ptr()->g_simulation_starttime, 1);
+         unsigned long long  elapsed_time = MAX(curr_time - gpgpu_ctx->the_gpgpusim->g_simulation_starttime, 1);
          if ( (elapsed_time - last_liveness_message_time) >= m_config.liveness_message_freq && DTRACE(LIVENESS) ) {
             days    = elapsed_time/(3600*24);
             hrs     = elapsed_time/3600 - 24*days;
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index 76c7a06..fba770d 100644
--- a/src/gpgpu-sim/gpu-sim.h
+++ b/src/gpgpu-sim/gpu-sim.h
@@ -368,8 +368,8 @@ private:
     double l2_period;
 
     // GPGPU-Sim timing model options
-    unsigned gpu_max_cycle_opt;
-    unsigned gpu_max_insn_opt;
+    unsigned long long gpu_max_cycle_opt;
+    unsigned long long gpu_max_insn_opt;
     unsigned gpu_max_cta_opt;
     char *gpgpu_runtime_stat;
     bool  gpgpu_flush_l1_cache;
diff --git a/src/gpgpu-sim/l2cache.cc b/src/gpgpu-sim/l2cache.cc
index 39a5812..fb4ce32 100644
--- a/src/gpgpu-sim/l2cache.cc
+++ b/src/gpgpu-sim/l2cache.cc
@@ -82,7 +82,7 @@ void memory_partition_unit::handle_memcpy_to_gpu( size_t addr, unsigned global_s
     unsigned p = global_sub_partition_id_to_local_id(global_subpart_id);
     std::string mystring =
         mask.to_string<char,std::string::traits_type,std::string::allocator_type>();
-    MEMPART_DPRINTF("Copy Engine Request Received For Address=%llx, local_subpart=%u, global_subpart=%u, sector_mask=%s \n", addr, p, global_subpart_id, mystring.c_str()); 
+    MEMPART_DPRINTF("Copy Engine Request Received For Address=%zx, local_subpart=%u, global_subpart=%u, sector_mask=%s \n", addr, p, global_subpart_id, mystring.c_str()); 
     m_sub_partition[p]->force_l2_tag_update(addr,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle, mask);
 }
 
@@ -681,7 +681,7 @@ std::vector<mem_fetch*> memory_sub_partition::breakdown_request_to_sector_reques
 			}
 		} else
 			{
-			    printf("Invalid sector received, address = 0x%06x, sector mask = %s, data size = %d",
+			    printf("Invalid sector received, address = 0x%06llx, sector mask = %s, data size = %d",
 			    		mf->get_addr(), mf->get_access_sector_mask(), mf->get_data_size());
 				assert(0 && "Undefined sector mask is received");
 			}
@@ -716,7 +716,7 @@ std::vector<mem_fetch*> memory_sub_partition::breakdown_request_to_sector_reques
 			 byte_sector_mask <<= SECTOR_SIZE;
 		}
 	} else {
-		 printf("Invalid sector received, address = 0x%06x, sector mask = %d, byte mask = , data size = %d",
+		 printf("Invalid sector received, address = 0x%06llx, sector mask = %d, byte mask = , data size = %u",
 					    		mf->get_addr(), mf->get_access_sector_mask().count(), mf->get_data_size());
 		 assert(0 && "Undefined data size is received");
 	}
diff --git a/src/gpgpu-sim/local_interconnect.cc b/src/gpgpu-sim/local_interconnect.cc
index da8a65c..c70477c 100644
--- a/src/gpgpu-sim/local_interconnect.cc
+++ b/src/gpgpu-sim/local_interconnect.cc
@@ -253,7 +253,7 @@ LocalInterconnect::LocalInterconnect(const struct inct_config& m_localinct_confi
 }
 
 LocalInterconnect::~LocalInterconnect(){
-	for (int i=0; i<m_inct_config.subnets; ++i) {
+	for (unsigned i = 0; i < m_inct_config.subnets; ++i) {
 		delete net[i];
 	}
 }
diff --git a/src/gpgpu-sim/scoreboard.cc b/src/gpgpu-sim/scoreboard.cc
index 80f95c6..1017e75 100644
--- a/src/gpgpu-sim/scoreboard.cc
+++ b/src/gpgpu-sim/scoreboard.cc
@@ -140,10 +140,10 @@ bool Scoreboard::checkCollision( unsigned wid, const class inst_t *inst ) const
 	// Get list of all input and output registers
 	std::set<int> inst_regs;
 
-	for(int iii=0;iii<inst->outcount;iii++)
+	for(unsigned iii=0; iii < inst->outcount; iii++)
 		inst_regs.insert(inst->out[iii]);
 
-	for(int jjj=0;jjj<inst->incount;jjj++)
+	for(unsigned jjj=0;jjj<inst->incount;jjj++)
 		inst_regs.insert(inst->in[jjj]);
 
         if(inst->pred > 0) inst_regs.insert(inst->pred);
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index 0514a77..86508f6 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -87,7 +87,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu,
                                   shader_core_stats *stats )
    : core_t( gpu, NULL, config->warp_size, config->n_thread_per_shader ),
      m_barriers( this, config->max_warps_per_shader, config->max_cta_per_core, config->max_barriers_per_cta, config->warp_size ),
-     m_dynamic_warp_id(0), m_active_warps(0)
+     m_active_warps(0), m_dynamic_warp_id(0)
 {
     m_cluster = cluster;
     m_config = config;
@@ -164,7 +164,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu,
                                          NUM_CONCRETE_SCHEDULERS;
     assert ( scheduler != NUM_CONCRETE_SCHEDULERS );
     
-    for (int i = 0; i < m_config->gpgpu_num_sched_per_core; i++) {
+    for (unsigned i = 0; i < m_config->gpgpu_num_sched_per_core; i++) {
         switch( scheduler )
         {
             case CONCRETE_SCHEDULER_LRR:
@@ -263,7 +263,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu,
         //distribute i's evenly though schedulers;
         schedulers[i%m_config->gpgpu_num_sched_per_core]->add_supervised_warp_id(i);
     }
-    for ( int i = 0; i < m_config->gpgpu_num_sched_per_core; ++i ) {
+    for ( unsigned i = 0; i < m_config->gpgpu_num_sched_per_core; ++i ) {
         schedulers[i]->done_adding_supervised_warps();
     }
     
@@ -474,7 +474,7 @@ void shader_core_ctx::init_warps( unsigned cta_id, unsigned start_thread, unsign
             }
             m_simt_stack[i]->launch(start_pc,active_threads);
 
-              if(m_gpu->resume_option==1 && kernel_id==m_gpu->resume_kernel && ctaid>=m_gpu->resume_CTA && ctaid<m_gpu->checkpoint_CTA_t )
+              if(m_gpu->resume_option == 1 && kernel_id == m_gpu->resume_kernel && ctaid >= m_gpu->resume_CTA && ctaid < m_gpu->checkpoint_CTA_t )
                { 
                 char fname[2048];
                 snprintf(fname,2048,"checkpoint_files/warp_%d_%d_simt.txt",i%warp_per_cta,ctaid );
@@ -868,7 +868,7 @@ void shader_core_ctx::func_exec_inst( warp_inst_t &inst )
 
 void shader_core_ctx::issue_warp( register_set& pipe_reg_set, const warp_inst_t* next_inst, const active_mask_t &active_mask, unsigned warp_id, unsigned sch_id )
 {
-	warp_inst_t** pipe_reg = pipe_reg = pipe_reg_set.get_free(m_config->sub_core_model, sch_id);
+	warp_inst_t** pipe_reg = pipe_reg_set.get_free(m_config->sub_core_model, sch_id);
     assert(pipe_reg);
 
     m_warp[warp_id].ibuffer_free();
@@ -2147,7 +2147,7 @@ ldst_unit::ldst_unit( mem_fetch_interface *icnt,
         l1_latency_queue.resize(m_config->m_L1D_config.l1_banks);
         assert(m_config->m_L1D_config.l1_latency > 0);
 
-		for(int j=0; j<m_config->m_L1D_config.l1_banks; j++ )
+		for(unsigned j = 0; j < m_config->m_L1D_config.l1_banks; j++ )
 			l1_latency_queue[j].resize(m_config->m_L1D_config.l1_latency,(mem_fetch*)NULL);
 
     }
@@ -2459,7 +2459,7 @@ void shader_core_ctx::register_cta_thread_exit( unsigned cta_num, kernel_info_t
       m_barriers.deallocate_barrier(cta_num);
       shader_CTA_count_unlog(m_sid, 1);
 
-     SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Finished CTA #%d (%lld,%lld), %u CTAs running\n",
+     SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Finished CTA #%u (%lld,%lld), %u CTAs running\n",
         cta_num, m_gpu->gpu_sim_cycle, m_gpu->gpu_tot_sim_cycle, m_n_active_cta);
 
       if( m_n_active_cta == 0 ) {
diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h
index 62e0e42..667cb2d 100644
--- a/src/gpgpu-sim/shader.h
+++ b/src/gpgpu-sim/shader.h
@@ -1392,9 +1392,9 @@ class shader_core_config : public core_config
 	 	If we won't remove it, old regression will be broken.
 		So to support the legacy config files it's best to handle in this way.
          */  
-	int num_config_to_read=N_PIPELINE_STAGES-2*(!gpgpu_tensor_core_avail);
+	int num_config_to_read= N_PIPELINE_STAGES - 2 * (!gpgpu_tensor_core_avail);
 
-        for (unsigned i = 0; i <num_config_to_read; i++) { 
+        for (int i = 0; i < num_config_to_read; i++) { 
 	    assert(toks);
 	    ntok = sscanf(toks,"%d", &pipe_widths[i]);
 	    assert(ntok == 1); 
@@ -1455,7 +1455,7 @@ class shader_core_config : public core_config
 
     bool gpgpu_dwf_reg_bankconflict;
 
-    int gpgpu_num_sched_per_core;
+    unsigned gpgpu_num_sched_per_core;
     int gpgpu_max_insn_issue_per_warp;
     bool gpgpu_dual_issue_diff_exec_units;
 
diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc
index 816159f..846773d 100644
--- a/src/gpgpusim_entrypoint.cc
+++ b/src/gpgpusim_entrypoint.cc
@@ -43,46 +43,28 @@ static int sg_argc = 3;
 static const char *sg_argv[] = {"", "-config","gpgpusim.config"};
 
 
-GPGPUsim_ctx* the_gpgpusim =  NULL;
-
-GPGPUsim_ctx* GPGPUsim_ctx_ptr(){
-	if(the_gpgpusim == NULL)
-		the_gpgpusim = GPGPU_Context()->the_gpgpusim;
-
-	return the_gpgpusim;
-}
-
-class gpgpu_sim* g_the_gpu() {
-	return GPGPUsim_ctx_ptr()->g_the_gpu;
-}
-
-class stream_manager* g_stream_manager()  {
-	return GPGPUsim_ctx_ptr()->g_stream_manager;
-}
-
-static void print_simulation_time();
-
-void *gpgpu_sim_thread_sequential(void*)
+void * gpgpu_sim_thread_sequential(void * ctx_ptr)
 {
+    gpgpu_context * ctx = (gpgpu_context *)ctx_ptr;
    // at most one kernel running at a time
    bool done;
    do {
-      sem_wait(&(GPGPUsim_ctx_ptr()->g_sim_signal_start));
+      sem_wait(&(ctx->the_gpgpusim->g_sim_signal_start));
       done = true;
-      if( GPGPUsim_ctx_ptr()->g_the_gpu->get_more_cta_left() ) {
+      if( ctx->the_gpgpusim->g_the_gpu->get_more_cta_left() ) {
           done = false;
-          GPGPUsim_ctx_ptr()->g_the_gpu->init();
-          while( GPGPUsim_ctx_ptr()->g_the_gpu->active() ) {
-        	  GPGPUsim_ctx_ptr()->g_the_gpu->cycle();
-        	  GPGPUsim_ctx_ptr()->g_the_gpu->deadlock_check();
+          ctx->the_gpgpusim->g_the_gpu->init();
+          while( ctx->the_gpgpusim->g_the_gpu->active() ) {
+        	  ctx->the_gpgpusim->g_the_gpu->cycle();
+        	  ctx->the_gpgpusim->g_the_gpu->deadlock_check();
           }
-          GPGPUsim_ctx_ptr()->g_the_gpu->print_stats();
-          GPGPUsim_ctx_ptr()->g_the_gpu->update_stats();
-          print_simulation_time();
+          ctx->the_gpgpusim->g_the_gpu->print_stats();
+          ctx->the_gpgpusim->g_the_gpu->update_stats();
+          ctx->print_simulation_time();
       }
-      sem_post(&(GPGPUsim_ctx_ptr()->g_sim_signal_finish));
+      sem_post(&(ctx->the_gpgpusim->g_sim_signal_finish));
    } while(!done);
-   sem_post(&(GPGPUsim_ctx_ptr()->g_sim_signal_exit));
+   sem_post(&(ctx->the_gpgpusim->g_sim_signal_exit));
    return NULL;
 }
 
@@ -94,8 +76,9 @@ static void termination_callback()
     fflush(stdout);
 }
 
-void *gpgpu_sim_thread_concurrent(void*)
+void *gpgpu_sim_thread_concurrent(void * ctx_ptr)
 {
+    gpgpu_context * ctx = (gpgpu_context *)ctx_ptr;
     atexit(termination_callback);
     // concurrent kernel execution simulation thread
     do {
@@ -103,19 +86,19 @@ void *gpgpu_sim_thread_concurrent(void*)
            printf("GPGPU-Sim: *** simulation thread starting and spinning waiting for work ***\n");
            fflush(stdout);
         }
-        while( GPGPUsim_ctx_ptr()->g_stream_manager->empty_protected() && !GPGPUsim_ctx_ptr()->g_sim_done )
+        while( ctx->the_gpgpusim->g_stream_manager->empty_protected() && !ctx->the_gpgpusim->g_sim_done )
             ;
         if(g_debug_execution >= 3) {
            printf("GPGPU-Sim: ** START simulation thread (detected work) **\n");
-           GPGPUsim_ctx_ptr()->g_stream_manager->print(stdout);
+           ctx->the_gpgpusim->g_stream_manager->print(stdout);
            fflush(stdout);
         }
-        pthread_mutex_lock(&(GPGPUsim_ctx_ptr()->g_sim_lock));
-        GPGPUsim_ctx_ptr()->g_sim_active = true;
-        pthread_mutex_unlock(&(GPGPUsim_ctx_ptr()->g_sim_lock));
+        pthread_mutex_lock(&(ctx->the_gpgpusim->g_sim_lock));
+        ctx->the_gpgpusim->g_sim_active = true;
+        pthread_mutex_unlock(&(ctx->the_gpgpusim->g_sim_lock));
         bool active = false;
         bool sim_cycles = false;
-        GPGPUsim_ctx_ptr()->g_the_gpu->init();
+        ctx->the_gpgpusim->g_the_gpu->init();
         do {
             // check if a kernel has completed
             // launch operation on device if one is pending and can be run
@@ -127,82 +110,82 @@ void *gpgpu_sim_thread_concurrent(void*)
             // another kernel, the gpu is not re-initialized and the inter-kernel
             // behaviour may be incorrect. Check that a kernel has finished and
             // no other kernel is currently running.
-            if(GPGPUsim_ctx_ptr()->g_stream_manager->operation(&sim_cycles) && !GPGPUsim_ctx_ptr()->g_the_gpu->active())
+            if(ctx->the_gpgpusim->g_stream_manager->operation(&sim_cycles) && !ctx->the_gpgpusim->g_the_gpu->active())
                 break;
 
             //functional simulation
-            if( GPGPUsim_ctx_ptr()->g_the_gpu->is_functional_sim()) {
-                kernel_info_t * kernel = GPGPUsim_ctx_ptr()->g_the_gpu->get_functional_kernel();
+            if( ctx->the_gpgpusim->g_the_gpu->is_functional_sim()) {
+                kernel_info_t * kernel = ctx->the_gpgpusim->g_the_gpu->get_functional_kernel();
                 assert(kernel);
-                GPGPUsim_ctx_ptr()->gpgpu_ctx->func_sim->gpgpu_cuda_ptx_sim_main_func(*kernel);
-                GPGPUsim_ctx_ptr()->g_the_gpu->finish_functional_sim(kernel);
+                ctx->the_gpgpusim->gpgpu_ctx->func_sim->gpgpu_cuda_ptx_sim_main_func(*kernel);
+                ctx->the_gpgpusim->g_the_gpu->finish_functional_sim(kernel);
             }
 
             //performance simulation
-            if( GPGPUsim_ctx_ptr()->g_the_gpu->active() ) {
-            	GPGPUsim_ctx_ptr()->g_the_gpu->cycle();
+            if( ctx->the_gpgpusim->g_the_gpu->active() ) {
+            	ctx->the_gpgpusim->g_the_gpu->cycle();
             	sim_cycles = true;
-            	GPGPUsim_ctx_ptr()->g_the_gpu->deadlock_check();
+            	ctx->the_gpgpusim->g_the_gpu->deadlock_check();
             }else {
-                if(GPGPUsim_ctx_ptr()->g_the_gpu->cycle_insn_cta_max_hit()){
-                	GPGPUsim_ctx_ptr()->g_stream_manager->stop_all_running_kernels();
-                	GPGPUsim_ctx_ptr()->g_sim_done = true;
-                	GPGPUsim_ctx_ptr()->break_limit = true;
+                if(ctx->the_gpgpusim->g_the_gpu->cycle_insn_cta_max_hit()){
+                	ctx->the_gpgpusim->g_stream_manager->stop_all_running_kernels();
+                	ctx->the_gpgpusim->g_sim_done = true;
+                	ctx->the_gpgpusim->break_limit = true;
                 }
             }
 
-            active=GPGPUsim_ctx_ptr()->g_the_gpu->active() || !(GPGPUsim_ctx_ptr()->g_stream_manager->empty_protected());
+            active=ctx->the_gpgpusim->g_the_gpu->active() || !(ctx->the_gpgpusim->g_stream_manager->empty_protected());
 
-        } while( active && !GPGPUsim_ctx_ptr()->g_sim_done);
+        } while( active && !ctx->the_gpgpusim->g_sim_done);
         if(g_debug_execution >= 3) {
            printf("GPGPU-Sim: ** STOP simulation thread (no work) **\n");
            fflush(stdout);
         }
         if(sim_cycles) {
-        	GPGPUsim_ctx_ptr()->g_the_gpu->print_stats();
-        	GPGPUsim_ctx_ptr()->g_the_gpu->update_stats();
-            print_simulation_time();
+        	ctx->the_gpgpusim->g_the_gpu->print_stats();
+        	ctx->the_gpgpusim->g_the_gpu->update_stats();
+            ctx->print_simulation_time();
         }
-        pthread_mutex_lock(&(GPGPUsim_ctx_ptr()->g_sim_lock));
-        GPGPUsim_ctx_ptr()->g_sim_active = false;
-        pthread_mutex_unlock(&(GPGPUsim_ctx_ptr()->g_sim_lock));
-    } while( !GPGPUsim_ctx_ptr()->g_sim_done );
+        pthread_mutex_lock(&(ctx->the_gpgpusim->g_sim_lock));
+        ctx->the_gpgpusim->g_sim_active = false;
+        pthread_mutex_unlock(&(ctx->the_gpgpusim->g_sim_lock));
+    } while( !ctx->the_gpgpusim->g_sim_done );
 
     printf("GPGPU-Sim: *** simulation thread exiting ***\n");
     fflush(stdout);
 
-    if(GPGPUsim_ctx_ptr()->break_limit) {
+    if(ctx->the_gpgpusim->break_limit) {
     	printf("GPGPU-Sim: ** break due to reaching the maximum cycles (or instructions) **\n");
     	exit(1);
     }
 
-    sem_post(&(GPGPUsim_ctx_ptr()->g_sim_signal_exit));
+    sem_post(&(ctx->the_gpgpusim->g_sim_signal_exit));
     return NULL;
 }
 
-void synchronize()
+void gpgpu_context::synchronize()
 {
     printf("GPGPU-Sim: synchronize waiting for inactive GPU simulation\n");
-    GPGPUsim_ctx_ptr()->g_stream_manager->print(stdout);
+    the_gpgpusim->g_stream_manager->print(stdout);
     fflush(stdout);
 //    sem_wait(&g_sim_signal_finish);
     bool done = false;
     do {
-        pthread_mutex_lock(&(GPGPUsim_ctx_ptr()->g_sim_lock));
-        done = ( GPGPUsim_ctx_ptr()->g_stream_manager->empty() && !GPGPUsim_ctx_ptr()->g_sim_active ) || GPGPUsim_ctx_ptr()->g_sim_done;
-        pthread_mutex_unlock(&(GPGPUsim_ctx_ptr()->g_sim_lock));
+        pthread_mutex_lock(&(the_gpgpusim->g_sim_lock));
+        done = ( the_gpgpusim->g_stream_manager->empty() && !the_gpgpusim->g_sim_active ) || the_gpgpusim->g_sim_done;
+        pthread_mutex_unlock(&(the_gpgpusim->g_sim_lock));
     } while (!done);
     printf("GPGPU-Sim: detected inactive GPU simulation thread\n");
     fflush(stdout);
 //    sem_post(&g_sim_signal_start);
 }
 
-void exit_simulation()
+void gpgpu_context::exit_simulation()
 {
-	GPGPUsim_ctx_ptr()->g_sim_done=true;
+	the_gpgpusim->g_sim_done=true;
     printf("GPGPU-Sim: exit_simulation called\n");
     fflush(stdout);
-    sem_wait(&(GPGPUsim_ctx_ptr()->g_sim_signal_exit));
+    sem_wait(&(the_gpgpusim->g_sim_signal_exit));
     printf("GPGPU-Sim: simulation thread signaled exit\n");
     fflush(stdout);
 }
@@ -219,8 +202,8 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf()
    func_sim->ptx_opcocde_latency_options(opp);
 
    icnt_reg_options(opp);
-   GPGPUsim_ctx_ptr()->g_the_gpu_config = new gpgpu_sim_config(this);
-   GPGPUsim_ctx_ptr()->g_the_gpu_config->reg_options(opp); // register GPU microrachitecture options
+   the_gpgpusim->g_the_gpu_config = new gpgpu_sim_config(this);
+   the_gpgpusim->g_the_gpu_config->reg_options(opp); // register GPU microrachitecture options
 
    option_parser_cmdline(opp, sg_argc, sg_argv); // parse configuration options
    fprintf(stdout, "GPGPU-Sim: Configuration options:\n\n");
@@ -228,37 +211,37 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf()
    // Set the Numeric locale to a standard locale where a decimal point is a "dot" not a "comma"
    // so it does the parsing correctly independent of the system environment variables
    assert(setlocale(LC_NUMERIC,"C"));
-   GPGPUsim_ctx_ptr()->g_the_gpu_config->init();
+   the_gpgpusim->g_the_gpu_config->init();
 
-   GPGPUsim_ctx_ptr()->g_the_gpu = new gpgpu_sim(*(GPGPUsim_ctx_ptr()->g_the_gpu_config), this);
-   GPGPUsim_ctx_ptr()->g_stream_manager = new stream_manager((GPGPUsim_ctx_ptr()->g_the_gpu), func_sim->g_cuda_launch_blocking);
+   the_gpgpusim->g_the_gpu = new gpgpu_sim(*(the_gpgpusim->g_the_gpu_config), this);
+   the_gpgpusim->g_stream_manager = new stream_manager((the_gpgpusim->g_the_gpu), func_sim->g_cuda_launch_blocking);
 
-   GPGPUsim_ctx_ptr()->g_simulation_starttime = time((time_t *)NULL);
+   the_gpgpusim->g_simulation_starttime = time((time_t *)NULL);
 
-   sem_init(&(GPGPUsim_ctx_ptr()->g_sim_signal_start),0,0);
-   sem_init(&(GPGPUsim_ctx_ptr()->g_sim_signal_finish),0,0);
-   sem_init(&(GPGPUsim_ctx_ptr()->g_sim_signal_exit),0,0);
+   sem_init(&(the_gpgpusim->g_sim_signal_start),0,0);
+   sem_init(&(the_gpgpusim->g_sim_signal_finish),0,0);
+   sem_init(&(the_gpgpusim->g_sim_signal_exit),0,0);
 
-   return GPGPUsim_ctx_ptr()->g_the_gpu;
+   return the_gpgpusim->g_the_gpu;
 }
 
-void start_sim_thread(int api)
+void gpgpu_context::start_sim_thread(int api)
 {
-    if( GPGPUsim_ctx_ptr()->g_sim_done ) {
-    	GPGPUsim_ctx_ptr()->g_sim_done = false;
+    if( the_gpgpusim->g_sim_done ) {
+    	the_gpgpusim->g_sim_done = false;
         if( api == 1 ) {
-           pthread_create(&(GPGPUsim_ctx_ptr()->g_simulation_thread),NULL,gpgpu_sim_thread_concurrent,NULL);
+           pthread_create(&(the_gpgpusim->g_simulation_thread),NULL,gpgpu_sim_thread_concurrent,(void *)this);
         } else {
-           pthread_create(&(GPGPUsim_ctx_ptr()->g_simulation_thread),NULL,gpgpu_sim_thread_sequential,NULL);
+           pthread_create(&(the_gpgpusim->g_simulation_thread),NULL,gpgpu_sim_thread_sequential,(void *)this);
         }
     }
 }
 
-void print_simulation_time()
+void gpgpu_context::print_simulation_time()
 {
    time_t current_time, difference, d, h, m, s;
    current_time = time((time_t *)NULL);
-   difference = MAX(current_time - GPGPUsim_ctx_ptr()->g_simulation_starttime, 1);
+   difference = MAX(current_time - the_gpgpusim->g_simulation_starttime, 1);
 
    d = difference/(3600*24);
    h = difference/3600 - 24*d;
@@ -268,16 +251,18 @@ void print_simulation_time()
    fflush(stderr);
    printf("\n\ngpgpu_simulation_time = %u days, %u hrs, %u min, %u sec (%u sec)\n",
           (unsigned)d, (unsigned)h, (unsigned)m, (unsigned)s, (unsigned)difference );
-   printf("gpgpu_simulation_rate = %u (inst/sec)\n", (unsigned)(GPGPUsim_ctx_ptr()->g_the_gpu->gpu_tot_sim_insn / difference) );
-   printf("gpgpu_simulation_rate = %u (cycle/sec)\n", (unsigned)(GPGPUsim_ctx_ptr()->g_the_gpu->gpu_tot_sim_cycle / difference) );
+   printf("gpgpu_simulation_rate = %u (inst/sec)\n", (unsigned)(the_gpgpusim->g_the_gpu->gpu_tot_sim_insn / difference) );
+   const unsigned cycles_per_sec = (unsigned)(the_gpgpusim->g_the_gpu->gpu_tot_sim_cycle / difference);
+   printf("gpgpu_simulation_rate = %u (cycle/sec)\n", cycles_per_sec );
+   printf("gpgpu_silicon_slowdown = %ux\n", the_gpgpusim->g_the_gpu->shader_clock() * 1000 / cycles_per_sec);
    fflush(stdout);
 }
 
-int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid )
+int gpgpu_context::gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid )
 {
-   GPGPUsim_ctx_ptr()->g_the_gpu->launch(grid);
-   sem_post(&(GPGPUsim_ctx_ptr()->g_sim_signal_start));
-   sem_wait(&(GPGPUsim_ctx_ptr()->g_sim_signal_finish));
+   the_gpgpusim->g_the_gpu->launch(grid);
+   sem_post(&(the_gpgpusim->g_sim_signal_start));
+   sem_wait(&(the_gpgpusim->g_sim_signal_finish));
    return 0;
 }
 
diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h
index 887b3c8..9f408df 100644
--- a/src/gpgpusim_entrypoint.h
+++ b/src/gpgpusim_entrypoint.h
@@ -76,12 +76,4 @@ class GPGPUsim_ctx {
 
 };
 
-void start_sim_thread(int api);
-
-class gpgpu_sim* g_the_gpu();
-struct GPGPUsim_ctx* GPGPUsim_ctx_ptr();
-class stream_manager* g_stream_manager();
-
-int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid );
-
 #endif