Refactor GPGPUSim_Context and GPGPUSim_Init

Signed-off-by: Mengchi Zhang <[email protected]>

1 files changed, 689 insertions, 488 deletions

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 716e297..175cbc5 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -199,8 +199,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();
 
@@ -252,15 +251,14 @@ struct _cuda_device_id *gpgpu_context::GPGPUSim_Init()
 	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 +274,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 +287,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();
@@ -585,7 +582,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 +729,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 +760,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 +869,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));
@@ -960,7 +957,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 +991,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 +1031,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 +1051,164 @@ 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){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	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(ctx);
+	(*array) = (struct cudaArray*) malloc(sizeof(struct cudaArray));
+	(*array)->desc = *desc;
+	(*array)->width = width;
+	(*array)->height = height;
+	(*array)->size = size;
+	(*array)->dimensions = 2;
+	((*array)->devPtr32)= (int) (long long)context->get_device()->get_gpgpu()->gpu_mallocarray(size);
+	printf("GPGPU-Sim PTX: cudaMallocArray: devPtr32 = %d\n", ((*array)->devPtr32));
+	((*array)->devPtr) = (void*) (long long) ((*array)->devPtr32);
+	if ( ((*array)->devPtr) ) {
+		return g_last_cudaError = cudaSuccess;
+	} else {
+		return g_last_cudaError = cudaErrorMemoryAllocation;
+	}
+}
+
+__host__ cudaError_t CUDARTAPI 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(ctx);
+	gpgpu_t *gpu = context->get_device()->get_gpgpu();
+	size_t size = count;
+	printf("GPGPU-Sim PTX: cudaMemcpyToArray\n");
+	if( kind == cudaMemcpyHostToDevice )
+		gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size);
+	else if( kind == cudaMemcpyDeviceToHost )
+		gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size);
+	else if( kind == cudaMemcpyDeviceToDevice )
+		gpu->memcpy_gpu_to_gpu( (size_t)(dst->devPtr), (size_t)src, size);
+	else {
+		printf("GPGPU-Sim PTX: cudaMemcpyToArray - ERROR : unsupported cudaMemcpyKind\n");
+		abort();
+	}
+	dst->devPtr32 = (unsigned) (size_t)(dst->devPtr);
+	return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI 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)
+{
+    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();
+	size_t size = spitch*height;
+	gpgpusim_ptx_assert( (dpitch==spitch), "different src and dst pitch not supported yet" );
+	if( kind == cudaMemcpyHostToDevice )
+		gpu->memcpy_to_gpu( (size_t)dst, src, size );
+	else if( kind == cudaMemcpyDeviceToHost )
+		gpu->memcpy_from_gpu( dst, (size_t)src, size );
+	else if( kind == cudaMemcpyDeviceToDevice )
+		gpu->memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, size);
+	else {
+		printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n");
+		abort();
+	}
+	return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI 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(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;
+	gpgpusim_ptx_assert( ((channel_size%8) == 0), "none byte multiple destination channel size not supported (sz=%u)", channel_size );
+	unsigned elem_size = channel_size/8;
+	gpgpusim_ptx_assert( (dst->dimensions==2), "copy to none 2D array not supported" );
+	gpgpusim_ptx_assert( (wOffset==0), "non-zero wOffset not yet supported" );
+	gpgpusim_ptx_assert( (hOffset==0), "non-zero hOffset not yet supported" );
+	gpgpusim_ptx_assert( (dst->height == (int)height), "partial copy not supported" );
+	gpgpusim_ptx_assert( (elem_size*dst->width == width), "partial copy not supported" );
+	gpgpusim_ptx_assert( (spitch == width), "spitch != width not supported" );
+	if( kind == cudaMemcpyHostToDevice )
+		gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size);
+	else if( kind == cudaMemcpyDeviceToHost )
+		gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size);
+	else if( kind == cudaMemcpyDeviceToDevice )
+		gpu->memcpy_gpu_to_gpu( (size_t)dst->devPtr, (size_t)src, size);
+	else {
+		printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n");
+		abort();
+	}
+	dst->devPtr32 = (unsigned) (size_t)(dst->devPtr);
+	return g_last_cudaError = cudaSuccess;
+}
+
+#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)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	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 cudaMemsetInternal(void *mem, int c, size_t count, gpgpu_context* gpgpu_ctx = NULL)
 {
     gpgpu_context *ctx;
     if (gpgpu_ctx){
@@ -1042,12 +1219,45 @@ cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, gpgpu
 	if(g_debug_execution >= 3){
 	    announce_call(__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;
+}
+
+//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__);
+    }
+	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;
+}
+
+cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, 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();
+	}
 	GLint buffer_size=0;
