From 9e52c143a883f682c02d81149748cdf8aa5508f7 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 12 Jun 2019 01:42:52 -0400 Subject: Move some function from ptx_loader to gpgpu_context Signed-off-by: Mengchi Zhang --- libopencl/opencl_runtime_api.cc | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) (limited to 'libopencl') diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 97a54d8..0ec635e 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -84,6 +84,7 @@ #include "../src/gpgpusim_entrypoint.h" #include "../src/gpgpu-sim/gpu-sim.h" #include "../src/gpgpu-sim/shader.h" +#include "../libcuda/gpgpu_context.h" //# define __my_func__ __PRETTY_FUNCTION__ # if defined __cplusplus ? __GNUC_PREREQ (2, 6) : __GNUC_PREREQ (2, 4) @@ -424,6 +425,8 @@ void register_ptx_function( const char *name, function_info *impl ) void _cl_program::Build(const char *options) { + gpgpu_context *ctx; + ctx = GPGPU_Context(); printf("GPGPU-Sim OpenCL API: compiling OpenCL kernels...\n"); std::map::iterator i; for( i = m_pgm.begin(); i!= m_pgm.end(); i++ ) { @@ -576,7 +579,7 @@ void _cl_program::Build(const char *options) } } info.m_asm = tmp; - info.m_symtab = gpgpu_ptx_sim_load_ptx_from_string( tmp, source_num ); + info.m_symtab = ctx->gpgpu_ptx_sim_load_ptx_from_string( tmp, source_num ); gpgpu_ptxinfo_load_from_string( tmp, source_num ); free(tmp); } -- cgit v1.3 From 7d02cbb061485db38ed8e5f6bf06c9b2fa40eed2 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 12 Jun 2019 13:19:08 -0400 Subject: Integrate ptxinfo into gpgpu_context Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 2 +- libcuda/gpgpu_context.h | 5 +++++ libopencl/opencl_runtime_api.cc | 2 +- src/cuda-sim/ptx_loader.cc | 43 +++++++++++++++++++---------------------- src/cuda-sim/ptx_loader.h | 2 -- 5 files changed, 27 insertions(+), 27 deletions(-) (limited to 'libopencl') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 14e3329..0a6eabb 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -640,7 +640,7 @@ void** cudaRegisterFatBinaryInternal( void *fatCubin, gpgpu_context* gpgpu_ctx = } else { symtab=ctx->gpgpu_ptx_sim_load_ptx_from_string(ptx,source_num); context->add_binary(symtab,fat_cubin_handle); - gpgpu_ptxinfo_load_from_string( ptx, source_num, max_capability, context->no_of_ptx ); + ctx->gpgpu_ptxinfo_load_from_string( ptx, source_num, max_capability, context->no_of_ptx ); } source_num++; ctx->api->load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 16626eb..d3db1ad 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -1,19 +1,24 @@ #ifndef __gpgpu_context_h__ #define __gpgpu_context_h__ #include "cuda_api_object.h" +#include "../src/cuda-sim/ptx_loader.h" class gpgpu_context { public: gpgpu_context() { api = new cuda_runtime_api(); + ptxinfo = new ptxinfo_data(); } // global list // objects pointers for each file cuda_runtime_api* api; + ptxinfo_data* ptxinfo; // member function list 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 ); + void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version); + void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version=20, int no_of_ptx=0 ); }; gpgpu_context* GPGPU_Context(); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 0ec635e..50a02fa 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -580,7 +580,7 @@ void _cl_program::Build(const char *options) } info.m_asm = tmp; info.m_symtab = ctx->gpgpu_ptx_sim_load_ptx_from_string( tmp, source_num ); - gpgpu_ptxinfo_load_from_string( tmp, source_num ); + ctx->gpgpu_ptxinfo_load_from_string( tmp, source_num ); free(tmp); } printf("GPGPU-Sim OpenCL API: finished compiling OpenCL kernels.\n"); diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 9857741..2a6d930 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -332,7 +332,7 @@ char* get_app_binary_name(){ return self_exe_path; } -void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version) +void gpgpu_context::gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version) { std::string ptxas_filename(std::string(filename) + "as"); char buff[1024], extra_flags[1024]; @@ -352,17 +352,16 @@ void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_versio } FILE *ptxinfo_in; - ptxinfo_data ptxinfo; - ptxinfo.g_ptxinfo_filename = strdup(ptxas_filename.c_str()); - ptxinfo_in = fopen(ptxinfo.g_ptxinfo_filename,"r"); - ptxinfo_lex_init(&(ptxinfo.scanner)); - ptxinfo_set_in(ptxinfo_in, ptxinfo.scanner); - ptxinfo_parse(ptxinfo.scanner, &ptxinfo); - ptxinfo_lex_destroy(ptxinfo.scanner); + ptxinfo->g_ptxinfo_filename = strdup(ptxas_filename.c_str()); + ptxinfo_in = fopen(ptxinfo->g_ptxinfo_filename,"r"); + ptxinfo_lex_init(&(ptxinfo->scanner)); + ptxinfo_set_in(ptxinfo_in, ptxinfo->scanner); + ptxinfo_parse(ptxinfo->scanner, ptxinfo); + ptxinfo_lex_destroy(ptxinfo->scanner); fclose(ptxinfo_in); } -void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version, int no_of_ptx ) +void gpgpu_context::gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version, int no_of_ptx ) { //do ptxas for individual files instead of one big embedded ptx. This prevents the duplicate defs and declarations. char ptx_file[1000]; @@ -420,14 +419,13 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num if( result != 0 ) { // 65280 = duplicate errors if (result == 65280) { - ptxinfo_data ptxinfo; FILE *ptxinfo_in; ptxinfo_in = fopen(tempfile_ptxinfo,"r"); - ptxinfo.g_ptxinfo_filename = tempfile_ptxinfo; - ptxinfo_lex_init(&(ptxinfo.scanner)); - ptxinfo_set_in(ptxinfo_in, ptxinfo.scanner); - ptxinfo_parse(ptxinfo.scanner, &ptxinfo); - ptxinfo_lex_destroy(ptxinfo.scanner); + ptxinfo->g_ptxinfo_filename = tempfile_ptxinfo; + ptxinfo_lex_init(&(ptxinfo->scanner)); + ptxinfo_set_in(ptxinfo_in, ptxinfo->scanner); + ptxinfo_parse(ptxinfo->scanner, ptxinfo); + ptxinfo_lex_destroy(ptxinfo->scanner); fclose(ptxinfo_in); fix_duplicate_errors(fname2); @@ -507,18 +505,17 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num } } - ptxinfo_data ptxinfo; if(no_of_ptx>0) - ptxinfo.g_ptxinfo_filename = final_tempfile_ptxinfo; + ptxinfo->g_ptxinfo_filename = final_tempfile_ptxinfo; else - ptxinfo.g_ptxinfo_filename = tempfile_ptxinfo; + ptxinfo->g_ptxinfo_filename = tempfile_ptxinfo; FILE *ptxinfo_in; - ptxinfo_in = fopen(ptxinfo.g_ptxinfo_filename,"r"); + ptxinfo_in = fopen(ptxinfo->g_ptxinfo_filename,"r"); - ptxinfo_lex_init(&(ptxinfo.scanner)); - ptxinfo_set_in(ptxinfo_in, ptxinfo.scanner); - ptxinfo_parse(ptxinfo.scanner, &ptxinfo); - ptxinfo_lex_destroy(ptxinfo.scanner); + ptxinfo_lex_init(&(ptxinfo->scanner)); + ptxinfo_set_in(ptxinfo_in, ptxinfo->scanner); + ptxinfo_parse(ptxinfo->scanner, ptxinfo); + ptxinfo_lex_destroy(ptxinfo->scanner); fclose(ptxinfo_in); snprintf(commandline,1024,"rm -f *info"); diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index c3ce888..3637bff 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -44,8 +44,6 @@ class ptxinfo_data{ extern bool g_override_embedded_ptx; extern int no_of_ptx; //counter to track number of ptx files to be extracted in an application. -void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version=20, int no_of_ptx=0 ); -void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version ); char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str); bool keep_intermediate_files(); -- cgit v1.3 From cb678c3670de4a435a3260ed80dc476da3860082 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 1 Jul 2019 16:23:33 -0400 Subject: Move g_debug_ir_generation and GPGPUSim_Init Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 674 ++++++++++++++++++++++------------------ libcuda/gpgpu_context.h | 2 + libopencl/opencl_runtime_api.cc | 4 +- src/cuda-sim/ptx_parser.cc | 3 +- src/cuda-sim/ptx_parser.h | 3 + src/gpgpusim_entrypoint.cc | 5 +- src/gpgpusim_entrypoint.h | 1 - 7 files changed, 375 insertions(+), 317 deletions(-) (limited to 'libopencl') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 61f8415..a34727f 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -197,7 +197,7 @@ void register_ptx_function( const char *name, function_info *impl ) # endif #endif -struct _cuda_device_id *GPGPUSim_Init() +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; @@ -255,9 +255,10 @@ struct _cuda_device_id *GPGPUSim_Init() static CUctx_st* GPGPUSim_Context() { //static CUctx_st *the_context = NULL; + gpgpu_context *cur_ctx = GPGPU_Context(); CUctx_st *the_context = GPGPUsim_ctx_ptr()->the_context; if( the_context == NULL ) { - _cuda_device_id *the_gpu = GPGPUSim_Init(); + _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; } @@ -496,7 +497,7 @@ cudaError_t cudaSetDeviceInternal(int device, gpgpu_context* gpgpu_ctx = NULL) announce_call(__my_func__); } //set the active device to run cuda - if ( device <= GPGPUSim_Init()->num_devices() ) { + if ( device <= ctx->GPGPUSim_Init()->num_devices() ) { ctx->api->g_active_device = device; return g_last_cudaError = cudaSuccess; } else { @@ -519,6 +520,55 @@ cudaError_t cudaGetDeviceInternal(int *device, gpgpu_context* gpgpu_ctx = NULL) return g_last_cudaError = cudaSuccess; } +__host__ cudaError_t CUDARTAPI cudaDeviceGetLimitInternal( size_t* pValue, cudaLimit limit, 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_device_id *dev = ctx->GPGPUSim_Init(); + const struct cudaDeviceProp *prop = dev->get_prop(); + const gpgpu_sim_config& config=dev->get_gpgpu()->get_config(); + switch(limit) { + case 0: // cudaLimitStackSize + *pValue=config.stack_limit(); + break; + case 2: // cudaLimitMallocHeapSize + *pValue=config.heap_limit(); + break; +#if (CUDART_VERSION > 5050) + case 3: // cudaLimitDevRuntimeSyncDepth + if(prop->major > 2){ + *pValue=config.sync_depth_limit(); + break; + } + else{ + printf("ERROR:Limit %s is not supported on this architecture \n",limit); + abort(); + } + case 4: // cudaLimitDevRuntimePendingLaunchCount + if(prop->major > 2){ + *pValue=config.pending_launch_count_limit(); + break; + } + else{ + printf("ERROR:Limit %s is not supported on this architecture \n",limit); + abort(); + } +#endif + default: + printf("ERROR:Limit %s unimplemented \n",limit); + abort(); + } + return g_last_cudaError = cudaSuccess; + +} + void** cudaRegisterFatBinaryInternal( void *fatCubin, gpgpu_context* gpgpu_ctx = NULL) { @@ -739,6 +789,260 @@ cudaError_t cudaConfigureCallInternal(dim3 gridDim, dim3 blockDim, size_t shared return g_last_cudaError = cudaSuccess; } +__host__ cudaError_t CUDARTAPI cudaGetDeviceCountInternal(int *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_device_id *dev = ctx->GPGPUSim_Init(); + *count = dev->num_devices(); + return g_last_cudaError = cudaSuccess; +} + +__host__ cudaError_t CUDARTAPI cudaGetDevicePropertiesInternal(struct cudaDeviceProp *prop, int device, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + _cuda_device_id *dev = ctx->GPGPUSim_Init(); + if (device <= dev->num_devices() ) { + *prop= *dev->get_prop(); + return g_last_cudaError = cudaSuccess; + } else { + return g_last_cudaError = cudaErrorInvalidDevice; + } +} + +#if (CUDART_VERSION > 5000) +__host__ cudaError_t CUDARTAPI cudaDeviceGetAttributeInternal(int *value, enum cudaDeviceAttr attr, int device, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + const struct cudaDeviceProp *prop; + _cuda_device_id *dev = ctx->GPGPUSim_Init(); + if (device <= dev->num_devices() ) { + prop = dev->get_prop(); + switch (attr) { + case 1: + *value= prop->maxThreadsDim[0] * prop->maxThreadsDim[1] * prop->maxThreadsDim[2] * prop->maxGridSize[0] * prop->maxGridSize[1] * prop->maxGridSize[2]; + break; + case 2: + *value= prop->maxThreadsDim[0]; + break; + case 3: + *value= prop->maxThreadsDim[1]; + break; + case 4: + *value= prop->maxThreadsDim[2]; + break; + case 5: + *value= prop->maxGridSize[0]; + break; + case 6: + *value= prop->maxGridSize[1]; + break; + case 7: + *value= prop->maxGridSize[2]; + break; + case 8: + *value= prop->sharedMemPerBlock; + break; + case 9: + *value= prop->totalConstMem; + break; + case 10: + *value= prop->warpSize; + break; + case 11: + *value= 16;//dummy value + break; + case 12: + *value= prop->regsPerBlock; + break; + case 13: + *value= 1480000;//for 1080ti + break; + case 14: + *value= prop->textureAlignment ; + break; + case 15: + *value = 0; + break; + case 16: + *value= prop->multiProcessorCount ; + break; + case 17: + case 18: + case 19: + *value = 0; + break; + case 21: + case 22: + case 23: + case 24: + case 25: + case 26: + case 27: + case 28: + case 42: + case 45: + case 46: + case 47: + case 48: + case 49: + case 52: + case 53: + case 55: + case 56: + case 57: + case 58: + case 59: + case 60: + case 61: + case 62: + case 63: + case 64: + case 66: + case 67: + case 69: + case 70: + case 71: + case 73: + case 74: + case 77: + *value = 1000;//dummy value + break; + case 29: + case 43: + case 54: + case 65: + case 68: + case 72: + *value = 10;//dummy value + break; + case 30: + case 51: + *value = 128;//dummy value + break; + case 31: + *value = 1; + break; + case 32: + *value = 0; + break; + case 33: + case 50: + *value = 0;//dummy value + break; + case 34: + *value= 0; + break; + case 35: + *value = 0; + break; + case 36: + *value = 1250000;//CK value for 1080ti + break; + case 37: + *value = 352;//value for 1080ti + break; + case 38: + *value = 3000000;//value for 1080ti + break; + case 39: + *value= dev->get_gpgpu()->threads_per_core(); + break; + case 40: + *value= 0; + break; + case 41: + *value= 0; + break; + case 75://cudaDevAttrComputeCapabilityMajor + *value= prop->major ; + break; + case 76://cudaDevAttrComputeCapabilityMinor + *value= prop->minor ; + break; + case 78: + *value= 0 ; //TODO: as of now, we dont support stream priorities. + break; + case 79: + *value= 0; + break; + case 80: + *value= 0; + break; + #if (CUDART_VERSION > 5050) + case 81: + *value= prop->sharedMemPerMultiprocessor; + break; + case 82: + *value= prop->regsPerMultiprocessor; + break; + #endif + case 83: + case 84: + case 85: + case 86: + *value= 0; + break; + case 87: + *value= 4;//dummy value + break; + case 88: + case 89: + case 90: + case 91: + case 95: + *value= 0; + break; + default: + printf("ERROR: Attribute number %d unimplemented \n",attr); + abort(); + } + return g_last_cudaError = cudaSuccess; + } else { + return g_last_cudaError = cudaErrorInvalidDevice; + } +} +#endif + +__host__ cudaError_t CUDARTAPI cudaChooseDeviceInternal(int *device, const struct cudaDeviceProp *prop, 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_device_id *dev = ctx->GPGPUSim_Init(); + *device = dev->get_id(); + return g_last_cudaError = cudaSuccess; +} + cudaError_t cudaSetupArgumentInternal(const void *arg, size_t size, size_t offset, gpgpu_context* gpgpu_ctx = NULL) { gpgpu_context *ctx; @@ -1058,6 +1362,52 @@ cudaError_t cudaHostAllocInternal(void **pHost, size_t bytes, unsigned int flag #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 ((prop.regsPerBlock / attr->numRegs) < max) + max = prop.regsPerBlock / attr->numRegs; + + 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; +} + + /******************************************************************************* * * * * @@ -1503,235 +1853,24 @@ __host__ cudaError_t CUDARTAPI cudaGetSymbolSize(size_t *size, const char *symbo *******************************************************************************/ __host__ cudaError_t CUDARTAPI cudaGetDeviceCount(int *count) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - _cuda_device_id *dev = GPGPUSim_Init(); - *count = dev->num_devices(); - return g_last_cudaError = cudaSuccess; + return cudaGetDeviceCountInternal(count); } - -__host__ cudaError_t CUDARTAPI cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device) -{ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - _cuda_device_id *dev = GPGPUSim_Init(); - if (device <= dev->num_devices() ) { - *prop= *dev->get_prop(); - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorInvalidDevice; - } + +__host__ cudaError_t CUDARTAPI cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device) +{ + return cudaGetDevicePropertiesInternal(prop, device); } #if (CUDART_VERSION > 5000) __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDeviceAttr attr, int device) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - const struct cudaDeviceProp *prop; - _cuda_device_id *dev = GPGPUSim_Init(); - if (device <= dev->num_devices() ) { - prop = dev->get_prop(); - switch (attr) { - case 1: - *value= prop->maxThreadsDim[0] * prop->maxThreadsDim[1] * prop->maxThreadsDim[2] * prop->maxGridSize[0] * prop->maxGridSize[1] * prop->maxGridSize[2]; - break; - case 2: - *value= prop->maxThreadsDim[0]; - break; - case 3: - *value= prop->maxThreadsDim[1]; - break; - case 4: - *value= prop->maxThreadsDim[2]; - break; - case 5: - *value= prop->maxGridSize[0]; - break; - case 6: - *value= prop->maxGridSize[1]; - break; - case 7: - *value= prop->maxGridSize[2]; - break; - case 8: - *value= prop->sharedMemPerBlock; - break; - case 9: - *value= prop->totalConstMem; - break; - case 10: - *value= prop->warpSize; - break; - case 11: - *value= 16;//dummy value - break; - case 12: - *value= prop->regsPerBlock; - break; - case 13: - *value= 1480000;//for 1080ti - break; - case 14: - *value= prop->textureAlignment ; - break; - case 15: - *value = 0; - break; - case 16: - *value= prop->multiProcessorCount ; - break; - case 17: - case 18: - case 19: - *value = 0; - break; - case 21: - case 22: - case 23: - case 24: - case 25: - case 26: - case 27: - case 28: - case 42: - case 45: - case 46: - case 47: - case 48: - case 49: - case 52: - case 53: - case 55: - case 56: - case 57: - case 58: - case 59: - case 60: - case 61: - case 62: - case 63: - case 64: - case 66: - case 67: - case 69: - case 70: - case 71: - case 73: - case 74: - case 77: - *value = 1000;//dummy value - break; - case 29: - case 43: - case 54: - case 65: - case 68: - case 72: - *value = 10;//dummy value - break; - case 30: - case 51: - *value = 128;//dummy value - break; - case 31: - *value = 1; - break; - case 32: - *value = 0; - break; - case 33: - case 50: - *value = 0;//dummy value - break; - case 34: - *value= 0; - break; - case 35: - *value = 0; - break; - case 36: - *value = 1250000;//CK value for 1080ti - break; - case 37: - *value = 352;//value for 1080ti - break; - case 38: - *value = 3000000;//value for 1080ti - break; - case 39: - *value= dev->get_gpgpu()->threads_per_core(); - break; - case 40: - *value= 0; - break; - case 41: - *value= 0; - break; - case 75://cudaDevAttrComputeCapabilityMajor - *value= prop->major ; - break; - case 76://cudaDevAttrComputeCapabilityMinor - *value= prop->minor ; - break; - case 78: - *value= 0 ; //TODO: as of now, we dont support stream priorities. - break; - case 79: - *value= 0; - break; - case 80: - *value= 0; - break; - #if (CUDART_VERSION > 5050) - case 81: - *value= prop->sharedMemPerMultiprocessor; - break; - case 82: - *value= prop->regsPerMultiprocessor; - break; - #endif - case 83: - case 84: - case 85: - case 86: - *value= 0; - break; - case 87: - *value= 4;//dummy value - break; - case 88: - case 89: - case 90: - case 91: - case 92: - case 93: - case 94: - case 95: - *value= 0; - break; - default: - printf("ERROR: Attribute number %d unimplemented \n",attr); - abort(); - } - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorInvalidDevice; - } + return cudaDeviceGetAttributeInternal(value, attr, device); } #endif __host__ cudaError_t CUDARTAPI cudaChooseDevice(int *device, const struct cudaDeviceProp *prop) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - _cuda_device_id *dev = GPGPUSim_Init(); - *device = dev->get_id(); - return g_last_cudaError = cudaSuccess; + return cudaChooseDeviceInternal(device, prop); } __host__ cudaError_t CUDARTAPI cudaSetDevice(int device) @@ -1744,47 +1883,9 @@ __host__ cudaError_t CUDARTAPI cudaGetDevice(int *device) return cudaGetDeviceInternal(device); } -__host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit limit ) +__host__ cudaError_t CUDARTAPI cudaDeviceGetLimit( size_t* pValue, cudaLimit limit ) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - _cuda_device_id *dev = GPGPUSim_Init(); - const struct cudaDeviceProp *prop = dev->get_prop(); - const gpgpu_sim_config& config=dev->get_gpgpu()->get_config(); - switch(limit) { - case 0: // cudaLimitStackSize - *pValue=config.stack_limit(); - break; - case 2: // cudaLimitMallocHeapSize - *pValue=config.heap_limit(); - break; -#if (CUDART_VERSION > 5050) - case 3: // cudaLimitDevRuntimeSyncDepth - if(prop->major > 2){ - *pValue=config.sync_depth_limit(); - break; - } - else{ - printf("ERROR:Limit %s is not supported on this architecture \n",limit); - abort(); - } - case 4: // cudaLimitDevRuntimePendingLaunchCount - if(prop->major > 2){ - *pValue=config.pending_launch_count_limit(); - break; - } - else{ - printf("ERROR:Limit %s is not supported on this architecture \n",limit); - abort(); - } -#endif - default: - printf("ERROR:Limit %s unimplemented \n",limit); - abort(); - } - return g_last_cudaError = cudaSuccess; - + return cudaDeviceGetLimitInternal( pValue, limit ); } __host__ cudaError_t CUDARTAPI cudaStreamGetPriority ( cudaStream_t hStream, int* priority ) @@ -3167,58 +3268,9 @@ cudaError_t CUDARTAPI cudaSetDeviceFlags( int flags ) } } -size_t getMaxThreadsPerBlock(struct cudaFuncAttributes *attr) { - _cuda_device_id *dev = GPGPUSim_Init(); - struct cudaDeviceProp prop; - - prop = *dev->get_prop(); - - size_t max = prop.maxThreadsPerBlock; - - if ((prop.regsPerBlock / attr->numRegs) < max) - max = prop.regsPerBlock / attr->numRegs; - - return max; -} - cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *attr, const char *hostFun ) { - 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); -#if CUDART_VERSION >= 3000 - attr->ptxVersion = kinfo->ptx_version; - attr->binaryVersion = kinfo->sm_target; -#endif - } - return g_last_cudaError = cudaSuccess; -} - -cudaError_t CUDARTAPI cudaEventCreateWithFlags(cudaEvent_t *event, int flags) -{ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - CUevent_st *e = new CUevent_st(flags==cudaEventBlockingSync); - g_timer_events[e->get_uid()] = e; -#if CUDART_VERSION >= 3000 - *event = e; -#else - *event = e->get_uid(); -#endif - return g_last_cudaError = cudaSuccess; + return cudaFuncGetAttributesInternal(attr, hostFun ); } cudaError_t CUDARTAPI cudaDriverGetVersion(int *driverVersion) diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index fa5d02b..dd8b51d 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -26,6 +26,8 @@ class gpgpu_context { void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version=20, int no_of_ptx=0 ); 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(); + struct _cuda_device_id *GPGPUSim_Init(); }; gpgpu_context* GPGPU_Context(); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 50a02fa..e91e2e0 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -648,7 +648,9 @@ class _cl_device_id *GPGPUSim_Init() { static _cl_device_id *the_device = NULL; if( !the_device ) { - gpgpu_sim *the_gpu = gpgpu_ptx_sim_init_perf(); + 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); diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 2872b84..9094ec3 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -48,7 +48,6 @@ void set_ptx_warp_size(const struct core_config * warp_size) g_shader_core_config=warp_size; } -static bool g_debug_ir_generation=false; const char *g_filename; // the program intermediate representation... @@ -70,7 +69,7 @@ const char *decode_token( int type ) return g_ptx_token_decode[type].c_str(); } -void read_parser_environment_variables() +void ptx_recognizer::read_parser_environment_variables() { g_filename = getenv("PTX_SIM_KERNELFILE"); char *dbg_level = getenv("PTX_SIM_DEBUG"); diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h index 25c01fe..bc9a872 100644 --- a/src/cuda-sim/ptx_parser.h +++ b/src/cuda-sim/ptx_parser.h @@ -62,6 +62,7 @@ class ptx_recognizer { g_error_detected = 0; g_entry_func_param_index=0; g_func_info = NULL; + g_debug_ir_generation=false; } // global list yyscan_t scanner; @@ -103,6 +104,7 @@ class ptx_recognizer { unsigned g_entry_func_param_index; function_info *g_func_info; operand_info g_return_var; + bool g_debug_ir_generation; // member function list void init_directive_state(); @@ -169,6 +171,7 @@ class ptx_recognizer { void start_inst_group(); void end_inst_group(); bool check_for_duplicates( const char *identifier ); + void read_parser_environment_variables(); }; diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index b09674a..476a9d4 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -35,6 +35,7 @@ #include "gpgpu-sim/gpu-sim.h" #include "gpgpu-sim/icnt_wrapper.h" #include "stream_manager.h" +#include "../libcuda/gpgpu_context.h" #define MAX(a,b) (((a)>(b))?(a):(b)) @@ -208,12 +209,12 @@ void exit_simulation() extern bool g_cuda_launch_blocking; -gpgpu_sim *gpgpu_ptx_sim_init_perf() +gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf() { srand(1); print_splash(); read_sim_environment_variables(); - read_parser_environment_variables(); + ptx_parser->read_parser_environment_variables(); option_parser_t opp = option_parser_create(); ptx_reg_options(opp); diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h index 9ce7fef..a443151 100644 --- a/src/gpgpusim_entrypoint.h +++ b/src/gpgpusim_entrypoint.h @@ -74,7 +74,6 @@ struct GPGPUsim_ctx { }; -class gpgpu_sim *gpgpu_ptx_sim_init_perf(); void start_sim_thread(int api); class gpgpu_sim* g_the_gpu(); -- cgit v1.3 From 7e9591a24ffbe645cfaa7a84b33a82edb9fca0bc Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Sun, 7 Jul 2019 23:56:21 -0400 Subject: g_keep_intermediate_files Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 1 + libopencl/opencl_runtime_api.cc | 2 +- src/cuda-sim/ptx_ir.h | 3 --- src/cuda-sim/ptx_loader.cc | 7 +++---- src/cuda-sim/ptx_loader.h | 3 ++- 5 files changed, 7 insertions(+), 9 deletions(-) (limited to 'libopencl') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index d819559..a8b60f4 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -32,6 +32,7 @@ class gpgpu_context { class symbol_table* init_parser(const char*); class gpgpu_sim *gpgpu_ptx_sim_init_perf(); struct _cuda_device_id *GPGPUSim_Init(); + void ptx_reg_options(option_parser_t opp); }; gpgpu_context* GPGPU_Context(); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index e91e2e0..03ec80c 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -537,7 +537,7 @@ void _cl_program::Build(const char *options) exit(1); } } - if( !g_keep_intermediate_files ) { + if( !ctx->ptxinfo->g_keep_intermediate_files ) { // clean up files... snprintf(commandline,1024,"rm -f %s", cl_fname ); int result = system(commandline); diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 1af85de..fd869c6 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1575,11 +1575,8 @@ struct textureInfo { extern std::map g_sym_name_to_symbol_table; -extern bool g_keep_intermediate_files; - void gpgpu_ptx_assemble( std::string kname, void *kinfo ); #include "../option_parser.h" -void ptx_reg_options(option_parser_t opp); unsigned ptx_kernel_shmem_size( void *kernel_impl ); unsigned ptx_kernel_nregs( void *kernel_impl ); diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 921d1e6..e0d9a11 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -58,17 +58,16 @@ extern int ptxinfo_lex_destroy(yyscan_t scanner); static bool g_save_embedded_ptx; static int g_occupancy_sm_number; -bool g_keep_intermediate_files; bool m_ptx_save_converted_ptxplus; -bool keep_intermediate_files() {return g_keep_intermediate_files;} +bool ptxinfo_data::keep_intermediate_files() {return g_keep_intermediate_files;} -void ptx_reg_options(option_parser_t opp) +void gpgpu_context::ptx_reg_options(option_parser_t opp) { option_parser_register(opp, "-save_embedded_ptx", OPT_BOOL, &g_save_embedded_ptx, "saves ptx files embedded in binary as .ptx", "0"); - option_parser_register(opp, "-keep", OPT_BOOL, &g_keep_intermediate_files, + option_parser_register(opp, "-keep", OPT_BOOL, &(ptxinfo->g_keep_intermediate_files), "keep intermediate files created by GPGPU-Sim when interfacing with external programs", "0"); option_parser_register(opp, "-gpgpu_ptx_save_converted_ptxplus", OPT_BOOL, diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index c214b95..aa3e2f3 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -42,7 +42,9 @@ class ptxinfo_data{ unsigned col; const char *g_ptxinfo_filename; class gpgpu_context* gpgpu_ctx; + bool g_keep_intermediate_files; void ptxinfo_addinfo(); + bool keep_intermediate_files(); }; @@ -50,6 +52,5 @@ extern bool g_override_embedded_ptx; extern int no_of_ptx; //counter to track number of ptx files to be extracted in an application. char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str); -bool keep_intermediate_files(); #endif -- cgit v1.3 From 98bb2046a70715c8599561ca78425a7ff99ce53b Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 8 Jul 2019 14:21:08 -0400 Subject: Move cp_count Signed-off-by: Mengchi Zhang --- libopencl/opencl_runtime_api.cc | 4 +++- src/cuda-sim/cuda-sim.cc | 5 ++--- src/cuda-sim/cuda-sim.h | 4 +++- src/gpgpu-sim/gpu-sim.h | 4 ++-- src/gpgpusim_entrypoint.cc | 4 ++-- src/gpgpusim_entrypoint.h | 1 - 6 files changed, 12 insertions(+), 10 deletions(-) (limited to 'libopencl') diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 03ec80c..d302ff8 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -877,6 +877,8 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, const cl_event * event_wait_list, cl_event * event) CL_API_SUFFIX__VERSION_1_0 { + gpgpu_context *ctx; + ctx = GPGPU_Context(); int _global_size[3]; int zeros[3] = { 0, 0, 0}; printf("\n\n\n"); @@ -956,7 +958,7 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, } kernel_info_t *grid = gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu); if ( g_ptx_sim_mode ) - gpgpu_opencl_ptx_sim_main_func( grid ); + ctx->func_sim->gpgpu_opencl_ptx_sim_main_func( grid ); else gpgpu_opencl_ptx_sim_main_perf( grid ); return CL_SUCCESS; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index df0bbd7..12f38af 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -61,7 +61,6 @@ int g_debug_execution = 0; int g_debug_thread_uid = 0; addr_t g_debug_pc = 0xBEEF1518; // Output debug information to file options -int cp_count; int cp_cta_resume; unsigned g_ptx_sim_num_insn = 0; @@ -2146,7 +2145,7 @@ unsigned max_cta (const struct gpgpu_ptx_sim_info *kernel_info, unsigned threads This function simulates the CUDA code functionally, it takes a kernel_info_t parameter which holds the data for the CUDA kernel to be executed !*/ -void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) +void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) { printf("GPGPU-Sim: Performing Functional Simulation, executing kernel %s...\n",kernel.name().c_str()); @@ -2312,7 +2311,7 @@ void functionalCoreSim::createWarp(unsigned warpId) void functionalCoreSim::execute(int inst_count, unsigned ctaid_cp) { - cp_count= m_gpu->checkpoint_insn_Y; + m_gpu->gpgpu_ctx->func_sim->cp_count= m_gpu->checkpoint_insn_Y; cp_cta_resume= m_gpu->checkpoint_CTA_t; initializeCTA(ctaid_cp); diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 96d34f5..4c8d953 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -55,7 +55,6 @@ extern class kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info * struct dim3 gridDim, struct dim3 blockDim, class gpgpu_t *gpu ); -extern void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL = false ); extern void print_splash(); extern void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size ); extern void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ); @@ -137,8 +136,11 @@ class cuda_sim { public: //global variables char *opcode_latency_int; + int cp_count; //global functions void ptx_opcocde_latency_options (option_parser_t opp); + void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL = false ); + int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid ); }; #endif diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 7eeb7dd..b47ab16 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -494,6 +494,8 @@ public: */ simt_core_cluster * getSIMTCluster(); + // backward pointer + class gpgpu_context* gpgpu_ctx; private: // clocks @@ -511,8 +513,6 @@ private: void gpgpu_debug(); ///// data ///// -// backward pointer - class gpgpu_context* gpgpu_ctx; class simt_core_cluster **m_cluster; class memory_partition_unit **m_memory_partition_unit; diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 5018305..d9d1023 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -134,7 +134,7 @@ void *gpgpu_sim_thread_concurrent(void*) if( GPGPUsim_ctx_ptr()->g_the_gpu->is_functional_sim()) { kernel_info_t * kernel = GPGPUsim_ctx_ptr()->g_the_gpu->get_functional_kernel(); assert(kernel); - gpgpu_cuda_ptx_sim_main_func(*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); } @@ -287,7 +287,7 @@ int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ) /*! * This function call the CUDA PTX functional simulator */ -int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid ) +int cuda_sim::gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid ) { //calling the CUDA PTX simulator, sending the kernel by reference and a flag set to true, //the flag used by the function to distinguish OpenCL calls from the CUDA simulation calls which diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h index dfb82d0..887b3c8 100644 --- a/src/gpgpusim_entrypoint.h +++ b/src/gpgpusim_entrypoint.h @@ -83,6 +83,5 @@ struct GPGPUsim_ctx* GPGPUsim_ctx_ptr(); class stream_manager* g_stream_manager(); int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ); -int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid ); #endif -- cgit v1.3 From cda7a145b9e28eff0f3e9ac8197c2b6215755fc8 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 9 Jul 2019 14:32:24 -0400 Subject: Move g_ptx_kernel_count Signed-off-by: Mengchi Zhang --- libcuda/cuda_api_object.h | 5 ++++- libcuda/cuda_runtime_api.cc | 2 +- libcuda/gpgpu_context.h | 2 +- libopencl/opencl_runtime_api.cc | 2 +- src/cuda-sim/cuda-sim.cc | 13 ++++++------- src/cuda-sim/cuda-sim.h | 14 ++++++++------ src/gpgpu-sim/gpu-sim.cc | 4 ++-- 7 files changed, 23 insertions(+), 19 deletions(-) (limited to 'libopencl') diff --git a/libcuda/cuda_api_object.h b/libcuda/cuda_api_object.h index 0054697..db5e6a4 100644 --- a/libcuda/cuda_api_object.h +++ b/libcuda/cuda_api_object.h @@ -169,9 +169,10 @@ private: class cuda_runtime_api { public: - cuda_runtime_api() { + cuda_runtime_api( gpgpu_context* ctx ) { g_glbmap = NULL; g_active_device = 0; //active gpu that runs the code + gpgpu_ctx = ctx; } // global list std::list cuobjdumpSectionList; @@ -187,6 +188,8 @@ class cuda_runtime_api { std::map pinned_memory_size; glbmap_entry_t* g_glbmap; int g_active_device; //active gpu that runs the code + // backward pointer + class gpgpu_context* gpgpu_ctx; // member function list void cuobjdumpInit(); void extract_code_using_cuobjdump(); diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 43c5e1f..43c8bae 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -3545,7 +3545,7 @@ kernel_info_t * cuda_runtime_api::gpgpu_cuda_ptx_sim_init_grid( const char *host } entry->finalize(result->get_param_memory()); - g_ptx_kernel_count++; + gpgpu_ctx->func_sim->g_ptx_kernel_count++; fflush(stdout); if(g_debug_execution >= 4){ diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 2e21009..a2ae7b6 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -10,7 +10,7 @@ class gpgpu_context { public: gpgpu_context() { g_global_allfiles_symbol_table = NULL; - api = new cuda_runtime_api(); + api = new cuda_runtime_api(this); ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); the_gpgpusim = new GPGPUsim_ctx(this); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index d302ff8..0a6eb3e 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -956,7 +956,7 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); } - kernel_info_t *grid = gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu); + kernel_info_t *grid = ctx->func_sim->gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu); if ( g_ptx_sim_mode ) ctx->func_sim->gpgpu_opencl_ptx_sim_main_func( grid ); else diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index e86395d..0ed125a 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -56,7 +56,6 @@ typedef void * yyscan_t; int gpgpu_ptx_instruction_classification; void ** g_inst_classification_stat = NULL; void ** g_inst_op_classification_stat= NULL; -int g_ptx_kernel_count = -1; // used for classification stat collection purposes int g_debug_execution = 0; int g_debug_thread_uid = 0; addr_t g_debug_pc = 0xBEEF1518; @@ -1480,7 +1479,7 @@ bool ptx_debug_exec_dump_cond(int thd_uid, addr_t pc) return false; } -void init_inst_classification_stat() +void cuda_sim::init_inst_classification_stat() { static std::set init; if( init.find(g_ptx_kernel_count) != init.end() ) @@ -1690,7 +1689,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) ptx_file_line_stats_add_exec_count(pI); if ( gpgpu_ptx_instruction_classification ) { - init_inst_classification_stat(); + m_gpu->gpgpu_ctx->func_sim->init_inst_classification_stat(); unsigned space_type=0; switch ( pI->get_space().get_type() ) { case global_space: space_type = 10; break; @@ -1706,9 +1705,9 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) space_type = 0 ; break; } - StatAddSample( g_inst_classification_stat[g_ptx_kernel_count], op_classification); - if (space_type) StatAddSample( g_inst_classification_stat[g_ptx_kernel_count], ( int )space_type); - StatAddSample( g_inst_op_classification_stat[g_ptx_kernel_count], (int) pI->get_opcode() ); + StatAddSample( g_inst_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], op_classification); + if (space_type) StatAddSample( g_inst_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], ( int )space_type); + StatAddSample( g_inst_op_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], (int) pI->get_opcode() ); } if ( (m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn % 100000) == 0 ) { dim3 ctaid = get_ctaid(); @@ -1917,7 +1916,7 @@ size_t get_kernel_code_size( class function_info *entry ) } -kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, +kernel_info_t *cuda_sim::gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, gpgpu_ptx_sim_arg_list_t args, struct dim3 gridDim, struct dim3 blockDim, diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 16ee46e..e259f1f 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -46,13 +46,7 @@ extern int g_debug_execution; extern int g_debug_thread_uid; extern void ** g_inst_classification_stat; extern void ** g_inst_op_classification_stat; -extern int g_ptx_kernel_count; // used for classification stat collection purposes -extern class kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, - gpgpu_ptx_sim_arg_list_t args, - struct dim3 gridDim, - struct dim3 blockDim, - class gpgpu_t *gpu ); extern void print_splash(); extern void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size ); extern void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ); @@ -134,6 +128,7 @@ class cuda_sim { public: cuda_sim() { g_ptx_sim_num_insn = 0; + g_ptx_kernel_count = -1; // used for classification stat collection purposes } //global variables char *opcode_latency_int; @@ -151,10 +146,17 @@ class cuda_sim { int g_ptxinfo_error_detected; unsigned g_ptx_sim_num_insn; char *cdp_latency_str; + int g_ptx_kernel_count; // used for classification stat collection purposes //global functions void ptx_opcocde_latency_options (option_parser_t opp); void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL = false ); int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid ); + void init_inst_classification_stat(); + kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, + gpgpu_ptx_sim_arg_list_t args, + struct dim3 gridDim, + struct dim3 blockDim, + gpgpu_t *gpu ); }; #endif diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 4f9ccbf..9f47067 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1201,8 +1201,8 @@ void gpgpu_sim::gpu_print_stat() insn_warp_occ_print(stdout); } if ( gpgpu_ptx_instruction_classification ) { - StatDisp( g_inst_classification_stat[g_ptx_kernel_count]); - StatDisp( g_inst_op_classification_stat[g_ptx_kernel_count]); + StatDisp( g_inst_classification_stat[gpgpu_ctx->func_sim->g_ptx_kernel_count]); + StatDisp( g_inst_op_classification_stat[gpgpu_ctx->func_sim->g_ptx_kernel_count]); } #ifdef GPGPUSIM_POWER_MODEL -- cgit v1.3 From 4d4d5938d715d2b79a617c32583184426b4a642d Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 9 Jul 2019 23:16:17 -0400 Subject: Move g_ptx_sim_mode Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 6 +++--- libcuda/gpgpu_context.h | 9 ++++++++- libopencl/opencl_runtime_api.cc | 6 +++--- src/cuda-sim/cuda-sim.cc | 6 ++---- src/cuda-sim/cuda-sim.h | 10 +++++++--- src/cuda-sim/cuda_device_runtime.cc | 7 ++++--- src/cuda-sim/cuda_device_runtime.h | 18 ++++++++++++++++-- src/gpgpu-sim/gpu-sim.cc | 4 ++-- src/gpgpusim_entrypoint.cc | 2 +- 9 files changed, 46 insertions(+), 22 deletions(-) (limited to 'libopencl') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index bbbaf23..59d2a60 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1076,7 +1076,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx = CUctx_st* context = GPGPUSim_Context(); char *mode = getenv("PTX_SIM_MODE_FUNC"); if( mode ) - sscanf(mode,"%u", &g_ptx_sim_mode); + sscanf(mode,"%u", &(ctx->func_sim->g_ptx_sim_mode)); gpgpusim_ptx_assert( !ctx->api->g_cuda_launch_stack.empty(), "empty launch stack" ); kernel_config config = ctx->api->g_cuda_launch_stack.back(); { @@ -1092,7 +1092,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx = } struct CUstream_st *stream = config.get_stream(); printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun, - g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 ); + (ctx->func_sim->g_ptx_sim_mode)?"functional simulation":"performance simulation", stream?stream->get_uid():0 ); kernel_info_t *grid = ctx->api->gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context); //do dynamic PDOM analysis for performance simulation scenario std::string kname = grid->name(); @@ -1143,7 +1143,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,g_ptx_sim_mode,stream); + stream_operation op(grid,ctx->func_sim->g_ptx_sim_mode,stream); g_stream_manager()->push(op); ctx->api->g_cuda_launch_stack.pop_back(); return g_last_cudaError = cudaSuccess; diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index a2ae7b6..3c9f87c 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -5,6 +5,7 @@ #include "../src/cuda-sim/ptx_parser.h" #include "../src/gpgpusim_entrypoint.h" #include "../src/cuda-sim/cuda-sim.h" +#include "../src/cuda-sim/cuda_device_runtime.h" class gpgpu_context { public: @@ -14,7 +15,10 @@ class gpgpu_context { ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); the_gpgpusim = new GPGPUsim_ctx(this); - func_sim = new cuda_sim(); + func_sim = new cuda_sim(this); +#if (CUDART_VERSION >= 5000) + device_runtime = new cuda_device_runtime(this); +#endif } // global list symbol_table *g_global_allfiles_symbol_table; @@ -25,6 +29,9 @@ class gpgpu_context { ptx_recognizer* ptx_parser; GPGPUsim_ctx* the_gpgpusim; cuda_sim* func_sim; +#if (CUDART_VERSION >= 5000) + cuda_device_runtime* device_runtime; +#endif // member function list void cuobjdumpParseBinary(unsigned int handle); class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 0a6eb3e..aaaec4f 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -884,9 +884,9 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, printf("\n\n\n"); char *mode = getenv("PTX_SIM_MODE_FUNC"); if ( mode ) - sscanf(mode,"%u", &g_ptx_sim_mode); + sscanf(mode,"%u", &(ctx->func_sim->g_ptx_sim_mode)); printf("GPGPU-Sim OpenCL API: clEnqueueNDRangeKernel '%s' (mode=%s)\n", kernel->name().c_str(), - g_ptx_sim_mode?"functional simulation":"performance simulation"); + (ctx->func_sim->g_ptx_sim_mode)?"functional simulation":"performance simulation"); if ( !work_dim || work_dim > 3 ) return CL_INVALID_WORK_DIMENSION; size_t _local_size[3]; if( local_work_size != NULL ) { @@ -957,7 +957,7 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); } kernel_info_t *grid = ctx->func_sim->gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu); - if ( g_ptx_sim_mode ) + 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 ); diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index b3e2965..7a7d205 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -2033,13 +2033,11 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co fflush(stdout); } -int g_ptx_sim_mode; // if non-zero run functional simulation only (i.e., no notion of a clock cycle) - extern int ptx_debug; bool g_cuda_launch_blocking = false; -void read_sim_environment_variables() +void cuda_sim::read_sim_environment_variables() { ptx_debug = 0; g_debug_execution = 0; @@ -2185,7 +2183,7 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL cta.execute(cp_count,temp); #if (CUDART_VERSION >= 5000) - launch_all_device_kernels(); + gpgpu_ctx->device_runtime->launch_all_device_kernels(); #endif } else diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 25ebf7b..3c4336d 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -36,12 +36,12 @@ #include #include"ptx_sim.h" +class gpgpu_context; class memory_space; class function_info; class symbol_table; extern const char *g_gpgpusim_version_string; -extern int g_ptx_sim_mode; extern int g_debug_execution; extern int g_debug_thread_uid; extern void ** g_inst_classification_stat; @@ -50,7 +50,6 @@ extern void ** g_inst_op_classification_stat; extern void print_splash(); extern void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu ); -extern void read_sim_environment_variables(); extern void ptxinfo_opencl_addinfo( std::map &kernels ); unsigned ptx_sim_init_thread( kernel_info_t &kernel, class ptx_thread_info** thread_info, @@ -124,9 +123,10 @@ struct gpgpu_ptx_sim_info get_ptxinfo(); class cuda_sim { public: - cuda_sim() { + cuda_sim( gpgpu_context* ctx ) { g_ptx_sim_num_insn = 0; g_ptx_kernel_count = -1; // used for classification stat collection purposes + gpgpu_ctx = ctx; } //global variables char *opcode_latency_int; @@ -147,6 +147,9 @@ class cuda_sim { int g_ptx_kernel_count; // used for classification stat collection purposes std::map g_global_name_lookup; // indexed by hostVar std::map g_const_name_lookup; // indexed by hostVar + int g_ptx_sim_mode; // if non-zero run functional simulation only (i.e., no notion of a clock cycle) + // backward pointer + class gpgpu_context* gpgpu_ctx; //global functions void ptx_opcocde_latency_options (option_parser_t opp); void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL = false ); @@ -159,6 +162,7 @@ class cuda_sim { gpgpu_t *gpu ); void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ); void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size ); + void read_sim_environment_variables(); }; #endif diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index be8369f..354fa79 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -20,6 +20,7 @@ unsigned long long g_max_total_param_size = 0; #include "../stream_manager.h" #include "../gpgpusim_entrypoint.h" #include "cuda_device_runtime.h" +#include "../../libcuda/gpgpu_context.h" #define DEV_RUNTIME_REPORT(a) \ if( g_debug_execution ) { \ @@ -318,17 +319,17 @@ void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_ } -void launch_one_device_kernel() { +void cuda_device_runtime::launch_one_device_kernel() { if(!g_cuda_device_launch_op.empty()) { device_launch_operation_t &op = g_cuda_device_launch_op.front(); - stream_operation stream_op = stream_operation(op.grid, g_ptx_sim_mode, op.stream); + stream_operation stream_op = stream_operation(op.grid, gpgpu_ctx->func_sim->g_ptx_sim_mode, op.stream); g_stream_manager()->push(stream_op); g_cuda_device_launch_op.pop_front(); } } -void launch_all_device_kernels() { +void cuda_device_runtime::launch_all_device_kernels() { while(!g_cuda_device_launch_op.empty()) { launch_one_device_kernel(); } diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h index 6dbcd71..851fed2 100644 --- a/src/cuda-sim/cuda_device_runtime.h +++ b/src/cuda-sim/cuda_device_runtime.h @@ -6,6 +6,20 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); -void launch_all_device_kernels(); -void launch_one_device_kernel(); +#endif +#if (CUDART_VERSION >= 5000) + +class gpgpu_context; + +class cuda_device_runtime { + public: + cuda_device_runtime( gpgpu_context* ctx ) { + gpgpu_ctx = ctx; + } + // backward pointer + class gpgpu_context* gpgpu_ctx; + void launch_all_device_kernels(); + void launch_one_device_kernel(); +}; + #endif diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 9f47067..bbcc078 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -503,7 +503,7 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) &gpgpu_ptx_instruction_classification, "if enabled will classify ptx instruction types per kernel (Max 255 kernels now)", "0"); - option_parser_register(opp, "-gpgpu_ptx_sim_mode", OPT_INT32, &g_ptx_sim_mode, + option_parser_register(opp, "-gpgpu_ptx_sim_mode", OPT_INT32, &(gpgpu_ctx->func_sim->g_ptx_sim_mode), "Select between Performance (default) or Functional simulation (1)", "0"); option_parser_register(opp, "-gpgpu_clock_domains", OPT_CSTR, &gpgpu_clock_domains, @@ -1753,7 +1753,7 @@ void gpgpu_sim::cycle() #if (CUDART_VERSION >= 5000) //launch device kernel - launch_one_device_kernel(); + gpgpu_ctx->device_runtime->launch_one_device_kernel(); #endif } } diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index d9d1023..683a695 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -213,7 +213,7 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf() { srand(1); print_splash(); - read_sim_environment_variables(); + func_sim->read_sim_environment_variables(); ptx_parser->read_parser_environment_variables(); option_parser_t opp = option_parser_create(); -- cgit v1.3 From c58d52ccde1eea6c45d0f8d2d5411bcf1ebab4f9 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 10 Jul 2019 00:45:39 -0400 Subject: Move g_globals and g_constants Signed-off-by: Mengchi Zhang --- libopencl/opencl_runtime_api.cc | 10 +++++----- src/cuda-sim/cuda-sim.cc | 5 +---- src/cuda-sim/cuda-sim.h | 4 +++- src/cuda-sim/ptx_parser.cc | 7 ++----- src/stream_manager.cc | 5 +++-- 5 files changed, 14 insertions(+), 17 deletions(-) (limited to 'libopencl') diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index aaaec4f..7f029c7 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -950,11 +950,11 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, gpgpu_t *gpu = command_queue->get_device()->the_device(); if (kernel->get_implementation()->get_ptx_version().ver() <3.0){ - gpgpu_ptx_sim_memcpy_symbol( "%_global_size", _global_size, 3 * sizeof(int), 0, 1, gpu ); - gpgpu_ptx_sim_memcpy_symbol( "%_work_dim", &work_dim, 1 * sizeof(int), 0, 1, gpu ); - gpgpu_ptx_sim_memcpy_symbol( "%_global_num_groups", &GridDim, 3 * sizeof(int), 0, 1, gpu ); - gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); - gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); + ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_size", _global_size, 3 * sizeof(int), 0, 1, gpu ); + ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_work_dim", &work_dim, 1 * sizeof(int), 0, 1, gpu ); + ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_num_groups", &GridDim, 3 * sizeof(int), 0, 1, gpu ); + ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); + ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); } kernel_info_t *grid = ctx->func_sim->gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu); if ( ctx->func_sim->g_ptx_sim_mode ) diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 7e8aab9..3d1da62 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1942,9 +1942,6 @@ void print_splash() } } -std::set g_globals; -std::set g_constants; - void cuda_sim::gpgpu_ptx_sim_register_const_variable(void *hostVar, const char *deviceName, size_t size ) { printf("GPGPU-Sim PTX registering constant %s (%zu bytes) to name mapping\n", deviceName, size ); @@ -1957,7 +1954,7 @@ void cuda_sim::gpgpu_ptx_sim_register_global_variable(void *hostVar, const char g_global_name_lookup[hostVar] = deviceName; } -void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu ) +void cuda_sim::gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu ) { printf("GPGPU-Sim PTX: starting gpgpu_ptx_sim_memcpy_symbol with hostVar 0x%p\n", hostVar); bool found_sym = false; diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 628e434..5c95100 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -46,7 +46,6 @@ extern int g_debug_execution; extern int g_debug_thread_uid; extern void print_splash(); -extern void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu ); extern void ptxinfo_opencl_addinfo( std::map &kernels ); unsigned ptx_sim_init_thread( kernel_info_t &kernel, @@ -161,6 +160,8 @@ class cuda_sim { bool g_cuda_launch_blocking; void ** g_inst_classification_stat; void ** g_inst_op_classification_stat; + std::set g_globals; + std::set g_constants; // backward pointer class gpgpu_context* gpgpu_ctx; //global functions @@ -179,6 +180,7 @@ class cuda_sim { void set_param_gpgpu_num_shaders(int num_shaders); struct rec_pts find_reconvergence_points( function_info *finfo ); address_type get_converge_point( address_type pc ); + void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu ); }; #endif diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 0139534..269ec4d 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -322,9 +322,6 @@ bool ptx_recognizer::check_for_duplicates( const char *identifier ) return ( s != NULL ); } -extern std::set g_globals; -extern std::set g_constants; - // Returns padding that needs to be inserted ahead of address to make it aligned to min(size, maxalign) /* * @param address the address in bytes @@ -453,7 +450,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi g_current_symbol_table->alloc_global( num_bits/8 + addr_pad ); } if( g_current_symbol_table == g_global_symbol_table ) { - g_constants.insert( identifier ); + gpgpu_ctx->func_sim->g_constants.insert( identifier ); } assert( g_current_symbol_table != NULL ); g_sym_name_to_symbol_table[ identifier ] = g_current_symbol_table; @@ -471,7 +468,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi fflush(stdout); g_last_symbol->set_address( addr+addr_pad ); g_current_symbol_table->alloc_global( num_bits/8 + addr_pad ); - g_globals.insert( identifier ); + gpgpu_ctx->func_sim->g_globals.insert( identifier ); assert( g_current_symbol_table != NULL ); g_sym_name_to_symbol_table[ identifier ] = g_current_symbol_table; break; diff --git a/src/stream_manager.cc b/src/stream_manager.cc index e07f4e4..be3dd71 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -29,6 +29,7 @@ #include "gpgpusim_entrypoint.h" #include "cuda-sim/cuda-sim.h" #include "gpgpu-sim/gpu-sim.h" +#include "../libcuda/gpgpu_context.h" unsigned CUstream_st::sm_next_stream_uid = 0; @@ -150,13 +151,13 @@ bool stream_operation::do_operation( gpgpu_sim *gpu ) case stream_memcpy_to_symbol: if(g_debug_execution >= 3) printf("memcpy to symbol\n"); - gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_src,m_cnt,m_offset,1,gpu); + gpu->gpgpu_ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_src,m_cnt,m_offset,1,gpu); m_stream->record_next_done(); break; case stream_memcpy_from_symbol: if(g_debug_execution >= 3) printf("memcpy from symbol\n"); - gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_dst,m_cnt,m_offset,0,gpu); + gpu->gpgpu_ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_dst,m_cnt,m_offset,0,gpu); m_stream->record_next_done(); break; case stream_kernel_launch: -- cgit v1.3