From 90e85664154a7b8eda8ddbf0429eb962c9be78e1 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Sun, 8 Sep 2019 00:56:20 -0400 Subject: Refactor GPGPUSim_Context and GPGPUSim_Init Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 1531 ++++++++++++++++++++++++------------------- 1 file changed, 866 insertions(+), 665 deletions(-) (limited to 'libcuda') 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::const_iterator i = ctx->api->pinned_memory_size.find(pHost); assert(i != ctx->api->pinned_memory_size.end()); @@ -1031,7 +1051,182 @@ 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){ + 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(); + 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){ @@ -1042,12 +1237,27 @@ cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, gpgpu 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,72 +1420,479 @@ cudaError_t CUDARTAPI cudaFuncGetAttributesInternal(struct cudaFuncAttributes *a 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) +#if (CUDART_VERSION > 5000) +__host__ cudaError_t CUDARTAPI cudaDeviceGetAttributeInternal(int *value, enum cudaDeviceAttr attr, int device, gpgpu_context* gpgpu_ctx = NULL) { - return cudaMallocInternal(devPtr, size); -} + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } -__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; - } -} + const struct cudaDeviceProp *prop; + _cuda_device_id *dev = ctx->GPGPUSim_Init(); -__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; + 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 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 { - return g_last_cudaError = cudaErrorMemoryAllocation; + 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 p = entry->get_param_config(i); + cudaSetupArgumentInternal(kernelParams[i], p.first, p.second, ctx); + } + cudaLaunchInternal(hostFun, ctx); + return CUDA_SUCCESS; +} +#endif /* CUDART_VERSION >= 4000 */ + +/******************************************************************************* + * * + * * + * * + *******************************************************************************/ + +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 __dv(1)); } __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr) @@ -1344,213 +1961,18 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t cou } else { printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported cudaMemcpyKind\n"); - abort(); - } - return g_last_cudaError = cudaSuccess; -} - -__host__ cudaError_t CUDARTAPI cudaMemcpyToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind) -{ - 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; -} - - -__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) -{ - 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; -} - - -__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; -} - - -__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)) -{ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - //CUctx_st *context = GPGPUSim_Context(); - assert(kind == cudaMemcpyHostToDevice); - printf("GPGPU-Sim PTX: cudaMemcpyToSymbol: symbol = %p\n", symbol); - //stream_operation( const char *symbol, const void *src, size_t count, size_t offset ) - g_stream_manager()->push( stream_operation(src,symbol,count,offset,0) ); - //gpgpu_ptx_sim_memcpy_symbol(symbol,src,count,offset,1,context->get_device()->get_gpgpu()); - return g_last_cudaError = cudaSuccess; -} - - -__host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbol(void *dst, const char *symbol, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost)) -{ - 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) ); - //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) -{ - 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; - default: - abort(); - } - return g_last_cudaError = cudaSuccess; -} - - -__host__ cudaError_t CUDARTAPI cudaMemcpyToArrayAsync(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream) -{ - 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; + abort(); + } + return g_last_cudaError = cudaSuccess; +} + +__host__ cudaError_t CUDARTAPI cudaMemcpyToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind) +{ + return cudaMemcpyToArrayInternal(dst, wOffset, hOffset, src, count, kind); } -__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) +__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__); @@ -1560,7 +1982,7 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2DAsync(void *dst, size_t dpitch, const } -__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) +__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__); @@ -1570,275 +1992,171 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2DToArrayAsync(struct cudaArray *dst, s } -__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) +__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__); - } - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + return cudaMemcpy2DInternal(dst, dpitch, src, spitch, width, height, kind); } -#if (CUDART_VERSION >= 8000) -cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks, const char *hostFunc, int blockSize, size_t dynamicSMemSize, unsigned int flags) +__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) { - 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 cudaMemcpy2DToArrayInternal(dst, wOffset, hOffset, src, spitch, width, height, kind); } -#endif - - - -/******************************************************************************* - * * - * * - * * - *******************************************************************************/ -__host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count) +__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__); } - 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; + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; } -#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; - } +__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; } -#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 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__); } - 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); + //CUctx_st *context = GPGPUSim_Context(); + assert(kind == cudaMemcpyHostToDevice); + printf("GPGPU-Sim PTX: cudaMemcpyToSymbol: symbol = %p\n", symbol); + //stream_operation( const char *symbol, const void *src, size_t count, size_t offset ) + g_stream_manager()->push( stream_operation(src,symbol,count,offset,0) ); + //gpgpu_ptx_sim_memcpy_symbol(symbol,src,count,offset,1,context->get_device()->get_gpgpu()); + return g_last_cudaError = cudaSuccess; +} + + +__host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbol(void *dst, const char *symbol, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost)) +{ + 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) ); + //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) +{ + 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; + default: + abort(); + } return g_last_cudaError = cudaSuccess; } + +__host__ cudaError_t CUDARTAPI cudaMemcpyToArrayAsync(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream) +{ + 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 +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 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 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 */ -- cgit v1.3 From 4671c7b3b3f1792f6f4b6c98e2bcffe7d072bac0 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 11 Sep 2019 21:06:37 -0400 Subject: Fix init bug in cudaMallocArray Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 175cbc5..8930d61 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1892,7 +1892,7 @@ __host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr, size_t *pitch, siz __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 __dv(1)); + return cudaMallocArrayInternal(array, desc, width, height); } __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr) -- cgit v1.3 From 5181fe19601cf10ce976c5a5965b993cf0213260 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 11 Sep 2019 23:09:12 -0400 Subject: Remove g_stream_manager() Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 348 ++++++++++++++++++++++-------------- src/abstract_hardware_model.cc | 8 +- src/cuda-sim/cuda-sim.cc | 2 +- src/cuda-sim/cuda_device_runtime.cc | 3 +- src/gpgpusim_entrypoint.cc | 4 - src/gpgpusim_entrypoint.h | 1 - 6 files changed, 224 insertions(+), 142 deletions(-) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 8930d61..9459b64 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -176,8 +176,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 @@ -941,7 +939,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; } @@ -1080,6 +1078,50 @@ __host__ cudaError_t CUDARTAPI cudaMallocArrayInternal(struct cudaArray **array, } } +__host__ cudaError_t CUDARTAPI cudaMemcpyInternal(void *dst, 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(); + //gpgpu_t *gpu = context->get_device()->get_gpgpu(); + if(g_debug_execution >= 3) + printf("GPGPU-Sim PTX: cudaMemcpy(): devPtr = %p\n", dst); + if( kind == cudaMemcpyHostToDevice ) + ctx->the_gpgpusim->g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) ); + else if( kind == cudaMemcpyDeviceToHost ) + ctx->the_gpgpusim->g_stream_manager->push( stream_operation((size_t)src,dst,count,0) ); + else if( kind == cudaMemcpyDeviceToDevice ) + 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) + ctx->the_gpgpusim->g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) ); // device to device + else + 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) + 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(); + } + } + } + else { + printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported cudaMemcpyKind\n"); + abort(); + } + return g_last_cudaError = cudaSuccess; +} + __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; @@ -1174,6 +1216,69 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2DToArrayInternal(struct cudaArray *dst return g_last_cudaError = cudaSuccess; } +__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) +{ + 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 == cudaMemcpyHostToDevice); + printf("GPGPU-Sim PTX: cudaMemcpyToSymbol: symbol = %p\n", symbol); + //stream_operation( const char *symbol, const void *src, size_t count, size_t offset ) + 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 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); + 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 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: 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(); + } + 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) { @@ -1729,6 +1834,48 @@ __host__ cudaError_t CUDARTAPI cudaLaunchKernelInternal( const char* hostFun, di 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; +} + void __cudaRegisterTextureInternal( void **fatCubinHandle, const struct textureReference *hostVar, @@ -1858,6 +2005,66 @@ CUresult CUDAAPI cuLaunchKernelInternal(CUfunction f, } #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; +} + /******************************************************************************* * * * * @@ -1930,40 +2137,7 @@ __host__ cudaError_t CUDARTAPI cudaFreeArray(struct cudaArray *array) __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__); - } - //CUctx_st *context = GPGPUSim_Context(); - //gpgpu_t *gpu = context->get_device()->get_gpgpu(); - if(g_debug_execution >= 3) - printf("GPGPU-Sim PTX: cudaMemcpy(): devPtr = %p\n", dst); - if( kind == cudaMemcpyHostToDevice ) - g_stream_manager()->push( stream_operation(src,(size_t)dst,count,0) ); - else if( kind == cudaMemcpyDeviceToHost ) - g_stream_manager()->push( stream_operation((size_t)src,dst,count,0) ); - else if( kind == cudaMemcpyDeviceToDevice ) - g_stream_manager()->push( stream_operation((size_t)src,(size_t)dst,count,0) ); - else if ( kind == cudaMemcpyDefault ) { - if ((size_t)src >= GLOBAL_HEAP_START) { - if ((size_t)dst >= GLOBAL_HEAP_START) - g_stream_manager()->push( stream_operation((size_t)src,(size_t)dst,count,0) ); // device to device - else - g_stream_manager()->push( stream_operation((size_t)src,dst,count,0) ); // device to host - } - else { - if ((size_t)dst >= GLOBAL_HEAP_START) - g_stream_manager()->push( stream_operation(src,(size_t)dst,count,0) ); - else { - printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported transfer: host to host\n"); - abort(); - } - } - } - else { - printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported cudaMemcpyKind\n"); - abort(); - } - return g_last_cudaError = cudaSuccess; + 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) @@ -2022,30 +2196,13 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2DArrayToArray(struct cudaArray *dst, s __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__); - } - //CUctx_st *context = GPGPUSim_Context(); - assert(kind == cudaMemcpyHostToDevice); - printf("GPGPU-Sim PTX: cudaMemcpyToSymbol: symbol = %p\n", symbol); - //stream_operation( const char *symbol, const void *src, size_t count, size_t offset ) - g_stream_manager()->push( stream_operation(src,symbol,count,offset,0) ); - //gpgpu_ptx_sim_memcpy_symbol(symbol,src,count,offset,1,context->get_device()->get_gpgpu()); - return g_last_cudaError = cudaSuccess; + 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)) { - 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) ); - //gpgpu_ptx_sim_memcpy_symbol(symbol,dst,count,offset,0,context->get_device()->get_gpgpu()); - return g_last_cudaError = cudaSuccess; + return cudaMemcpyFromSymbolInternal(dst, symbol, count, offset, kind); } __host__ cudaError_t CUDARTAPI cudaMemGetInfo (size_t *free, size_t *total){ @@ -2067,18 +2224,7 @@ __host__ cudaError_t CUDARTAPI cudaMemGetInfo (size_t *free, size_t *total){ __host__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream) { - 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; - default: - abort(); - } - return g_last_cudaError = cudaSuccess; + return cudaMemcpyAsyncInternal(dst, src, count, kind, stream); } @@ -2414,18 +2560,7 @@ __host__ cudaError_t CUDARTAPI cudaLaunchKernel( const char* hostFun, dim3 gridD __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 @@ -2452,16 +2587,7 @@ __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) @@ -2516,52 +2642,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) diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 9a91818..07232ee 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 * 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/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 67d7f9b..fc1ea8e 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -2188,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******* 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/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 19a525e..5ce40c6 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -52,10 +52,6 @@ GPGPUsim_ctx* GPGPUsim_ctx_ptr(){ return the_gpgpusim; } -class stream_manager* g_stream_manager() { - return GPGPUsim_ctx_ptr()->g_stream_manager; -} - static void print_simulation_time(); void *gpgpu_sim_thread_sequential(void*) diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h index ba3aac9..b2e22e6 100644 --- a/src/gpgpusim_entrypoint.h +++ b/src/gpgpusim_entrypoint.h @@ -79,7 +79,6 @@ class GPGPUsim_ctx { void start_sim_thread(int api); struct GPGPUsim_ctx* GPGPUsim_ctx_ptr(); -class stream_manager* g_stream_manager(); int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ); -- cgit v1.3 From 48887cfe0261abde6de23fc5d1d76694426b7e8e Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Thu, 12 Sep 2019 02:01:01 -0400 Subject: Remove GPGPUsim_ctx_ptr() Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 116 ++++++++++++++++++++--------- libcuda/gpgpu_context.h | 5 ++ libopencl/opencl_runtime_api.cc | 8 +- src/cuda-sim/cuda-sim.cc | 2 +- src/gpgpu-sim/gpu-sim.cc | 4 +- src/gpgpusim_entrypoint.cc | 161 +++++++++++++++++++--------------------- src/gpgpusim_entrypoint.h | 6 -- 7 files changed, 167 insertions(+), 135 deletions(-) (limited to 'libcuda') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 9459b64..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 { @@ -242,8 +239,8 @@ 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; @@ -1876,6 +1873,28 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroyInternal(cudaStream_t stream, gp 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, @@ -2065,6 +2084,53 @@ __host__ cudaError_t CUDARTAPI cudaStreamWaitEventInternal(cudaStream_t stream, 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; +} + /******************************************************************************* * * * * @@ -2592,18 +2658,7 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t 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) @@ -2722,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*) { @@ -3441,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, 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/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index fc1ea8e..7a130ea 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -2203,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 diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index d236c74..56ea8c4 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -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 ); @@ -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/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 5ce40c6..846773d 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -43,38 +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; -} - -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; } @@ -86,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 { @@ -95,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 @@ -119,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); } @@ -211,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"); @@ -220,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; @@ -260,18 +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) ); - const unsigned cycles_per_sec = (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", GPGPUsim_ctx_ptr()->g_the_gpu->shader_clock() * 1000 / 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 b2e22e6..9f408df 100644 --- a/src/gpgpusim_entrypoint.h +++ b/src/gpgpusim_entrypoint.h @@ -76,10 +76,4 @@ class GPGPUsim_ctx { }; -void start_sim_thread(int api); - -struct GPGPUsim_ctx* GPGPUsim_ctx_ptr(); - -int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ); - #endif -- cgit v1.3