-	CUctx_st* context = GPGPUSim_Context();
+	CUctx_st* context = GPGPUSim_Context(ctx);
 
 	glbmap_entry_t *p = ctx->api->g_glbmap;
 	while ( p && p->m_bufferObj != bufferObj )
@@ -1118,7 +1328,7 @@ cuLinkAddFileInternal(CUlinkState state, CUjitInputType type, const char *path,
 
     //blocking
     assert(type==CU_JIT_INPUT_PTX);
-	CUctx_st *context = GPGPUSim_Context();
+	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");
@@ -1190,7 +1400,7 @@ cudaError_t CUDARTAPI cudaFuncGetAttributesInternal(struct cudaFuncAttributes *a
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	CUctx_st *context = GPGPUSim_Context();
+	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();
@@ -1210,6 +1420,443 @@ cudaError_t CUDARTAPI cudaFuncGetAttributesInternal(struct cudaFuncAttributes *a
 	return g_last_cudaError = cudaSuccess;
 }
 
+#if (CUDART_VERSION > 5000)
+__host__ cudaError_t CUDARTAPI cudaDeviceGetAttributeInternal(int *value, enum cudaDeviceAttr attr, int device, 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__);
+		}
+
+        const struct cudaDeviceProp *prop;
+        _cuda_device_id *dev = ctx->GPGPUSim_Init();
+
+        if (device <= dev->num_devices() )  {
+                prop = dev->get_prop();
+                switch (attr) {
+                case 1:
+                        *value= prop->maxThreadsPerBlock;
+                        break;
+                case 2:
+                        *value= prop->maxThreadsDim[0];
+                        break;
+                case 3:
+                        *value= prop->maxThreadsDim[1];
+                        break;
+                case 4:
+                        *value= prop->maxThreadsDim[2];
+                        break; 
+		case 5:
+                        *value= prop->maxGridSize[0];
+                        break;
+                case 6:
+                        *value= prop->maxGridSize[1];
+                        break;
+                case 7:
+                        *value= prop->maxGridSize[2];
+                        break;
+		case 8:
+                        *value= prop->sharedMemPerBlock;
+                        break;
+		case 9:
+                        *value= prop->totalConstMem;
+                        break;
+                case 10:
+                        *value= prop->warpSize;
+                        break;
+                case 11:
+                        *value= 16;//dummy value
+                        break;
+                case 12:
+                        *value= prop->regsPerBlock;
+                        break;
+                case 13:
+                        *value= 1480000;//for 1080ti
+                        break;
+                case 14:
+                        *value= prop->textureAlignment ;
+                        break;
+                case 15:
+                        *value = 0;
+                        break;
+                case 16:
+                        *value= prop->multiProcessorCount ;
+                        break;
+                case 17:
+                case 18:
+                case 19:
+                        *value = 0;
+                        break;
+                case 21:
+                case 22:
+                case 23:
+                case 24:
+                case 25:
+                case 26:
+                case 27:
+                case 28:
+                case 42:
+                case 45:
+                case 46:
+                case 47:
+                case 48:
+                case 49:
+                case 52:
+                case 53:
+                case 55:
+                case 56:
+                case 57:
+                case 58:
+                case 59:
+                case 60:
+                case 61:
+                case 62:
+                case 63:
+                case 64:
+                case 66:
+                case 67:
+                case 69:
+                case 70:
+                case 71:
+                case 73:
+                case 74:
+                case 77:
+                        *value = 1000;//dummy value
+                        break;
+                case 29:
+                case 43:
+                case 54:
+                case 65:
+                case 68:
+                case 72:
+                        *value = 10;//dummy value
+                        break;
+                case 30:
+                case 51:
+                        *value = 128;//dummy value
+                        break;
+                case 31:
+                        *value = 1;
+                        break;
+                case 32:
+                        *value = 0;
+                        break;
+                case 33:
+                case 50:
+                        *value = 0;//dummy value
+                        break;
+		        case 34:
+                        *value= 0;
+                        break;
+                case 35:
+                        *value = 0;
+                        break;
+                case 36:
+                        *value = 1250000;//CK value for 1080ti
+                        break;
+                case 37:
+                        *value = 352;//value for 1080ti
+                        break;
+                case 38:
+                        *value = 3000000;//value for 1080ti
+                        break;
+                case 39:
+                        *value= dev->get_gpgpu()->threads_per_core();
+                        break;
+                case 40:
+                        *value= 0;
+                        break;
+                case 41:
+                        *value= 0;
+                        break;
+                case 75://cudaDevAttrComputeCapabilityMajor
+                        *value= prop->major ;
+                        break;
+                case 76://cudaDevAttrComputeCapabilityMinor
+                        *value= prop->minor ;
+                        break;
+                case 78:
+                        *value= 0 ; //TODO: as of now, we dont support stream priorities.
+                        break;
+                case 79:
+                        *value= 0;
+                        break;
+                case 80:
+                        *value= 0;
+                        break;
+		#if (CUDART_VERSION > 5050)
+		case 81:
+                        *value= prop->sharedMemPerMultiprocessor;
+                        break;
+                case 82:
+                        *value= prop->regsPerMultiprocessor;
+                        break;
+		#endif
+                case 83:
+                case 84:
+                case 85:
+                case 86:
+                        *value= 0;
+                        break;
+                case 87:
+                        *value= 4;//dummy value
+                        break;
+                case 88:
+                case 89:
+                        *value= 0;
+                        break;
+		default:
+			printf("ERROR: Attribute number %d unimplemented \n",attr);
+			abort();
+                }
+                return g_last_cudaError = cudaSuccess;
+        } else {
+                return g_last_cudaError = cudaErrorInvalidDevice;
+        }
+}
+#endif
+
+__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__);
+    }
+	CUctx_st *context = GPGPUSim_Context(ctx);
+	gpgpu_t *gpu = context->get_device()->get_gpgpu();
+	printf("GPGPU-Sim PTX: in cudaBindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference));
+	struct cudaArray *array;
+	array = (struct cudaArray*) malloc(sizeof(struct cudaArray));
+	array->desc = *desc;
+	array->size = size;
+	array->width = size;
+	array->height = 1;
+	array->dimensions = 1;
+	array->devPtr = (void*)devPtr;
+	array->devPtr32 = (int)(long long)devPtr;
+	offset = 0;
+	printf("GPGPU-Sim PTX:   size = %zu\n", size);
+	printf("GPGPU-Sim PTX:   texref = %p, array = %p\n", texref, array);
+	printf("GPGPU-Sim PTX:   devPtr32 = %x\n", array->devPtr32);
+	printf("GPGPU-Sim PTX:   Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref));
+	printf("GPGPU-Sim PTX:   ChannelFormatDesc: x=%d, y=%d, z=%d, w=%d\n", desc->x, desc->y, desc->z, desc->w);
+	printf("GPGPU-Sim PTX:   Texture Normalized? = %d\n", texref->normalized);
+	gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array);
+	devPtr = (void*)(long long)array->devPtr32;
+	printf("GPGPU-Sim PTX: devPtr = %p\n", devPtr);
+	return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI 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;
+}
+
+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 */
 
 /*******************************************************************************
  *                                                                              *
@@ -1240,42 +1887,12 @@ __host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t 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;
-	}
+    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))
 {
-	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();
-	(*array) = (struct cudaArray*) malloc(sizeof(struct cudaArray));
-	(*array)->desc = *desc;
-	(*array)->width = width;
-	(*array)->height = height;
-	(*array)->size = size;
-	(*array)->dimensions = 2;
-	((*array)->devPtr32)= (int) (long long)context->get_device()->get_gpgpu()->gpu_mallocarray(size);
-	printf("GPGPU-Sim PTX: cudaMallocArray: devPtr32 = %d\n", ((*array)->devPtr32));
-	((*array)->devPtr) = (void*) (long long) ((*array)->devPtr32);
-	if ( ((*array)->devPtr) ) {
-		return g_last_cudaError = cudaSuccess;
-	} else {
-		return g_last_cudaError = cudaErrorMemoryAllocation;
-	}
+    return cudaMallocArrayInternal(array, desc, width, height __dv(1));
 }
 
 __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr)
@@ -1351,25 +1968,7 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t cou
 
 __host__ cudaError_t CUDARTAPI cudaMemcpyToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind)
 { 
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	CUctx_st *context = GPGPUSim_Context();
-	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	size_t size = count;
-	printf("GPGPU-Sim PTX: cudaMemcpyToArray\n");
-	if( kind == cudaMemcpyHostToDevice )
-		gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size);
-	else if( kind == cudaMemcpyDeviceToHost )
-		gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size);
-	else if( kind == cudaMemcpyDeviceToDevice )
-		gpu->memcpy_gpu_to_gpu( (size_t)(dst->devPtr), (size_t)src, size);
-	else {
-		printf("GPGPU-Sim PTX: cudaMemcpyToArray - ERROR : unsupported cudaMemcpyKind\n");
-		abort();
-	}
-	dst->devPtr32 = (unsigned) (size_t)(dst->devPtr);
-	return g_last_cudaError = cudaSuccess;
+    return cudaMemcpyToArrayInternal(dst, wOffset, hOffset, src, count, kind);
 }
 
 
@@ -1395,59 +1994,14 @@ __host__ cudaError_t CUDARTAPI cudaMemcpyArrayToArray(struct cudaArray *dst, siz
 
 __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();
-	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	size_t size = spitch*height;
-	gpgpusim_ptx_assert( (dpitch==spitch), "different src and dst pitch not supported yet" );
-	if( kind == cudaMemcpyHostToDevice )
-		gpu->memcpy_to_gpu( (size_t)dst, src, size );
-	else if( kind == cudaMemcpyDeviceToHost )
-		gpu->memcpy_from_gpu( dst, (size_t)src, size );
-	else if( kind == cudaMemcpyDeviceToDevice )
-		gpu->memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, size);
-	else {
-		printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n");
-		abort();
-	}
-	return g_last_cudaError = cudaSuccess;
+    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)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	CUctx_st *context = GPGPUSim_Context();
-	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	size_t size = spitch*height;
-	size_t channel_size = dst->desc.w+dst->desc.x+dst->desc.y+dst->desc.z;
-	gpgpusim_ptx_assert( ((channel_size%8) == 0), "none byte multiple destination channel size not supported (sz=%u)", channel_size );
-	unsigned elem_size = channel_size/8;
-	gpgpusim_ptx_assert( (dst->dimensions==2), "copy to none 2D array not supported" );
-	gpgpusim_ptx_assert( (wOffset==0), "non-zero wOffset not yet supported" );
-	gpgpusim_ptx_assert( (hOffset==0), "non-zero hOffset not yet supported" );
-	gpgpusim_ptx_assert( (dst->height == (int)height), "partial copy not supported" );
-	gpgpusim_ptx_assert( (elem_size*dst->width == width), "partial copy not supported" );
-	gpgpusim_ptx_assert( (spitch == width), "spitch != width not supported" );
-	if( kind == cudaMemcpyHostToDevice )
-		gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size);
-	else if( kind == cudaMemcpyDeviceToHost )
-		gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size);
-	else if( kind == cudaMemcpyDeviceToDevice )
-		gpu->memcpy_gpu_to_gpu( (size_t)dst->devPtr, (size_t)src, size);
-	else {
-		printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n");
-		abort();
-	}
-	dst->devPtr32 = (unsigned) (size_t)(dst->devPtr);
-	return g_last_cudaError = cudaSuccess;
+    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){
@@ -1457,7 +2011,6 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArray(void *dst, size_t dpitch, c
 	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){
@@ -1467,7 +2020,6 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2DArrayToArray(struct cudaArray *dst, s
 	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){
@@ -1582,27 +2134,7 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArrayAsync(void *dst, size_t dpit
 #if (CUDART_VERSION >= 8000)
 cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks, const char *hostFunc, int blockSize, size_t dynamicSMemSize, unsigned int flags)
 {
-	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);
-		return g_last_cudaError = cudaSuccess;
-	} else {
-		cuda_not_implemented(__my_func__,__LINE__);
-		return g_last_cudaError = cudaErrorUnknown;
-	}
+    return cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlagsInternal(numBlocks, hostFunc, blockSize, dynamicSMemSize, flags);
 }
 
 #endif
@@ -1616,227 +2148,13 @@ cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int
  *******************************************************************************/
 __host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count)
 {
-	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);
-	return g_last_cudaError = cudaSuccess;
+    return cudaMemsetInternal(mem, c, count);
 }
 
