diff options
28 files changed, 1199 insertions, 908 deletions
diff --git a/Jenkinsfile b/Jenkinsfile index a345bc0..8a3ad19 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -13,11 +13,11 @@ pipeline { parallel "4.2": { sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/4.2_env_setup.sh &&\ source `pwd`/setup_environment &&\ - make -j' + make -j 10' }, "10.1" : { sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/10.1_env_setup.sh &&\ source `pwd`/setup_environment &&\ - make -j' + make -j 10' } } } @@ -32,13 +32,13 @@ pipeline { source `pwd`/setup_environment &&\ cd gpgpu-sim_simulations && \ source ./benchmarks/src/setup_environment && \ - make -j -C ./benchmarks/src rodinia_2.0-ft sdk-4.2 && \ + make -j 10 -C ./benchmarks/src rodinia_2.0-ft sdk-4.2 && \ make -C ./benchmarks/src data' sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/10.1_env_setup.sh &&\ source `pwd`/setup_environment &&\ cd gpgpu-sim_simulations && \ source ./benchmarks/src/setup_environment && \ - make -j -C ./benchmarks/src/ rodinia_2.0-ft sdk-4.2 && \ + make -j 10 -C ./benchmarks/src/ rodinia_2.0-ft sdk-4.2 && \ make -C ./benchmarks/src data' } } @@ -77,8 +77,8 @@ pipeline { ./gpgpu-sim_simulations/util/plotting/plot-get-stats.py -c per-app-merge-10.1.csv -P cuda-10.1 &&\ ./gpgpu-sim_simulations/util/plotting/merge-stats.py -c ./gpgpu-sim-results-repo/${JOB_NAME}/stats-per-kernel-4.2-ptxplus.csv,./stats-per-kernel-4.2-ptxplus.csv -R > per-kernel-merge-4.2-ptxplus.csv &&\ ./gpgpu-sim_simulations/util/plotting/merge-stats.py -c ./gpgpu-sim-results-repo/${JOB_NAME}/stats-per-kernel-10.1.csv,./stats-per-kernel-10.1.csv -R > per-kernel-merge-10.1.csv &&\ - ./gpgpu-sim_simulations/util/plotting/plot-correlation.py -c per-kernel-merge-4.2-ptxplus.csv -p cuda-4.2 | grep "Correl=" | tee correl.4.2.txt &&\ - ./gpgpu-sim_simulations/util/plotting/plot-correlation.py -c per-kernel-merge-10.1.csv -p cuda-10.1 | grep "Correl=" | tee correl.10.1.txt &&\ + ./gpgpu-sim_simulations/util/plotting/plot-correlation.py -c per-kernel-merge-4.2-ptxplus.csv -p cuda-4.2 | grep -B 1 "Correl=" | tee correl.4.2.txt &&\ + ./gpgpu-sim_simulations/util/plotting/plot-correlation.py -c per-kernel-merge-10.1.csv -p cuda-10.1 | grep -B 1 "Correl=" | tee correl.10.1.txt &&\ mkdir -p ./gpgpu-sim-results-repo/${JOB_NAME}/ && cp stats-per-*.csv ./gpgpu-sim-results-repo/${JOB_NAME}/ &&\ cd ./gpgpu-sim-results-repo &&\ git diff --quiet && git diff --staged --quiet || git commit -am "Jenkins automated checkin ${JOB_NAME} Build:${BUILD_NUMBER}" &&\ diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index e71db4c..5cefe60 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -150,9 +150,6 @@ #endif -extern void synchronize(); -extern void exit_simulation(); - /*DEVICE_BUILTIN*/ struct cudaArray { @@ -176,8 +173,6 @@ struct cudaArray cudaError_t g_last_cudaError = cudaSuccess; -//extern stream_manager *g_stream_manager(); - void register_ptx_function( const char *name, function_info *impl ) { // no longer need this @@ -199,8 +194,7 @@ void register_ptx_function( const char *name, function_info *impl ) struct _cuda_device_id *gpgpu_context::GPGPUSim_Init() { - //static _cuda_device_id *the_device = NULL; - _cuda_device_id *the_device = GPGPUsim_ctx_ptr()->the_cude_device; + _cuda_device_id *the_device = the_gpgpusim->the_cude_device; if( !the_device ) { gpgpu_sim *the_gpu = gpgpu_ptx_sim_init_perf(); @@ -245,22 +239,21 @@ struct _cuda_device_id *gpgpu_context::GPGPUSim_Init() prop->maxThreadsPerMultiProcessor = the_gpu->threads_per_core(); #endif the_gpu->set_prop(prop); - GPGPUsim_ctx_ptr()->the_cude_device = new _cuda_device_id(the_gpu); - the_device = GPGPUsim_ctx_ptr()->the_cude_device; + the_gpgpusim->the_cude_device = new _cuda_device_id(the_gpu); + the_device = the_gpgpusim->the_cude_device; } start_sim_thread(1); return the_device; } -static CUctx_st* GPGPUSim_Context() +CUctx_st* GPGPUSim_Context(gpgpu_context * ctx) { //static CUctx_st *the_context = NULL; - gpgpu_context *cur_ctx = GPGPU_Context(); - CUctx_st *the_context = GPGPUsim_ctx_ptr()->the_context; + CUctx_st *the_context = ctx->the_gpgpusim->the_context; if( the_context == NULL ) { - _cuda_device_id *the_gpu = cur_ctx->GPGPUSim_Init(); - GPGPUsim_ctx_ptr()->the_context = new CUctx_st(the_gpu); - the_context = GPGPUsim_ctx_ptr()->the_context; + _cuda_device_id *the_gpu = ctx->GPGPUSim_Init(); + ctx->the_gpgpusim->the_context = new CUctx_st(the_gpu); + the_context = ctx->the_gpgpusim->the_context; } return the_context; } @@ -276,9 +269,9 @@ gpgpu_context* GPGPU_Context() void ptxinfo_data::ptxinfo_addinfo() { + CUctx_st *context = GPGPUSim_Context(gpgpu_ctx); if(!get_ptxinfo_kname()){ /* This info is not per kernel (since CUDA 5.0 some info (e.g. gmem, and cmem) is added at the beginning for the whole binary ) */ - CUctx_st *context = GPGPUSim_Context(); print_ptxinfo(); context->add_ptxinfo(get_ptxinfo()); clear_ptxinfo(); @@ -289,7 +282,6 @@ gpgpu_context* GPGPU_Context() clear_ptxinfo(); return; } - CUctx_st *context = GPGPUSim_Context(); print_ptxinfo(); context->add_ptxinfo( get_ptxinfo_kname(), get_ptxinfo() ); clear_ptxinfo(); @@ -548,7 +540,7 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimitInternal( size_t* pValue, cudaL break; } else{ - printf("ERROR:Limit %s is not supported on this architecture \n",limit); + printf("ERROR:Limit %d is not supported on this architecture \n", limit); abort(); } case 4: // cudaLimitDevRuntimePendingLaunchCount @@ -557,12 +549,12 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimitInternal( size_t* pValue, cudaL break; } else{ - printf("ERROR:Limit %s is not supported on this architecture \n",limit); + printf("ERROR:Limit %d is not supported on this architecture \n",limit); abort(); } #endif default: - printf("ERROR:Limit %s unimplemented \n",limit); + printf("ERROR:Limit %d unimplemented \n",limit); abort(); } return g_last_cudaError = cudaSuccess; @@ -585,7 +577,7 @@ void** cudaRegisterFatBinaryInternal( void *fatCubin, gpgpu_context* gpgpu_ctx = printf("GPGPU-Sim PTX: ERROR ** this version of GPGPU-Sim requires CUDA 2.1 or higher\n"); exit(1); #endif - CUctx_st *context = GPGPUSim_Context(); + CUctx_st *context = GPGPUSim_Context(ctx); static unsigned next_fat_bin_handle = 1; if(context->get_device()->get_gpgpu()->get_config().use_cuobjdump()) { // The following workaround has only been verified on 64-bit systems. @@ -732,7 +724,7 @@ void cudaRegisterFunctionInternal( if(g_debug_execution >= 3){ announce_call(__my_func__); } - CUctx_st *context = GPGPUSim_Context(); + CUctx_st *context = GPGPUSim_Context(ctx); unsigned fat_cubin_handle = (unsigned)(unsigned long long)fatCubinHandle; printf("GPGPU-Sim PTX: __cudaRegisterFunction %s : hostFun 0x%p, fat_cubin_handle = %u\n", deviceFun, hostFun, fat_cubin_handle); @@ -763,7 +755,7 @@ void cudaRegisterVarInternal( } printf("GPGPU-Sim PTX: __cudaRegisterVar: hostVar = %p; deviceAddress = %s; deviceName = %s\n", hostVar, deviceAddress, deviceName); printf("GPGPU-Sim PTX: __cudaRegisterVar: Registering const memory space of %d bytes\n", size); - if(GPGPUSim_Context()->get_device()->get_gpgpu()->get_config().use_cuobjdump()) + if(GPGPUSim_Context(ctx)->get_device()->get_gpgpu()->get_config().use_cuobjdump()) ctx->cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle); fflush(stdout); if ( constant && !global && !ext ) { @@ -872,7 +864,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx = if(g_debug_execution >= 3){ announce_call(__my_func__); } - CUctx_st* context = GPGPUSim_Context(); + CUctx_st* context = GPGPUSim_Context(ctx); char *mode = getenv("PTX_SIM_MODE_FUNC"); if( mode ) sscanf(mode,"%u", &(ctx->func_sim->g_ptx_sim_mode)); @@ -944,7 +936,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx = printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n", kname.c_str(), stream?stream->get_uid():0, gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z ); stream_operation op(grid,ctx->func_sim->g_ptx_sim_mode,stream); - g_stream_manager()->push(op); + ctx->the_gpgpusim->g_stream_manager->push(op); ctx->api->g_cuda_launch_stack.pop_back(); return g_last_cudaError = cudaSuccess; } @@ -960,7 +952,7 @@ cudaError_t cudaMallocInternal(void **devPtr, size_t size, gpgpu_context* gpgpu_ if(g_debug_execution >= 3){ announce_call(__my_func__); } - CUctx_st* context = GPGPUSim_Context(); + CUctx_st* context = GPGPUSim_Context(ctx); *devPtr = context->get_device()->get_gpgpu()->gpu_malloc(size); if(g_debug_execution >= 3){ printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *devPtr); @@ -994,6 +986,29 @@ cudaError_t cudaMallocHostInternal(void **ptr, size_t size, gpgpu_context* gpgpu } } +__host__ cudaError_t CUDARTAPI cudaMallocPitchInternal(void **devPtr, size_t *pitch, size_t width, size_t height, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + unsigned malloc_width_inbytes = width; + printf("GPGPU-Sim PTX: cudaMallocPitch (width = %d)\n", malloc_width_inbytes); + CUctx_st* context = GPGPUSim_Context(ctx); + *devPtr = context->get_device()->get_gpgpu()->gpu_malloc(malloc_width_inbytes*height); + pitch[0] = malloc_width_inbytes; + if ( *devPtr ) { + return g_last_cudaError = cudaSuccess; + } else { + return g_last_cudaError = cudaErrorMemoryAllocation; + } +} + cudaError_t cudaHostGetDevicePointerInternal(void **pDevice, void *pHost, unsigned int flags, gpgpu_context* gpgpu_ctx = NULL) { gpgpu_context *ctx; @@ -1011,7 +1026,7 @@ cudaError_t cudaHostGetDevicePointerInternal(void **pDevice, void *pHost, unsign //only cpu memory allocation happens in cudaHostAlloc. Linking with device pointer to pinned memory happens here. //TODO: once kernel is executed, the contents in global pointer of GPU must be copied back to CPU host pointer! flags=0; - CUctx_st* context = GPGPUSim_Context(); + CUctx_st* context = GPGPUSim_Context(ctx); gpgpu_t *gpu = context->get_device()->get_gpgpu(); std::map<void *, size_t>::const_iterator i = ctx->api->pinned_memory_size.find(pHost); assert(i != ctx->api->pinned_memory_size.end()); @@ -1031,7 +1046,7 @@ cudaError_t cudaHostGetDevicePointerInternal(void **pDevice, void *pHost, unsign } } -cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, gpgpu_context* gpgpu_ctx = NULL) +__host__ cudaError_t CUDARTAPI cudaMallocArrayInternal(struct cudaArray **array, const struct cudaChannelFormatDesc *desc, size_t width, size_t height __dv(1), gpgpu_context* gpgpu_ctx = NULL) { gpgpu_context *ctx; if (gpgpu_ctx){ @@ -1042,226 +1057,8 @@ cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, gpgpu if(g_debug_execution >= 3){ announce_call(__my_func__); } - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } -#ifdef OPENGL_SUPPORT - GLint buffer_size=0; - CUctx_st* context = GPGPUSim_Context(); - - glbmap_entry_t *p = ctx->api->g_glbmap; - while ( p && p->m_bufferObj != bufferObj ) - p = p->m_next; - if ( p == NULL ) { - glBindBuffer(GL_ARRAY_BUFFER,bufferObj); - glGetBufferParameteriv(GL_ARRAY_BUFFER,GL_BUFFER_SIZE,&buffer_size); - assert( buffer_size != 0 ); - *devPtr = context->get_device()->get_gpgpu()->gpu_malloc(buffer_size); - - // create entry and insert to front of list - glbmap_entry_t *n = (glbmap_entry_t *) calloc(1,sizeof(glbmap_entry_t)); - n->m_next = ctx->api->g_glbmap; - ctx->api->g_glbmap = n; - - // initialize entry - n->m_bufferObj = bufferObj; - n->m_devPtr = *devPtr; - n->m_size = buffer_size; - - p = n; - } else { - buffer_size = p->m_size; - *devPtr = p->m_devPtr; - } - - if ( *devPtr ) { - char *data = (char *) calloc(p->m_size,1); - glGetBufferSubData(GL_ARRAY_BUFFER,0,buffer_size,data); - memcpy_to_gpu( (size_t) *devPtr, data, buffer_size ); - free(data); - printf("GPGPU-Sim PTX: cudaGLMapBufferObject %zu bytes starting at 0x%llx..\n", (size_t)buffer_size, - (unsigned long long) *devPtr); - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorMemoryAllocation; - } - - return g_last_cudaError = cudaSuccess; -#else - fflush(stdout); - fflush(stderr); - printf("GPGPU-Sim PTX: GPGPU-Sim support for OpenGL integration disabled -- exiting\n"); - fflush(stdout); - exit(50); -#endif -} - -#if CUDART_VERSION >= 6050 -CUresult -cuLinkAddFileInternal(CUlinkState state, CUjitInputType type, const char *path, - unsigned int numOptions, CUjit_option *options, void **optionValues, gpgpu_context* gpgpu_ctx = NULL) -{ - gpgpu_context *ctx; - if (gpgpu_ctx){ - ctx = gpgpu_ctx; - } else { - ctx = GPGPU_Context(); - } - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - static bool addedFile = false; - if (addedFile){ - printf("GPGPU-Sim PTX: ERROR: cuLinkAddFile does not support multiple files\n"); - abort(); - } - - //blocking - assert(type==CU_JIT_INPUT_PTX); - CUctx_st *context = GPGPUSim_Context(); - char *file = getenv("PTX_JIT_PATH"); - if(file==NULL){ - printf("GPGPU-Sim PTX: ERROR: PTX_JIT_PATH has not been set\n"); - abort(); - } - strcat(file,"/"); - strcat(file,path); - symbol_table *symtab = ctx->gpgpu_ptx_sim_load_ptx_from_filename( file ); - std::string fname(path); - ctx->api->name_symtab[fname] = symtab; - context->add_binary(symtab, 1); - ctx->api->load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); - ctx->api->load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); - addedFile = true; - return CUDA_SUCCESS; -} -#endif - -#if (CUDART_VERSION >= 2010) - -cudaError_t cudaHostAllocInternal(void **pHost, size_t bytes, unsigned int flags, gpgpu_context* gpgpu_ctx = NULL) -{ - gpgpu_context *ctx; - if (gpgpu_ctx){ - ctx = gpgpu_ctx; - } else { - ctx = GPGPU_Context(); - } - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - *pHost = malloc(bytes); - //need to track the size allocated so that cudaHostGetDevicePointer() can function properly. - //TODO: vary this function behavior based on flags value (following nvidia documentation) - ctx->api->pinned_memory_size[*pHost]=bytes; - if( *pHost ) - return g_last_cudaError = cudaSuccess; - else - return g_last_cudaError = cudaErrorMemoryAllocation; -} - -#endif - -size_t getMaxThreadsPerBlock(struct cudaFuncAttributes *attr, gpgpu_context *ctx) { - _cuda_device_id *dev = ctx->GPGPUSim_Init(); - struct cudaDeviceProp prop; - - prop = *dev->get_prop(); - - size_t max = prop.maxThreadsPerBlock; - - if (attr->numRegs && (prop.regsPerBlock / attr->numRegs) < max) - max = prop.regsPerBlock / attr->numRegs; - - if (attr->sharedSizeBytes && (prop.sharedMemPerBlock / attr->sharedSizeBytes) < max) - max = prop.sharedMemPerBlock / attr->sharedSizeBytes; - - return max; -} - -cudaError_t CUDARTAPI cudaFuncGetAttributesInternal(struct cudaFuncAttributes *attr, const char *hostFun, gpgpu_context* gpgpu_ctx = NULL ) -{ - gpgpu_context *ctx; - if (gpgpu_ctx){ - ctx = gpgpu_ctx; - } else { - ctx = GPGPU_Context(); - } - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - CUctx_st *context = GPGPUSim_Context(); - function_info *entry = context->get_kernel(hostFun); - if( entry ) { - const struct gpgpu_ptx_sim_info *kinfo = entry->get_kernel_info(); - attr->sharedSizeBytes = kinfo->smem; - attr->constSizeBytes = kinfo->cmem; - attr->localSizeBytes = kinfo->lmem; - attr->numRegs = kinfo->regs; - if(kinfo->maxthreads > 0) - attr->maxThreadsPerBlock = kinfo->maxthreads; - else - attr->maxThreadsPerBlock = getMaxThreadsPerBlock(attr, ctx); -#if CUDART_VERSION >= 3000 - attr->ptxVersion = kinfo->ptx_version; - attr->binaryVersion = kinfo->sm_target; -#endif - } - return g_last_cudaError = cudaSuccess; -} - - -/******************************************************************************* - * * - * * - * * - *******************************************************************************/ - -extern "C" { - -/******************************************************************************* - * * - * * - * * - *******************************************************************************/ -cudaError_t cudaPeekAtLastError(void) -{ - return g_last_cudaError; -} - -__host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) -{ - return cudaMallocInternal(devPtr, size); -} - -__host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size) -{ - return cudaMallocHostInternal(ptr, size); -} -__host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr, size_t *pitch, size_t width, size_t height) -{ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - unsigned malloc_width_inbytes = width; - printf("GPGPU-Sim PTX: cudaMallocPitch (width = %d)\n", malloc_width_inbytes); - CUctx_st* ctx = GPGPUSim_Context(); - *devPtr = ctx->get_device()->get_gpgpu()->gpu_malloc(malloc_width_inbytes*height); - pitch[0] = malloc_width_inbytes; - if ( *devPtr ) { - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorMemoryAllocation; - } -} - -__host__ cudaError_t CUDARTAPI cudaMallocArray(struct cudaArray **array, const struct cudaChannelFormatDesc *desc, size_t width, size_t height __dv(1)) -{ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } unsigned size = width * height * ((desc->x + desc->y + desc->z + desc->w)/8); - CUctx_st* context = GPGPUSim_Context(); + CUctx_st* context = GPGPUSim_Context(ctx); (*array) = (struct cudaArray*) malloc(sizeof(struct cudaArray)); (*array)->desc = *desc; (*array)->width = width; @@ -1278,41 +1075,14 @@ __host__ cudaError_t CUDARTAPI cudaMallocArray(struct cudaArray **array, const s } } -__host__ cudaError_t CUDARTAPI cudaFree(void *devPtr) -{ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - // TODO... manage g_global_mem space? - return g_last_cudaError = cudaSuccess; -} -__host__ cudaError_t CUDARTAPI cudaFreeHost(void *ptr) +__host__ cudaError_t CUDARTAPI cudaMemcpyInternal(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, gpgpu_context* gpgpu_ctx = NULL) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); } - free (ptr); // this will crash the system if called twice - return g_last_cudaError = cudaSuccess; -} - -__host__ cudaError_t CUDARTAPI cudaFreeArray(struct cudaArray *array) -{ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - // TODO... manage g_global_mem space? - return g_last_cudaError = cudaSuccess; -}; - - -/******************************************************************************* - * * - * * - * * - *******************************************************************************/ - -__host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind) -{ if(g_debug_execution >= 3){ announce_call(__my_func__); } @@ -1321,21 +1091,21 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t cou if(g_debug_execution >= 3) printf("GPGPU-Sim PTX: cudaMemcpy(): devPtr = %p\n", dst); if( kind == cudaMemcpyHostToDevice ) - g_stream_manager()->push( stream_operation(src,(size_t)dst,count,0) ); + ctx->the_gpgpusim->g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) ); else if( kind == cudaMemcpyDeviceToHost ) - g_stream_manager()->push( stream_operation((size_t)src,dst,count,0) ); + ctx->the_gpgpusim->g_stream_manager->push( stream_operation((size_t)src,dst,count,0) ); else if( kind == cudaMemcpyDeviceToDevice ) - g_stream_manager()->push( stream_operation((size_t)src,(size_t)dst,count,0) ); + ctx->the_gpgpusim->g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) ); else if ( kind == cudaMemcpyDefault ) { if ((size_t)src >= GLOBAL_HEAP_START) { if ((size_t)dst >= GLOBAL_HEAP_START) - g_stream_manager()->push( stream_operation((size_t)src,(size_t)dst,count,0) ); // device to device + ctx->the_gpgpusim->g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) ); // device to device else - g_stream_manager()->push( stream_operation((size_t)src,dst,count,0) ); // device to host + ctx->the_gpgpusim->g_stream_manager->push( stream_operation((size_t)src,dst,count,0) ); // device to host } else { if ((size_t)dst >= GLOBAL_HEAP_START) - g_stream_manager()->push( stream_operation(src,(size_t)dst,count,0) ); + ctx->the_gpgpusim->g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) ); else { printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported transfer: host to host\n"); abort(); @@ -1349,12 +1119,18 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t cou return g_last_cudaError = cudaSuccess; } -__host__ cudaError_t CUDARTAPI cudaMemcpyToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind) +__host__ cudaError_t CUDARTAPI cudaMemcpyToArrayInternal(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind, gpgpu_context* gpgpu_ctx = NULL) { + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } if(g_debug_execution >= 3){ announce_call(__my_func__); } - CUctx_st *context = GPGPUSim_Context(); + CUctx_st *context = GPGPUSim_Context(ctx); gpgpu_t *gpu = context->get_device()->get_gpgpu(); size_t size = count; printf("GPGPU-Sim PTX: cudaMemcpyToArray\n"); @@ -1372,33 +1148,18 @@ __host__ cudaError_t CUDARTAPI cudaMemcpyToArray(struct cudaArray *dst, size_t w return g_last_cudaError = cudaSuccess; } - -__host__ cudaError_t CUDARTAPI cudaMemcpyFromArray(void *dst, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t count, enum cudaMemcpyKind kind) -{ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; -} - - -__host__ cudaError_t CUDARTAPI cudaMemcpyArrayToArray(struct cudaArray *dst, size_t wOffsetDst, size_t hOffsetDst, const struct cudaArray *src, size_t wOffsetSrc, size_t hOffsetSrc, size_t count, enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToDevice)) +__host__ cudaError_t CUDARTAPI cudaMemcpy2DInternal(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, gpgpu_context* gpgpu_ctx = NULL) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); } - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; -} - - -__host__ cudaError_t CUDARTAPI cudaMemcpy2D(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind) -{ if(g_debug_execution >= 3){ announce_call(__my_func__); } - CUctx_st *context = GPGPUSim_Context(); + CUctx_st *context = GPGPUSim_Context(ctx); gpgpu_t *gpu = context->get_device()->get_gpgpu(); size_t size = spitch*height; gpgpusim_ptx_assert( (dpitch==spitch), "different src and dst pitch not supported yet" ); @@ -1415,13 +1176,18 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2D(void *dst, size_t dpitch, const void return g_last_cudaError = cudaSuccess; } - -__host__ cudaError_t CUDARTAPI cudaMemcpy2DToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind) +__host__ cudaError_t CUDARTAPI cudaMemcpy2DToArrayInternal(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, gpgpu_context* gpgpu_ctx = NULL) { + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } if(g_debug_execution >= 3){ announce_call(__my_func__); } - CUctx_st *context = GPGPUSim_Context(); + CUctx_st *context = GPGPUSim_Context(ctx); gpgpu_t *gpu = context->get_device()->get_gpgpu(); size_t size = spitch*height; size_t channel_size = dst->desc.w+dst->desc.x+dst->desc.y+dst->desc.z; @@ -1447,29 +1213,14 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2DToArray(struct cudaArray *dst, size_t return g_last_cudaError = cudaSuccess; } - -__host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArray(void *dst, size_t dpitch, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t width, size_t height, enum cudaMemcpyKind kind) -{ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; -} - - -__host__ cudaError_t CUDARTAPI cudaMemcpy2DArrayToArray(struct cudaArray *dst, size_t wOffsetDst, size_t hOffsetDst, const struct cudaArray *src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToDevice)) +__host__ cudaError_t CUDARTAPI cudaMemcpyToSymbolInternal(const char *symbol, const void *src, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyHostToDevice), gpgpu_context* gpgpu_ctx = NULL) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); } - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; -} - - -__host__ cudaError_t CUDARTAPI cudaMemcpyToSymbol(const char *symbol, const void *src, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyHostToDevice)) -{ if(g_debug_execution >= 3){ announce_call(__my_func__); } @@ -1477,52 +1228,47 @@ __host__ cudaError_t CUDARTAPI cudaMemcpyToSymbol(const char *symbol, const void assert(kind == cudaMemcpyHostToDevice); printf("GPGPU-Sim PTX: cudaMemcpyToSymbol: symbol = %p\n", symbol); //stream_operation( const char *symbol, const void *src, size_t count, size_t offset ) - g_stream_manager()->push( stream_operation(src,symbol,count,offset,0) ); + ctx->the_gpgpusim->g_stream_manager->push( stream_operation(src,symbol,count,offset,0) ); //gpgpu_ptx_sim_memcpy_symbol(symbol,src,count,offset,1,context->get_device()->get_gpgpu()); return g_last_cudaError = cudaSuccess; } -__host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbol(void *dst, const char *symbol, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost)) +__host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbolInternal(void *dst, const char *symbol, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost), gpgpu_context* gpgpu_ctx = NULL) { + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } if(g_debug_execution >= 3){ announce_call(__my_func__); } //CUctx_st *context = GPGPUSim_Context(); assert(kind == cudaMemcpyDeviceToHost); printf("GPGPU-Sim PTX: cudaMemcpyFromSymbol: symbol = %p\n", symbol); - g_stream_manager()->push( stream_operation(symbol,dst,count,offset,0) ); + ctx->the_gpgpusim->g_stream_manager->push( stream_operation(symbol,dst,count,offset,0) ); //gpgpu_ptx_sim_memcpy_symbol(symbol,dst,count,offset,0,context->get_device()->get_gpgpu()); return g_last_cudaError = cudaSuccess; } -__host__ cudaError_t CUDARTAPI cudaMemGetInfo (size_t *free, size_t *total){ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - //placeholder; should interact with cudaMalloc and cudaFree? - *free = 10000000000; - *total = 10000000000; - - return g_last_cudaError = cudaSuccess; -} - -/******************************************************************************* - * * - * * - * * - *******************************************************************************/ - -__host__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream) +__host__ cudaError_t CUDARTAPI cudaMemcpyAsyncInternal(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream, gpgpu_context* gpgpu_ctx = NULL) { + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } if(g_debug_execution >= 3){ announce_call(__my_func__); } struct CUstream_st *s = (struct CUstream_st *)stream; switch( kind ) { - case cudaMemcpyHostToDevice: g_stream_manager()->push( stream_operation(src,(size_t)dst,count,s) ); break; - case cudaMemcpyDeviceToHost: g_stream_manager()->push( stream_operation((size_t)src,dst,count,s) ); break; - case cudaMemcpyDeviceToDevice: g_stream_manager()->push( stream_operation((size_t)src,(size_t)dst,count,s) ); break; + case cudaMemcpyHostToDevice: ctx->the_gpgpusim->g_stream_manager->push( stream_operation(src,(size_t)dst,count,s) ); break; + case cudaMemcpyDeviceToHost: ctx->the_gpgpusim->g_stream_manager->push( stream_operation((size_t)src,dst,count,s) ); break; + case cudaMemcpyDeviceToDevice: ctx->the_gpgpusim->g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,s) ); break; default: abort(); } @@ -1530,98 +1276,249 @@ __host__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_ } -__host__ cudaError_t CUDARTAPI cudaMemcpyToArrayAsync(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream) +#if (CUDART_VERSION >= 8000) +cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlagsInternal(int* numBlocks, const char *hostFunc, int blockSize, size_t dynamicSMemSize, unsigned int flags, gpgpu_context* gpgpu_ctx = NULL) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); } - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + printf("GPGPU-Sim PTX: cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags %p\n", hostFunc); + CUctx_st *context = GPGPUSim_Context(ctx); + function_info *entry = context->get_kernel(hostFunc); + printf("Calculate Maxium Active Block with function ptr=%p, blockSize=%d, SMemSize=%d\n", hostFunc, blockSize, dynamicSMemSize); + if (flags == cudaOccupancyDefault) { + //create kernel_info based on entry + dim3 gridDim(context->get_device()->get_gpgpu()->max_cta_per_core() + * context->get_device()->get_gpgpu()->get_config().num_shader()); + dim3 blockDim(blockSize); + kernel_info_t result(gridDim, blockDim, entry); + //if(entry == NULL){ + // *numBlocks = 1; + // return g_last_cudaError = cudaErrorUnknown; + //} + *numBlocks = context->get_device()->get_gpgpu()->get_max_cta(result); + printf("Maximum size is %d with gridDim %d and blockDim %d\n", *numBlocks, gridDim.x, blockDim.x); + return g_last_cudaError = cudaSuccess; + } else { + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; + } } +#endif -__host__ cudaError_t CUDARTAPI cudaMemcpyFromArrayAsync(void *dst, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream) +__host__ cudaError_t CUDARTAPI cudaMemsetInternal(void *mem, int c, size_t count, gpgpu_context* gpgpu_ctx = NULL) { + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } if(g_debug_execution >= 3){ announce_call(__my_func__); } - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + CUctx_st *context = GPGPUSim_Context(ctx); + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + gpu->gpu_memset((size_t)mem, c, count); + return g_last_cudaError = cudaSuccess; } - -__host__ cudaError_t CUDARTAPI cudaMemcpy2DAsync(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream) +//memset operation is done but i think its not async? +__host__ cudaError_t CUDARTAPI cudaMemsetAsyncInternal(void *mem, int c, size_t count, cudaStream_t stream=0, gpgpu_context* gpgpu_ctx = NULL) { + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } if(g_debug_execution >= 3){ announce_call(__my_func__); } - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + printf("GPGPU-Sim PTX: WARNING: Asynchronous memset not supported (%s)\n", __my_func__); + CUctx_st *context = GPGPUSim_Context(ctx); + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + gpu->gpu_memset((size_t)mem, c, count); + return g_last_cudaError = cudaSuccess; } - -__host__ cudaError_t CUDARTAPI cudaMemcpy2DToArrayAsync(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream) +cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, gpgpu_context* gpgpu_ctx = NULL) { if(g_debug_execution >= 3){ announce_call(__my_func__); } - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; -} +#ifdef OPENGL_SUPPORT + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + GLint buffer_size=0; + CUctx_st* context = GPGPUSim_Context(ctx); + + glbmap_entry_t *p = ctx->api->g_glbmap; + while ( p && p->m_bufferObj != bufferObj ) + p = p->m_next; + if ( p == NULL ) { + glBindBuffer(GL_ARRAY_BUFFER,bufferObj); + glGetBufferParameteriv(GL_ARRAY_BUFFER,GL_BUFFER_SIZE,&buffer_size); + assert( buffer_size != 0 ); + *devPtr = context->get_device()->get_gpgpu()->gpu_malloc(buffer_size); + // create entry and insert to front of list + glbmap_entry_t *n = (glbmap_entry_t *) calloc(1,sizeof(glbmap_entry_t)); + n->m_next = ctx->api->g_glbmap; + ctx->api->g_glbmap = n; -__host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArrayAsync(void *dst, size_t dpitch, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream) + // initialize entry + n->m_bufferObj = bufferObj; + n->m_devPtr = *devPtr; + n->m_size = buffer_size; + + p = n; + } else { + buffer_size = p->m_size; + *devPtr = p->m_devPtr; + } + + if ( *devPtr ) { + char *data = (char *) calloc(p->m_size,1); + glGetBufferSubData(GL_ARRAY_BUFFER,0,buffer_size,data); + memcpy_to_gpu( (size_t) *devPtr, data, buffer_size ); + free(data); + printf("GPGPU-Sim PTX: cudaGLMapBufferObject %zu bytes starting at 0x%llx..\n", (size_t)buffer_size, + (unsigned long long) *devPtr); + return g_last_cudaError = cudaSuccess; + } else { + return g_last_cudaError = cudaErrorMemoryAllocation; + } + + return g_last_cudaError = cudaSuccess; +#else + fflush(stdout); + fflush(stderr); + printf("GPGPU-Sim PTX: GPGPU-Sim support for OpenGL integration disabled -- exiting\n"); + fflush(stdout); + exit(50); +#endif +} + +#if CUDART_VERSION >= 6050 +CUresult +cuLinkAddFileInternal(CUlinkState state, CUjitInputType type, const char *path, + unsigned int numOptions, CUjit_option *options, void **optionValues, gpgpu_context* gpgpu_ctx = NULL) { + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } if(g_debug_execution >= 3){ announce_call(__my_func__); } - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; + static bool addedFile = false; + if (addedFile){ + printf("GPGPU-Sim PTX: ERROR: cuLinkAddFile does not support multiple files\n"); + abort(); + } + + //blocking + assert(type==CU_JIT_INPUT_PTX); + CUctx_st *context = GPGPUSim_Context(ctx); + char *file = getenv("PTX_JIT_PATH"); + if(file==NULL){ + printf("GPGPU-Sim PTX: ERROR: PTX_JIT_PATH has not been set\n"); + abort(); + } + strcat(file,"/"); + strcat(file,path); + symbol_table *symtab = ctx->gpgpu_ptx_sim_load_ptx_from_filename( file ); + std::string fname(path); + ctx->api->name_symtab[fname] = symtab; + context->add_binary(symtab, 1); + ctx->api->load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); + ctx->api->load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); + addedFile = true; + return CUDA_SUCCESS; } +#endif -#if (CUDART_VERSION >= 8000) -cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks, const char *hostFunc, int blockSize, size_t dynamicSMemSize, unsigned int flags) +#if (CUDART_VERSION >= 2010) + +cudaError_t cudaHostAllocInternal(void **pHost, size_t bytes, unsigned int flags, gpgpu_context* gpgpu_ctx = NULL) { - printf("GPGPU-Sim PTX: cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags %p\n", hostFunc); - CUctx_st *context = GPGPUSim_Context(); - function_info *entry = context->get_kernel(hostFunc); - printf("Calculate Maxium Active Block with function ptr=%p, blockSize=%d, SMemSize=%d\n", hostFunc, blockSize, dynamicSMemSize); - if (flags == cudaOccupancyDefault) { - //create kernel_info based on entry - dim3 gridDim(context->get_device()->get_gpgpu()->max_cta_per_core() - * context->get_device()->get_gpgpu()->get_config().num_shader()); - dim3 blockDim(blockSize); - kernel_info_t result(gridDim, blockDim, entry); - //if(entry == NULL){ - // *numBlocks = 1; - // return g_last_cudaError = cudaErrorUnknown; - //} - *numBlocks = context->get_device()->get_gpgpu()->get_max_cta(result); - printf("Maximum size is %d with gridDim %d and blockDim %d\n", *numBlocks, gridDim.x, blockDim.x); + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + *pHost = malloc(bytes); + //need to track the size allocated so that cudaHostGetDevicePointer() can function properly. + //TODO: vary this function behavior based on flags value (following nvidia documentation) + ctx->api->pinned_memory_size[*pHost]=bytes; + if( *pHost ) return g_last_cudaError = cudaSuccess; - } else { - cuda_not_implemented(__my_func__,__LINE__); - return g_last_cudaError = cudaErrorUnknown; - } + else + return g_last_cudaError = cudaErrorMemoryAllocation; } #endif +size_t getMaxThreadsPerBlock(struct cudaFuncAttributes *attr, gpgpu_context *ctx) { + _cuda_device_id *dev = ctx->GPGPUSim_Init(); + struct cudaDeviceProp prop; + prop = *dev->get_prop(); -/******************************************************************************* - * * - * * - * * - *******************************************************************************/ -__host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count) + size_t max = prop.maxThreadsPerBlock; + + if (attr->numRegs && (prop.regsPerBlock / attr->numRegs) < max) + max = prop.regsPerBlock / attr->numRegs; + + if (attr->sharedSizeBytes && (prop.sharedMemPerBlock / attr->sharedSizeBytes) < max) + max = prop.sharedMemPerBlock / attr->sharedSizeBytes; + + return max; +} + +cudaError_t CUDARTAPI cudaFuncGetAttributesInternal(struct cudaFuncAttributes *attr, const char *hostFun, gpgpu_context* gpgpu_ctx = NULL ) { + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } if(g_debug_execution >= 3){ announce_call(__my_func__); } - CUctx_st *context = GPGPUSim_Context(); - gpgpu_t *gpu = context->get_device()->get_gpgpu(); - gpu->gpu_memset((size_t)mem, c, count); + CUctx_st *context = GPGPUSim_Context(ctx); + function_info *entry = context->get_kernel(hostFun); + if( entry ) { + const struct gpgpu_ptx_sim_info *kinfo = entry->get_kernel_info(); + attr->sharedSizeBytes = kinfo->smem; + attr->constSizeBytes = kinfo->cmem; + attr->localSizeBytes = kinfo->lmem; + attr->numRegs = kinfo->regs; + if(kinfo->maxthreads > 0) + attr->maxThreadsPerBlock = kinfo->maxthreads; + else + attr->maxThreadsPerBlock = getMaxThreadsPerBlock(attr, ctx); +#if CUDART_VERSION >= 3000 + attr->ptxVersion = kinfo->ptx_version; + attr->binaryVersion = kinfo->sm_target; +#endif + } return g_last_cudaError = cudaSuccess; } @@ -1826,19 +1723,652 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttributeInternal(int *value, enum c } #endif -//memset operation is done but i think its not async? -__host__ cudaError_t CUDARTAPI cudaMemsetAsync(void *mem, int c, size_t count, cudaStream_t stream=0) +__host__ cudaError_t CUDARTAPI cudaBindTextureInternal(size_t *offset, + const struct textureReference *texref, + const void *devPtr, + const struct cudaChannelFormatDesc *desc, + size_t size __dv(UINT_MAX), + gpgpu_context* gpgpu_ctx = NULL) { + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } if(g_debug_execution >= 3){ announce_call(__my_func__); } - printf("GPGPU-Sim PTX: WARNING: Asynchronous memset not supported (%s)\n", __my_func__); - CUctx_st *context = GPGPUSim_Context(); + CUctx_st *context = GPGPUSim_Context(ctx); gpgpu_t *gpu = context->get_device()->get_gpgpu(); - gpu->gpu_memset((size_t)mem, c, count); + printf("GPGPU-Sim PTX: in cudaBindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference)); + struct cudaArray *array; + array = (struct cudaArray*) malloc(sizeof(struct cudaArray)); + array->desc = *desc; + array->size = size; + array->width = size; + array->height = 1; + array->dimensions = 1; + array->devPtr = (void*)devPtr; + array->devPtr32 = (int)(long long)devPtr; + offset = 0; + printf("GPGPU-Sim PTX: size = %zu\n", size); + printf("GPGPU-Sim PTX: texref = %p, array = %p\n", texref, array); + printf("GPGPU-Sim PTX: devPtr32 = %x\n", array->devPtr32); + printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref)); + printf("GPGPU-Sim PTX: ChannelFormatDesc: x=%d, y=%d, z=%d, w=%d\n", desc->x, desc->y, desc->z, desc->w); + printf("GPGPU-Sim PTX: Texture Normalized? = %d\n", texref->normalized); + gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array); + devPtr = (void*)(long long)array->devPtr32; + printf("GPGPU-Sim PTX: devPtr = %p\n", devPtr); return g_last_cudaError = cudaSuccess; } +__host__ cudaError_t CUDARTAPI cudaBindTextureToArrayInternal(const struct textureReference *texref, const struct cudaArray *array, const struct cudaChannelFormatDesc *desc, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + CUctx_st *context = GPGPUSim_Context(ctx); + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + printf("GPGPU-Sim PTX: in cudaBindTextureToArray: %p %p\n", texref, array); + printf("GPGPU-Sim PTX: devPtr32 = %x\n", array->devPtr32); + printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref)); + printf("GPGPU-Sim PTX: Texture Normalized? = %d\n", texref->normalized); + gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array); + return g_last_cudaError = cudaSuccess; +} + +__host__ cudaError_t CUDARTAPI cudaUnbindTextureInternal(const struct textureReference *texref, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + CUctx_st *context = GPGPUSim_Context(ctx); + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + printf("GPGPU-Sim PTX: in cudaUnbindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference)); + printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref)); + + gpu->gpgpu_ptx_sim_unbindTexture(texref); + return g_last_cudaError = cudaSuccess; +} + +__host__ cudaError_t CUDARTAPI cudaLaunchKernelInternal( const char* hostFun, dim3 gridDim, dim3 blockDim, const void** args, size_t sharedMem, cudaStream_t stream, gpgpu_context* gpgpu_ctx = NULL ) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + CUctx_st *context = GPGPUSim_Context(ctx); + function_info *entry = context->get_kernel(hostFun); +#if CUDART_VERSION < 10000 + cudaConfigureCallInternal(gridDim, blockDim, sharedMem, stream, ctx); +#endif + for(unsigned i = 0; i < entry->num_args(); i++){ + std::pair<size_t, unsigned> p = entry->get_param_config(i); + cudaSetupArgumentInternal(args[i], p.first, p.second); + } + + cudaLaunchInternal(hostFun); + return g_last_cudaError = cudaSuccess; +} + +__host__ cudaError_t CUDARTAPI cudaStreamCreateInternal(cudaStream_t *stream, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + printf("GPGPU-Sim PTX: cudaStreamCreate\n"); +#if (CUDART_VERSION >= 3000) + *stream = new struct CUstream_st(); + ctx->the_gpgpusim->g_stream_manager->add_stream(*stream); +#else + *stream = 0; + printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); +#endif + return g_last_cudaError = cudaSuccess; +} + +__host__ cudaError_t CUDARTAPI cudaStreamDestroyInternal(cudaStream_t stream, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } +#if (CUDART_VERSION >= 3000) + //per-stream synchronization required for application using external libraries without explicit synchronization in the code to + //avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. + stream->synchronize(); + ctx->the_gpgpusim->g_stream_manager->destroy_stream(stream); +#endif + return g_last_cudaError = cudaSuccess; +} + +__host__ cudaError_t CUDARTAPI cudaStreamSynchronizeInternal(cudaStream_t stream, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } +#if (CUDART_VERSION >= 3000) + if( stream == NULL ) + ctx->synchronize(); + return g_last_cudaError = cudaSuccess; + stream->synchronize(); +#else + printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); +#endif + return g_last_cudaError = cudaSuccess; +} + +void __cudaRegisterTextureInternal( + void **fatCubinHandle, + const struct textureReference *hostVar, + const void **deviceAddress, + const char *deviceName, + int dim, + int norm, + int ext, + gpgpu_context* gpgpu_ctx = NULL +) //passes in a newly created textureReference +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + std::string devStr (deviceName); + #if (CUDART_VERSION > 4020) + if (devStr.size() > 2 && devStr.data()[0] == ':' && devStr.data()[1] == ':') + devStr = devStr.replace(0, 2, ""); + #endif + CUctx_st *context = GPGPUSim_Context(ctx); + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + printf("GPGPU-Sim PTX: in __cudaRegisterTexture:\n"); + gpu->gpgpu_ptx_sim_bindNameToTexture(devStr.data(), hostVar, dim, norm, ext); + printf("GPGPU-Sim PTX: int dim = %d\n", dim); + printf("GPGPU-Sim PTX: int norm = %d\n", norm); + printf("GPGPU-Sim PTX: int ext = %d\n", ext); + printf("GPGPU-Sim PTX: Execution warning: Not finished implementing \"%s\"\n", __my_func__ ); +} + +cudaError_t cudaGLUnmapBufferObjectInternal(GLuint bufferObj, gpgpu_context* gpgpu_ctx = NULL) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } +#ifdef OPENGL_SUPPORT + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + CUctx_st* ctx = GPGPUSim_Context(ctx); + glbmap_entry_t *p = ctx->api->g_glbmap; + while ( p && p->m_bufferObj != bufferObj ) + p = p->m_next; + if ( p == NULL ) + return g_last_cudaError = cudaErrorUnknown; + + char *data = (char *) calloc(p->m_size,1); + memcpy_from_gpu( data,(size_t)p->m_devPtr,p->m_size ); + glBufferSubData(GL_ARRAY_BUFFER,0,p->m_size,data); + free(data); + + return g_last_cudaError = cudaSuccess; +#else + fflush(stdout); + fflush(stderr); + printf("GPGPU-Sim PTX: support for OpenGL integration disabled -- exiting\n"); + fflush(stdout); + exit(50); +#endif +} + +#if CUDART_VERSION >= 3000 + +__host__ cudaError_t CUDARTAPI cudaFuncSetCacheConfigInternal(const char *func, enum cudaFuncCache cacheConfig, gpgpu_context* gpgpu_ctx = NULL ) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + CUctx_st *context = GPGPUSim_Context(ctx); + context->get_device()->get_gpgpu()->set_cache_config(context->get_kernel(func)->get_name(), (FuncCache)cacheConfig); + return g_last_cudaError = cudaSuccess; +} + +#endif + +#if CUDART_VERSION >= 4000 +CUresult CUDAAPI cuLaunchKernelInternal(CUfunction f, + unsigned int gridDimX, + unsigned int gridDimY, + unsigned int gridDimZ, + unsigned int blockDimX, + unsigned int blockDimY, + unsigned int blockDimZ, + unsigned int sharedMemBytes, + CUstream hStream, + void **kernelParams, + void **extra, + gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + if (extra!=NULL){ + printf("GPGPU-Sim CUDA DRIVER API: ERROR: Currently do not support void** extra.\n"); + abort(); + } + const char *hostFun = (const char*) f; + CUctx_st *context = GPGPUSim_Context(ctx); + function_info *entry = context->get_kernel(hostFun); + cudaConfigureCallInternal(dim3(gridDimX, gridDimY, gridDimZ), dim3(blockDimX, blockDimY, blockDimZ), sharedMemBytes, (cudaStream_t) hStream, ctx); + for(unsigned i = 0; i < entry->num_args(); i++){ + std::pair<size_t, unsigned> p = entry->get_param_config(i); + cudaSetupArgumentInternal(kernelParams[i], p.first, p.second, ctx); + } + cudaLaunchInternal(hostFun, ctx); + return CUDA_SUCCESS; +} +#endif /* CUDART_VERSION >= 4000 */ + +CUevent_st *get_event(cudaEvent_t event) +{ + unsigned event_uid; +#if CUDART_VERSION >= 3000 + event_uid = event->get_uid(); +#else + event_uid = event; +#endif + event_tracker_t::iterator e = g_timer_events.find(event_uid); + if( e == g_timer_events.end() ) + return NULL; + return e->second; +} + +__host__ cudaError_t CUDARTAPI cudaEventRecordInternal(cudaEvent_t event, cudaStream_t stream, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + CUevent_st *e = get_event(event); + if( !e ) return g_last_cudaError = cudaErrorUnknown; + struct CUstream_st *s = (struct CUstream_st *)stream; + stream_operation op(e,s); + ctx->the_gpgpusim->g_stream_manager->push(op); + return g_last_cudaError = cudaSuccess; +} + +__host__ cudaError_t CUDARTAPI cudaStreamWaitEventInternal(cudaStream_t stream, cudaEvent_t event, unsigned int flags, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + //reference: https://www.cs.cmu.edu/afs/cs/academic/class/15668-s11/www/cuda-doc/html/group__CUDART__STREAM_gfe68d207dc965685d92d3f03d77b0876.html + CUevent_st *e = get_event(event); + if( !e ){ + printf("GPGPU-Sim API: Warning: cudaEventRecord has not been called on event before calling cudaStreamWaitEvent.\nNothing to be done.\n"); + return g_last_cudaError = cudaSuccess; + } + if (!stream){ + ctx->the_gpgpusim->g_stream_manager->pushCudaStreamWaitEventToAllStreams(e, flags); + } else { + struct CUstream_st *s = (struct CUstream_st *)stream; + stream_operation op(s,e,flags); + ctx->the_gpgpusim->g_stream_manager->push(op); + } + return g_last_cudaError = cudaSuccess; +} + +__host__ cudaError_t CUDARTAPI cudaThreadExitInternal(gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + ctx->exit_simulation(); + return g_last_cudaError = cudaSuccess; +} + +__host__ cudaError_t CUDARTAPI cudaThreadSynchronizeInternal(gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + //Called on host side + ctx->synchronize(); + return g_last_cudaError = cudaSuccess; +} + +cudaError_t CUDARTAPI cudaDeviceSynchronizeInternal(gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + //Blocks until the device has completed all preceding requested tasks + ctx->synchronize(); + return g_last_cudaError = cudaSuccess; +} + +/******************************************************************************* + * * + * * + * * + *******************************************************************************/ + +extern "C" { + +/******************************************************************************* + * * + * * + * * + *******************************************************************************/ +cudaError_t cudaPeekAtLastError(void) +{ + return g_last_cudaError; +} + +__host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) +{ + return cudaMallocInternal(devPtr, size); +} + +__host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size) +{ + return cudaMallocHostInternal(ptr, size); +} +__host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr, size_t *pitch, size_t width, size_t height) +{ + return cudaMallocPitchInternal(devPtr, pitch, width, height); +} + +__host__ cudaError_t CUDARTAPI cudaMallocArray(struct cudaArray **array, const struct cudaChannelFormatDesc *desc, size_t width, size_t height __dv(1)) +{ + return cudaMallocArrayInternal(array, desc, width, height); +} + +__host__ cudaError_t CUDARTAPI cudaFree(void *devPtr) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + // TODO... manage g_global_mem space? + return g_last_cudaError = cudaSuccess; +} +__host__ cudaError_t CUDARTAPI cudaFreeHost(void *ptr) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + free (ptr); // this will crash the system if called twice + return g_last_cudaError = cudaSuccess; +} + +__host__ cudaError_t CUDARTAPI cudaFreeArray(struct cudaArray *array) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + // TODO... manage g_global_mem space? + return g_last_cudaError = cudaSuccess; +}; + + +/******************************************************************************* + * * + * * + * * + *******************************************************************************/ + +__host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind) +{ + return cudaMemcpyInternal(dst, src, count, kind); +} + +__host__ cudaError_t CUDARTAPI cudaMemcpyToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind) +{ + return cudaMemcpyToArrayInternal(dst, wOffset, hOffset, src, count, kind); +} + + +__host__ cudaError_t CUDARTAPI cudaMemcpyFromArray(void *dst, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t count, enum cudaMemcpyKind kind) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; +} + + +__host__ cudaError_t CUDARTAPI cudaMemcpyArrayToArray(struct cudaArray *dst, size_t wOffsetDst, size_t hOffsetDst, const struct cudaArray *src, size_t wOffsetSrc, size_t hOffsetSrc, size_t count, enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToDevice)) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; +} + + +__host__ cudaError_t CUDARTAPI cudaMemcpy2D(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind) +{ + return cudaMemcpy2DInternal(dst, dpitch, src, spitch, width, height, kind); +} + +__host__ cudaError_t CUDARTAPI cudaMemcpy2DToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind) +{ + return cudaMemcpy2DToArrayInternal(dst, wOffset, hOffset, src, spitch, width, height, kind); +} + +__host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArray(void *dst, size_t dpitch, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t width, size_t height, enum cudaMemcpyKind kind) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; +} + +__host__ cudaError_t CUDARTAPI cudaMemcpy2DArrayToArray(struct cudaArray *dst, size_t wOffsetDst, size_t hOffsetDst, const struct cudaArray *src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToDevice)) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; +} + +__host__ cudaError_t CUDARTAPI cudaMemcpyToSymbol(const char *symbol, const void *src, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyHostToDevice)) +{ + return cudaMemcpyToSymbolInternal(symbol, src, count, offset, kind); +} + + +__host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbol(void *dst, const char *symbol, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost)) +{ + return cudaMemcpyFromSymbolInternal(dst, symbol, count, offset, kind); +} + +__host__ cudaError_t CUDARTAPI cudaMemGetInfo (size_t *free, size_t *total){ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + //placeholder; should interact with cudaMalloc and cudaFree? + *free = 10000000000; + *total = 10000000000; + + return g_last_cudaError = cudaSuccess; +} + +/******************************************************************************* + * * + * * + * * + *******************************************************************************/ + +__host__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream) +{ + return cudaMemcpyAsyncInternal(dst, src, count, kind, stream); +} + + +__host__ cudaError_t CUDARTAPI cudaMemcpyToArrayAsync(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; +} + + +__host__ cudaError_t CUDARTAPI cudaMemcpyFromArrayAsync(void *dst, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; +} + + +__host__ cudaError_t CUDARTAPI cudaMemcpy2DAsync(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; +} + + +__host__ cudaError_t CUDARTAPI cudaMemcpy2DToArrayAsync(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; +} + + +__host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArrayAsync(void *dst, size_t dpitch, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream) +{ + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + cuda_not_implemented(__my_func__,__LINE__); + return g_last_cudaError = cudaErrorUnknown; +} + +#if (CUDART_VERSION >= 8000) +cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks, const char *hostFunc, int blockSize, size_t dynamicSMemSize, unsigned int flags) +{ + return cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlagsInternal(numBlocks, hostFunc, blockSize, dynamicSMemSize, flags); +} + +#endif + + + +/******************************************************************************* + * * + * * + * * + *******************************************************************************/ +__host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count) +{ + return cudaMemsetInternal(mem, c, count); +} + +//memset operation is done but i think its not async? +__host__ cudaError_t CUDARTAPI cudaMemsetAsync(void *mem, int c, size_t count, cudaStream_t stream=0) +{ + return cudaMemsetAsyncInternal(mem, c, count, stream=0); +} + __host__ cudaError_t CUDARTAPI cudaMemset2D(void *mem, size_t pitch, int c, size_t width, size_t height) { if(g_debug_execution >= 3){ @@ -1986,61 +2516,18 @@ __host__ cudaError_t CUDARTAPI cudaBindTexture(size_t *offset, const struct cudaChannelFormatDesc *desc, size_t size __dv(UINT_MAX)) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - CUctx_st *context = GPGPUSim_Context(); - gpgpu_t *gpu = context->get_device()->get_gpgpu(); - printf("GPGPU-Sim PTX: in cudaBindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference)); - struct cudaArray *array; - array = (struct cudaArray*) malloc(sizeof(struct cudaArray)); - array->desc = *desc; - array->size = size; - array->width = size; - array->height = 1; - array->dimensions = 1; - array->devPtr = (void*)devPtr; - array->devPtr32 = (int)(long long)devPtr; - offset = 0; - printf("GPGPU-Sim PTX: size = %zu\n", size); - printf("GPGPU-Sim PTX: texref = %p, array = %p\n", texref, array); - printf("GPGPU-Sim PTX: devPtr32 = %x\n", array->devPtr32); - printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref)); - printf("GPGPU-Sim PTX: ChannelFormatDesc: x=%d, y=%d, z=%d, w=%d\n", desc->x, desc->y, desc->z, desc->w); - printf("GPGPU-Sim PTX: Texture Normalized? = %d\n", texref->normalized); - gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array); - devPtr = (void*)(long long)array->devPtr32; - printf("GPGPU-Sim PTX: devPtr = %p\n", devPtr); - return g_last_cudaError = cudaSuccess; + return cudaBindTextureInternal(offset, texref, devPtr, desc, size __dv(UINT_MAX)); } __host__ cudaError_t CUDARTAPI cudaBindTextureToArray(const struct textureReference *texref, const struct cudaArray *array, const struct cudaChannelFormatDesc *desc) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - CUctx_st *context = GPGPUSim_Context(); - gpgpu_t *gpu = context->get_device()->get_gpgpu(); - printf("GPGPU-Sim PTX: in cudaBindTextureToArray: %p %p\n", texref, array); - printf("GPGPU-Sim PTX: devPtr32 = %x\n", array->devPtr32); - printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref)); - printf("GPGPU-Sim PTX: Texture Normalized? = %d\n", texref->normalized); - gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array); - return g_last_cudaError = cudaSuccess; + return cudaBindTextureToArrayInternal(texref, array, desc); } -__host__ cudaError_t CUDARTAPI cudaUnbindTexture(const struct textureReference *texref){ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - CUctx_st *context = GPGPUSim_Context(); - gpgpu_t *gpu = context->get_device()->get_gpgpu(); - printf("GPGPU-Sim PTX: in cudaUnbindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference)); - printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref)); - - gpu->gpgpu_ptx_sim_unbindTexture(texref); - return g_last_cudaError = cudaSuccess; +__host__ cudaError_t CUDARTAPI cudaUnbindTexture(const struct textureReference *texref) +{ + return cudaUnbindTextureInternal(texref); } __host__ cudaError_t CUDARTAPI cudaGetTextureAlignmentOffset(size_t *offset, const struct textureReference *texref) @@ -2125,24 +2612,9 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) return cudaLaunchInternal( hostFun ); } -__host__ cudaError_t CUDARTAPI cudaLaunchKernel ( const char* hostFun, dim3 gridDim, dim3 blockDim, const void** args, size_t sharedMem, cudaStream_t stream ) +__host__ cudaError_t CUDARTAPI cudaLaunchKernel( const char* hostFun, dim3 gridDim, dim3 blockDim, const void** args, size_t sharedMem, cudaStream_t stream ) { - - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - CUctx_st *context = GPGPUSim_Context(); - function_info *entry = context->get_kernel(hostFun); -#if CUDART_VERSION < 10000 - cudaConfigureCallInternal(gridDim, blockDim, sharedMem, stream); -#endif - for(unsigned i = 0; i < entry->num_args(); i++){ - std::pair<size_t, unsigned> p = entry->get_param_config(i); - cudaSetupArgumentInternal(args[i], p.first, p.second); - } - - cudaLaunchInternal(hostFun); - return g_last_cudaError = cudaSuccess; + return cudaLaunchKernelInternal(hostFun, gridDim, blockDim, args, sharedMem, stream); } @@ -2154,18 +2626,7 @@ __host__ cudaError_t CUDARTAPI cudaLaunchKernel ( const char* hostFun, dim3 grid __host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - printf("GPGPU-Sim PTX: cudaStreamCreate\n"); -#if (CUDART_VERSION >= 3000) - *stream = new struct CUstream_st(); - g_stream_manager()->add_stream(*stream); -#else - *stream = 0; - printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); -#endif - return g_last_cudaError = cudaSuccess; + return cudaStreamCreateInternal(stream); } //TODO: introduce priorities @@ -2192,32 +2653,12 @@ __host__ __device__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } -#if (CUDART_VERSION >= 3000) - //per-stream synchronization required for application using external libraries without explicit synchronization in the code to - //avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. - stream->synchronize(); - g_stream_manager()->destroy_stream(stream); -#endif - return g_last_cudaError = cudaSuccess; + return cudaStreamDestroyInternal(stream); } __host__ cudaError_t CUDARTAPI cudaStreamSynchronize(cudaStream_t stream) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } -#if (CUDART_VERSION >= 3000) - if( stream == NULL ) - synchronize(); - return g_last_cudaError = cudaSuccess; - stream->synchronize(); -#else - printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); -#endif - return g_last_cudaError = cudaSuccess; + return cudaStreamSynchronizeInternal(stream); } __host__ cudaError_t CUDARTAPI cudaStreamQuery(cudaStream_t stream) @@ -2256,52 +2697,14 @@ __host__ cudaError_t CUDARTAPI cudaEventCreate(cudaEvent_t *event) return g_last_cudaError = cudaSuccess; } -CUevent_st *get_event(cudaEvent_t event) -{ - unsigned event_uid; -#if CUDART_VERSION >= 3000 - event_uid = event->get_uid(); -#else - event_uid = event; -#endif - event_tracker_t::iterator e = g_timer_events.find(event_uid); - if( e == g_timer_events.end() ) - return NULL; - return e->second; -} - __host__ cudaError_t CUDARTAPI cudaEventRecord(cudaEvent_t event, cudaStream_t stream) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - CUevent_st *e = get_event(event); - if( !e ) return g_last_cudaError = cudaErrorUnknown; - struct CUstream_st *s = (struct CUstream_st *)stream; - stream_operation op(e,s); - g_stream_manager()->push(op); - return g_last_cudaError = cudaSuccess; + return cudaEventRecordInternal(event, stream); } __host__ cudaError_t CUDARTAPI cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event, unsigned int flags) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - //reference: https://www.cs.cmu.edu/afs/cs/academic/class/15668-s11/www/cuda-doc/html/group__CUDART__STREAM_gfe68d207dc965685d92d3f03d77b0876.html - CUevent_st *e = get_event(event); - if( !e ){ - printf("GPGPU-Sim API: Warning: cudaEventRecord has not been called on event before calling cudaStreamWaitEvent.\nNothing to be done.\n"); - return g_last_cudaError = cudaSuccess; - } - if (!stream){ - g_stream_manager()->pushCudaStreamWaitEventToAllStreams(e, flags); - } else { - struct CUstream_st *s = (struct CUstream_st *)stream; - stream_operation op(s,e,flags); - g_stream_manager()->push(op); - } - return g_last_cudaError = cudaSuccess; + return cudaStreamWaitEventInternal(stream, event, flags); } __host__ cudaError_t CUDARTAPI cudaEventQuery(cudaEvent_t event) @@ -2374,22 +2777,13 @@ __host__ cudaError_t CUDARTAPI cudaEventElapsedTime(float *ms, cudaEvent_t start __host__ cudaError_t CUDARTAPI cudaThreadExit(void) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - exit_simulation(); - return g_last_cudaError = cudaSuccess; + return cudaThreadExitInternal(); } __host__ cudaError_t CUDARTAPI cudaThreadSynchronize(void) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - //Called on host side - synchronize(); - return g_last_cudaError = cudaSuccess; -}; + return cudaThreadSynchronizeInternal(); +} int CUDARTAPI __cudaSynchronizeThreads(void**, void*) { @@ -2502,7 +2896,6 @@ void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context){ while (std::getline(infile, line)) { //int pos = line.find(std::string(get_app_binary_name(app_binary))); - const char *ptx_file = line.c_str(); int pos1 = line.find("sm_"); int pos2 = line.find_last_of("."); if (pos1==std::string::npos&&pos2==std::string::npos){ @@ -2529,12 +2922,11 @@ void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context){ * enabled * */ void cuda_runtime_api::extract_code_using_cuobjdump(){ - CUctx_st *context = GPGPUSim_Context(); - unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability(); + CUctx_st *context = GPGPUSim_Context(gpgpu_ctx); //prevent the dumping by cuobjdump everytime we execute the code! const char *override_cuobjdump = getenv("CUOBJDUMP_SIM_FILE"); - char command[1000], ptx_file[1000]; + char command[1000]; std::string app_binary = get_app_binary(); //Running cuobjdump using dynamic link to current process snprintf(command,1000,"md5sum %s ", app_binary.c_str()); @@ -2888,7 +3280,7 @@ cuobjdumpPTXSection* cuda_runtime_api::findPTXSection(const std::string identifi //! Extract the code using cuobjdump and remove unnecessary sections void cuda_runtime_api::cuobjdumpInit(){ - CUctx_st *context = GPGPUSim_Context(); + CUctx_st *context = GPGPUSim_Context(gpgpu_ctx); extract_code_using_cuobjdump(); //extract all the output of cuobjdump to _cuobjdump_*.* const char* pre_load = getenv("CUOBJDUMP_SIM_FILE"); if (pre_load ==NULL || strlen(pre_load)==0){ @@ -2901,7 +3293,7 @@ void cuda_runtime_api::cuobjdumpInit(){ //! Either submit PTX for simulation or convert SASS to PTXPlus and submit it void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){ - CUctx_st *context = GPGPUSim_Context(); + CUctx_st *context = GPGPUSim_Context(this); if(api->fatbin_registered[handle]) return; api->fatbin_registered[handle] = true; std::string fname = api->fatbinmap[handle]; @@ -3095,15 +3487,11 @@ cudaError_t cudaDeviceReset ( void ) { } return g_last_cudaError = cudaSuccess; } -cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - //Blocks until the device has completed all preceding requested tasks - synchronize(); - return g_last_cudaError = cudaSuccess; -} +cudaError_t CUDARTAPI cudaDeviceSynchronize(void) +{ + return cudaDeviceSynchronizeInternal(); +} void __cudaRegisterShared( void **fatCubinHandle, @@ -3142,22 +3530,7 @@ void __cudaRegisterTexture( int ext ) //passes in a newly created textureReference { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - std::string devStr (deviceName); - #if (CUDART_VERSION > 4020) - if (devStr.size() > 2 && devStr.data()[0] == ':' && devStr.data()[1] == ':') - devStr = devStr.replace(0, 2, ""); - #endif - CUctx_st *context = GPGPUSim_Context(); - gpgpu_t *gpu = context->get_device()->get_gpgpu(); - printf("GPGPU-Sim PTX: in __cudaRegisterTexture:\n"); - gpu->gpgpu_ptx_sim_bindNameToTexture(devStr.data(), hostVar, dim, norm, ext); - printf("GPGPU-Sim PTX: int dim = %d\n", dim); - printf("GPGPU-Sim PTX: int norm = %d\n", norm); - printf("GPGPU-Sim PTX: int ext = %d\n", ext); - printf("GPGPU-Sim PTX: Execution warning: Not finished implementing \"%s\"\n", __my_func__ ); + __cudaRegisterTextureInternal(fatCubinHandle, hostVar, deviceAddress, deviceName, dim, norm, ext); } @@ -3189,30 +3562,7 @@ cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj) cudaError_t cudaGLUnmapBufferObject(GLuint bufferObj) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } -#ifdef OPENGL_SUPPORT - CUctx_st* ctx = GPGPUSim_Context(); - glbmap_entry_t *p = ctx->api->g_glbmap; - while ( p && p->m_bufferObj != bufferObj ) - p = p->m_next; - if ( p == NULL ) - return g_last_cudaError = cudaErrorUnknown; - - char *data = (char *) calloc(p->m_size,1); - memcpy_from_gpu( data,(size_t)p->m_devPtr,p->m_size ); - glBufferSubData(GL_ARRAY_BUFFER,0,p->m_size,data); - free(data); - - return g_last_cudaError = cudaSuccess; -#else - fflush(stdout); - fflush(stderr); - printf("GPGPU-Sim PTX: support for OpenGL integration disabled -- exiting\n"); - fflush(stdout); - exit(50); -#endif + return cudaGLUnmapBufferObjectInternal(bufferObj); } cudaError_t cudaGLUnregisterBufferObject(GLuint bufferObj) @@ -3324,12 +3674,7 @@ cudaError_t CUDARTAPI cudaRuntimeGetVersion(int *runtimeVersion) #if CUDART_VERSION >= 3000 __host__ cudaError_t CUDARTAPI cudaFuncSetCacheConfig(const char *func, enum cudaFuncCache cacheConfig ) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - CUctx_st *context = GPGPUSim_Context(); - context->get_device()->get_gpgpu()->set_cache_config(context->get_kernel(func)->get_name(), (FuncCache)cacheConfig); - return g_last_cudaError = cudaSuccess; + return cudaFuncSetCacheConfigInternal(func, cacheConfig); } //Jin: hack for cdp @@ -5035,23 +5380,7 @@ CUresult CUDAAPI cuLaunchKernel(CUfunction f, void **kernelParams, void **extra) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - if (extra!=NULL){ - printf("GPGPU-Sim CUDA DRIVER API: ERROR: Currently do not support void** extra.\n"); - abort(); - } - const char *hostFun = (const char*) f; - CUctx_st *context = GPGPUSim_Context(); - function_info *entry = context->get_kernel(hostFun); - cudaConfigureCallInternal(dim3(gridDimX, gridDimY, gridDimZ), dim3(blockDimX, blockDimY, blockDimZ), sharedMemBytes, (cudaStream_t) hStream); - for(unsigned i = 0; i < entry->num_args(); i++){ - std::pair<size_t, unsigned> p = entry->get_param_config(i); - cudaSetupArgument(kernelParams[i], p.first, p.second); - } - cudaLaunchInternal(hostFun); - return CUDA_SUCCESS; + return cuLaunchKernelInternal(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra); } #endif /* CUDART_VERSION >= 4000 */ diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 45c5cdd..61d7507 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -52,6 +52,10 @@ class gpgpu_context { cuda_device_runtime* device_runtime; ptx_stats* stats; // member function list + void synchronize(); + void exit_simulation(); + void print_simulation_time(); + int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ); void cuobjdumpParseBinary(unsigned int handle); class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); class symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename ); @@ -60,6 +64,7 @@ class gpgpu_context { void print_ptx_file( const char *p, unsigned source_num, const char *filename ); class symbol_table* init_parser(const char*); class gpgpu_sim *gpgpu_ptx_sim_init_perf(); + void start_sim_thread(int api); struct _cuda_device_id *GPGPUSim_Init(); void ptx_reg_options(option_parser_t opp); const ptx_instruction* pc_to_instruction(unsigned pc); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 7f029c7..b032c05 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -647,13 +647,13 @@ unsigned _cl_kernel::sm_context_uid = 0; class _cl_device_id *GPGPUSim_Init() { static _cl_device_id *the_device = NULL; + gpgpu_context *ctx; + ctx = GPGPU_Context(); if( !the_device ) { - gpgpu_context *ctx; - ctx = GPGPU_Context(); gpgpu_sim *the_gpu = ctx->gpgpu_ptx_sim_init_perf(); the_device = new _cl_device_id(the_gpu); } - start_sim_thread(2); + ctx->start_sim_thread(2); return the_device; } @@ -960,7 +960,7 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, if ( ctx->func_sim->g_ptx_sim_mode ) ctx->func_sim->gpgpu_opencl_ptx_sim_main_func( grid ); else - gpgpu_opencl_ptx_sim_main_perf( grid ); + ctx->gpgpu_opencl_ptx_sim_main_perf( grid ); return CL_SUCCESS; } diff --git a/nightly.jenkinsfile b/nightly.jenkinsfile index e5ffa57..c841fb0 100644 --- a/nightly.jenkinsfile +++ b/nightly.jenkinsfile @@ -17,7 +17,7 @@ pipeline { steps { sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\ source `pwd`/setup_environment &&\ - make -j' + make -j 10' } } stage('nightly-simulations-build'){ @@ -34,7 +34,7 @@ pipeline { source `pwd`/setup_environment &&\ cd gpgpu-sim_simulations && \ source ./benchmarks/src/setup_environment && \ - make -i -j -C ./benchmarks/src/ all && \ + make -i -j 10 -C ./benchmarks/src/ all && \ make -C ./benchmarks/src data' } } @@ -46,7 +46,7 @@ pipeline { sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\ source `pwd`/setup_environment &&\ ./gpgpu-sim_simulations/util/job_launching/run_simulations.py \ - -B `cat ./gpgpu-sim_simulations/util/job_launching/apps/all-apps.list` -C TITANV-2B -N nightly-$$ && \ + -B `cat ./gpgpu-sim_simulations/util/job_launching/apps/correlation-apps.list` -C QV100-2B -N nightly-$$ && \ ./gpgpu-sim_simulations/util/job_launching/monitor_func_test.py -I -S 1800 -v \ -s stats-per-app-9.1.csv -T 12 -K -N nightly-$$ && \ ./gpgpu-sim_simulations/util/plotting/plot-get-stats.py -c stats-per-app-9.1.csv -P cuda-9.1.nightly' @@ -65,8 +65,8 @@ pipeline { sh 'source /home/tgrogers-raid/a/common/gpgpu-sim-setup/9.1_env_setup.sh &&\ ./gpgpu-sim_simulations/util/job_launching/get_stats.py -R -K -k \ - -B `cat ./gpgpu-sim_simulations/util/job_launching/apps/all-apps.list` \ - -C TITANV-2B > stats-per-kernel-9.1.csv &&\ + -B `cat ./gpgpu-sim_simulations/util/job_launching/apps/correlation-apps.list` \ + -C QV100-2B > stats-per-kernel-9.1.csv &&\ ./gpgpu-sim_simulations/util/plotting/merge-stats.py \ -c ./gpgpu-sim-results-repo/${JOB_NAME}/stats-per-app-9.1.csv,./stats-per-app-9.1.csv -R \ > per-app-merge-9.1.csv &&\ @@ -76,7 +76,7 @@ pipeline { source `pwd`/setup_environment &&\ ./gpgpu-sim_simulations/util/plotting/plot-correlation.py -c per-kernel-merge-9.1.csv \ -p cuda-9.1.nightly -b ./gpgpu-sim_simulations/util/plotting/known.correlation.outliers.list \ - | grep "Correl=" | tee correl.9.1.txt &&\ + | grep "Correl=" -B 1 | tee correl.9.1.txt &&\ cp stats-per-*.csv ./gpgpu-sim-results-repo/${JOB_NAME}/ &&\ cd ./gpgpu-sim-results-repo &&\ git pull &&\ diff --git a/setup_environment b/setup_environment index 58f4713..ca60d6b 100644 --- a/setup_environment +++ b/setup_environment @@ -7,13 +7,10 @@ export GPGPUSIM_ROOT="$( cd "$( dirname "$BASH_SOURCE" )" && pwd )" GPGPUSIM_VERSION_STRING=`cat $GPGPUSIM_ROOT/version | awk '/Version/ {print $8}'` #Detect Git branch and commit # -CURRENT_PWD=`pwd` -cd $GPGPUSIM_ROOT -GIT_COMMIT=`git log -n 1 | head -1 | sed -re 's/commit (.*)/\1/'` -GIT_FILES_CHANGED=`git diff --numstat | wc | sed -re 's/^\s+([0-9]+).*/\1./'` -GIT_FILES_CHANGED+=`git diff --numstat --cached | wc | sed -re 's/^\s+([0-9]+).*/\1/'` +GIT_COMMIT=`git --git-dir=$GPGPUSIM_ROOT/.git log -n 1 | head -1 | sed -re 's/commit (.*)/\1/'` +GIT_FILES_CHANGED=`git --git-dir=$GPGPUSIM_ROOT/.git diff --numstat | wc | sed -re 's/^\s+([0-9]+).*/\1./'` +GIT_FILES_CHANGED+=`git --git-dir=$GPGPUSIM_ROOT/.git diff --numstat --cached | wc | sed -re 's/^\s+([0-9]+).*/\1/'` GPGPUSIM_BUILD_STRING="gpgpu-sim_git-commit-$GIT_COMMIT-modified_$GIT_FILES_CHANGED" -cd $CURRENT_PWD echo -n "GPGPU-Sim version $GPGPUSIM_VERSION_STRING (build $GPGPUSIM_BUILD_STRING) "; diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 9aa1d73..e8e4b51 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -808,14 +808,14 @@ void kernel_info_t::notify_parent_finished() { if(m_parent_kernel) { m_kernel_entry->gpgpu_ctx->device_runtime->g_total_param_size -= ((m_kernel_entry->get_args_aligned_size() + 255)/256*256); m_parent_kernel->remove_child(this); - g_stream_manager()->register_finished_kernel(m_parent_kernel->get_uid()); + m_kernel_entry->gpgpu_ctx->the_gpgpusim->g_stream_manager->register_finished_kernel(m_parent_kernel->get_uid()); } } CUstream_st * kernel_info_t::create_stream_cta(dim3 ctaid) { assert(get_default_stream_cta(ctaid)); CUstream_st * stream = new CUstream_st(); - g_stream_manager()->add_stream(stream); + m_kernel_entry->gpgpu_ctx->the_gpgpusim->g_stream_manager->add_stream(stream); assert(m_cta_streams.find(ctaid) != m_cta_streams.end()); assert(m_cta_streams[ctaid].size() >= 1); //must have default stream m_cta_streams[ctaid].push_back(stream); @@ -831,7 +831,7 @@ CUstream_st * kernel_info_t::get_default_stream_cta(dim3 ctaid) { else { m_cta_streams[ctaid] = std::list<CUstream_st *>(); CUstream_st * stream = new CUstream_st(); - g_stream_manager()->add_stream(stream); + m_kernel_entry->gpgpu_ctx->the_gpgpusim->g_stream_manager->add_stream(stream); m_cta_streams[ctaid].push_back(stream); return stream; } @@ -863,7 +863,7 @@ void kernel_info_t::destroy_cta_streams() { for(auto s = m_cta_streams.begin(); s != m_cta_streams.end(); s++) { stream_size += s->second.size(); for(auto ss = s->second.begin(); ss != s->second.end(); ss++) - g_stream_manager()->destroy_stream(*ss); + m_kernel_entry->gpgpu_ctx->the_gpgpusim->g_stream_manager->destroy_stream(*ss); s->second.clear(); } printf("size %lu\n", stream_size); diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 60f7f68..341e44c 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -447,19 +447,25 @@ protected: class gpgpu_sim * m_gpu; }; -#define GLOBAL_HEAP_START 0xC0000000 - // start allocating from this address (lower values used for allocating globals in .ptx file) -#define SHARED_MEM_SIZE_MAX (96*1024) -#define LOCAL_MEM_SIZE_MAX (16*1024) -#define MAX_STREAMING_MULTIPROCESSORS 80 //scale it to Volta -#define MAX_THREAD_PER_SM 2048 -#define MAX_WARP_PER_SM 64 -#define TOTAL_LOCAL_MEM_PER_SM (MAX_THREAD_PER_SM*LOCAL_MEM_SIZE_MAX) -#define TOTAL_SHARED_MEM (MAX_STREAMING_MULTIPROCESSORS*SHARED_MEM_SIZE_MAX) -#define TOTAL_LOCAL_MEM (MAX_STREAMING_MULTIPROCESSORS*MAX_THREAD_PER_SM*LOCAL_MEM_SIZE_MAX) -#define SHARED_GENERIC_START (GLOBAL_HEAP_START-TOTAL_SHARED_MEM) -#define LOCAL_GENERIC_START (SHARED_GENERIC_START-TOTAL_LOCAL_MEM) -#define STATIC_ALLOC_LIMIT (GLOBAL_HEAP_START - (TOTAL_LOCAL_MEM+TOTAL_SHARED_MEM)) +// Let's just upgrade to C++11 so we can use constexpr here... +// start allocating from this address (lower values used for allocating globals in .ptx file) +const unsigned long long GLOBAL_HEAP_START = 0xC0000000; +// Volta max shmem size is 96kB +const unsigned long long SHARED_MEM_SIZE_MAX = 96 * (1 << 10); +// Volta max local mem is 16kB +const unsigned long long LOCAL_MEM_SIZE_MAX = 1 << 14; +// Volta Titan V has 80 SMs +const unsigned MAX_STREAMING_MULTIPROCESSORS = 80; +// Max 2048 threads / SM +const unsigned MAX_THREAD_PER_SM = 1 << 11; +// MAX 64 warps / SM +const unsigned MAX_WARP_PER_SM = 1 << 6; +const unsigned long long TOTAL_LOCAL_MEM_PER_SM = MAX_THREAD_PER_SM * LOCAL_MEM_SIZE_MAX; +const unsigned long long TOTAL_SHARED_MEM = MAX_STREAMING_MULTIPROCESSORS * SHARED_MEM_SIZE_MAX; +const unsigned long long TOTAL_LOCAL_MEM = MAX_STREAMING_MULTIPROCESSORS * MAX_THREAD_PER_SM * LOCAL_MEM_SIZE_MAX; +const unsigned long long SHARED_GENERIC_START = GLOBAL_HEAP_START - TOTAL_SHARED_MEM; +const unsigned long long LOCAL_GENERIC_START = SHARED_GENERIC_START - TOTAL_LOCAL_MEM; +const unsigned long long STATIC_ALLOC_LIMIT = GLOBAL_HEAP_START - (TOTAL_LOCAL_MEM + TOTAL_SHARED_MEM); #if !defined(__CUDA_RUNTIME_API_H__) @@ -525,10 +531,10 @@ private: int checkpoint_option; int checkpoint_kernel; int checkpoint_CTA; - int resume_option; - int resume_kernel; - int resume_CTA; - int checkpoint_CTA_t; + unsigned resume_option; + unsigned resume_kernel; + unsigned resume_CTA; + unsigned checkpoint_CTA_t; int checkpoint_insn_Y; int g_ptx_inst_debug_to_file; char* g_ptx_inst_debug_file; @@ -546,10 +552,10 @@ public: int checkpoint_option; int checkpoint_kernel; int checkpoint_CTA; - int resume_option; - int resume_kernel; - int resume_CTA; - int checkpoint_CTA_t; + unsigned resume_option; + unsigned resume_kernel; + unsigned resume_CTA; + unsigned checkpoint_CTA_t; int checkpoint_insn_Y; //Move some cycle core stats here instead of being global @@ -998,7 +1004,7 @@ public: printf("Printing mem access generated\n"); std::list<mem_access_t>::iterator it; for (it = m_accessq.begin(); it != m_accessq.end(); ++it){ - printf("MEM_TXN_GEN:%s:%x, Size:%d \n",mem_access_type_str(it->get_type()), it->get_addr(),it->get_size()); + printf("MEM_TXN_GEN:%s:%llx, Size:%d \n",mem_access_type_str(it->get_type()), it->get_addr(),it->get_size()); } } } diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 1b0e841..7a130ea 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -429,7 +429,7 @@ void gpgpu_t::memcpy_to_gpu( size_t dst_start_addr, const void *src, size_t coun // Copy into the performance model. //extern gpgpu_sim* g_the_gpu; - g_the_gpu()->perf_memcpy_to_gpu(dst_start_addr, count); + gpgpu_ctx->the_gpgpusim->g_the_gpu->perf_memcpy_to_gpu(dst_start_addr, count); if(g_debug_execution >= 3) { printf( " done.\n"); fflush(stdout); @@ -448,7 +448,7 @@ void gpgpu_t::memcpy_from_gpu( void *dst, size_t src_start_addr, size_t count ) // Copy into the performance model. //extern gpgpu_sim* g_the_gpu; - g_the_gpu()->perf_memcpy_to_gpu(src_start_addr, count); + gpgpu_ctx->the_gpgpusim->g_the_gpu->perf_memcpy_to_gpu(src_start_addr, count); if(g_debug_execution >= 3) { printf( " done.\n"); fflush(stdout); @@ -1254,7 +1254,7 @@ void function_info::param_to_shared( memory_space *shared_mem, symbol_table *sym { // TODO: call this only for PTXPlus with GT200 models //extern gpgpu_sim* g_the_gpu; - if (not g_the_gpu()->get_config().convert_to_ptxplus()) return; + if (not gpgpu_ctx->the_gpgpusim->g_the_gpu->get_config().convert_to_ptxplus()) return; // copies parameters into simulated shared memory for( std::map<unsigned,param_info>::iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) { @@ -1309,7 +1309,7 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt char buff[1024]; std::string filename_c(filename+"_c"); snprintf(buff,1024,"c++filt %s > %s", get_name().c_str(), filename_c.c_str()); - system(buff); + assert(system(buff) != NULL); FILE *fp = fopen(filename_c.c_str(), "r"); fgets(buff, 1024, fp); fclose(fp); @@ -1432,13 +1432,13 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt fout = fopen(ptx_config_fn.c_str(), "a"); assert(fout!=NULL); for (unsigned i = 0; i<line_number; i++){ - fgets(buff, 1024, fin); + assert(fgets(buff, 1024, fin) != NULL); assert(!feof(fin)); } fprintf(fout, "\n\n"); do{ fprintf(fout, "%s", buff); - fgets(buff, 1024, fin); + assert(fgets(buff, 1024, fin) != NULL); if(feof(fin)){ break; } @@ -1491,7 +1491,6 @@ static unsigned get_tex_datasize( const ptx_instruction *pI, ptx_thread_info *th const operand_info &src1 = pI->src1(); //the name of the texture std::string texname = src1.name(); - gpgpu_t *gpu = thread->get_gpu(); /* For programs with many streams, textures can be bound and unbound asynchronously. This means we need to use the kernel's "snapshot" of @@ -1577,7 +1576,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) } //Tensorcore is warp synchronous operation. So these instructions needs to be executed only once. To make the simulation faster removing the redundant tensorcore operation - if(!tensorcore_op(inst_opcode)||(tensorcore_op(inst_opcode))&&(lane_id==0)){ + if(!tensorcore_op(inst_opcode)||((tensorcore_op(inst_opcode))&&(lane_id==0))){ switch ( inst_opcode ) { #define OP_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,this); op_classification = CLASSIFICATION; break; #define OP_W_DEF(OP,FUNC,STR,DST,CLASSIFICATION) case OP: FUNC(pI,get_core(),inst); op_classification = CLASSIFICATION; break; @@ -2138,19 +2137,13 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL kernel_func_info->set_pdom(); } - unsigned max_cta_tot = max_cta(kernel_info,kernel.threads_per_cta(), g_the_gpu()->getShaderCoreConfig()->warp_size, g_the_gpu()->getShaderCoreConfig()->n_thread_per_shader, g_the_gpu()->getShaderCoreConfig()->gpgpu_shmem_size, g_the_gpu()->getShaderCoreConfig()->gpgpu_shader_registers, g_the_gpu()->getShaderCoreConfig()->max_cta_per_core); + unsigned max_cta_tot = max_cta(kernel_info,kernel.threads_per_cta(), gpgpu_ctx->the_gpgpusim->g_the_gpu->getShaderCoreConfig()->warp_size, gpgpu_ctx->the_gpgpusim->g_the_gpu->getShaderCoreConfig()->n_thread_per_shader, gpgpu_ctx->the_gpgpusim->g_the_gpu->getShaderCoreConfig()->gpgpu_shmem_size, gpgpu_ctx->the_gpgpusim->g_the_gpu->getShaderCoreConfig()->gpgpu_shader_registers, gpgpu_ctx->the_gpgpusim->g_the_gpu->getShaderCoreConfig()->max_cta_per_core); printf("Max CTA : %d\n",max_cta_tot); - - - - - int inst_count=50; - int cp_op= g_the_gpu()->checkpoint_option; - int cp_CTA = g_the_gpu()->checkpoint_CTA; - int cp_kernel= g_the_gpu()->checkpoint_kernel; - cp_count= g_the_gpu()->checkpoint_insn_Y; - cp_cta_resume= g_the_gpu()->checkpoint_CTA_t; + int cp_op= gpgpu_ctx->the_gpgpusim->g_the_gpu->checkpoint_option; + int cp_kernel= gpgpu_ctx->the_gpgpusim->g_the_gpu->checkpoint_kernel; + cp_count= gpgpu_ctx->the_gpgpusim->g_the_gpu->checkpoint_insn_Y; + cp_cta_resume= gpgpu_ctx->the_gpgpusim->g_the_gpu->checkpoint_CTA_t; int cta_launched =0; //we excute the kernel one CTA (Block) at the time, as synchronization functions work block wise @@ -2162,8 +2155,8 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL { functionalCoreSim cta( &kernel, - g_the_gpu(), - g_the_gpu()->getShaderCoreConfig()->warp_size + gpgpu_ctx->the_gpgpusim->g_the_gpu, + gpgpu_ctx->the_gpgpusim->g_the_gpu->getShaderCoreConfig()->warp_size ); cta.execute(cp_count,temp); @@ -2184,7 +2177,7 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL { char f1name[2048]; snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", kernel.get_uid() ); - g_checkpoint->store_global_mem(g_the_gpu()->get_global_memory(), f1name , "%08x"); + g_checkpoint->store_global_mem(gpgpu_ctx->the_gpgpusim->g_the_gpu->get_global_memory(), f1name , (char *)"%08x"); } @@ -2195,7 +2188,7 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL //openCL kernel simulation calls don't register the kernel so we don't register its exit if(!openCL) { //extern stream_manager *g_stream_manager; - g_stream_manager()->register_finished_kernel(kernel.get_uid()); + gpgpu_ctx->the_gpgpusim->g_stream_manager->register_finished_kernel(kernel.get_uid()); } //******PRINTING******* @@ -2210,7 +2203,7 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL //g_simulation_starttime is initilized by gpgpu_ptx_sim_init_perf() in gpgpusim_entrypoint.cc upon starting gpgpu-sim time_t end_time, elapsed_time, days, hrs, minutes, sec; end_time = time((time_t *)NULL); - elapsed_time = MAX(end_time - GPGPUsim_ctx_ptr()->g_simulation_starttime, 1); + elapsed_time = MAX(end_time - gpgpu_ctx->the_gpgpusim->g_simulation_starttime, 1); //calculating and printing simulation time in terms of days, hours, minutes and seconds @@ -2312,18 +2305,15 @@ void functionalCoreSim::execute(int inst_count, unsigned ctaid_cp) checkpoint *g_checkpoint; g_checkpoint = new checkpoint(); - symbol * sym; ptx_reg_t regval; regval.u64= 123; - symbol_table * symtab= m_kernel->entry()->get_symtab(); - unsigned ctaid =m_kernel->get_next_cta_id_single(); if(m_gpu->checkpoint_option==1 && (m_kernel->get_uid()==m_gpu->checkpoint_kernel) && (ctaid_cp>=m_gpu->checkpoint_CTA) && (ctaid_cp<m_gpu->checkpoint_CTA_t)) { char fname[2048]; snprintf(fname,2048,"checkpoint_files/shared_mem_%d.txt",ctaid-1 ); - g_checkpoint->store_global_mem(m_thread[0]->m_shared_mem, fname , "%08x"); + g_checkpoint->store_global_mem(m_thread[0]->m_shared_mem, fname , (char *)"%08x"); for(int i=0; i<32*m_warp_count;i++) { char fname[2048]; @@ -2331,7 +2321,7 @@ void functionalCoreSim::execute(int inst_count, unsigned ctaid_cp) m_thread[i]->print_reg_thread(fname); char f1name[2048]; snprintf(f1name,2048,"checkpoint_files/local_mem_thread_%d_%d_reg.txt",i,ctaid-1 ); - g_checkpoint->store_global_mem(m_thread[i]->m_local_mem, f1name , "%08x"); + g_checkpoint->store_global_mem(m_thread[i]->m_local_mem, f1name , (char *)"%08x"); m_thread[i]->set_done(); m_thread[i]->exitCore(); m_thread[i]->registerExit(); diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index dc3adc3..4baced5 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -27,7 +27,6 @@ } -//extern stream_manager *g_stream_manager(); //Handling device runtime api: //void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize) @@ -285,7 +284,7 @@ void cuda_device_runtime::launch_one_device_kernel() { device_launch_operation_t &op = g_cuda_device_launch_op.front(); stream_operation stream_op = stream_operation(op.grid, gpgpu_ctx->func_sim->g_ptx_sim_mode, op.stream); - g_stream_manager()->push(stream_op); + gpgpu_ctx->the_gpgpusim->g_stream_manager->push(stream_op); g_cuda_device_launch_op.pop_front(); } } diff --git a/src/cuda-sim/half.h b/src/cuda-sim/half.h index 8f1a8eb..9f74bb7 100644 --- a/src/cuda-sim/half.h +++ b/src/cuda-sim/half.h @@ -642,10 +642,10 @@ namespace half_float if(exp > 16)
{
if(R == std::round_toward_infinity)
- return hbits | 0x7C00 - (hbits>>15);
+ return hbits | (0x7C00 - (hbits>>15));
else if(R == std::round_toward_neg_infinity)
- return hbits | 0x7BFF + (hbits>>15);
- return hbits | 0x7BFF + (R!=std::round_toward_zero);
+ return hbits | (0x7BFF + (hbits>>15));
+ return hbits | (0x7BFF + (R!=std::round_toward_zero));
}
if(exp < -13)
value = std::ldexp(value, 24);
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 58a077e..014e588 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -203,7 +203,7 @@ void ptx_thread_info::print_reg_thread(char * fname) const std::string &name = it->first->name(); const std::string &dec= it->first->decl_location(); unsigned size = it->first->get_size_in_bytes(); - fprintf(fp,"%s %llu %s %d\n",name.c_str(),it->second, dec.c_str(),size ); + fprintf(fp,"%s %llu %s %d\n", name.c_str(), it->second, dec.c_str(), size); } //m_regs.pop_back(); @@ -224,7 +224,6 @@ void ptx_thread_info::resume_reg_thread(char * fname, symbol_table * symtab) { symbol *reg; char * pch; - unsigned size; pch = strtok (line," "); char * name =pch; reg= symtab->lookup(name); @@ -232,11 +231,7 @@ void ptx_thread_info::resume_reg_thread(char * fname, symbol_table * symtab) pch = strtok (NULL," "); data = atoi(pch); pch = strtok (NULL," "); - char * decl= pch; pch = strtok (NULL," "); - size = atoi(pch); - - m_regs.back()[reg] = data; } fclose ( fp2 ); @@ -1819,9 +1814,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) ptx_reg_t matrix_d[16][16]; ptx_reg_t src_data; ptx_thread_info *thread; - int stride; - unsigned wmma_type = pI->get_wmma_type(); unsigned a_layout = pI->get_wmma_layout(0); unsigned b_layout = pI->get_wmma_layout(1); unsigned type = pI->get_type(); @@ -1833,7 +1826,6 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) tid= inst.warp_id_func()*core->get_warp_size(); else tid= inst.warp_id()*core->get_warp_size(); - unsigned thread_group_index; float temp; half temp2; @@ -1847,9 +1839,9 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) ptx_reg_t v[8]; thread->get_vector_operand_values( src_a, v, nelem ); if(core->get_gpu()->gpgpu_ctx->debug_tensorcore){ - printf("Thread%d_Iteration=%d\n:",thrd,operand_num); - for(k=0;k<nelem;k++){ - printf("%x ",v[k].u64); + printf("Thread%d_Iteration=%d\n:", thrd, operand_num); + for(k = 0; k < nelem; k++){ + printf("%llx ",v[k].u64); } printf("\n"); } @@ -2027,7 +2019,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) printf("thread%d:",thrd); for(k=0;k<8;k++){ - printf("%x ",matrix_d[row_t[k]][col_t[k]].f16); + printf("%x ", (unsigned int)matrix_d[row_t[k]][col_t[k]].f16); } printf("\n"); } @@ -2038,7 +2030,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) nw_data4.s64=((matrix_d[row_t[6]][col_t[6]].s64 & 0xffff))|((matrix_d[row_t[7]][col_t[7]].s64&0xffff)<<16); thread->set_vector_operand_values(dst,nw_data1,nw_data2,nw_data3,nw_data4); if(core->get_gpu()->gpgpu_ctx->debug_tensorcore) - printf("thread%d=%x,%x,%x,%x",thrd,nw_data1.s64,nw_data2.s64,nw_data3.s64,nw_data4.s64); + printf("thread%d=%llx,%llx,%llx,%llx", thrd, nw_data1.s64, nw_data2.s64, nw_data3.s64, nw_data4.s64); } else{ @@ -2298,9 +2290,8 @@ unsigned int saturatei(unsigned int a, unsigned int max) ptx_reg_t f2x( ptx_reg_t x, unsigned from_width, unsigned to_width, int to_sign, int rounding_mode, int saturation_mode ) { - half mytemp; - float myfloat; - half_float::half tmp_h; + half mytemp; + half_float::half tmp_h; //assert( from_width == 32); enum cudaRoundMode mode = cudaRoundZero; @@ -3085,7 +3076,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) size_t size; unsigned smid; int t; - int thrd,odd,inx,k; + int thrd, k; ptx_thread_info *thread; const operand_info &src = pI->operand_lookup(1); @@ -3105,15 +3096,13 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) _memory_op_t insn_memory_op = pI->has_memory_read() ? memory_load : memory_store; for (thrd=0; thrd < core->get_warp_size(); thrd++) { thread = core->get_thread_info()[tid+thrd]; - odd=thrd%2; - inx=thrd/2; - ptx_reg_t addr_reg = thread->get_operand_value(src1, src, type, thread, 1); + ptx_reg_t addr_reg = thread->get_operand_value(src1, src, type, thread, 1); ptx_reg_t src2_data = thread->get_operand_value(src2, src, type, thread, 1); const operand_info &src_a= pI->operand_lookup(1); unsigned nelem = src_a.get_vect_nelem(); ptx_reg_t* v= new ptx_reg_t[8]; thread->get_vector_operand_values( src_a, v, nelem ); - stride=src2_data.u32; + stride = src2_data.u32; memory_space_t space = pI->get_space(); @@ -3130,9 +3119,9 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) } decode_space(space,thread,src1,mem,addr); - type_info_key::type_decode(type,size,t); + type_info_key::type_decode(type, size, t); if(core->get_gpu()->gpgpu_ctx->debug_tensorcore) - printf("mma_st: thrd=%d,addr=%x, fp(size=%d), stride=%d\n",thrd,addr_reg.u32,size,src2_data.u32); + printf("mma_st: thrd=%d, addr=%x, fp(size=%zu), stride=%d\n", thrd, addr_reg.u32, size, src2_data.u32); addr_t new_addr = addr+thread_group_offset(thrd,wmma_type,wmma_layout,type,stride)*size/8; addr_t push_addr; @@ -3152,7 +3141,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) mem_txn_addr[num_mem_txn++]=push_addr; if(core->get_gpu()->gpgpu_ctx->debug_tensorcore){ - printf("wmma:store:thread%d=%x,%x,%x,%x,%x,%x,%x,%x\n",thrd,v[0].s64,v[1].s64,v[2].s64,v[3].s64,v[4].s64,v[5].s64,v[6].s64,v[7].s64); + printf("wmma:store:thread%d=%llx,%llx,%llx,%llx,%llx,%llx,%llx,%llx\n",thrd,v[0].s64,v[1].s64,v[2].s64,v[3].s64,v[4].s64,v[5].s64,v[6].s64,v[7].s64); float temp; int l; printf("thread=%d:",thrd); @@ -3179,7 +3168,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) } if(core->get_gpu()->gpgpu_ctx->debug_tensorcore) - printf("wmma:store:thread%d=%x,%x,%x,%x,%x,%x,%x,%x\n",thrd,nw_v[0].s64,nw_v[1].s64,nw_v[2].s64,nw_v[3].s64,nw_v[4].s64,nw_v[5].s64,nw_v[6].s64,nw_v[7].s64); + printf("wmma:store:thread%d=%llx,%llx,%llx,%llx,%llx,%llx,%llx,%llx\n",thrd,nw_v[0].s64,nw_v[1].s64,nw_v[2].s64,nw_v[3].s64,nw_v[4].s64,nw_v[5].s64,nw_v[6].s64,nw_v[7].s64); } } @@ -3238,11 +3227,11 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) } decode_space(space,thread,src1,mem,addr); - type_info_key::type_decode(type,size,t); + type_info_key::type_decode(type, size, t); ptx_reg_t data[16]; if(core->get_gpu()->gpgpu_ctx->debug_tensorcore) - printf("mma_ld: thrd=%d,addr=%x, fpsize=%d, stride=%d\n",thrd,src1_data.u32,size,src2_data.u32); + printf("mma_ld: thrd=%d,addr=%x, fpsize=%zu, stride=%d\n", thrd, src1_data.u32, size, src2_data.u32); addr_t new_addr = addr+thread_group_offset(thrd,wmma_type,wmma_layout,type,stride)*size/8; addr_t fetch_addr; @@ -3341,7 +3330,7 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) if(type==F16_TYPE){ printf("\nmma_ld:thread%d= ",thrd); for(i=0;i<16;i++){ - printf("%x ",data[i].u64); + printf("%llx ",data[i].u64); } printf("\n"); @@ -3361,7 +3350,7 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) printf("\n"); printf("\nmma_ld:thread%d= ",thrd); for(i=0;i<8;i++){ - printf("%x ",data[i].u64); + printf("%llx ",data[i].u64); } printf("\n"); } @@ -3388,15 +3377,15 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) else thread->set_wmma_vector_operand_values(dst,nw_data[0],nw_data[1],nw_data[2],nw_data[3],nw_data[4],nw_data[5],nw_data[6],nw_data[7]); if(core->get_gpu()->gpgpu_ctx->debug_tensorcore){ - printf("mma_ld:data[0].s64=%x,data[1].s64=%x,new_data[0].s64=%x\n",data[0].u64,data[1].u64,nw_data[0].u64); - printf("mma_ld:data[2].s64=%x,data[3].s64=%x,new_data[1].s64=%x\n",data[2].u64,data[3].u64,nw_data[1].u64); - printf("mma_ld:data[4].s64=%x,data[5].s64=%x,new_data[2].s64=%x\n",data[4].u64,data[5].u64,nw_data[2].u64); - printf("mma_ld:data[6].s64=%x,data[7].s64=%x,new_data[3].s64=%x\n",data[6].u64,data[7].u64,nw_data[3].u64); + printf("mma_ld:data[0].s64=%llx,data[1].s64=%llx,new_data[0].s64=%llx\n",data[0].u64,data[1].u64,nw_data[0].u64); + printf("mma_ld:data[2].s64=%llx,data[3].s64=%llx,new_data[1].s64=%llx\n",data[2].u64,data[3].u64,nw_data[1].u64); + printf("mma_ld:data[4].s64=%llx,data[5].s64=%llx,new_data[2].s64=%llx\n",data[4].u64,data[5].u64,nw_data[2].u64); + printf("mma_ld:data[6].s64=%llx,data[7].s64=%llx,new_data[3].s64=%llx\n",data[6].u64,data[7].u64,nw_data[3].u64); if(wmma_type!=LOAD_C){ - printf("mma_ld:data[8].s64=%x,data[9].s64=%x,new_data[4].s64=%x\n",data[8].u64,data[9].u64,nw_data[4].s64); - printf("mma_ld:data[10].s64=%x,data[11].s64=%x,new_data[5].s64=%x\n",data[10].u64,data[11].u64,nw_data[5].u64); - printf("mma_ld:data[12].s64=%x,data[13].s64=%x,new_data[6].s64=%x\n",data[12].u64,data[13].u64,nw_data[6].u64); - printf("mma_ld:data[14].s64=%x,data[15].s64=%x,new_data[7].s64=%x\n",data[14].u64,data[15].u64,nw_data[3].u64); + printf("mma_ld:data[8].s64=%llx,data[9].s64=%llx,new_data[4].s64=%llx\n",data[8].u64,data[9].u64,nw_data[4].s64); + printf("mma_ld:data[10].s64=%llx,data[11].s64=%llx,new_data[5].s64=%llx\n",data[10].u64,data[11].u64,nw_data[5].u64); + printf("mma_ld:data[12].s64=%llx,data[13].s64=%llx,new_data[6].s64=%llx\n",data[12].u64,data[13].u64,nw_data[6].u64); + printf("mma_ld:data[14].s64=%llx,data[15].s64=%llx,new_data[7].s64=%llx\n",data[14].u64,data[15].u64,nw_data[3].u64); } } } @@ -4132,9 +4121,9 @@ int prmt_mode_present(int mode) } return returnval; } -int read_byte(int mode,int control,int d_sel_index,signed long long value){ +int read_byte(int mode, int control, int d_sel_index, signed long long value){ - int returnval; + int returnval = 0; int prmt_f4e_mode[4][4]={{0,1,2,3},{1,2,3,4},{2,3,4,5},{3,4,5,6}}; int prmt_b4e_mode[4][4]={{0,7,6,5},{1,0,7,6},{2,1,0,7},{3,2,1,0}}; int prmt_rc8_mode[4][4]={{0,0,0,0},{1,1,1,1},{2,2,2,2},{3,3,3,3}}; @@ -4157,11 +4146,12 @@ int read_byte(int mode,int control,int d_sel_index,signed long long value){ case PRMT_RC8_MODE: returnval=prmt_rc8_mode[control][d_sel_index];break; case PRMT_ECL_MODE: returnval=prmt_ecl_mode[control][d_sel_index];break; case PRMT_ECR_MODE: returnval=prmt_ecr_mode[control][d_sel_index];break; - case PRMT_RC16_MODE: returnval=prmt_rc16_mode[control][d_sel_index];break; - default: printf("ERROR\n");break; + case PRMT_RC16_MODE: returnval=prmt_rc16_mode[control][d_sel_index];break; + // Change the default from printing "ERROR" to just asserting + default: assert(false); } } - return (returnval<<8*d_sel_index); + return (returnval << 8 * d_sel_index); } void prmt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 6978cc1..d8943d2 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1415,7 +1415,7 @@ unsigned function_info::print_insn( unsigned pc, FILE * fp ) const snprintf(command,1024,"c++filt -p %s",m_name.c_str()); FILE *p = popen(command,"r"); buffer[0]=0; - fgets(buffer, 1023, p); + assert(fgets(buffer, 1023, p) != NULL); // Remove trailing "\n" in buffer char *c; if ((c=strchr(buffer, '\n')) != NULL) *c = '\0'; diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index dca3cec..33bcf45 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -214,7 +214,8 @@ void fix_duplicate_errors(char fname2[1024]) { long filesize = ftell(ptxsource); rewind(ptxsource); char *ptxdata = (char*)malloc((filesize+1)*sizeof(char)); - fread(ptxdata, filesize, 1, ptxsource); + // Fail if we do not read the file + assert(fread(ptxdata, filesize, 1, ptxsource) == 1); fclose(ptxsource); FILE *ptxdest = fopen(fname2,"w"); diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 81b70af..a4f4a0c 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -421,7 +421,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi assert( (num_bits%8) == 0 ); addr = g_current_symbol_table->get_sstarr_next(); addr_pad = pad_address(addr, num_bits/8, 128); - printf("from 0x%x to 0x%lx (sstarr memory space)\n", + printf("from 0x%llx to 0x%llx (sstarr memory space)\n", addr+addr_pad, addr+addr_pad + num_bits/8); fflush(stdout); diff --git a/src/gpgpu-sim/addrdec.cc b/src/gpgpu-sim/addrdec.cc index 3262456..f225933 100644 --- a/src/gpgpu-sim/addrdec.cc +++ b/src/gpgpu-sim/addrdec.cc @@ -184,7 +184,7 @@ void linear_to_raw_address_translation::addrdec_tlx(new_addr_type addr, addrdec_ } assert(tlx->chip < m_n_channel); - assert(tlx->sub_partition < m_n_channel*m_n_sub_partition_in_channel); + assert(tlx->sub_partition < m_n_channel * m_n_sub_partition_in_channel); return; break; } diff --git a/src/gpgpu-sim/addrdec.h b/src/gpgpu-sim/addrdec.h index a5333fb..c9a1420 100644 --- a/src/gpgpu-sim/addrdec.h +++ b/src/gpgpu-sim/addrdec.h @@ -92,7 +92,7 @@ private: new_addr_type sub_partition_id_mask; unsigned int gap; - int m_n_channel; + unsigned m_n_channel; int m_n_sub_partition_in_channel; }; diff --git a/src/gpgpu-sim/dram.cc b/src/gpgpu-sim/dram.cc index d443d79..9c33822 100644 --- a/src/gpgpu-sim/dram.cc +++ b/src/gpgpu-sim/dram.cc @@ -482,7 +482,6 @@ void dram_t::cycle() bool memory_pending_rw_found=false; for (unsigned j=0;j<m_config->nbk;j++) { - unsigned grp = get_bankgrp_number(j); if (bk[j]->mrq && (((bk[j]->curr_row == bk[j]->mrq->row) && (bk[j]->mrq->rw == READ) && (bk[j]->state == BANK_ACTIVE)) @@ -817,10 +816,10 @@ void dram_t::visualize() const void dram_t::print_stat( FILE* simFile ) { - fprintf(simFile,"DRAM (%llu): n_cmd=%llu n_nop=%llu n_act=%llu n_pre=%llu n_ref=%llu n_req=%llu n_rd=%llu n_write=%llu bw_util=%.4g ", + fprintf(simFile,"DRAM (%u): n_cmd=%llu n_nop=%llu n_act=%llu n_pre=%llu n_ref=%llu n_req=%llu n_rd=%llu n_write=%llu bw_util=%.4g ", id, n_cmd, n_nop, n_act, n_pre, n_ref, n_req, n_rd, n_wr, (float)bwutil/n_cmd); - fprintf(simFile, "mrqq: %d %.4g mrqsmax=%d ", max_mrqs, (float)ave_mrqs/n_cmd, max_mrqs_temp); + fprintf(simFile, "mrqq: %d %.4g mrqsmax=%llu ", max_mrqs, (float)ave_mrqs/n_cmd, max_mrqs_temp); fprintf(simFile, "\n"); fprintf(simFile, "dram_util_bins:"); for (unsigned i=0;i<10;i++) fprintf(simFile, " %d", dram_util_bins[i]); @@ -899,10 +898,10 @@ void dram_t::set_dram_power_stats( unsigned &cmd, unsigned dram_t::get_bankgrp_number(unsigned i) { if(m_config->dram_bnkgrp_indexing_policy == HIGHER_BITS) { //higher bits - return i>>m_config->bk_tag_length; + return i >> m_config->bk_tag_length; } else if (m_config->dram_bnkgrp_indexing_policy == LOWER_BITS) { //lower bits - return i&((m_config->nbkgrp-1)); + return i & ((m_config->nbkgrp - 1)); } else { assert(1); diff --git a/src/gpgpu-sim/gpu-cache.cc b/src/gpgpu-sim/gpu-cache.cc index b7d0ac3..65f2f6d 100644 --- a/src/gpgpu-sim/gpu-cache.cc +++ b/src/gpgpu-sim/gpu-cache.cc @@ -790,7 +790,7 @@ void cache_stats::print_stats(FILE *fout, const char *cache_name) const{ } for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { if(total_access[type] > 0) - fprintf(fout, "\t%s[%s][%s] = %llu\n", + fprintf(fout, "\t%s[%s][%s] = %u\n", m_cache_name.c_str(), mem_access_type_str((enum mem_access_type)type), "TOTAL_ACCESS", @@ -803,7 +803,7 @@ void cache_stats::print_fail_stats(FILE *fout, const char *cache_name) const{ for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { for (unsigned fail = 0; fail < NUM_CACHE_RESERVATION_FAIL_STATUS; ++fail) { if(m_fail_stats[type][fail] > 0){ - fprintf(fout, "\t%s[%s][%s] = %u\n", + fprintf(fout, "\t%s[%s][%s] = %llu\n", m_cache_name.c_str(), mem_access_type_str((enum mem_access_type)type), cache_fail_status_str((enum cache_reservation_fail_reason)fail), @@ -1430,8 +1430,6 @@ data_cache::wr_miss_wa_lazy_fetch_on_read( new_addr_type addr, { new_addr_type block_addr = m_config.block_addr(addr); - new_addr_type mshr_addr = m_config.mshr_addr(mf->get_addr()); - //if the request writes to the whole cache line/sector, then, write and set cache line Modified. //and no need to send read request to memory or reserve mshr diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index a151379..ce5c4cb 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -475,10 +475,10 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) m_shader_config.reg_options(opp); m_memory_config.reg_options(opp); power_config::reg_options(opp); - option_parser_register(opp, "-gpgpu_max_cycle", OPT_INT32, &gpu_max_cycle_opt, + option_parser_register(opp, "-gpgpu_max_cycle", OPT_INT64, &gpu_max_cycle_opt, "terminates gpu simulation early (0 = no limit)", "0"); - option_parser_register(opp, "-gpgpu_max_insn", OPT_INT32, &gpu_max_insn_opt, + option_parser_register(opp, "-gpgpu_max_insn", OPT_INT64, &gpu_max_insn_opt, "terminates gpu simulation early (0 = no limit)", "0"); option_parser_register(opp, "-gpgpu_max_cta", OPT_INT32, &gpu_max_cta_opt, @@ -1117,8 +1117,8 @@ void gpgpu_sim::gpu_print_stat() printf("gpu_tot_sim_insn = %lld\n", gpu_tot_sim_insn+gpu_sim_insn); printf("gpu_tot_ipc = %12.4f\n", (float)(gpu_tot_sim_insn+gpu_sim_insn) / (gpu_tot_sim_cycle+gpu_sim_cycle)); printf("gpu_tot_issued_cta = %lld\n", gpu_tot_issued_cta + m_total_cta_launched); - printf("gpu_occupancy = %.4f\% \n", gpu_occupancy.get_occ_fraction() * 100); - printf("gpu_tot_occupancy = %.4f\% \n", (gpu_occupancy + gpu_tot_occupancy).get_occ_fraction() * 100); + printf("gpu_occupancy = %.4f%% \n", gpu_occupancy.get_occ_fraction() * 100); + printf("gpu_tot_occupancy = %.4f%% \n", (gpu_occupancy + gpu_tot_occupancy).get_occ_fraction() * 100); fprintf(statfout, "max_total_param_size = %llu\n", gpgpu_ctx->device_runtime->g_max_total_param_size); @@ -1144,7 +1144,7 @@ void gpgpu_sim::gpu_print_stat() time_t curr_time; time(&curr_time); - unsigned long long elapsed_time = MAX( curr_time - GPGPUsim_ctx_ptr()->g_simulation_starttime, 1 ); + unsigned long long elapsed_time = MAX( curr_time - gpgpu_ctx->the_gpgpusim->g_simulation_starttime, 1 ); printf( "gpu_total_sim_rate=%u\n", (unsigned)( ( gpu_tot_sim_insn + gpu_sim_insn ) / elapsed_time ) ); //shader_print_l1_miss_stat( stdout ); @@ -1359,7 +1359,7 @@ bool shader_core_ctx::occupy_shader_resource_1block(kernel_info_t & k, bool occu m_occupied_regs += (padded_cta_size * ((kernel_info->regs+3)&~3)); m_occupied_ctas++; - SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Occupied %d threads, %d shared mem, %d registers, %d ctas\n", + SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Occupied %u threads, %u shared mem, %u registers, %u ctas\n", m_occupied_n_threads, m_occupied_shmem, m_occupied_regs, m_occupied_ctas); } @@ -1476,7 +1476,7 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel ) nthreads_in_block += ptx_sim_init_thread(kernel,&m_thread[i],m_sid,i,cta_size-(i-start_thread),m_config->n_thread_per_shader,this,free_cta_hw_id,warp_id,m_cluster->get_gpu()); m_threadState[i].m_active = true; // load thread local memory and register file - if(m_gpu->resume_option==1 && kernel.get_uid()==m_gpu->resume_kernel && ctaid>=m_gpu->resume_CTA && ctaid<m_gpu->checkpoint_CTA_t ) + if(m_gpu->resume_option == 1 && kernel.get_uid() == m_gpu->resume_kernel && ctaid >= m_gpu->resume_CTA && ctaid < m_gpu->checkpoint_CTA_t ) { char fname[2048]; snprintf(fname,2048,"checkpoint_files/thread_%d_%d_reg.txt",i%cta_size,ctaid ); @@ -1491,7 +1491,7 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel ) assert( nthreads_in_block > 0 && nthreads_in_block <= m_config->n_thread_per_shader); // should be at least one, but less than max m_cta_status[free_cta_hw_id]=nthreads_in_block; - if(m_gpu->resume_option==1 && kernel.get_uid()==m_gpu->resume_kernel && ctaid>=m_gpu->resume_CTA && ctaid<m_gpu->checkpoint_CTA_t ) + if(m_gpu->resume_option == 1 && kernel.get_uid() == m_gpu->resume_kernel && ctaid >= m_gpu->resume_CTA && ctaid < m_gpu->checkpoint_CTA_t ) { char f1name[2048]; snprintf(f1name,2048,"checkpoint_files/shared_mem_%d.txt", ctaid); @@ -1717,7 +1717,7 @@ void gpgpu_sim::cycle() time_t days, hrs, minutes, sec; time_t curr_time; time(&curr_time); - unsigned long long elapsed_time = MAX(curr_time - GPGPUsim_ctx_ptr()->g_simulation_starttime, 1); + unsigned long long elapsed_time = MAX(curr_time - gpgpu_ctx->the_gpgpusim->g_simulation_starttime, 1); if ( (elapsed_time - last_liveness_message_time) >= m_config.liveness_message_freq && DTRACE(LIVENESS) ) { days = elapsed_time/(3600*24); hrs = elapsed_time/3600 - 24*days; diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 76c7a06..fba770d 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -368,8 +368,8 @@ private: double l2_period; // GPGPU-Sim timing model options - unsigned gpu_max_cycle_opt; - unsigned gpu_max_insn_opt; + unsigned long long gpu_max_cycle_opt; + unsigned long long gpu_max_insn_opt; unsigned gpu_max_cta_opt; char *gpgpu_runtime_stat; bool gpgpu_flush_l1_cache; diff --git a/src/gpgpu-sim/l2cache.cc b/src/gpgpu-sim/l2cache.cc index 39a5812..fb4ce32 100644 --- a/src/gpgpu-sim/l2cache.cc +++ b/src/gpgpu-sim/l2cache.cc @@ -82,7 +82,7 @@ void memory_partition_unit::handle_memcpy_to_gpu( size_t addr, unsigned global_s unsigned p = global_sub_partition_id_to_local_id(global_subpart_id); std::string mystring = mask.to_string<char,std::string::traits_type,std::string::allocator_type>(); - MEMPART_DPRINTF("Copy Engine Request Received For Address=%llx, local_subpart=%u, global_subpart=%u, sector_mask=%s \n", addr, p, global_subpart_id, mystring.c_str()); + MEMPART_DPRINTF("Copy Engine Request Received For Address=%zx, local_subpart=%u, global_subpart=%u, sector_mask=%s \n", addr, p, global_subpart_id, mystring.c_str()); m_sub_partition[p]->force_l2_tag_update(addr,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle, mask); } @@ -681,7 +681,7 @@ std::vector<mem_fetch*> memory_sub_partition::breakdown_request_to_sector_reques } } else { - printf("Invalid sector received, address = 0x%06x, sector mask = %s, data size = %d", + printf("Invalid sector received, address = 0x%06llx, sector mask = %s, data size = %d", mf->get_addr(), mf->get_access_sector_mask(), mf->get_data_size()); assert(0 && "Undefined sector mask is received"); } @@ -716,7 +716,7 @@ std::vector<mem_fetch*> memory_sub_partition::breakdown_request_to_sector_reques byte_sector_mask <<= SECTOR_SIZE; } } else { - printf("Invalid sector received, address = 0x%06x, sector mask = %d, byte mask = , data size = %d", + printf("Invalid sector received, address = 0x%06llx, sector mask = %d, byte mask = , data size = %u", mf->get_addr(), mf->get_access_sector_mask().count(), mf->get_data_size()); assert(0 && "Undefined data size is received"); } diff --git a/src/gpgpu-sim/local_interconnect.cc b/src/gpgpu-sim/local_interconnect.cc index da8a65c..c70477c 100644 --- a/src/gpgpu-sim/local_interconnect.cc +++ b/src/gpgpu-sim/local_interconnect.cc @@ -253,7 +253,7 @@ LocalInterconnect::LocalInterconnect(const struct inct_config& m_localinct_confi } LocalInterconnect::~LocalInterconnect(){ - for (int i=0; i<m_inct_config.subnets; ++i) { + for (unsigned i = 0; i < m_inct_config.subnets; ++i) { delete net[i]; } } diff --git a/src/gpgpu-sim/scoreboard.cc b/src/gpgpu-sim/scoreboard.cc index 80f95c6..1017e75 100644 --- a/src/gpgpu-sim/scoreboard.cc +++ b/src/gpgpu-sim/scoreboard.cc @@ -140,10 +140,10 @@ bool Scoreboard::checkCollision( unsigned wid, const class inst_t *inst ) const // Get list of all input and output registers std::set<int> inst_regs; - for(int iii=0;iii<inst->outcount;iii++) + for(unsigned iii=0; iii < inst->outcount; iii++) inst_regs.insert(inst->out[iii]); - for(int jjj=0;jjj<inst->incount;jjj++) + for(unsigned jjj=0;jjj<inst->incount;jjj++) inst_regs.insert(inst->in[jjj]); if(inst->pred > 0) inst_regs.insert(inst->pred); diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 0514a77..86508f6 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -87,7 +87,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, shader_core_stats *stats ) : core_t( gpu, NULL, config->warp_size, config->n_thread_per_shader ), m_barriers( this, config->max_warps_per_shader, config->max_cta_per_core, config->max_barriers_per_cta, config->warp_size ), - m_dynamic_warp_id(0), m_active_warps(0) + m_active_warps(0), m_dynamic_warp_id(0) { m_cluster = cluster; m_config = config; @@ -164,7 +164,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, NUM_CONCRETE_SCHEDULERS; assert ( scheduler != NUM_CONCRETE_SCHEDULERS ); - for (int i = 0; i < m_config->gpgpu_num_sched_per_core; i++) { + for (unsigned i = 0; i < m_config->gpgpu_num_sched_per_core; i++) { switch( scheduler ) { case CONCRETE_SCHEDULER_LRR: @@ -263,7 +263,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, //distribute i's evenly though schedulers; schedulers[i%m_config->gpgpu_num_sched_per_core]->add_supervised_warp_id(i); } - for ( int i = 0; i < m_config->gpgpu_num_sched_per_core; ++i ) { + for ( unsigned i = 0; i < m_config->gpgpu_num_sched_per_core; ++i ) { schedulers[i]->done_adding_supervised_warps(); } @@ -474,7 +474,7 @@ void shader_core_ctx::init_warps( unsigned cta_id, unsigned start_thread, unsign } m_simt_stack[i]->launch(start_pc,active_threads); - if(m_gpu->resume_option==1 && kernel_id==m_gpu->resume_kernel && ctaid>=m_gpu->resume_CTA && ctaid<m_gpu->checkpoint_CTA_t ) + if(m_gpu->resume_option == 1 && kernel_id == m_gpu->resume_kernel && ctaid >= m_gpu->resume_CTA && ctaid < m_gpu->checkpoint_CTA_t ) { char fname[2048]; snprintf(fname,2048,"checkpoint_files/warp_%d_%d_simt.txt",i%warp_per_cta,ctaid ); @@ -868,7 +868,7 @@ void shader_core_ctx::func_exec_inst( warp_inst_t &inst ) void shader_core_ctx::issue_warp( register_set& pipe_reg_set, const warp_inst_t* next_inst, const active_mask_t &active_mask, unsigned warp_id, unsigned sch_id ) { - warp_inst_t** pipe_reg = pipe_reg = pipe_reg_set.get_free(m_config->sub_core_model, sch_id); + warp_inst_t** pipe_reg = pipe_reg_set.get_free(m_config->sub_core_model, sch_id); assert(pipe_reg); m_warp[warp_id].ibuffer_free(); @@ -2147,7 +2147,7 @@ ldst_unit::ldst_unit( mem_fetch_interface *icnt, l1_latency_queue.resize(m_config->m_L1D_config.l1_banks); assert(m_config->m_L1D_config.l1_latency > 0); - for(int j=0; j<m_config->m_L1D_config.l1_banks; j++ ) + for(unsigned j = 0; j < m_config->m_L1D_config.l1_banks; j++ ) l1_latency_queue[j].resize(m_config->m_L1D_config.l1_latency,(mem_fetch*)NULL); } @@ -2459,7 +2459,7 @@ void shader_core_ctx::register_cta_thread_exit( unsigned cta_num, kernel_info_t m_barriers.deallocate_barrier(cta_num); shader_CTA_count_unlog(m_sid, 1); - SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Finished CTA #%d (%lld,%lld), %u CTAs running\n", + SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Finished CTA #%u (%lld,%lld), %u CTAs running\n", cta_num, m_gpu->gpu_sim_cycle, m_gpu->gpu_tot_sim_cycle, m_n_active_cta); if( m_n_active_cta == 0 ) { diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index 62e0e42..667cb2d 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1392,9 +1392,9 @@ class shader_core_config : public core_config If we won't remove it, old regression will be broken. So to support the legacy config files it's best to handle in this way. */ - int num_config_to_read=N_PIPELINE_STAGES-2*(!gpgpu_tensor_core_avail); + int num_config_to_read= N_PIPELINE_STAGES - 2 * (!gpgpu_tensor_core_avail); - for (unsigned i = 0; i <num_config_to_read; i++) { + for (int i = 0; i < num_config_to_read; i++) { assert(toks); ntok = sscanf(toks,"%d", &pipe_widths[i]); assert(ntok == 1); @@ -1455,7 +1455,7 @@ class shader_core_config : public core_config bool gpgpu_dwf_reg_bankconflict; - int gpgpu_num_sched_per_core; + unsigned gpgpu_num_sched_per_core; int gpgpu_max_insn_issue_per_warp; bool gpgpu_dual_issue_diff_exec_units; diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 816159f..846773d 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -43,46 +43,28 @@ static int sg_argc = 3; static const char *sg_argv[] = {"", "-config","gpgpusim.config"}; -GPGPUsim_ctx* the_gpgpusim = NULL; - -GPGPUsim_ctx* GPGPUsim_ctx_ptr(){ - if(the_gpgpusim == NULL) - the_gpgpusim = GPGPU_Context()->the_gpgpusim; - - return the_gpgpusim; -} - -class gpgpu_sim* g_the_gpu() { - return GPGPUsim_ctx_ptr()->g_the_gpu; -} - -class stream_manager* g_stream_manager() { - return GPGPUsim_ctx_ptr()->g_stream_manager; -} - -static void print_simulation_time(); - -void *gpgpu_sim_thread_sequential(void*) +void * gpgpu_sim_thread_sequential(void * ctx_ptr) { + gpgpu_context * ctx = (gpgpu_context *)ctx_ptr; // at most one kernel running at a time bool done; do { - sem_wait(&(GPGPUsim_ctx_ptr()->g_sim_signal_start)); + sem_wait(&(ctx->the_gpgpusim->g_sim_signal_start)); done = true; - if( GPGPUsim_ctx_ptr()->g_the_gpu->get_more_cta_left() ) { + if( ctx->the_gpgpusim->g_the_gpu->get_more_cta_left() ) { done = false; - GPGPUsim_ctx_ptr()->g_the_gpu->init(); - while( GPGPUsim_ctx_ptr()->g_the_gpu->active() ) { - GPGPUsim_ctx_ptr()->g_the_gpu->cycle(); - GPGPUsim_ctx_ptr()->g_the_gpu->deadlock_check(); + ctx->the_gpgpusim->g_the_gpu->init(); + while( ctx->the_gpgpusim->g_the_gpu->active() ) { + ctx->the_gpgpusim->g_the_gpu->cycle(); + ctx->the_gpgpusim->g_the_gpu->deadlock_check(); } - GPGPUsim_ctx_ptr()->g_the_gpu->print_stats(); - GPGPUsim_ctx_ptr()->g_the_gpu->update_stats(); - print_simulation_time(); + ctx->the_gpgpusim->g_the_gpu->print_stats(); + ctx->the_gpgpusim->g_the_gpu->update_stats(); + ctx->print_simulation_time(); } - sem_post(&(GPGPUsim_ctx_ptr()->g_sim_signal_finish)); + sem_post(&(ctx->the_gpgpusim->g_sim_signal_finish)); } while(!done); - sem_post(&(GPGPUsim_ctx_ptr()->g_sim_signal_exit)); + sem_post(&(ctx->the_gpgpusim->g_sim_signal_exit)); return NULL; } @@ -94,8 +76,9 @@ static void termination_callback() fflush(stdout); } -void *gpgpu_sim_thread_concurrent(void*) +void *gpgpu_sim_thread_concurrent(void * ctx_ptr) { + gpgpu_context * ctx = (gpgpu_context *)ctx_ptr; atexit(termination_callback); // concurrent kernel execution simulation thread do { @@ -103,19 +86,19 @@ void *gpgpu_sim_thread_concurrent(void*) printf("GPGPU-Sim: *** simulation thread starting and spinning waiting for work ***\n"); fflush(stdout); } - while( GPGPUsim_ctx_ptr()->g_stream_manager->empty_protected() && !GPGPUsim_ctx_ptr()->g_sim_done ) + while( ctx->the_gpgpusim->g_stream_manager->empty_protected() && !ctx->the_gpgpusim->g_sim_done ) ; if(g_debug_execution >= 3) { printf("GPGPU-Sim: ** START simulation thread (detected work) **\n"); - GPGPUsim_ctx_ptr()->g_stream_manager->print(stdout); + ctx->the_gpgpusim->g_stream_manager->print(stdout); fflush(stdout); } - pthread_mutex_lock(&(GPGPUsim_ctx_ptr()->g_sim_lock)); - GPGPUsim_ctx_ptr()->g_sim_active = true; - pthread_mutex_unlock(&(GPGPUsim_ctx_ptr()->g_sim_lock)); + pthread_mutex_lock(&(ctx->the_gpgpusim->g_sim_lock)); + ctx->the_gpgpusim->g_sim_active = true; + pthread_mutex_unlock(&(ctx->the_gpgpusim->g_sim_lock)); bool active = false; bool sim_cycles = false; - GPGPUsim_ctx_ptr()->g_the_gpu->init(); + ctx->the_gpgpusim->g_the_gpu->init(); do { // check if a kernel has completed // launch operation on device if one is pending and can be run @@ -127,82 +110,82 @@ void *gpgpu_sim_thread_concurrent(void*) // another kernel, the gpu is not re-initialized and the inter-kernel // behaviour may be incorrect. Check that a kernel has finished and // no other kernel is currently running. - if(GPGPUsim_ctx_ptr()->g_stream_manager->operation(&sim_cycles) && !GPGPUsim_ctx_ptr()->g_the_gpu->active()) + if(ctx->the_gpgpusim->g_stream_manager->operation(&sim_cycles) && !ctx->the_gpgpusim->g_the_gpu->active()) break; //functional simulation - if( GPGPUsim_ctx_ptr()->g_the_gpu->is_functional_sim()) { - kernel_info_t * kernel = GPGPUsim_ctx_ptr()->g_the_gpu->get_functional_kernel(); + if( ctx->the_gpgpusim->g_the_gpu->is_functional_sim()) { + kernel_info_t * kernel = ctx->the_gpgpusim->g_the_gpu->get_functional_kernel(); assert(kernel); - GPGPUsim_ctx_ptr()->gpgpu_ctx->func_sim->gpgpu_cuda_ptx_sim_main_func(*kernel); - GPGPUsim_ctx_ptr()->g_the_gpu->finish_functional_sim(kernel); + ctx->the_gpgpusim->gpgpu_ctx->func_sim->gpgpu_cuda_ptx_sim_main_func(*kernel); + ctx->the_gpgpusim->g_the_gpu->finish_functional_sim(kernel); } //performance simulation - if( GPGPUsim_ctx_ptr()->g_the_gpu->active() ) { - GPGPUsim_ctx_ptr()->g_the_gpu->cycle(); + if( ctx->the_gpgpusim->g_the_gpu->active() ) { + ctx->the_gpgpusim->g_the_gpu->cycle(); sim_cycles = true; - GPGPUsim_ctx_ptr()->g_the_gpu->deadlock_check(); + ctx->the_gpgpusim->g_the_gpu->deadlock_check(); }else { - if(GPGPUsim_ctx_ptr()->g_the_gpu->cycle_insn_cta_max_hit()){ - GPGPUsim_ctx_ptr()->g_stream_manager->stop_all_running_kernels(); - GPGPUsim_ctx_ptr()->g_sim_done = true; - GPGPUsim_ctx_ptr()->break_limit = true; + if(ctx->the_gpgpusim->g_the_gpu->cycle_insn_cta_max_hit()){ + ctx->the_gpgpusim->g_stream_manager->stop_all_running_kernels(); + ctx->the_gpgpusim->g_sim_done = true; + ctx->the_gpgpusim->break_limit = true; } } - active=GPGPUsim_ctx_ptr()->g_the_gpu->active() || !(GPGPUsim_ctx_ptr()->g_stream_manager->empty_protected()); + active=ctx->the_gpgpusim->g_the_gpu->active() || !(ctx->the_gpgpusim->g_stream_manager->empty_protected()); - } while( active && !GPGPUsim_ctx_ptr()->g_sim_done); + } while( active && !ctx->the_gpgpusim->g_sim_done); if(g_debug_execution >= 3) { printf("GPGPU-Sim: ** STOP simulation thread (no work) **\n"); fflush(stdout); } if(sim_cycles) { - GPGPUsim_ctx_ptr()->g_the_gpu->print_stats(); - GPGPUsim_ctx_ptr()->g_the_gpu->update_stats(); - print_simulation_time(); + ctx->the_gpgpusim->g_the_gpu->print_stats(); + ctx->the_gpgpusim->g_the_gpu->update_stats(); + ctx->print_simulation_time(); } - pthread_mutex_lock(&(GPGPUsim_ctx_ptr()->g_sim_lock)); - GPGPUsim_ctx_ptr()->g_sim_active = false; - pthread_mutex_unlock(&(GPGPUsim_ctx_ptr()->g_sim_lock)); - } while( !GPGPUsim_ctx_ptr()->g_sim_done ); + pthread_mutex_lock(&(ctx->the_gpgpusim->g_sim_lock)); + ctx->the_gpgpusim->g_sim_active = false; + pthread_mutex_unlock(&(ctx->the_gpgpusim->g_sim_lock)); + } while( !ctx->the_gpgpusim->g_sim_done ); printf("GPGPU-Sim: *** simulation thread exiting ***\n"); fflush(stdout); - if(GPGPUsim_ctx_ptr()->break_limit) { + if(ctx->the_gpgpusim->break_limit) { printf("GPGPU-Sim: ** break due to reaching the maximum cycles (or instructions) **\n"); exit(1); } - sem_post(&(GPGPUsim_ctx_ptr()->g_sim_signal_exit)); + sem_post(&(ctx->the_gpgpusim->g_sim_signal_exit)); return NULL; } -void synchronize() +void gpgpu_context::synchronize() { printf("GPGPU-Sim: synchronize waiting for inactive GPU simulation\n"); - GPGPUsim_ctx_ptr()->g_stream_manager->print(stdout); + the_gpgpusim->g_stream_manager->print(stdout); fflush(stdout); // sem_wait(&g_sim_signal_finish); bool done = false; do { - pthread_mutex_lock(&(GPGPUsim_ctx_ptr()->g_sim_lock)); - done = ( GPGPUsim_ctx_ptr()->g_stream_manager->empty() && !GPGPUsim_ctx_ptr()->g_sim_active ) || GPGPUsim_ctx_ptr()->g_sim_done; - pthread_mutex_unlock(&(GPGPUsim_ctx_ptr()->g_sim_lock)); + pthread_mutex_lock(&(the_gpgpusim->g_sim_lock)); + done = ( the_gpgpusim->g_stream_manager->empty() && !the_gpgpusim->g_sim_active ) || the_gpgpusim->g_sim_done; + pthread_mutex_unlock(&(the_gpgpusim->g_sim_lock)); } while (!done); printf("GPGPU-Sim: detected inactive GPU simulation thread\n"); fflush(stdout); // sem_post(&g_sim_signal_start); } -void exit_simulation() +void gpgpu_context::exit_simulation() { - GPGPUsim_ctx_ptr()->g_sim_done=true; + the_gpgpusim->g_sim_done=true; printf("GPGPU-Sim: exit_simulation called\n"); fflush(stdout); - sem_wait(&(GPGPUsim_ctx_ptr()->g_sim_signal_exit)); + sem_wait(&(the_gpgpusim->g_sim_signal_exit)); printf("GPGPU-Sim: simulation thread signaled exit\n"); fflush(stdout); } @@ -219,8 +202,8 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf() func_sim->ptx_opcocde_latency_options(opp); icnt_reg_options(opp); - GPGPUsim_ctx_ptr()->g_the_gpu_config = new gpgpu_sim_config(this); - GPGPUsim_ctx_ptr()->g_the_gpu_config->reg_options(opp); // register GPU microrachitecture options + the_gpgpusim->g_the_gpu_config = new gpgpu_sim_config(this); + the_gpgpusim->g_the_gpu_config->reg_options(opp); // register GPU microrachitecture options option_parser_cmdline(opp, sg_argc, sg_argv); // parse configuration options fprintf(stdout, "GPGPU-Sim: Configuration options:\n\n"); @@ -228,37 +211,37 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf() // Set the Numeric locale to a standard locale where a decimal point is a "dot" not a "comma" // so it does the parsing correctly independent of the system environment variables assert(setlocale(LC_NUMERIC,"C")); - GPGPUsim_ctx_ptr()->g_the_gpu_config->init(); + the_gpgpusim->g_the_gpu_config->init(); - GPGPUsim_ctx_ptr()->g_the_gpu = new gpgpu_sim(*(GPGPUsim_ctx_ptr()->g_the_gpu_config), this); - GPGPUsim_ctx_ptr()->g_stream_manager = new stream_manager((GPGPUsim_ctx_ptr()->g_the_gpu), func_sim->g_cuda_launch_blocking); + the_gpgpusim->g_the_gpu = new gpgpu_sim(*(the_gpgpusim->g_the_gpu_config), this); + the_gpgpusim->g_stream_manager = new stream_manager((the_gpgpusim->g_the_gpu), func_sim->g_cuda_launch_blocking); - GPGPUsim_ctx_ptr()->g_simulation_starttime = time((time_t *)NULL); + the_gpgpusim->g_simulation_starttime = time((time_t *)NULL); - sem_init(&(GPGPUsim_ctx_ptr()->g_sim_signal_start),0,0); - sem_init(&(GPGPUsim_ctx_ptr()->g_sim_signal_finish),0,0); - sem_init(&(GPGPUsim_ctx_ptr()->g_sim_signal_exit),0,0); + sem_init(&(the_gpgpusim->g_sim_signal_start),0,0); + sem_init(&(the_gpgpusim->g_sim_signal_finish),0,0); + sem_init(&(the_gpgpusim->g_sim_signal_exit),0,0); - return GPGPUsim_ctx_ptr()->g_the_gpu; + return the_gpgpusim->g_the_gpu; } -void start_sim_thread(int api) +void gpgpu_context::start_sim_thread(int api) { - if( GPGPUsim_ctx_ptr()->g_sim_done ) { - GPGPUsim_ctx_ptr()->g_sim_done = false; + if( the_gpgpusim->g_sim_done ) { + the_gpgpusim->g_sim_done = false; if( api == 1 ) { - pthread_create(&(GPGPUsim_ctx_ptr()->g_simulation_thread),NULL,gpgpu_sim_thread_concurrent,NULL); + pthread_create(&(the_gpgpusim->g_simulation_thread),NULL,gpgpu_sim_thread_concurrent,(void *)this); } else { - pthread_create(&(GPGPUsim_ctx_ptr()->g_simulation_thread),NULL,gpgpu_sim_thread_sequential,NULL); + pthread_create(&(the_gpgpusim->g_simulation_thread),NULL,gpgpu_sim_thread_sequential,(void *)this); } } } -void print_simulation_time() +void gpgpu_context::print_simulation_time() { time_t current_time, difference, d, h, m, s; current_time = time((time_t *)NULL); - difference = MAX(current_time - GPGPUsim_ctx_ptr()->g_simulation_starttime, 1); + difference = MAX(current_time - the_gpgpusim->g_simulation_starttime, 1); d = difference/(3600*24); h = difference/3600 - 24*d; @@ -268,16 +251,18 @@ void print_simulation_time() fflush(stderr); printf("\n\ngpgpu_simulation_time = %u days, %u hrs, %u min, %u sec (%u sec)\n", (unsigned)d, (unsigned)h, (unsigned)m, (unsigned)s, (unsigned)difference ); - printf("gpgpu_simulation_rate = %u (inst/sec)\n", (unsigned)(GPGPUsim_ctx_ptr()->g_the_gpu->gpu_tot_sim_insn / difference) ); - printf("gpgpu_simulation_rate = %u (cycle/sec)\n", (unsigned)(GPGPUsim_ctx_ptr()->g_the_gpu->gpu_tot_sim_cycle / difference) ); + printf("gpgpu_simulation_rate = %u (inst/sec)\n", (unsigned)(the_gpgpusim->g_the_gpu->gpu_tot_sim_insn / difference) ); + const unsigned cycles_per_sec = (unsigned)(the_gpgpusim->g_the_gpu->gpu_tot_sim_cycle / difference); + printf("gpgpu_simulation_rate = %u (cycle/sec)\n", cycles_per_sec ); + printf("gpgpu_silicon_slowdown = %ux\n", the_gpgpusim->g_the_gpu->shader_clock() * 1000 / cycles_per_sec); fflush(stdout); } -int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ) +int gpgpu_context::gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ) { - GPGPUsim_ctx_ptr()->g_the_gpu->launch(grid); - sem_post(&(GPGPUsim_ctx_ptr()->g_sim_signal_start)); - sem_wait(&(GPGPUsim_ctx_ptr()->g_sim_signal_finish)); + the_gpgpusim->g_the_gpu->launch(grid); + sem_post(&(the_gpgpusim->g_sim_signal_start)); + sem_wait(&(the_gpgpusim->g_sim_signal_finish)); return 0; } diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h index 887b3c8..9f408df 100644 --- a/src/gpgpusim_entrypoint.h +++ b/src/gpgpusim_entrypoint.h @@ -76,12 +76,4 @@ class GPGPUsim_ctx { }; -void start_sim_thread(int api); - -class gpgpu_sim* g_the_gpu(); -struct GPGPUsim_ctx* GPGPUsim_ctx_ptr(); -class stream_manager* g_stream_manager(); - -int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ); - #endif |