-#if (CUDART_VERSION > 5000)
-__host__ cudaError_t CUDARTAPI cudaDeviceGetAttributeInternal(int *value, enum cudaDeviceAttr attr, int device, 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__);
-		}
-
-        const struct cudaDeviceProp *prop;
-        _cuda_device_id *dev = ctx->GPGPUSim_Init();
-
-        if (device <= dev->num_devices() )  {
-                prop = dev->get_prop();
-                switch (attr) {
-                case 1:
-                        *value= prop->maxThreadsPerBlock;
-                        break;
-                case 2:
-                        *value= prop->maxThreadsDim[0];
-                        break;
-                case 3:
-                        *value= prop->maxThreadsDim[1];
-                        break;
-                case 4:
-                        *value= prop->maxThreadsDim[2];
-                        break; 
-		case 5:
-                        *value= prop->maxGridSize[0];
-                        break;
-                case 6:
-                        *value= prop->maxGridSize[1];
-                        break;
-                case 7:
-                        *value= prop->maxGridSize[2];
-                        break;
-		case 8:
-                        *value= prop->sharedMemPerBlock;
-                        break;
-		case 9:
-                        *value= prop->totalConstMem;
-                        break;
-                case 10:
-                        *value= prop->warpSize;
-                        break;
-                case 11:
-                        *value= 16;//dummy value
-                        break;
-                case 12:
-                        *value= prop->regsPerBlock;
-                        break;
-                case 13:
-                        *value= 1480000;//for 1080ti
-                        break;
-                case 14:
-                        *value= prop->textureAlignment ;
-                        break;
-                case 15:
-                        *value = 0;
-                        break;
-                case 16:
-                        *value= prop->multiProcessorCount ;
-                        break;
-                case 17:
-                case 18:
-                case 19:
-                        *value = 0;
-                        break;
-                case 21:
-                case 22:
-                case 23:
-                case 24:
-                case 25:
-                case 26:
-                case 27:
-                case 28:
-                case 42:
-                case 45:
-                case 46:
-                case 47:
-                case 48:
-                case 49:
-                case 52:
-                case 53:
-                case 55:
-                case 56:
-                case 57:
-                case 58:
-                case 59:
-                case 60:
-                case 61:
-                case 62:
-                case 63:
-                case 64:
-                case 66:
-                case 67:
-                case 69:
-                case 70:
-                case 71:
-                case 73:
-                case 74:
-                case 77:
-                        *value = 1000;//dummy value
-                        break;
-                case 29:
-                case 43:
-                case 54:
-                case 65:
-                case 68:
-                case 72:
-                        *value = 10;//dummy value
-                        break;
-                case 30:
-                case 51:
-                        *value = 128;//dummy value
-                        break;
-                case 31:
-                        *value = 1;
-                        break;
-                case 32:
-                        *value = 0;
-                        break;
-                case 33:
-                case 50:
-                        *value = 0;//dummy value
-                        break;
-		        case 34:
-                        *value= 0;
-                        break;
-                case 35:
-                        *value = 0;
-                        break;
-                case 36:
-                        *value = 1250000;//CK value for 1080ti
-                        break;
-                case 37:
-                        *value = 352;//value for 1080ti
-                        break;
-                case 38:
-                        *value = 3000000;//value for 1080ti
-                        break;
-                case 39:
-                        *value= dev->get_gpgpu()->threads_per_core();
-                        break;
-                case 40:
-                        *value= 0;
-                        break;
-                case 41:
-                        *value= 0;
-                        break;
-                case 75://cudaDevAttrComputeCapabilityMajor
-                        *value= prop->major ;
-                        break;
-                case 76://cudaDevAttrComputeCapabilityMinor
-                        *value= prop->minor ;
-                        break;
-                case 78:
-                        *value= 0 ; //TODO: as of now, we dont support stream priorities.
-                        break;
-                case 79:
-                        *value= 0;
-                        break;
-                case 80:
-                        *value= 0;
-                        break;
-		#if (CUDART_VERSION > 5050)
-		case 81:
-                        *value= prop->sharedMemPerMultiprocessor;
-                        break;
-                case 82:
-                        *value= prop->regsPerMultiprocessor;
-                        break;
-		#endif
-                case 83:
-                case 84:
-                case 85:
-                case 86:
-                        *value= 0;
-                        break;
-                case 87:
-                        *value= 4;//dummy value
-                        break;
-                case 88:
-                case 89:
-                        *value= 0;
-                        break;
-		default:
-			printf("ERROR: Attribute number %d unimplemented \n",attr);
-			abort();
-                }
-                return g_last_cudaError = cudaSuccess;
-        } else {
-                return g_last_cudaError = cudaErrorInvalidDevice;
-        }
-}
-#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)
 {
-	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();
-	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	gpu->gpu_memset((size_t)mem, c, count);
-	return g_last_cudaError = cudaSuccess;
+    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)
@@ -1986,61 +2304,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 +2400,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);
 }
 
 
@@ -2528,7 +2788,7 @@ 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();
+    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"); 
@@ -2886,7 +3146,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){
@@ -2899,7 +3159,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];
@@ -3140,22 +3400,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);
 }
 
 
@@ -3187,30 +3432,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) 
@@ -3322,12 +3544,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
@@ -5033,23 +5250,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 */