From 8ad9dbc77947b4abd51ebb55ef4bbe80be01caaa Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 4 Jun 2019 14:26:42 -0400 Subject: Add gpgpu_context as the top level class Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) create mode 100644 libcuda/gpgpu_context.h (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h new file mode 100644 index 0000000..d6e564b --- /dev/null +++ b/libcuda/gpgpu_context.h @@ -0,0 +1,18 @@ +#include + +class cuobjdumpSection; +class cuobjdumpELFSection; +class cuobjdumpPTXSection; + +class gpgpu_context { + public: + std::list cuobjdumpSectionList; + void cuobjdumpInit(); + void cuobjdumpParseBinary(unsigned int handle); + void extract_code_using_cuobjdump(); + std::list pruneSectionList(CUctx_st *context); + std::list mergeMatchingSections(std::string identifier); + std::list mergeSections(); + cuobjdumpELFSection* findELFSection(const std::string identifier, std::list &libSectionList); + cuobjdumpPTXSection* findPTXSection(const std::string identifier, std::list &libSectionList); +}; -- cgit v1.3 From d553124832aca461dda4dd7d503748d22b44bbe2 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 5 Jun 2019 23:28:12 -0400 Subject: Add cuda_api_object and move some var Signed-off-by: Mengchi Zhang --- libcuda/cuda_api_object.h | 11 +++++++++++ libcuda/cuda_runtime_api.cc | 30 ++++++++++++++---------------- libcuda/gpgpu_context.h | 19 +++++++++++++++++-- 3 files changed, 42 insertions(+), 18 deletions(-) create mode 100644 libcuda/cuda_api_object.h (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/cuda_api_object.h b/libcuda/cuda_api_object.h new file mode 100644 index 0000000..86ffa98 --- /dev/null +++ b/libcuda/cuda_api_object.h @@ -0,0 +1,11 @@ +#ifndef __cuda_api_object_h__ +#define __cuda_api_object_h__ +class cuobjdumpSection; + +class cuda_runtime_api { + public: + // global list + std::list libSectionList; + // member function list +}; +#endif /* __cuda_api_object_h__ */ diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index c39571c..4221634 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -132,6 +132,7 @@ #include "__cudaFatFormat.h" #endif #include "gpgpu_context.h" +#include "cuda_api_object.h" #include "../src/gpgpu-sim/gpu-sim.h" #include "../src/cuda-sim/ptx_loader.h" #include "../src/cuda-sim/cuda-sim.h" @@ -314,9 +315,6 @@ struct CUctx_st { return i->second; } - std::list libSectionList; - //maps sm version number to set of filenames - std::map > version_filename; std::list g_cuda_launch_stack; std::mapfatbin_registered; std::map fatbinmap; @@ -2005,7 +2003,7 @@ char* get_app_binary_name(std::string abs_path){ } //extracts all ptx files from binary and dumps into prog_name.unique_no.sm_<>.ptx files -void extract_ptx_files_using_cuobjdump(CUctx_st *context){ +void gpgpu_context::extract_ptx_files_using_cuobjdump(CUctx_st *context){ extern bool g_cdp_enabled; char command[1000]; char *pytorch_bin = getenv("PYTORCH_BIN"); @@ -2066,10 +2064,10 @@ void extract_ptx_files_using_cuobjdump(CUctx_st *context){ } std::string vstr = line.substr(pos1+3,pos2-pos1-3); int version = atoi(vstr.c_str()); - if (context->version_filename.find(version)==context->version_filename.end()){ - context->version_filename[version] = std::set(); + if (version_filename.find(version)==version_filename.end()){ + version_filename[version] = std::set(); } - context->version_filename[version].insert(line); + version_filename[version].insert(line); } } @@ -2229,7 +2227,7 @@ void gpgpu_context::extract_code_using_cuobjdump(){ fclose(cuobjdump_in); std::getline(libsf, line); } - context->libSectionList = cuobjdumpSectionList; + api->libSectionList = cuobjdumpSectionList; //Restore the original section list cuobjdumpSectionList = tmpsl; @@ -2418,10 +2416,10 @@ cuobjdumpELFSection* findELFSectionInList(std::list sectionli } //! Find an ELF section in all the known lists -cuobjdumpELFSection* gpgpu_context::findELFSection(const std::string identifier, std::list &libSectionList){ +cuobjdumpELFSection* gpgpu_context::findELFSection(const std::string identifier){ cuobjdumpELFSection* sec = findELFSectionInList(cuobjdumpSectionList, identifier); if (sec!=NULL)return sec; - sec = findELFSectionInList(libSectionList, identifier); + sec = findELFSectionInList(api->libSectionList, identifier); if (sec!=NULL)return sec; std::cout << "Could not find " << identifier << std::endl; assert(0 && "Could not find the required ELF section"); @@ -2453,10 +2451,10 @@ cuobjdumpPTXSection* findPTXSectionInList(std::list §ionl } //! Find an PTX section in all the known lists -cuobjdumpPTXSection* gpgpu_context::findPTXSection(const std::string identifier, std::list &libSectionList){ +cuobjdumpPTXSection* gpgpu_context::findPTXSection(const std::string identifier){ cuobjdumpPTXSection* sec = findPTXSectionInList(cuobjdumpSectionList, identifier); if (sec!=NULL)return sec; - sec = findPTXSectionInList(libSectionList, identifier); + sec = findPTXSectionInList(api->libSectionList, identifier); if (sec!=NULL)return sec; std::cout << "Could not find " << identifier << std::endl; assert(0 && "Could not find the required PTX section"); @@ -2500,7 +2498,7 @@ void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){ #if (CUDART_VERSION >= 6000) //loops through all ptx files from smallest sm version to largest std::map >::iterator itr_m; - for (itr_m = context->version_filename.begin(); itr_m!=context->version_filename.end(); itr_m++){ + for (itr_m = version_filename.begin(); itr_m!=version_filename.end(); itr_m++){ std::set::iterator itr_s; for (itr_s = itr_m->second.begin(); itr_s!=itr_m->second.end(); itr_s++){ std::string ptx_filename = *itr_s; @@ -2512,7 +2510,7 @@ void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){ context->add_binary(symtab, handle); load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); - for (itr_m = context->version_filename.begin(); itr_m!=context->version_filename.end(); itr_m++){ + for (itr_m = version_filename.begin(); itr_m!=version_filename.end(); itr_m++){ std::set::iterator itr_s; for (itr_s = itr_m->second.begin(); itr_s!=itr_m->second.end(); itr_s++){ std::string ptx_filename = *itr_s; @@ -2536,7 +2534,7 @@ void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){ cuobjdumpPTXSection* ptx = NULL; const char* pre_load = getenv("CUOBJDUMP_SIM_FILE"); if(pre_load==NULL || strlen(pre_load)==0) - ptx = findPTXSection(fname, context->libSectionList); + ptx = findPTXSection(fname); char *ptxcode; const char *override_ptx_name = getenv("PTX_SIM_KERNELFILE"); if (override_ptx_name == NULL or getenv("PTX_SIM_USE_PTX_FILE") == NULL or strlen(getenv("PTX_SIM_USE_PTX_FILE"))==0) { @@ -2546,7 +2544,7 @@ void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){ ptxcode = readfile(override_ptx_name); } if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) { - cuobjdumpELFSection* elfsection = findELFSection(ptx->getIdentifier(), context->libSectionList); + cuobjdumpELFSection* elfsection = findELFSection(ptx->getIdentifier()); assert (elfsection!= NULL); char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus( ptx->getPTXfilename(), diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index d6e564b..576ec67 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -1,4 +1,9 @@ +#ifndef __gpgpu_context_h__ +#define __gpgpu_context_h__ #include +#include +#include +#include "cuda_api_object.h" class cuobjdumpSection; class cuobjdumpELFSection; @@ -6,13 +11,23 @@ class cuobjdumpPTXSection; class gpgpu_context { public: + gpgpu_context() { + api = new cuda_runtime_api(); + } + // global list std::list cuobjdumpSectionList; + //maps sm version number to set of filenames + std::map > version_filename; + cuda_runtime_api* api; + // member function list void cuobjdumpInit(); void cuobjdumpParseBinary(unsigned int handle); void extract_code_using_cuobjdump(); std::list pruneSectionList(CUctx_st *context); std::list mergeMatchingSections(std::string identifier); std::list mergeSections(); - cuobjdumpELFSection* findELFSection(const std::string identifier, std::list &libSectionList); - cuobjdumpPTXSection* findPTXSection(const std::string identifier, std::list &libSectionList); + cuobjdumpELFSection* findELFSection(const std::string identifier); + cuobjdumpPTXSection* findPTXSection(const std::string identifier); + void extract_ptx_files_using_cuobjdump(CUctx_st *context); }; +#endif /* __gpgpu_context_h__ */ -- cgit v1.3 From 570ce0a3d6e049ec3ee42329fbf25019fd4eb2e5 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Thu, 6 Jun 2019 01:57:40 -0400 Subject: Move 3 more var Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 92 +++++++++++++++++++++++++-------------------- libcuda/gpgpu_context.h | 6 +++ 2 files changed, 58 insertions(+), 40 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index ea96570..855f11d 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -313,10 +313,7 @@ struct CUctx_st { return i->second; } - std::mapfatbin_registered; - std::map fatbinmap; std::map g_mallocPtr_Size; - std::map name_symtab; std::map pinned_memory; //support for pinned memories added std::map pinned_memory_size; int no_of_ptx; @@ -647,8 +644,8 @@ static int get_app_cuda_version() { } //! Keep track of the association between filename and cubin handle -void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context){ - context->fatbinmap[handle] = filename; +void gpgpu_context::cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context){ + fatbinmap[handle] = filename; } /******************************************************************************* @@ -725,7 +722,7 @@ void** cudaRegisterFatBinaryInternal( void *fatCubin, gpgpu_context* gpgpu_ctx = */ assert(fat_cubin_handle >= 1); if (fat_cubin_handle==1) ctx->cuobjdumpInit(); - cuobjdumpRegisterFatBinary(fat_cubin_handle, filename, context); + ctx->cuobjdumpRegisterFatBinary(fat_cubin_handle, filename, context); return (void**)fat_cubin_handle; } @@ -969,6 +966,46 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx = return g_last_cudaError = cudaSuccess; } +#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 = gpgpu_ptx_sim_load_ptx_from_filename( file ); + std::string fname(path); + ctx->name_symtab[fname] = symtab; + context->add_binary(symtab, 1); + load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); + load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); + addedFile = true; + return CUDA_SUCCESS; +} +#endif /******************************************************************************* * * @@ -2715,12 +2752,12 @@ void gpgpu_context::cuobjdumpInit(){ void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){ CUctx_st *context = GPGPUSim_Context(); - if(context->fatbin_registered[handle]) return; - context->fatbin_registered[handle] = true; - std::string fname = context->fatbinmap[handle]; + if(fatbin_registered[handle]) return; + fatbin_registered[handle] = true; + std::string fname = fatbinmap[handle]; - if (context->name_symtab.find(fname) != context->name_symtab.end()) { - symbol_table *symtab = context->name_symtab[fname]; + if (name_symtab.find(fname) != name_symtab.end()) { + symbol_table *symtab = name_symtab[fname]; context->add_binary(symtab, handle); return; } @@ -2737,7 +2774,7 @@ void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){ symtab = gpgpu_ptx_sim_load_ptx_from_filename( ptx_filename.c_str() ); } } - context->name_symtab[fname] = symtab; + name_symtab[fname] = symtab; context->add_binary(symtab, handle); load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); @@ -2795,7 +2832,7 @@ void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){ } load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); - context->name_symtab[fname] = symtab; + name_symtab[fname] = symtab; //TODO: Remove temporarily files as per configurations } @@ -3909,33 +3946,8 @@ CUresult CUDAAPI cuLinkAddFile(CUlinkState state, CUjitInputType type, const char *path, unsigned int numOptions, CUjit_option *options, void **optionValues) { - 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 = gpgpu_ptx_sim_load_ptx_from_filename( file ); - std::string fname(path); - context->name_symtab[fname] = symtab; - context->add_binary(symtab, 1); - load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); - load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); - addedFile = true; - return CUDA_SUCCESS; + return cuLinkAddFileInternal(state, type, path, + numOptions, options, optionValues) } #endif diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 576ec67..0543ff8 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -3,11 +3,13 @@ #include #include #include +#include #include "cuda_api_object.h" class cuobjdumpSection; class cuobjdumpELFSection; class cuobjdumpPTXSection; +class symbol_table; class gpgpu_context { public: @@ -16,6 +18,9 @@ class gpgpu_context { } // global list std::list cuobjdumpSectionList; + std::mapfatbin_registered; + std::map fatbinmap; + std::map name_symtab; //maps sm version number to set of filenames std::map > version_filename; cuda_runtime_api* api; @@ -29,5 +34,6 @@ class gpgpu_context { cuobjdumpELFSection* findELFSection(const std::string identifier); cuobjdumpPTXSection* findPTXSection(const std::string identifier); void extract_ptx_files_using_cuobjdump(CUctx_st *context); + void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context); }; #endif /* __gpgpu_context_h__ */ -- cgit v1.3 From bc8d1e10507f043c916ee051dc9a687adc6d9b4b Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Thu, 6 Jun 2019 02:47:07 -0400 Subject: Move more vars Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 117 ++++++++++++++++++++++++++------------------ libcuda/gpgpu_context.h | 11 +++++ 2 files changed, 81 insertions(+), 47 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 66562aa..02e2b2e 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -158,11 +158,11 @@ extern void exit_simulation(); static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr, gpgpu_t *gpu ); static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu ); -static kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, - gpgpu_ptx_sim_arg_list_t args, - struct dim3 gridDim, - struct dim3 blockDim, - struct CUctx_st* context ); +//static kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, +// gpgpu_ptx_sim_arg_list_t args, +// struct dim3 gridDim, +// struct dim3 blockDim, +// struct CUctx_st* context ); /*DEVICE_BUILTIN*/ struct cudaArray @@ -313,7 +313,6 @@ struct CUctx_st { return i->second; } - std::map g_mallocPtr_Size; std::map pinned_memory; //support for pinned memories added std::map pinned_memory_size; int no_of_ptx; @@ -910,7 +909,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 ); - kernel_info_t *grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context); + kernel_info_t *grid = ctx->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(); function_info *kernel_func_info = grid->entry(); @@ -966,6 +965,66 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx = return g_last_cudaError = cudaSuccess; } +cudaError_t cudaMallocInternal(void **devPtr, size_t size, 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(); + *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); + ctx->g_mallocPtr_Size[(unsigned long long)*devPtr] = size; + } + 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; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + //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(); + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + std::map::const_iterator i = context->pinned_memory_size.find(pHost); + assert(i != context->pinned_memory_size.end()); + size_t size = i->second; + *pDevice = gpu->gpu_malloc(size); + if(g_debug_execution >= 3){ + printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *pDevice); + ctx->g_mallocPtr_Size[(unsigned long long)*pDevice] = size; + } + if ( *pDevice ) { + context->pinned_memory[pHost]=pDevice; + //Copy contents in cpu to gpu + gpu->memcpy_to_gpu((size_t)*pDevice,pHost,size); + return g_last_cudaError = cudaSuccess; + } else { + return g_last_cudaError = cudaErrorMemoryAllocation; + } +} #if CUDART_VERSION >= 6050 CUresult cuLinkAddFileInternal(CUlinkState state, CUjitInputType type, const char *path, @@ -1027,20 +1086,7 @@ cudaError_t cudaPeekAtLastError(void) __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - CUctx_st* context = GPGPUSim_Context(); - *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); - context->g_mallocPtr_Size[(unsigned long long)*devPtr] = size; - } - if ( *devPtr ) { - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorMemoryAllocation; - } + return cudaMallocInternal(devPtr, size); } __host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size) @@ -3111,30 +3157,7 @@ cudaError_t CUDARTAPI cudaHostAlloc(void **pHost, size_t bytes, unsigned int fl cudaError_t CUDARTAPI cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - //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(); - gpgpu_t *gpu = context->get_device()->get_gpgpu(); - std::map::const_iterator i = context->pinned_memory_size.find(pHost); - assert(i != context->pinned_memory_size.end()); - size_t size = i->second; - *pDevice = gpu->gpu_malloc(size); - if(g_debug_execution >= 3){ - printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *pDevice); - context->g_mallocPtr_Size[(unsigned long long)*pDevice] = size; - } - if ( *pDevice ) { - context->pinned_memory[pHost]=pDevice; - //Copy contents in cpu to gpu - gpu->memcpy_to_gpu((size_t)*pDevice,pHost,size); - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorMemoryAllocation; - } + return cudaHostGetDevicePointerInternal(pDevice, pHost, flags); } __host__ cudaError_t CUDARTAPI cudaPointerGetAttributes( @@ -3450,7 +3473,7 @@ static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu return nc_bytes; } -kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, +kernel_info_t * gpgpu_context::gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, gpgpu_ptx_sim_arg_list_t args, struct dim3 gridDim, struct dim3 blockDim, @@ -3482,7 +3505,7 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, fflush(stdout); if(g_debug_execution >= 4){ - entry->ptx_jit_config(context->g_mallocPtr_Size, result->get_param_memory(), (gpgpu_t *) context->get_device()->get_gpgpu(), gridDim, blockDim); + entry->ptx_jit_config(g_mallocPtr_Size, result->get_param_memory(), (gpgpu_t *) context->get_device()->get_gpgpu(), gridDim, blockDim); } return result; diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 0543ff8..f29e2e0 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -10,6 +10,10 @@ class cuobjdumpSection; class cuobjdumpELFSection; class cuobjdumpPTXSection; class symbol_table; +class gpgpu_ptx_sim_arg; +class kernel_info_t; + +typedef std::list gpgpu_ptx_sim_arg_list_t; class gpgpu_context { public: @@ -21,6 +25,7 @@ class gpgpu_context { std::mapfatbin_registered; std::map fatbinmap; std::map name_symtab; + std::map g_mallocPtr_Size; //maps sm version number to set of filenames std::map > version_filename; cuda_runtime_api* api; @@ -35,5 +40,11 @@ class gpgpu_context { cuobjdumpPTXSection* findPTXSection(const std::string identifier); void extract_ptx_files_using_cuobjdump(CUctx_st *context); void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context); + kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, + gpgpu_ptx_sim_arg_list_t args, + struct dim3 gridDim, + struct dim3 blockDim, + struct CUctx_st* context ); + }; #endif /* __gpgpu_context_h__ */ -- cgit v1.3 From 7e286ec0ffc963d307551caced7f4a52241dced4 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Thu, 6 Jun 2019 12:53:34 -0400 Subject: Move pinned_memory etc Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 80 +++++++++++++++++++++++++++++---------------- libcuda/gpgpu_context.h | 3 ++ 2 files changed, 54 insertions(+), 29 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 02e2b2e..41f27bf 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -313,8 +313,6 @@ struct CUctx_st { return i->second; } - std::map pinned_memory; //support for pinned memories added - std::map pinned_memory_size; int no_of_ptx; typedef struct glbmap_entry glbmap_entry_t; @@ -989,6 +987,27 @@ cudaError_t cudaMallocInternal(void **devPtr, size_t size, gpgpu_context* gpgpu_ } } +cudaError_t cudaMallocHostInternal(void **ptr, size_t size, 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__); + } + *ptr = malloc(size); + if ( *ptr ) { + //track pinned memory size allocated in the host so that same amount of memory is also allocated in GPU. + ctx->pinned_memory_size[*ptr]=size; + 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; @@ -1008,8 +1027,8 @@ cudaError_t cudaHostGetDevicePointerInternal(void **pDevice, void *pHost, unsign flags=0; CUctx_st* context = GPGPUSim_Context(); gpgpu_t *gpu = context->get_device()->get_gpgpu(); - std::map::const_iterator i = context->pinned_memory_size.find(pHost); - assert(i != context->pinned_memory_size.end()); + std::map::const_iterator i = ctx->pinned_memory_size.find(pHost); + assert(i != ctx->pinned_memory_size.end()); size_t size = i->second; *pDevice = gpu->gpu_malloc(size); if(g_debug_execution >= 3){ @@ -1017,7 +1036,7 @@ cudaError_t cudaHostGetDevicePointerInternal(void **pDevice, void *pHost, unsign ctx->g_mallocPtr_Size[(unsigned long long)*pDevice] = size; } if ( *pDevice ) { - context->pinned_memory[pHost]=pDevice; + ctx->pinned_memory[pHost]=pDevice; //Copy contents in cpu to gpu gpu->memcpy_to_gpu((size_t)*pDevice,pHost,size); return g_last_cudaError = cudaSuccess; @@ -1066,6 +1085,31 @@ cuLinkAddFileInternal(CUlinkState state, CUjitInputType type, const char *path, } #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->pinned_memory_size[*pHost]=bytes; + if( *pHost ) + return g_last_cudaError = cudaSuccess; + else + return g_last_cudaError = cudaErrorMemoryAllocation; +} + +#endif + /******************************************************************************* * * * * @@ -1091,18 +1135,7 @@ __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) __host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - CUctx_st* context = GPGPUSim_Context(); - *ptr = malloc(size); - if ( *ptr ) { - //track pinned memory size allocated in the host so that same amount of memory is also allocated in GPU. - context->pinned_memory_size[*ptr]=size; - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorMemoryAllocation; - } + return cudaMallocHostInternal(ptr, size); } __host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr, size_t *pitch, size_t width, size_t height) { @@ -3141,18 +3174,7 @@ cudaError_t cudaGLUnregisterBufferObject(GLuint bufferObj) cudaError_t CUDARTAPI cudaHostAlloc(void **pHost, size_t bytes, unsigned int flags) { - 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) - CUctx_st* context = GPGPUSim_Context(); - context->pinned_memory_size[*pHost]=bytes; - if( *pHost ) - return g_last_cudaError = cudaSuccess; - else - return g_last_cudaError = cudaErrorMemoryAllocation; + return cudaHostAllocInternal(pHost, bytes, flags); } cudaError_t CUDARTAPI cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags) diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index f29e2e0..6878d5c 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -28,6 +28,9 @@ class gpgpu_context { std::map g_mallocPtr_Size; //maps sm version number to set of filenames std::map > version_filename; + std::map pinned_memory; //support for pinned memories added + std::map pinned_memory_size; + // objects pointers for each file cuda_runtime_api* api; // member function list void cuobjdumpInit(); -- cgit v1.3 From ba8374c72558e4b89a0bddc973bcc87a10e2ab5f Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Thu, 6 Jun 2019 13:20:29 -0400 Subject: Move g_glbmap Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 134 ++++++++++++++++++++++---------------------- libcuda/gpgpu_context.h | 15 +++++ 2 files changed, 82 insertions(+), 67 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 41f27bf..18a9abb 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -237,17 +237,6 @@ private: struct _cuda_device_id *m_next; }; -#ifndef OPENGL_SUPPORT -typedef unsigned long GLuint; -#endif - -struct glbmap_entry { - GLuint m_bufferObj; - void *m_devPtr; - size_t m_size; - struct glbmap_entry *m_next; -}; - struct CUctx_st { CUctx_st( _cuda_device_id *gpu ) { @@ -255,7 +244,6 @@ struct CUctx_st { m_binary_info.cmem = 0; m_binary_info.gmem = 0; no_of_ptx=0; - g_glbmap = NULL; } _cuda_device_id *get_device() { return m_gpu; } @@ -314,9 +302,6 @@ struct CUctx_st { } int no_of_ptx; - typedef struct glbmap_entry glbmap_entry_t; - - glbmap_entry_t* g_glbmap; private: _cuda_device_id *m_gpu; // selected gpu @@ -1044,6 +1029,72 @@ cudaError_t cudaHostGetDevicePointerInternal(void **pDevice, void *pHost, unsign return g_last_cudaError = cudaErrorMemoryAllocation; } } + +cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, 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(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->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->g_glbmap; + ctx->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, @@ -3079,58 +3130,7 @@ cudaError_t cudaGLRegisterBufferObject(GLuint bufferObj) cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } -#ifdef OPENGL_SUPPORT - GLint buffer_size=0; - CUctx_st* ctx = GPGPUSim_Context(); - - glbmap_entry_t *p = ctx->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 = ctx->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->g_glbmap; - ctx->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 + return cudaGLMapBufferObjectInternal(devPtr, bufferObj); } cudaError_t cudaGLUnmapBufferObject(GLuint bufferObj) diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 6878d5c..7569ea6 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -15,10 +15,24 @@ class kernel_info_t; typedef std::list gpgpu_ptx_sim_arg_list_t; +#ifndef OPENGL_SUPPORT +typedef unsigned long GLuint; +#endif + +struct glbmap_entry { + GLuint m_bufferObj; + void *m_devPtr; + size_t m_size; + struct glbmap_entry *m_next; +}; + +typedef struct glbmap_entry glbmap_entry_t; + class gpgpu_context { public: gpgpu_context() { api = new cuda_runtime_api(); + g_glbmap = NULL; } // global list std::list cuobjdumpSectionList; @@ -30,6 +44,7 @@ class gpgpu_context { std::map > version_filename; std::map pinned_memory; //support for pinned memories added std::map pinned_memory_size; + glbmap_entry_t* g_glbmap; // objects pointers for each file cuda_runtime_api* api; // member function list -- cgit v1.3 From 87409183125b863be0c3e0470b2417c3c92a748b Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Sun, 9 Jun 2019 02:47:39 -0400 Subject: Move some vars back to cuda_api_object Signed-off-by: Mengchi Zhang --- libcuda/cuda_api_object.h | 54 +++++++++++++++++++++++++++++++++++++++ libcuda/cuda_runtime_api.cc | 62 ++++++++++++++++++++++----------------------- libcuda/gpgpu_context.h | 53 -------------------------------------- 3 files changed, 85 insertions(+), 84 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/cuda_api_object.h b/libcuda/cuda_api_object.h index 2001f91..41337c6 100644 --- a/libcuda/cuda_api_object.h +++ b/libcuda/cuda_api_object.h @@ -1,14 +1,68 @@ #ifndef __cuda_api_object_h__ #define __cuda_api_object_h__ + +#include +#include +#include +#include + class cuobjdumpSection; +class cuobjdumpELFSection; +class cuobjdumpPTXSection; +class symbol_table; +class gpgpu_ptx_sim_arg; class kernel_config; +class kernel_info_t; + +typedef std::list gpgpu_ptx_sim_arg_list_t; + +#ifndef OPENGL_SUPPORT +typedef unsigned long GLuint; +#endif + +struct glbmap_entry { + GLuint m_bufferObj; + void *m_devPtr; + size_t m_size; + struct glbmap_entry *m_next; +}; + +typedef struct glbmap_entry glbmap_entry_t; class cuda_runtime_api { public: + cuda_runtime_api() { + g_glbmap = NULL; + } // global list + std::list cuobjdumpSectionList; std::list libSectionList; std::list g_cuda_launch_stack; + std::mapfatbin_registered; + std::map fatbinmap; + std::map name_symtab; + std::map g_mallocPtr_Size; + //maps sm version number to set of filenames + std::map > version_filename; + std::map pinned_memory; //support for pinned memories added + std::map pinned_memory_size; + glbmap_entry_t* g_glbmap; // member function list + void cuobjdumpInit(); + void extract_code_using_cuobjdump(); + void extract_ptx_files_using_cuobjdump(CUctx_st *context); + void cuobjdumpParseBinary(unsigned int handle); + std::list pruneSectionList(CUctx_st *context); + std::list mergeMatchingSections(std::string identifier); + std::list mergeSections(); + cuobjdumpELFSection* findELFSection(const std::string identifier); + cuobjdumpPTXSection* findPTXSection(const std::string identifier); + void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context); + kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, + gpgpu_ptx_sim_arg_list_t args, + struct dim3 gridDim, + struct dim3 blockDim, + struct CUctx_st* context ); }; #endif /* __cuda_api_object_h__ */ diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 18a9abb..72ab002 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -626,7 +626,7 @@ static int get_app_cuda_version() { } //! Keep track of the association between filename and cubin handle -void gpgpu_context::cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context){ +void cuda_runtime_api::cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context){ fatbinmap[handle] = filename; } @@ -703,8 +703,8 @@ void** cudaRegisterFatBinaryInternal( void *fatCubin, gpgpu_context* gpgpu_ctx = * then for next calls, only returns the appropriate number */ assert(fat_cubin_handle >= 1); - if (fat_cubin_handle==1) ctx->cuobjdumpInit(); - ctx->cuobjdumpRegisterFatBinary(fat_cubin_handle, filename, context); + if (fat_cubin_handle==1) ctx->api->cuobjdumpInit(); + ctx->api->cuobjdumpRegisterFatBinary(fat_cubin_handle, filename, context); return (void**)fat_cubin_handle; } @@ -801,7 +801,7 @@ void cudaRegisterFunctionInternal( printf("GPGPU-Sim PTX: __cudaRegisterFunction %s : hostFun 0x%p, fat_cubin_handle = %u\n", deviceFun, hostFun, fat_cubin_handle); if(context->get_device()->get_gpgpu()->get_config().use_cuobjdump()) - ctx->cuobjdumpParseBinary(fat_cubin_handle); + ctx->api->cuobjdumpParseBinary(fat_cubin_handle); context->register_function( fat_cubin_handle, hostFun, deviceFun ); } @@ -828,7 +828,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()) - ctx->cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle); + ctx->api->cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle); fflush(stdout); if ( constant && !global && !ext ) { gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size); @@ -892,7 +892,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 ); - kernel_info_t *grid = ctx->gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context); + 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(); function_info *kernel_func_info = grid->entry(); @@ -963,7 +963,7 @@ cudaError_t cudaMallocInternal(void **devPtr, size_t size, gpgpu_context* gpgpu_ *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); - ctx->g_mallocPtr_Size[(unsigned long long)*devPtr] = size; + ctx->api->g_mallocPtr_Size[(unsigned long long)*devPtr] = size; } if ( *devPtr ) { return g_last_cudaError = cudaSuccess; @@ -986,7 +986,7 @@ cudaError_t cudaMallocHostInternal(void **ptr, size_t size, gpgpu_context* gpgpu *ptr = malloc(size); if ( *ptr ) { //track pinned memory size allocated in the host so that same amount of memory is also allocated in GPU. - ctx->pinned_memory_size[*ptr]=size; + ctx->api->pinned_memory_size[*ptr]=size; return g_last_cudaError = cudaSuccess; } else { return g_last_cudaError = cudaErrorMemoryAllocation; @@ -1012,16 +1012,16 @@ cudaError_t cudaHostGetDevicePointerInternal(void **pDevice, void *pHost, unsign flags=0; CUctx_st* context = GPGPUSim_Context(); gpgpu_t *gpu = context->get_device()->get_gpgpu(); - std::map::const_iterator i = ctx->pinned_memory_size.find(pHost); - assert(i != ctx->pinned_memory_size.end()); + std::map::const_iterator i = ctx->api->pinned_memory_size.find(pHost); + assert(i != ctx->api->pinned_memory_size.end()); size_t size = i->second; *pDevice = gpu->gpu_malloc(size); if(g_debug_execution >= 3){ printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *pDevice); - ctx->g_mallocPtr_Size[(unsigned long long)*pDevice] = size; + ctx->api->g_mallocPtr_Size[(unsigned long long)*pDevice] = size; } if ( *pDevice ) { - ctx->pinned_memory[pHost]=pDevice; + ctx->api->pinned_memory[pHost]=pDevice; //Copy contents in cpu to gpu gpu->memcpy_to_gpu((size_t)*pDevice,pHost,size); return g_last_cudaError = cudaSuccess; @@ -1048,7 +1048,7 @@ cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, gpgpu GLint buffer_size=0; CUctx_st* context = GPGPUSim_Context(); - glbmap_entry_t *p = ctx->g_glbmap; + glbmap_entry_t *p = ctx->api->g_glbmap; while ( p && p->m_bufferObj != bufferObj ) p = p->m_next; if ( p == NULL ) { @@ -1059,8 +1059,8 @@ cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, gpgpu // 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->g_glbmap; - ctx->g_glbmap = n; + n->m_next = ctx->api->g_glbmap; + ctx->api->g_glbmap = n; // initialize entry n->m_bufferObj = bufferObj; @@ -1127,7 +1127,7 @@ cuLinkAddFileInternal(CUlinkState state, CUjitInputType type, const char *path, strcat(file,path); symbol_table *symtab = gpgpu_ptx_sim_load_ptx_from_filename( file ); std::string fname(path); - ctx->name_symtab[fname] = symtab; + ctx->api->name_symtab[fname] = symtab; context->add_binary(symtab, 1); load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); @@ -1152,7 +1152,7 @@ cudaError_t cudaHostAllocInternal(void **pHost, size_t bytes, unsigned int flag *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->pinned_memory_size[*pHost]=bytes; + ctx->api->pinned_memory_size[*pHost]=bytes; if( *pHost ) return g_last_cudaError = cudaSuccess; else @@ -2428,7 +2428,7 @@ __host__ cudaError_t CUDARTAPI cudaGetExportTable(const void **ppExportTable, co //#include "../../cuobjdump_to_ptxplus/cuobjdump_parser.h" //extracts all ptx files from binary and dumps into prog_name.unique_no.sm_<>.ptx files -void gpgpu_context::extract_ptx_files_using_cuobjdump(CUctx_st *context){ +void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context){ extern bool g_cdp_enabled; char command[1000]; char *pytorch_bin = getenv("PYTORCH_BIN"); @@ -2506,7 +2506,7 @@ void gpgpu_context::extract_ptx_files_using_cuobjdump(CUctx_st *context){ * It is also responsible for extracting the libraries linked to the binary if the option is * enabled * */ -void gpgpu_context::extract_code_using_cuobjdump(){ +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(); @@ -2630,7 +2630,7 @@ void gpgpu_context::extract_code_using_cuobjdump(){ fclose(cuobjdump_in); std::getline(libsf, line); } - api->libSectionList = cuobjdumpSectionList; + libSectionList = cuobjdumpSectionList; //Restore the original section list cuobjdumpSectionList = tmpsl; @@ -2676,7 +2676,7 @@ void printSectionList(std::list sl) { } //! Remove unecessary sm versions from the section list -std::list gpgpu_context::pruneSectionList(CUctx_st *context) { +std::list cuda_runtime_api::pruneSectionList(CUctx_st *context) { unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability(); //For ptxplus, force the max capability to 19 if it's higher or unspecified(0) @@ -2729,7 +2729,7 @@ std::list gpgpu_context::pruneSectionList(CUctx_st *context) } //! Merge all PTX sections that have a specific identifier into one file -std::list gpgpu_context::mergeMatchingSections(std::string identifier){ +std::list cuda_runtime_api::mergeMatchingSections(std::string identifier){ const char *ptxcode = ""; std::list::iterator old_iter; cuobjdumpPTXSection* old_ptxsection = NULL; @@ -2772,7 +2772,7 @@ std::list gpgpu_context::mergeMatchingSections(std::string id } //! Merge any PTX sections with matching identifiers -std::list gpgpu_context::mergeSections(){ +std::list cuda_runtime_api::mergeSections(){ std::vector identifier; cuobjdumpPTXSection* ptxsection; @@ -2819,10 +2819,10 @@ cuobjdumpELFSection* findELFSectionInList(std::list sectionli } //! Find an ELF section in all the known lists -cuobjdumpELFSection* gpgpu_context::findELFSection(const std::string identifier){ +cuobjdumpELFSection* cuda_runtime_api::findELFSection(const std::string identifier){ cuobjdumpELFSection* sec = findELFSectionInList(cuobjdumpSectionList, identifier); if (sec!=NULL)return sec; - sec = findELFSectionInList(api->libSectionList, identifier); + sec = findELFSectionInList(libSectionList, identifier); if (sec!=NULL)return sec; std::cout << "Could not find " << identifier << std::endl; assert(0 && "Could not find the required ELF section"); @@ -2854,10 +2854,10 @@ cuobjdumpPTXSection* findPTXSectionInList(std::list §ionl } //! Find an PTX section in all the known lists -cuobjdumpPTXSection* gpgpu_context::findPTXSection(const std::string identifier){ +cuobjdumpPTXSection* cuda_runtime_api::findPTXSection(const std::string identifier){ cuobjdumpPTXSection* sec = findPTXSectionInList(cuobjdumpSectionList, identifier); if (sec!=NULL)return sec; - sec = findPTXSectionInList(api->libSectionList, identifier); + sec = findPTXSectionInList(libSectionList, identifier); if (sec!=NULL)return sec; std::cout << "Could not find " << identifier << std::endl; assert(0 && "Could not find the required PTX section"); @@ -2867,7 +2867,7 @@ cuobjdumpPTXSection* gpgpu_context::findPTXSection(const std::string identifier) //! Extract the code using cuobjdump and remove unnecessary sections -void gpgpu_context::cuobjdumpInit(){ +void cuda_runtime_api::cuobjdumpInit(){ CUctx_st *context = GPGPUSim_Context(); extract_code_using_cuobjdump(); //extract all the output of cuobjdump to _cuobjdump_*.* const char* pre_load = getenv("CUOBJDUMP_SIM_FILE"); @@ -2879,7 +2879,7 @@ void gpgpu_context::cuobjdumpInit(){ //! Either submit PTX for simulation or convert SASS to PTXPlus and submit it -void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){ +void cuda_runtime_api::cuobjdumpParseBinary(unsigned int handle){ CUctx_st *context = GPGPUSim_Context(); if(fatbin_registered[handle]) return; @@ -3140,7 +3140,7 @@ cudaError_t cudaGLUnmapBufferObject(GLuint bufferObj) } #ifdef OPENGL_SUPPORT CUctx_st* ctx = GPGPUSim_Context(); - glbmap_entry_t *p = ctx->g_glbmap; + glbmap_entry_t *p = ctx->api->g_glbmap; while ( p && p->m_bufferObj != bufferObj ) p = p->m_next; if ( p == NULL ) @@ -3495,7 +3495,7 @@ static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu return nc_bytes; } -kernel_info_t * gpgpu_context::gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, +kernel_info_t * cuda_runtime_api::gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, gpgpu_ptx_sim_arg_list_t args, struct dim3 gridDim, struct dim3 blockDim, diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 7569ea6..4622f00 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -1,68 +1,15 @@ #ifndef __gpgpu_context_h__ #define __gpgpu_context_h__ -#include -#include -#include -#include #include "cuda_api_object.h" -class cuobjdumpSection; -class cuobjdumpELFSection; -class cuobjdumpPTXSection; -class symbol_table; -class gpgpu_ptx_sim_arg; -class kernel_info_t; - -typedef std::list gpgpu_ptx_sim_arg_list_t; - -#ifndef OPENGL_SUPPORT -typedef unsigned long GLuint; -#endif - -struct glbmap_entry { - GLuint m_bufferObj; - void *m_devPtr; - size_t m_size; - struct glbmap_entry *m_next; -}; - -typedef struct glbmap_entry glbmap_entry_t; - class gpgpu_context { public: gpgpu_context() { api = new cuda_runtime_api(); - g_glbmap = NULL; } // global list - std::list cuobjdumpSectionList; - std::mapfatbin_registered; - std::map fatbinmap; - std::map name_symtab; - std::map g_mallocPtr_Size; - //maps sm version number to set of filenames - std::map > version_filename; - std::map pinned_memory; //support for pinned memories added - std::map pinned_memory_size; - glbmap_entry_t* g_glbmap; // objects pointers for each file cuda_runtime_api* api; // member function list - void cuobjdumpInit(); - void cuobjdumpParseBinary(unsigned int handle); - void extract_code_using_cuobjdump(); - std::list pruneSectionList(CUctx_st *context); - std::list mergeMatchingSections(std::string identifier); - std::list mergeSections(); - cuobjdumpELFSection* findELFSection(const std::string identifier); - cuobjdumpPTXSection* findPTXSection(const std::string identifier); - void extract_ptx_files_using_cuobjdump(CUctx_st *context); - void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context); - kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, - gpgpu_ptx_sim_arg_list_t args, - struct dim3 gridDim, - struct dim3 blockDim, - struct CUctx_st* context ); - }; #endif /* __gpgpu_context_h__ */ -- cgit v1.3 From 269fbbdb72dacdf4dc3c482b213174336c3eeb1f Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 10 Jun 2019 22:42:27 -0400 Subject: Move getters for gpgpu_context to globals Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 2 +- libcuda/gpgpu_context.h | 2 ++ 2 files changed, 3 insertions(+), 1 deletion(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 2e2b50b..c5ebc20 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -273,7 +273,7 @@ static CUctx_st* GPGPUSim_Context() return the_context; } -static gpgpu_context* GPGPU_Context() +gpgpu_context* GPGPU_Context() { static gpgpu_context *gpgpu_ctx = NULL; if( gpgpu_ctx == NULL ) { diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 4622f00..6c8a293 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -12,4 +12,6 @@ class gpgpu_context { cuda_runtime_api* api; // member function list }; +gpgpu_context* GPGPU_Context(); + #endif /* __gpgpu_context_h__ */ -- cgit v1.3 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 --- libcuda/cuda_api_object.h | 1 - libcuda/cuda_runtime_api.cc | 42 ++++++++++++++++++++--------------------- libcuda/gpgpu_context.h | 3 +++ libopencl/opencl_runtime_api.cc | 5 ++++- src/cuda-sim/ptx_loader.cc | 5 +++-- src/cuda-sim/ptx_loader.h | 2 -- 6 files changed, 31 insertions(+), 27 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/cuda_api_object.h b/libcuda/cuda_api_object.h index f9a4fde..0054697 100644 --- a/libcuda/cuda_api_object.h +++ b/libcuda/cuda_api_object.h @@ -191,7 +191,6 @@ class cuda_runtime_api { void cuobjdumpInit(); void extract_code_using_cuobjdump(); void extract_ptx_files_using_cuobjdump(CUctx_st *context); - void cuobjdumpParseBinary(unsigned int handle); std::list pruneSectionList(CUctx_st *context); std::list mergeMatchingSections(std::string identifier); std::list mergeSections(); diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 09a13a7..25642a7 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -685,7 +685,7 @@ void cudaRegisterFunctionInternal( printf("GPGPU-Sim PTX: __cudaRegisterFunction %s : hostFun 0x%p, fat_cubin_handle = %u\n", deviceFun, hostFun, fat_cubin_handle); if(context->get_device()->get_gpgpu()->get_config().use_cuobjdump()) - ctx->api->cuobjdumpParseBinary(fat_cubin_handle); + ctx->cuobjdumpParseBinary(fat_cubin_handle); context->register_function( fat_cubin_handle, hostFun, deviceFun ); } @@ -712,7 +712,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()) - ctx->api->cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle); + ctx->cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle); fflush(stdout); if ( constant && !global && !ext ) { gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size); @@ -1009,7 +1009,7 @@ cuLinkAddFileInternal(CUlinkState state, CUjitInputType type, const char *path, } strcat(file,"/"); strcat(file,path); - symbol_table *symtab = gpgpu_ptx_sim_load_ptx_from_filename( file ); + 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); @@ -2750,15 +2750,15 @@ void cuda_runtime_api::cuobjdumpInit(){ //! Either submit PTX for simulation or convert SASS to PTXPlus and submit it -void cuda_runtime_api::cuobjdumpParseBinary(unsigned int handle){ +void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){ CUctx_st *context = GPGPUSim_Context(); - if(fatbin_registered[handle]) return; - fatbin_registered[handle] = true; - std::string fname = fatbinmap[handle]; + if(api->fatbin_registered[handle]) return; + api->fatbin_registered[handle] = true; + std::string fname = api->fatbinmap[handle]; - if (name_symtab.find(fname) != name_symtab.end()) { - symbol_table *symtab = name_symtab[fname]; + if (api->name_symtab.find(fname) != api->name_symtab.end()) { + symbol_table *symtab = api->name_symtab[fname]; context->add_binary(symtab, handle); return; } @@ -2767,7 +2767,7 @@ void cuda_runtime_api::cuobjdumpParseBinary(unsigned int handle){ #if (CUDART_VERSION >= 6000) //loops through all ptx files from smallest sm version to largest std::map >::iterator itr_m; - for (itr_m = version_filename.begin(); itr_m!=version_filename.end(); itr_m++){ + for (itr_m = api->version_filename.begin(); itr_m!=api->version_filename.end(); itr_m++){ std::set::iterator itr_s; for (itr_s = itr_m->second.begin(); itr_s!=itr_m->second.end(); itr_s++){ std::string ptx_filename = *itr_s; @@ -2775,11 +2775,11 @@ void cuda_runtime_api::cuobjdumpParseBinary(unsigned int handle){ symtab = gpgpu_ptx_sim_load_ptx_from_filename( ptx_filename.c_str() ); } } - name_symtab[fname] = symtab; + api->name_symtab[fname] = symtab; context->add_binary(symtab, handle); - load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); - load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); - for (itr_m = version_filename.begin(); itr_m!=version_filename.end(); itr_m++){ + api->load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); + api->load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); + for (itr_m = api->version_filename.begin(); itr_m!=api->version_filename.end(); itr_m++){ std::set::iterator itr_s; for (itr_s = itr_m->second.begin(); itr_s!=itr_m->second.end(); itr_s++){ std::string ptx_filename = *itr_s; @@ -2791,8 +2791,8 @@ void cuda_runtime_api::cuobjdumpParseBinary(unsigned int handle){ #endif unsigned max_capability = 0; - for ( std::list::iterator iter = cuobjdumpSectionList.begin(); - iter != cuobjdumpSectionList.end(); + for ( std::list::iterator iter = api->cuobjdumpSectionList.begin(); + iter != api->cuobjdumpSectionList.end(); iter++){ unsigned capability = (*iter)->getArch(); if (capability > max_capability) max_capability = capability; @@ -2803,7 +2803,7 @@ void cuda_runtime_api::cuobjdumpParseBinary(unsigned int handle){ cuobjdumpPTXSection* ptx = NULL; const char* pre_load = getenv("CUOBJDUMP_SIM_FILE"); if(pre_load==NULL || strlen(pre_load)==0) - ptx = findPTXSection(fname); + ptx = api->findPTXSection(fname); char *ptxcode; const char *override_ptx_name = getenv("PTX_SIM_KERNELFILE"); if (override_ptx_name == NULL or getenv("PTX_SIM_USE_PTX_FILE") == NULL or strlen(getenv("PTX_SIM_USE_PTX_FILE"))==0) { @@ -2813,7 +2813,7 @@ void cuda_runtime_api::cuobjdumpParseBinary(unsigned int handle){ ptxcode = readfile(override_ptx_name); } if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) { - cuobjdumpELFSection* elfsection = findELFSection(ptx->getIdentifier()); + cuobjdumpELFSection* elfsection = api->findELFSection(ptx->getIdentifier()); assert (elfsection!= NULL); char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus( ptx->getPTXfilename(), @@ -2831,9 +2831,9 @@ void cuda_runtime_api::cuobjdumpParseBinary(unsigned int handle){ context->add_binary(symtab, handle); gpgpu_ptxinfo_load_from_string( ptxcode, handle, max_capability, context->no_of_ptx ); } - load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); - load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); - name_symtab[fname] = symtab; + api->load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); + api->load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu()); + api->name_symtab[fname] = symtab; //TODO: Remove temporarily files as per configurations } diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 6c8a293..16626eb 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -11,6 +11,9 @@ class gpgpu_context { // objects pointers for each file cuda_runtime_api* api; // 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 ); }; gpgpu_context* GPGPU_Context(); 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); } diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index f037c34..d7e9b71 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -33,6 +33,7 @@ #include #include #include +#include "../../libcuda/gpgpu_context.h" /// globals @@ -165,7 +166,7 @@ char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptxfilenam } -symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ) +symbol_table *gpgpu_context::gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ) { char buf[1024]; snprintf(buf,1024,"_%u.ptx", source_num ); @@ -200,7 +201,7 @@ symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source return symtab; } -symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename ) +symbol_table *gpgpu_context::gpgpu_ptx_sim_load_ptx_from_filename( const char *filename ) { symbol_table *symtab=init_parser(filename); printf("GPGPU-Sim PTX: finished parsing EMBEDDED .ptx file %s\n",filename); diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index c4d8292..2af611a 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -42,8 +42,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. -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_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); -- 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 'libcuda/gpgpu_context.h') 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 7ac6034a99b52d09db7ef07bc008b8f6b039f929 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 12 Jun 2019 18:03:35 -0400 Subject: Move print_ptx_file Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 1 + src/cuda-sim/ptx_loader.cc | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index d3db1ad..dbc3140 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -19,6 +19,7 @@ class gpgpu_context { 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 ); + void print_ptx_file( const char *p, unsigned source_num, const char *filename ); }; gpgpu_context* GPGPU_Context(); diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 2a6d930..38597e0 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -81,7 +81,7 @@ void ptx_reg_options(option_parser_t opp) "0"); } -void print_ptx_file( const char *p, unsigned source_num, const char *filename ) +void gpgpu_context::print_ptx_file( const char *p, unsigned source_num, const char *filename ) { printf("\nGPGPU-Sim PTX: file _%u.ptx contents:\n\n", source_num ); char *s = strdup(p); -- cgit v1.3 From c146d3707216f235a1dde3bd978239fd55b111ed Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 12 Jun 2019 18:17:45 -0400 Subject: Move all ptx_parser inside gpgpu_context Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 4 ++++ src/cuda-sim/ptx_loader.cc | 9 ++++----- src/cuda-sim/ptx_parser.cc | 18 +++++++++--------- src/cuda-sim/ptx_parser.h | 1 - 4 files changed, 17 insertions(+), 15 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index dbc3140..b586e47 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -2,17 +2,20 @@ #define __gpgpu_context_h__ #include "cuda_api_object.h" #include "../src/cuda-sim/ptx_loader.h" +#include "../src/cuda-sim/ptx_parser.h" class gpgpu_context { public: gpgpu_context() { api = new cuda_runtime_api(); ptxinfo = new ptxinfo_data(); + ptx_parser = new ptx_recognizer(); } // global list // objects pointers for each file cuda_runtime_api* api; ptxinfo_data* ptxinfo; + ptx_recognizer* ptx_parser; // member function list void cuobjdumpParseBinary(unsigned int handle); class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); @@ -20,6 +23,7 @@ class gpgpu_context { 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 ); void print_ptx_file( const char *p, unsigned source_num, const char *filename ); + class symbol_table* init_parser(const char*); }; gpgpu_context* GPGPU_Context(); diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 38597e0..921d1e6 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -170,10 +170,9 @@ symbol_table *gpgpu_context::gpgpu_ptx_sim_load_ptx_from_string( const char *p, fclose(fp); } symbol_table *symtab=init_parser(buf); - ptx_recognizer recognizer; - ptx_lex_init(&(recognizer.scanner)); - ptx__scan_string(p, recognizer.scanner); - int errors = ptx_parse (recognizer.scanner, &recognizer); + ptx_lex_init(&(ptx_parser->scanner)); + ptx__scan_string(p, ptx_parser->scanner); + int errors = ptx_parse (ptx_parser->scanner, ptx_parser); if ( errors ) { char fname[1024]; snprintf(fname,1024,"_ptx_errors_XXXXXX"); @@ -186,7 +185,7 @@ symbol_table *gpgpu_context::gpgpu_ptx_sim_load_ptx_from_string( const char *p, abort(); exit(40); } - ptx_lex_destroy(recognizer.scanner); + ptx_lex_destroy(ptx_parser->scanner); if ( g_debug_execution >= 100 ) print_ptx_file(p,source_num,buf); diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 897beaf..7080ee9 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -27,6 +27,7 @@ #include "ptx_parser.h" #include "ptx_ir.h" +#include "../../libcuda/gpgpu_context.h" typedef void * yyscan_t; #include "ptx.tab.h" @@ -122,7 +123,7 @@ void ptx_recognizer::init_instruction_state() init_directive_state(); } -symbol_table *init_parser( const char *ptx_filename ) +symbol_table * gpgpu_context::init_parser( const char *ptx_filename ) { g_filename = strdup(ptx_filename); if (g_global_allfiles_symbol_table == NULL) { @@ -151,17 +152,16 @@ symbol_table *init_parser( const char *ptx_filename ) g_ptx_token_decode[generic_space] = "generic_space"; g_ptx_token_decode[instruction_space] = "instruction_space"; - ptx_recognizer recognizer; - ptx_lex_init(&(recognizer.scanner)); - recognizer.init_directive_state(); - recognizer.init_instruction_state(); + ptx_lex_init(&(ptx_parser->scanner)); + ptx_parser->init_directive_state(); + ptx_parser->init_instruction_state(); FILE *ptx_in; ptx_in = fopen(ptx_filename, "r"); - ptx_set_in(ptx_in, recognizer.scanner); - ptx_parse(recognizer.scanner, &recognizer); - ptx_in = ptx_get_in(recognizer.scanner); - ptx_lex_destroy(recognizer.scanner); + ptx_set_in(ptx_in, ptx_parser->scanner); + ptx_parse(ptx_parser->scanner, ptx_parser); + ptx_in = ptx_get_in(ptx_parser->scanner); + ptx_lex_destroy(ptx_parser->scanner); fclose(ptx_in); return g_global_symbol_table; } diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h index c52fe20..5293c2e 100644 --- a/src/cuda-sim/ptx_parser.h +++ b/src/cuda-sim/ptx_parser.h @@ -34,7 +34,6 @@ extern const char *g_filename; extern int g_error_detected; #ifdef __cplusplus -class symbol_table* init_parser(const char*); const class ptx_instruction *ptx_instruction_lookup( const char *filename, unsigned linenumber ); #endif -- cgit v1.3 From 5530b51c7c572d25717ad7f5d81196011dfcb540 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 12 Jun 2019 18:49:58 -0400 Subject: Move symbol table globals Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 2 ++ src/cuda-sim/ptx_parser.cc | 10 +++------- src/cuda-sim/ptx_parser.h | 7 +++++++ 3 files changed, 12 insertions(+), 7 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index b586e47..fa5d02b 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -7,11 +7,13 @@ class gpgpu_context { public: gpgpu_context() { + g_global_allfiles_symbol_table = NULL; api = new cuda_runtime_api(); ptxinfo = new ptxinfo_data(); ptx_parser = new ptx_recognizer(); } // global list + symbol_table *g_global_allfiles_symbol_table; // objects pointers for each file cuda_runtime_api* api; ptxinfo_data* ptxinfo; diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 7080ee9..32e12a1 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -52,12 +52,8 @@ static bool g_debug_ir_generation=false; const char *g_filename; // the program intermediate representation... -static symbol_table *g_global_allfiles_symbol_table = NULL; -static symbol_table *g_global_symbol_table = NULL; std::map g_sym_name_to_symbol_table; -static symbol_table *g_current_symbol_table = NULL; static std::list g_instructions; -static symbol *g_last_symbol = NULL; int g_error_detected = 0; @@ -128,7 +124,7 @@ symbol_table * gpgpu_context::init_parser( const char *ptx_filename ) g_filename = strdup(ptx_filename); if (g_global_allfiles_symbol_table == NULL) { g_global_allfiles_symbol_table = new symbol_table("global_allfiles", 0, NULL); - g_global_symbol_table = g_current_symbol_table = g_global_allfiles_symbol_table; + ptx_parser->g_global_symbol_table = ptx_parser->g_current_symbol_table = g_global_allfiles_symbol_table; } /*else { g_global_symbol_table = g_current_symbol_table = new symbol_table("global",0,g_global_allfiles_symbol_table); @@ -163,7 +159,7 @@ symbol_table * gpgpu_context::init_parser( const char *ptx_filename ) ptx_in = ptx_get_in(ptx_parser->scanner); ptx_lex_destroy(ptx_parser->scanner); fclose(ptx_in); - return g_global_symbol_table; + return ptx_parser->g_global_symbol_table; } static int g_entry_point; @@ -331,7 +327,7 @@ void ptx_recognizer::set_variable_type() g_extern_spec ); } -bool check_for_duplicates( const char *identifier ) +bool ptx_recognizer::check_for_duplicates( const char *identifier ) { const symbol *s = g_current_symbol_table->lookup(identifier); return ( s != NULL ); diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h index 5293c2e..cc8fe18 100644 --- a/src/cuda-sim/ptx_parser.h +++ b/src/cuda-sim/ptx_parser.h @@ -56,6 +56,9 @@ class ptx_recognizer { g_ident_add_uid = 0; g_const_alloc = 1; g_max_regs_per_thread = 0; + g_global_symbol_table = NULL; + g_current_symbol_table = NULL; + symbol *g_last_symbol = NULL; } // global list yyscan_t scanner; @@ -89,6 +92,9 @@ class ptx_recognizer { int g_ident_add_uid; unsigned g_const_alloc; unsigned g_max_regs_per_thread; + symbol_table *g_global_symbol_table; + symbol_table *g_current_symbol_table; + symbol *g_last_symbol; // member function list void init_directive_state(); @@ -154,6 +160,7 @@ class ptx_recognizer { //Jin: handle instructino group for cdp void start_inst_group(); void end_inst_group(); + bool check_for_duplicates( const char *identifier ); }; -- 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 'libcuda/gpgpu_context.h') 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 71eaa4e7b4aba49efef39391f3e1bb3a7e4e1529 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 3 Jul 2019 02:44:08 -0400 Subject: Connect gpgpu_context and GPGPUsim_ctx Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 3 +++ src/gpgpusim_entrypoint.cc | 6 +++--- src/gpgpusim_entrypoint.h | 10 ++++++---- 3 files changed, 12 insertions(+), 7 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index dd8b51d..69341be 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -3,6 +3,7 @@ #include "cuda_api_object.h" #include "../src/cuda-sim/ptx_loader.h" #include "../src/cuda-sim/ptx_parser.h" +#include "../src/gpgpusim_entrypoint.h" class gpgpu_context { public: @@ -11,6 +12,7 @@ class gpgpu_context { api = new cuda_runtime_api(); ptxinfo = new ptxinfo_data(); ptx_parser = new ptx_recognizer(); + the_gpgpusim = new GPGPUsim_ctx(this); } // global list symbol_table *g_global_allfiles_symbol_table; @@ -18,6 +20,7 @@ class gpgpu_context { cuda_runtime_api* api; ptxinfo_data* ptxinfo; ptx_recognizer* ptx_parser; + GPGPUsim_ctx* the_gpgpusim; // 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/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 476a9d4..c91a9d1 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -43,11 +43,11 @@ static int sg_argc = 3; static const char *sg_argv[] = {"", "-config","gpgpusim.config"}; -struct GPGPUsim_ctx* the_gpgpusim = NULL; +GPGPUsim_ctx* the_gpgpusim = NULL; -struct GPGPUsim_ctx* GPGPUsim_ctx_ptr(){ +GPGPUsim_ctx* GPGPUsim_ctx_ptr(){ if(the_gpgpusim == NULL) - the_gpgpusim = new GPGPUsim_ctx(); + the_gpgpusim = GPGPU_Context()->the_gpgpusim; return the_gpgpusim; } diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h index a443151..dfb82d0 100644 --- a/src/gpgpusim_entrypoint.h +++ b/src/gpgpusim_entrypoint.h @@ -34,11 +34,11 @@ #include //extern time_t g_simulation_starttime; +class gpgpu_context; - -struct GPGPUsim_ctx { - - GPGPUsim_ctx() { +class GPGPUsim_ctx { + public: + GPGPUsim_ctx(gpgpu_context* ctx) { g_sim_active = false; g_sim_done = true; break_limit = false; @@ -49,6 +49,7 @@ struct GPGPUsim_ctx { g_stream_manager=NULL; the_cude_device=NULL; the_context=NULL; + gpgpu_ctx = ctx; } //struct gpgpu_ptx_sim_arg *grid_params; @@ -65,6 +66,7 @@ struct GPGPUsim_ctx { struct _cuda_device_id *the_cude_device; struct CUctx_st* the_context; + gpgpu_context* gpgpu_ctx; pthread_mutex_t g_sim_lock; -- cgit v1.3 From 60c38b0f77378eb111e88f632702d19dc3746cc7 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 3 Jul 2019 14:54:16 -0400 Subject: Remove g_filename Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 5 +++-- src/cuda-sim/ptx.l | 5 ++--- src/cuda-sim/ptx.y | 5 ++--- src/cuda-sim/ptx_loader.h | 5 +++++ src/cuda-sim/ptx_parser.cc | 26 ++++++++++++-------------- src/cuda-sim/ptx_parser.h | 7 +++++-- src/cuda-sim/ptxinfo.l | 4 ++-- src/gpgpu-sim/gpu-sim.h | 1 + 8 files changed, 32 insertions(+), 26 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 69341be..d819559 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -10,12 +10,13 @@ class gpgpu_context { gpgpu_context() { g_global_allfiles_symbol_table = NULL; api = new cuda_runtime_api(); - ptxinfo = new ptxinfo_data(); - ptx_parser = new ptx_recognizer(); + ptxinfo = new ptxinfo_data(this); + ptx_parser = new ptx_recognizer(this); the_gpgpusim = new GPGPUsim_ctx(this); } // global list symbol_table *g_global_allfiles_symbol_table; + const char *g_filename; // objects pointers for each file cuda_runtime_api* api; ptxinfo_data* ptxinfo; diff --git a/src/cuda-sim/ptx.l b/src/cuda-sim/ptx.l index 18952a9..3a2a839 100644 --- a/src/cuda-sim/ptx.l +++ b/src/cuda-sim/ptx.l @@ -40,6 +40,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include "ptx_parser.h" #include "ptx.tab.h" #include +#include "../../libcuda/gpgpu_context.h" #define LINEBUF_SIZE (4*1024) #define TC recognizer->col+=strlen(yytext); @@ -452,8 +453,6 @@ breakaddr TC; yylval->int_value = BREAKADDR_OP; return OPCODE; . TC; ptx_error(yyscanner, recognizer, (const char*)NULL); %% -extern const char *g_filename; - int ptx_error( yyscan_t yyscanner, ptx_recognizer* recognizer, const char *s ) { struct yyguts_t * yyg = (struct yyguts_t*)yyscanner; @@ -461,7 +460,7 @@ int ptx_error( yyscan_t yyscanner, ptx_recognizer* recognizer, const char *s ) recognizer->g_error_detected = 1; fflush(stdout); if( s != NULL ) - printf("%s:%u Syntax error:\n\n", g_filename, yylineno ); + printf("%s:%u Syntax error:\n\n", recognizer->gpgpu_ctx->g_filename, yylineno ); printf(" %s\n", recognizer->linebuf ); printf(" "); for( i=0; i < recognizer->col-1; i++ ) { diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y index 837fbe9..a01c3c6 100644 --- a/src/cuda-sim/ptx.y +++ b/src/cuda-sim/ptx.y @@ -30,6 +30,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. %{ typedef void * yyscan_t; class ptx_recognizer; +#include "../../libcuda/gpgpu_context.h" %} %define api.pure full @@ -628,11 +629,9 @@ address_expression: IDENTIFIER { recognizer->add_address_operand($1,0); } %% -extern const char *g_filename; - void syntax_not_implemented(yyscan_t yyscanner, ptx_recognizer* recognizer) { - printf("Parse error (%s): this syntax is not (yet) implemented:\n",g_filename); + printf("Parse error (%s): this syntax is not (yet) implemented:\n", recognizer->gpgpu_ctx->g_filename); ptx_error(yyscanner, recognizer, NULL); abort(); } diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index 77f27c8..c214b95 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -30,13 +30,18 @@ #include #define PTXINFO_LINEBUF_SIZE 1024 +class gpgpu_context; typedef void * yyscan_t; class ptxinfo_data{ public: + ptxinfo_data(gpgpu_context* ctx) { + gpgpu_ctx = ctx; + } yyscan_t scanner; char linebuf[PTXINFO_LINEBUF_SIZE]; unsigned col; const char *g_ptxinfo_filename; + class gpgpu_context* gpgpu_ctx; void ptxinfo_addinfo(); }; diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 09bc15d..05fc618 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -47,14 +47,12 @@ void ptx_recognizer::set_ptx_warp_size(const struct core_config * warp_size) g_shader_core_config=warp_size; } -const char *g_filename; - // the program intermediate representation... std::map g_sym_name_to_symbol_table; #define PTX_PARSE_DPRINTF(...) \ if( g_debug_ir_generation ) { \ - printf(" %s:%u => ",g_filename,ptx_get_lineno(scanner)); \ + printf(" %s:%u => ",gpgpu_ctx->g_filename,ptx_get_lineno(scanner)); \ printf(" (%s:%u) ", __FILE__, __LINE__); \ printf(__VA_ARGS__); \ printf("\n"); \ @@ -70,7 +68,7 @@ const char *decode_token( int type ) void ptx_recognizer::read_parser_environment_variables() { - g_filename = getenv("PTX_SIM_KERNELFILE"); + gpgpu_ctx->g_filename = getenv("PTX_SIM_KERNELFILE"); char *dbg_level = getenv("PTX_SIM_DEBUG"); if ( dbg_level && strlen(dbg_level) ) { int debug_execution=0; @@ -180,7 +178,7 @@ void ptx_recognizer::add_function_name( const char *name ) if( prior_decl ) { g_func_info->remove_args(); } - g_global_symbol_table->add_function( g_func_info, g_filename, ptx_get_lineno(scanner) ); + g_global_symbol_table->add_function( g_func_info, gpgpu_ctx->g_filename, ptx_get_lineno(scanner) ); } //Jin: handle instruction group for cdp @@ -229,7 +227,7 @@ void ptx_recognizer::parse_error_impl( const char *file, unsigned line, const ch va_end(ap); g_error_detected = 1; - printf("%s:%u: Parse error: %s (%s:%u)\n\n", g_filename, ptx_get_lineno(scanner), buf, file, line); + printf("%s:%u: Parse error: %s (%s:%u)\n\n", gpgpu_ctx->g_filename, ptx_get_lineno(scanner), buf, file, line); ptx_error(scanner, NULL); abort(); exit(1); @@ -284,12 +282,12 @@ void ptx_recognizer::add_instruction() g_wmma_options, g_scalar_type, g_space_spec, - g_filename, + gpgpu_ctx->g_filename, ptx_get_lineno(scanner), linebuf, g_shader_core_config ); g_instructions.push_back(i); - g_inst_lookup[g_filename][ptx_get_lineno(scanner)] = i; + g_inst_lookup[gpgpu_ctx->g_filename][ptx_get_lineno(scanner)] = i; init_instruction_state(); } @@ -391,7 +389,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi default: break; } - g_last_symbol = g_current_symbol_table->add_variable(identifier,type,num_bits/8,g_filename,ptx_get_lineno(scanner)); + g_last_symbol = g_current_symbol_table->add_variable(identifier,type,num_bits/8,gpgpu_ctx->g_filename,ptx_get_lineno(scanner)); switch ( ti.get_memory_space().get_type() ) { case reg_space: { regnum = g_current_symbol_table->next_reg_num(); @@ -654,7 +652,7 @@ void ptx_recognizer::add_label( const char *identifier ) if ( s != NULL ) { g_label = s; } else { - g_label = g_current_symbol_table->add_variable(identifier,NULL,0,g_filename,ptx_get_lineno(scanner)); + g_label = g_current_symbol_table->add_variable(identifier,NULL,0,gpgpu_ctx->g_filename,ptx_get_lineno(scanner)); } } @@ -912,7 +910,7 @@ void ptx_recognizer::add_scalar_operand( const char *identifier ) if ( s == NULL ) { if ( g_opcode == BRA_OP || g_opcode == CALLP_OP) { // forward branch target... - s = g_current_symbol_table->add_variable(identifier,NULL,0,g_filename,ptx_get_lineno(scanner)); + s = g_current_symbol_table->add_variable(identifier,NULL,0,gpgpu_ctx->g_filename,ptx_get_lineno(scanner)); } else { std::string msg = std::string("operand \"") + identifier + "\" has no declaration."; parse_error( msg.c_str() ); @@ -926,7 +924,7 @@ void ptx_recognizer::add_neg_pred_operand( const char *identifier ) PTX_PARSE_DPRINTF("add_neg_pred_operand"); const symbol *s = g_current_symbol_table->lookup(identifier); if ( s == NULL ) { - s = g_current_symbol_table->add_variable(identifier,NULL,1,g_filename,ptx_get_lineno(scanner)); + s = g_current_symbol_table->add_variable(identifier,NULL,1,gpgpu_ctx->g_filename,ptx_get_lineno(scanner)); } operand_info op(s); op.set_neg_pred(); @@ -962,7 +960,7 @@ void ptx_recognizer::add_version_info( float ver, unsigned ext ) void ptx_recognizer::add_file( unsigned num, const char *filename ) { - if( g_filename == NULL ) { + if( gpgpu_ctx->g_filename == NULL ) { char *b = strdup(filename); char *l=b; char *n=b; @@ -978,7 +976,7 @@ void ptx_recognizer::add_file( unsigned num, const char *filename ) char *q = strtok(NULL,"."); if( q && !strcmp(q,"cu") ) { - g_filename = strdup(buf); + gpgpu_ctx->g_filename = strdup(buf); } free( b ); diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h index ec300a1..e09aab1 100644 --- a/src/cuda-sim/ptx_parser.h +++ b/src/cuda-sim/ptx_parser.h @@ -30,17 +30,17 @@ #include "../abstract_hardware_model.h" #include "ptx_ir.h" -extern const char *g_filename; extern int g_error_detected; #ifdef __cplusplus const class ptx_instruction *ptx_instruction_lookup( const char *filename, unsigned linenumber ); #endif +class gpgpu_context; typedef void * yyscan_t; class ptx_recognizer { public: - ptx_recognizer() { + ptx_recognizer( gpgpu_context* ctx ) { scanner = NULL; g_size = -1; g_add_identifier_cached__identifier = NULL; @@ -63,6 +63,7 @@ class ptx_recognizer { g_entry_func_param_index=0; g_func_info = NULL; g_debug_ir_generation=false; + gpgpu_ctx = ctx; } // global list yyscan_t scanner; @@ -107,6 +108,8 @@ class ptx_recognizer { bool g_debug_ir_generation; int g_entry_point; const struct core_config *g_shader_core_config; + // backward pointer + class gpgpu_context* gpgpu_ctx; // member function list void init_directive_state(); diff --git a/src/cuda-sim/ptxinfo.l b/src/cuda-sim/ptxinfo.l index b0ada09..3a152b0 100644 --- a/src/cuda-sim/ptxinfo.l +++ b/src/cuda-sim/ptxinfo.l @@ -39,6 +39,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include "ptx_loader.h" #include "ptxinfo.tab.h" #include +#include "../../libcuda/gpgpu_context.h" #define LINEBUF_SIZE 1024 #define TC if( (yylineno == 1) && (ptxinfo->col + strlen(yytext) < LINEBUF_SIZE) ) { \ @@ -88,7 +89,6 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. %% extern int g_ptxinfo_error_detected; -extern const char *g_filename; int ptxinfo_error(yyscan_t yyscanner, ptxinfo_data* ptxinfo, const char* msg) { @@ -98,7 +98,7 @@ int ptxinfo_error(yyscan_t yyscanner, ptxinfo_data* ptxinfo, const char* msg) fflush(stdout); printf("GPGPU-Sim: ERROR while parsing output of ptxas (used to capture resource usage information)\n"); if( msg != NULL ) - printf("GPGPU-Sim: %s (%s:%u) Syntax error:\n\n", g_filename, ptxinfo->g_ptxinfo_filename, yylineno ); + printf("GPGPU-Sim: %s (%s:%u) Syntax error:\n\n", ptxinfo->gpgpu_ctx->g_filename, ptxinfo->g_ptxinfo_filename, yylineno ); printf(" %s\n", ptxinfo->linebuf ); printf(" "); for( i=0; i < ptxinfo->col-1; i++ ) { diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index b1d2e45..e2c913a 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -506,6 +506,7 @@ private: void gpgpu_debug(); ///// data ///// +// backward pointer class gpgpu_context* gpgpu_ctx; class simt_core_cluster **m_cluster; -- 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 'libcuda/gpgpu_context.h') 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 b3e786e3d8d720217f36a214e9b5be9a19ab9dd2 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 8 Jul 2019 12:37:52 -0400 Subject: Move opcode_latency_int thus pass gpgpu_context into many classes Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 3 +++ src/cuda-sim/cuda-sim.cc | 7 ++++--- src/cuda-sim/cuda-sim.h | 11 +++++++++-- src/cuda-sim/ptx_ir.cc | 4 +++- src/cuda-sim/ptx_ir.h | 7 ++++++- src/cuda-sim/ptx_parser.cc | 3 ++- src/gpgpu-sim/gpu-sim.cc | 2 +- src/gpgpu-sim/gpu-sim.h | 11 ++++++++--- src/gpgpu-sim/mem_latency_stat.cc | 2 +- src/gpgpu-sim/mem_latency_stat.h | 4 ++-- src/gpgpu-sim/power_interface.cc | 2 +- src/gpgpu-sim/power_interface.h | 2 +- src/gpgpu-sim/power_stat.cc | 6 +++--- src/gpgpu-sim/power_stat.h | 6 +++--- src/gpgpu-sim/shader.cc | 7 ++++--- src/gpgpu-sim/shader.h | 18 ++++++++++++------ src/gpgpusim_entrypoint.cc | 4 ++-- 17 files changed, 65 insertions(+), 34 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index a8b60f4..2e21009 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -4,6 +4,7 @@ #include "../src/cuda-sim/ptx_loader.h" #include "../src/cuda-sim/ptx_parser.h" #include "../src/gpgpusim_entrypoint.h" +#include "../src/cuda-sim/cuda-sim.h" class gpgpu_context { public: @@ -13,6 +14,7 @@ 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(); } // global list symbol_table *g_global_allfiles_symbol_table; @@ -22,6 +24,7 @@ class gpgpu_context { ptxinfo_data* ptxinfo; ptx_recognizer* ptx_parser; GPGPUsim_ctx* the_gpgpusim; + cuda_sim* func_sim; // 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/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 261d605..df0bbd7 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -51,6 +51,7 @@ typedef void * yyscan_t; #include "decuda_pred_table/decuda_pred_table.h" #include "../stream_manager.h" #include "cuda_device_runtime.h" +#include "../../libcuda/gpgpu_context.h" int gpgpu_ptx_instruction_classification; void ** g_inst_classification_stat = NULL; @@ -66,12 +67,12 @@ int cp_cta_resume; unsigned g_ptx_sim_num_insn = 0; unsigned gpgpu_param_num_shaders = 0; -char *opcode_latency_int, *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor; +char *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor; char *opcode_initiation_int, *opcode_initiation_fp, *opcode_initiation_dp,*opcode_initiation_sfu,*opcode_initiation_tensor; char *cdp_latency_str; unsigned cdp_latency[5]; -void ptx_opcocde_latency_options (option_parser_t opp) { +void cuda_sim::ptx_opcocde_latency_options (option_parser_t opp) { option_parser_register(opp, "-ptx_opcode_latency_int", OPT_CSTR, &opcode_latency_int, "Opcode latencies for integers " "Default 1,1,19,25,145", @@ -667,7 +668,7 @@ void ptx_instruction::set_opcode_and_latency() * [3] MAD * [4] DIV */ - sscanf(opcode_latency_int, "%u,%u,%u,%u,%u", + sscanf(gpgpu_ctx->func_sim->opcode_latency_int, "%u,%u,%u,%u,%u", &int_latency[0],&int_latency[1],&int_latency[2], &int_latency[3],&int_latency[4]); sscanf(opcode_latency_fp, "%u,%u,%u,%u,%u", diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index e690356..96d34f5 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -47,10 +47,9 @@ 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 char *opcode_latency_int, *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor; +extern char *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor; -void ptx_opcocde_latency_options (option_parser_t opp); 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, @@ -134,4 +133,12 @@ void print_ptxinfo(); void clear_ptxinfo(); struct gpgpu_ptx_sim_info get_ptxinfo(); +class cuda_sim { + public: + //global variables + char *opcode_latency_int; + //global functions + void ptx_opcocde_latency_options (option_parser_t opp); +}; + #endif diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 8cedf79..1bd409e 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1095,8 +1095,10 @@ ptx_instruction::ptx_instruction( int opcode, const char *file, unsigned line, const char *source, - const core_config *config ) : warp_inst_t(config) + const core_config *config, + gpgpu_context* ctx ) : warp_inst_t(config) { + gpgpu_ctx = ctx; m_uid = ++g_num_ptx_inst_uid; m_PC = 0; m_opcode = opcode; diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index fd869c6..1604551 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -43,6 +43,8 @@ #include "memory.h" +class gpgpu_context; + class type_info_key { public: type_info_key() @@ -931,7 +933,8 @@ public: const char *file, unsigned line, const char *source, - const core_config *config ); + const core_config *config, + gpgpu_context* ctx); void print_insn() const; virtual void print_insn( FILE *fp ) const; @@ -1187,6 +1190,8 @@ private: virtual void pre_decode(); friend class function_info; static unsigned g_num_ptx_inst_uid; + // backward pointer + class gpgpu_context* gpgpu_ctx; }; class param_info { diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 05fc618..0139534 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -285,7 +285,8 @@ void ptx_recognizer::add_instruction() gpgpu_ctx->g_filename, ptx_get_lineno(scanner), linebuf, - g_shader_core_config ); + g_shader_core_config, + gpgpu_ctx ); g_instructions.push_back(i); g_inst_lookup[gpgpu_ctx->g_filename][ptx_get_lineno(scanner)] = i; init_instruction_state(); diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 2ff37d1..39acdd9 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1826,7 +1826,7 @@ void gpgpu_sim::dump_pipeline( int mask, int s, int m ) const fflush(stdout); } -const struct shader_core_config * gpgpu_sim::getShaderCoreConfig() +const shader_core_config * gpgpu_sim::getShaderCoreConfig() { return m_shader_config; } diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index e2c913a..7eeb7dd 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -295,7 +295,10 @@ extern bool g_interactive_debugger_enabled; class gpgpu_sim_config : public power_config, public gpgpu_functional_sim_config { public: - gpgpu_sim_config() { m_valid = false; } + gpgpu_sim_config(gpgpu_context* ctx): m_shader_config(ctx) { + m_valid = false; + gpgpu_ctx = ctx; + } void reg_options(class OptionParser * opp); void init() { @@ -341,6 +344,8 @@ private: void init_clock_domains(void ); + // backward pointer + class gpgpu_context* gpgpu_ctx; bool m_valid; shader_core_config m_shader_config; memory_config m_memory_config; @@ -473,7 +478,7 @@ public: /*! * Returning the configuration of the shader core, used by the functional simulation only so far */ - const struct shader_core_config * getShaderCoreConfig(); + const shader_core_config * getShaderCoreConfig(); //! Get shader core Memory Configuration @@ -537,7 +542,7 @@ private: const gpgpu_sim_config &m_config; const struct cudaDeviceProp *m_cuda_properties; - const struct shader_core_config *m_shader_config; + const shader_core_config *m_shader_config; const struct memory_config *m_memory_config; // stats diff --git a/src/gpgpu-sim/mem_latency_stat.cc b/src/gpgpu-sim/mem_latency_stat.cc index 04dc75b..d08ba39 100644 --- a/src/gpgpu-sim/mem_latency_stat.cc +++ b/src/gpgpu-sim/mem_latency_stat.cc @@ -42,7 +42,7 @@ #include #include -memory_stats_t::memory_stats_t( unsigned n_shader, const struct shader_core_config *shader_config, const struct memory_config *mem_config, const class gpgpu_sim* gpu ) +memory_stats_t::memory_stats_t( unsigned n_shader, const shader_core_config *shader_config, const struct memory_config *mem_config, const class gpgpu_sim* gpu ) { assert( mem_config->m_valid ); assert( shader_config->m_valid ); diff --git a/src/gpgpu-sim/mem_latency_stat.h b/src/gpgpu-sim/mem_latency_stat.h index b9285c1..6ce568d 100644 --- a/src/gpgpu-sim/mem_latency_stat.h +++ b/src/gpgpu-sim/mem_latency_stat.h @@ -35,7 +35,7 @@ class memory_stats_t { public: memory_stats_t( unsigned n_shader, - const struct shader_core_config *shader_config, + const class shader_core_config *shader_config, const struct memory_config *mem_config, const class gpgpu_sim* gpu); @@ -53,7 +53,7 @@ public: unsigned m_n_shader; - const struct shader_core_config *m_shader_config; + const shader_core_config *m_shader_config; const struct memory_config *m_memory_config; const class gpgpu_sim* m_gpu; diff --git a/src/gpgpu-sim/power_interface.cc b/src/gpgpu-sim/power_interface.cc index 3861b6a..0272aa6 100644 --- a/src/gpgpu-sim/power_interface.cc +++ b/src/gpgpu-sim/power_interface.cc @@ -38,7 +38,7 @@ void init_mcpat(const gpgpu_sim_config &config, class gpgpu_sim_wrapper *wrapper } -void mcpat_cycle(const gpgpu_sim_config &config, const struct shader_core_config *shdr_config, class gpgpu_sim_wrapper *wrapper, class power_stat_t *power_stats, unsigned stat_sample_freq, unsigned tot_cycle, unsigned cycle, unsigned tot_inst, unsigned inst){ +void mcpat_cycle(const gpgpu_sim_config &config, const shader_core_config *shdr_config, class gpgpu_sim_wrapper *wrapper, class power_stat_t *power_stats, unsigned stat_sample_freq, unsigned tot_cycle, unsigned cycle, unsigned tot_inst, unsigned inst){ static bool mcpat_init=true; diff --git a/src/gpgpu-sim/power_interface.h b/src/gpgpu-sim/power_interface.h index afac22b..a388c23 100644 --- a/src/gpgpu-sim/power_interface.h +++ b/src/gpgpu-sim/power_interface.h @@ -36,7 +36,7 @@ #include "gpgpu_sim_wrapper.h" void init_mcpat(const gpgpu_sim_config &config, class gpgpu_sim_wrapper *wrapper, unsigned stat_sample_freq, unsigned tot_inst, unsigned inst); -void mcpat_cycle(const gpgpu_sim_config &config, const struct shader_core_config *shdr_config, class gpgpu_sim_wrapper *wrapper, class power_stat_t *power_stats, +void mcpat_cycle(const gpgpu_sim_config &config, const shader_core_config *shdr_config, class gpgpu_sim_wrapper *wrapper, class power_stat_t *power_stats, unsigned stat_sample_freq, unsigned tot_cycle, unsigned cycle, unsigned tot_inst, unsigned inst); void mcpat_reset_perf_count(class gpgpu_sim_wrapper *wrapper); diff --git a/src/gpgpu-sim/power_stat.cc b/src/gpgpu-sim/power_stat.cc index 4c995e9..007b4c6 100644 --- a/src/gpgpu-sim/power_stat.cc +++ b/src/gpgpu-sim/power_stat.cc @@ -42,7 +42,7 @@ -power_mem_stat_t::power_mem_stat_t(const struct memory_config *mem_config, const struct shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats){ +power_mem_stat_t::power_mem_stat_t(const struct memory_config *mem_config, const shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats){ assert( mem_config->m_valid ); m_mem_stats = mem_stats; m_config = mem_config; @@ -125,7 +125,7 @@ void power_mem_stat_t::print (FILE *fout) const { } -power_core_stat_t::power_core_stat_t( const struct shader_core_config *shader_config, shader_core_stats *core_stats ) +power_core_stat_t::power_core_stat_t( const shader_core_config *shader_config, shader_core_stats *core_stats ) { assert( shader_config->m_valid ); m_config = shader_config; @@ -266,7 +266,7 @@ for(unsigned i=0; inum_shader(); ++i){ } } -power_stat_t::power_stat_t( const struct shader_core_config *shader_config,float * average_pipeline_duty_cycle,float *active_sms,shader_core_stats * shader_stats, const struct memory_config *mem_config,memory_stats_t * memory_stats) +power_stat_t::power_stat_t( const shader_core_config *shader_config,float * average_pipeline_duty_cycle,float *active_sms,shader_core_stats * shader_stats, const struct memory_config *mem_config,memory_stats_t * memory_stats) { assert( shader_config->m_valid ); assert( mem_config->m_valid ); diff --git a/src/gpgpu-sim/power_stat.h b/src/gpgpu-sim/power_stat.h index 20af2e5..91fade9 100644 --- a/src/gpgpu-sim/power_stat.h +++ b/src/gpgpu-sim/power_stat.h @@ -73,7 +73,7 @@ struct shader_core_power_stats_pod { class power_core_stat_t : public shader_core_power_stats_pod { public: - power_core_stat_t(const struct shader_core_config *shader_config, shader_core_stats *core_stats); + power_core_stat_t(const shader_core_config *shader_config, shader_core_stats *core_stats); void visualizer_print( gzFile visualizer_file ); void print (FILE *fout); void init(); @@ -113,7 +113,7 @@ struct mem_power_stats_pod{ class power_mem_stat_t : public mem_power_stats_pod{ public: - power_mem_stat_t(const struct memory_config *mem_config, const struct shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats); + power_mem_stat_t(const struct memory_config *mem_config, const shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats); void visualizer_print( gzFile visualizer_file ); void print (FILE *fout) const; void init(); @@ -128,7 +128,7 @@ private: class power_stat_t { public: - power_stat_t( const struct shader_core_config *shader_config,float * average_pipeline_duty_cycle,float * active_sms,shader_core_stats * shader_stats, const struct memory_config *mem_config,memory_stats_t * memory_stats); + power_stat_t( const shader_core_config *shader_config,float * average_pipeline_duty_cycle,float * active_sms,shader_core_stats * shader_stats, const struct memory_config *mem_config,memory_stats_t * memory_stats); void visualizer_print( gzFile visualizer_file ); void print (FILE *fout) const; void save_stats(){ diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 69b619a..4d12068 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -46,6 +46,7 @@ #include #include "traffic_breakdown.h" #include "shader_trace.h" +#include "../../libcuda/gpgpu_context.h" #define PRIORITIZE_MSHR_OVER_WB 1 #define MAX(a,b) (((a)>(b))?(a):(b)) @@ -69,7 +70,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, class simt_core_cluster *cluster, unsigned shader_id, unsigned tpc_id, - const struct shader_core_config *config, + const shader_core_config *config, const struct memory_config *mem_config, shader_core_stats *stats ) : core_t( gpu, NULL, config->warp_size, config->n_thread_per_shader ), @@ -3018,7 +3019,7 @@ void shader_core_config::set_pipeline_latency() { * [3] MAD * [4] DIV */ - sscanf(opcode_latency_int, "%u,%u,%u,%u,%u", + sscanf(gpgpu_ctx->func_sim->opcode_latency_int, "%u,%u,%u,%u,%u", &int_latency[0],&int_latency[1],&int_latency[2], &int_latency[3],&int_latency[4]); sscanf(opcode_latency_fp, "%u,%u,%u,%u,%u", @@ -3786,7 +3787,7 @@ void opndcoll_rfu_t::collector_unit_t::dispatch() simt_core_cluster::simt_core_cluster( class gpgpu_sim *gpu, unsigned cluster_id, - const struct shader_core_config *config, + const shader_core_config *config, const struct memory_config *mem_config, shader_core_stats *stats, class memory_stats_t *mstats ) diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index 25b9607..e0cefac 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -69,6 +69,8 @@ #define WRITE_MASK_SIZE 8 +class gpgpu_context; + enum exec_unit_type_t { NONE = 0, @@ -294,7 +296,7 @@ typedef std::bitset warp_set_t; int register_bank(int regnum, int wid, unsigned num_banks, unsigned bank_warp_shift, bool sub_core_model, unsigned banks_per_sched, unsigned sched_id ); class shader_core_ctx; -struct shader_core_config; +class shader_core_config; class shader_core_stats; enum scheduler_prioritization_type @@ -1032,7 +1034,7 @@ struct ifetch_buffer_t { unsigned m_warp_id; }; -struct shader_core_config; +class shader_core_config; class simd_function_unit { public: @@ -1362,10 +1364,12 @@ const char* const pipeline_stage_name_decode[] = { "N_PIPELINE_STAGES" }; -struct shader_core_config : public core_config +class shader_core_config : public core_config { - shader_core_config(){ + public: + shader_core_config(gpgpu_context* ctx){ pipeline_widths_string = NULL; + gpgpu_ctx = ctx; } void init() @@ -1425,6 +1429,8 @@ struct shader_core_config : public core_config unsigned cid_to_sid( unsigned cid, unsigned cluster_id ) const { return cluster_id*n_simt_cores_per_cluster + cid; } void set_pipeline_latency(); + // backward pointer + class gpgpu_context* gpgpu_ctx; // data char *gpgpu_shader_core_pipeline_opt; bool gpgpu_perfect_mem; @@ -1770,7 +1776,7 @@ public: class simt_core_cluster *cluster, unsigned shader_id, unsigned tpc_id, - const struct shader_core_config *config, + const shader_core_config *config, const struct memory_config *mem_config, shader_core_stats *stats ); @@ -2065,7 +2071,7 @@ class simt_core_cluster { public: simt_core_cluster( class gpgpu_sim *gpu, unsigned cluster_id, - const struct shader_core_config *config, + const shader_core_config *config, const struct memory_config *mem_config, shader_core_stats *stats, memory_stats_t *mstats ); diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index b54f20c..5018305 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -218,10 +218,10 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf() option_parser_t opp = option_parser_create(); ptx_reg_options(opp); - ptx_opcocde_latency_options(opp); + func_sim->ptx_opcocde_latency_options(opp); icnt_reg_options(opp); - GPGPUsim_ctx_ptr()->g_the_gpu_config = new gpgpu_sim_config(); + 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 option_parser_cmdline(opp, sg_argc, sg_argv); // parse configuration options -- 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 'libcuda/gpgpu_context.h') 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 'libcuda/gpgpu_context.h') 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 2f22ddb07752bc956063f6548d4bc26a6ae066f2 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Thu, 11 Jul 2019 16:56:29 -0400 Subject: Fix symbols Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 4 ---- src/cuda-sim/cuda_device_runtime.h | 10 ++++++---- 2 files changed, 6 insertions(+), 8 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 3c9f87c..c0b250a 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -16,9 +16,7 @@ class gpgpu_context { ptx_parser = new ptx_recognizer(this); the_gpgpusim = new GPGPUsim_ctx(this); 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; @@ -29,9 +27,7 @@ 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/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h index b49d0eb..a6303b9 100644 --- a/src/cuda-sim/cuda_device_runtime.h +++ b/src/cuda-sim/cuda_device_runtime.h @@ -1,8 +1,7 @@ +#ifndef __cuda_device_runtime_h__ +#define __cuda_device_runtime_h__ //Jin: cuda_device_runtime.h //Defines CUDA device runtime APIs for CDP support -#if (CUDART_VERSION >= 5000) -#pragma once - class device_launch_config_t { public: @@ -51,11 +50,14 @@ class cuda_device_runtime { std::list g_cuda_device_launch_op; // backward pointer class gpgpu_context* gpgpu_ctx; +#if (CUDART_VERSION >= 5000) +#pragma once 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 gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); void launch_all_device_kernels(); void launch_one_device_kernel(); +#endif }; -#endif +#endif /* __cuda_device_runtime_h__ */ -- cgit v1.3 From 23c0bb224295dde9651fd915d854e4f7eafdf88f Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Sun, 14 Jul 2019 11:25:42 -0400 Subject: Move sm_next_access_uid Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 2 ++ src/abstract_hardware_model.cc | 13 ++++++++++--- src/abstract_hardware_model.h | 22 +++++++++------------- src/gpgpu-sim/dram.cc | 2 +- src/gpgpu-sim/dram.h | 5 +++-- src/gpgpu-sim/gpu-cache.cc | 6 ++++-- src/gpgpu-sim/gpu-sim.cc | 2 +- src/gpgpu-sim/gpu-sim.h | 15 +++++++++------ src/gpgpu-sim/l2cache.cc | 11 ++++++----- src/gpgpu-sim/l2cache.h | 12 ++++++------ src/gpgpu-sim/mem_fetch.cc | 4 ++-- src/gpgpu-sim/mem_fetch.h | 5 +++-- src/gpgpu-sim/mem_latency_stat.cc | 2 +- src/gpgpu-sim/mem_latency_stat.h | 5 +++-- src/gpgpu-sim/power_stat.cc | 4 ++-- src/gpgpu-sim/power_stat.h | 6 +++--- src/gpgpu-sim/shader.cc | 20 ++++++++++++++++---- src/gpgpu-sim/shader.h | 20 ++++---------------- 18 files changed, 85 insertions(+), 71 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index c0b250a..07473be 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -11,6 +11,7 @@ class gpgpu_context { public: gpgpu_context() { g_global_allfiles_symbol_table = NULL; + sm_next_access_uid=0; api = new cuda_runtime_api(this); ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); @@ -21,6 +22,7 @@ class gpgpu_context { // global list symbol_table *g_global_allfiles_symbol_table; const char *g_filename; + unsigned sm_next_access_uid; // objects pointers for each file cuda_runtime_api* api; ptxinfo_data* ptxinfo; diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index fde7874..fe10daa 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -41,9 +41,16 @@ #include #include "../libcuda/gpgpu_context.h" -unsigned mem_access_t::sm_next_access_uid = 0; unsigned warp_inst_t::sm_next_uid = 0; +void mem_access_t::init(gpgpu_context* ctx) +{ + gpgpu_ctx = ctx; + m_uid=++(gpgpu_ctx->sm_next_access_uid); + m_addr=0; + m_req_size=0; +} + checkpoint::checkpoint() { @@ -449,7 +456,7 @@ void warp_inst_t::generate_mem_accesses() byte_mask.set(idx+i); } for( a=accesses.begin(); a != accesses.end(); ++a ) - m_accessq.push_back( mem_access_t(access_type,a->first,cache_block_size,is_write,a->second, byte_mask, mem_access_sector_mask_t())); + m_accessq.push_back( mem_access_t(access_type,a->first,cache_block_size,is_write,a->second, byte_mask, mem_access_sector_mask_t(), m_config->gpgpu_ctx)); } if ( space.get_type() == global_space ) { @@ -681,7 +688,7 @@ void warp_inst_t::memory_coalescing_arch_reduce_and_send( bool is_write, mem_acc assert(lower_half_used && upper_half_used); } } - m_accessq.push_back( mem_access_t(access_type,addr,size,is_write,info.active,info.bytes, info.chunks) ); + m_accessq.push_back( mem_access_t(access_type,addr,size,is_write,info.active,info.bytes, info.chunks,m_config->gpgpu_ctx) ); } void warp_inst_t::completed( unsigned long long cycle ) const diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 8ef8376..a22c8c3 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -732,13 +732,14 @@ enum cache_operator_type { class mem_access_t { public: - mem_access_t() { init(); } + mem_access_t(gpgpu_context* ctx) { init(ctx); } mem_access_t( mem_access_type type, new_addr_type address, unsigned size, - bool wr ) + bool wr, + gpgpu_context* ctx) { - init(); + init(ctx); m_type = type; m_addr = address; m_req_size = size; @@ -750,10 +751,11 @@ public: bool wr, const active_mask_t &active_mask, const mem_access_byte_mask_t &byte_mask, - const mem_access_sector_mask_t §or_mask) + const mem_access_sector_mask_t §or_mask, + gpgpu_context* ctx) : m_warp_mask(active_mask), m_byte_mask(byte_mask), m_sector_mask(sector_mask) { - init(); + init(ctx); m_type = type; m_addr = address; m_req_size = size; @@ -786,13 +788,9 @@ public: } } + gpgpu_context* gpgpu_ctx; private: - void init() - { - m_uid=++sm_next_access_uid; - m_addr=0; - m_req_size=0; - } + void init(gpgpu_context* ctx); unsigned m_uid; new_addr_type m_addr; // request address @@ -802,8 +800,6 @@ private: active_mask_t m_warp_mask; mem_access_byte_mask_t m_byte_mask; mem_access_sector_mask_t m_sector_mask; - - static unsigned sm_next_access_uid; }; class mem_fetch; diff --git a/src/gpgpu-sim/dram.cc b/src/gpgpu-sim/dram.cc index 5e36d4b..d443d79 100644 --- a/src/gpgpu-sim/dram.cc +++ b/src/gpgpu-sim/dram.cc @@ -41,7 +41,7 @@ int PRINT_CYCLE = 0; template class fifo_pipeline; template class fifo_pipeline; -dram_t::dram_t( unsigned int partition_id, const struct memory_config *config, memory_stats_t *stats, +dram_t::dram_t( unsigned int partition_id, const memory_config *config, memory_stats_t *stats, memory_partition_unit *mp, gpgpu_sim* gpu ) { id = partition_id; diff --git a/src/gpgpu-sim/dram.h b/src/gpgpu-sim/dram.h index 7a3a2da..0bd9725 100644 --- a/src/gpgpu-sim/dram.h +++ b/src/gpgpu-sim/dram.h @@ -106,11 +106,12 @@ enum bank_grp_bits_position{ }; class mem_fetch; +class memory_config; class dram_t { public: - dram_t( unsigned int parition_id, const struct memory_config *config, class memory_stats_t *stats, + dram_t( unsigned int parition_id, const memory_config *config, class memory_stats_t *stats, class memory_partition_unit *mp, class gpgpu_sim* gpu ); bool full(bool is_write) const; @@ -145,7 +146,7 @@ public: - const struct memory_config *m_config; + const memory_config *m_config; private: bankgrp_t **bkgrp; diff --git a/src/gpgpu-sim/gpu-cache.cc b/src/gpgpu-sim/gpu-cache.cc index f1f6e19..1705821 100644 --- a/src/gpgpu-sim/gpu-cache.cc +++ b/src/gpgpu-sim/gpu-cache.cc @@ -1251,7 +1251,8 @@ data_cache::wr_miss_wa_naive( new_addr_type addr, false, // Now performing a read mf->get_access_warp_mask(), mf->get_access_byte_mask(), - mf->get_access_sector_mask()); + mf->get_access_sector_mask(), + m_gpu->gpgpu_ctx); mem_fetch *n_mf = new mem_fetch( *ma, NULL, @@ -1365,7 +1366,8 @@ data_cache::wr_miss_wa_fetch_on_write( new_addr_type addr, false, // Now performing a read mf->get_access_warp_mask(), mf->get_access_byte_mask(), - mf->get_access_sector_mask()); + mf->get_access_sector_mask(), + m_gpu->gpgpu_ctx); mem_fetch *n_mf = new mem_fetch( *ma, NULL, diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 0481259..0644b44 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1828,7 +1828,7 @@ const shader_core_config * gpgpu_sim::getShaderCoreConfig() return m_shader_config; } -const struct memory_config * gpgpu_sim::getMemoryConfig() +const memory_config * gpgpu_sim::getMemoryConfig() { return m_memory_config; } diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 78f0505..19e1eb3 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -33,6 +33,7 @@ #include "../trace.h" #include "addrdec.h" #include "shader.h" +#include "gpu-cache.h" #include #include #include @@ -143,13 +144,14 @@ struct power_config { }; - -struct memory_config { - memory_config() +class memory_config { + public: + memory_config(gpgpu_context* ctx) { m_valid = false; gpgpu_dram_timing_opt=NULL; gpgpu_L2_queue_config=NULL; + gpgpu_ctx = ctx; } void init() { @@ -291,13 +293,14 @@ struct memory_config { unsigned write_high_watermark; unsigned write_low_watermark; bool m_perf_sim_memcpy; + gpgpu_context* gpgpu_ctx; }; extern bool g_interactive_debugger_enabled; class gpgpu_sim_config : public power_config, public gpgpu_functional_sim_config { public: - gpgpu_sim_config(gpgpu_context* ctx): m_shader_config(ctx) { + gpgpu_sim_config(gpgpu_context* ctx): m_shader_config(ctx), m_memory_config(ctx) { m_valid = false; gpgpu_ctx = ctx; } @@ -507,7 +510,7 @@ public: /*! * Returning the memory configuration of the shader core, used by the functional simulation only so far */ - const struct memory_config * getMemoryConfig(); + const memory_config * getMemoryConfig(); //! Get shader core SIMT cluster @@ -567,7 +570,7 @@ private: const struct cudaDeviceProp *m_cuda_properties; const shader_core_config *m_shader_config; - const struct memory_config *m_memory_config; + const memory_config *m_memory_config; // stats class shader_core_stats *m_shader_stats; diff --git a/src/gpgpu-sim/l2cache.cc b/src/gpgpu-sim/l2cache.cc index 62e70a7..6540b52 100644 --- a/src/gpgpu-sim/l2cache.cc +++ b/src/gpgpu-sim/l2cache.cc @@ -49,7 +49,7 @@ mem_fetch * partition_mf_allocator::alloc(new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const { assert( wr ); - mem_access_t access( type, addr, size, wr ); + mem_access_t access( type, addr, size, wr, m_memory_config->gpgpu_ctx ); mem_fetch *mf = new mem_fetch( access, NULL, WRITE_PACKET_SIZE, @@ -62,7 +62,7 @@ mem_fetch * partition_mf_allocator::alloc(new_addr_type addr, mem_access_type ty } memory_partition_unit::memory_partition_unit( unsigned partition_id, - const struct memory_config *config, + const memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu) : m_id(partition_id), m_config(config), m_stats(stats), m_arbitration_metadata(config), m_gpu(gpu) @@ -95,7 +95,7 @@ memory_partition_unit::~memory_partition_unit() delete[] m_sub_partition; } -memory_partition_unit::arbitration_metadata::arbitration_metadata(const struct memory_config *config) +memory_partition_unit::arbitration_metadata::arbitration_metadata(const memory_config *config) : m_last_borrower(config->m_n_sub_partition_per_memory_channel - 1), m_private_credit(config->m_n_sub_partition_per_memory_channel, 0), m_shared_credit(0) @@ -312,7 +312,7 @@ void memory_partition_unit::print( FILE *fp ) const } memory_sub_partition::memory_sub_partition( unsigned sub_partition_id, - const struct memory_config *config, + const memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu) { @@ -640,7 +640,8 @@ std::vector memory_sub_partition::breakdown_request_to_sector_reques mf->is_write(), mf->get_access_warp_mask(), mf->get_access_byte_mask() & byte_sector_mask, - std::bitset().set(j)); + std::bitset().set(j), + m_gpu->gpgpu_ctx); mem_fetch *n_mf = new mem_fetch( *ma, NULL, diff --git a/src/gpgpu-sim/l2cache.h b/src/gpgpu-sim/l2cache.h index 8ff2666..1f74c47 100644 --- a/src/gpgpu-sim/l2cache.h +++ b/src/gpgpu-sim/l2cache.h @@ -58,7 +58,7 @@ private: class memory_partition_unit { public: - memory_partition_unit( unsigned partition_id, const struct memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu ); + memory_partition_unit( unsigned partition_id, const memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu ); ~memory_partition_unit(); bool busy() const; @@ -98,7 +98,7 @@ public: private: unsigned m_id; - const struct memory_config *m_config; + const memory_config *m_config; class memory_stats_t *m_stats; class memory_sub_partition **m_sub_partition; class dram_t *m_dram; @@ -106,7 +106,7 @@ private: class arbitration_metadata { public: - arbitration_metadata(const struct memory_config *config); + arbitration_metadata(const memory_config *config); // check if a subpartition still has credit bool has_credits(int inner_sub_partition_id) const; @@ -130,7 +130,7 @@ private: std::vector m_private_credit; int m_shared_credit; }; - arbitration_metadata m_arbitration_metadata; + arbitration_metadata m_arbitration_metadata; // determine wheither a given subpartition can issue to DRAM bool can_issue_to_dram(int inner_sub_partition_id); @@ -149,7 +149,7 @@ private: class memory_sub_partition { public: - memory_sub_partition( unsigned sub_partition_id, const struct memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu ); + memory_sub_partition( unsigned sub_partition_id, const memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu ); ~memory_sub_partition(); unsigned get_id() const { return m_id; } @@ -197,7 +197,7 @@ public: private: // data unsigned m_id; //< the global sub partition ID - const struct memory_config *m_config; + const memory_config *m_config; class l2_cache *m_L2cache; class L2interface *m_L2interface; class gpgpu_sim* m_gpu; diff --git a/src/gpgpu-sim/mem_fetch.cc b/src/gpgpu-sim/mem_fetch.cc index c9b0484..6a00889 100644 --- a/src/gpgpu-sim/mem_fetch.cc +++ b/src/gpgpu-sim/mem_fetch.cc @@ -39,10 +39,10 @@ mem_fetch::mem_fetch( const mem_access_t &access, unsigned wid, unsigned sid, unsigned tpc, - const struct memory_config *config, + const memory_config *config, unsigned long long cycle, mem_fetch *m_original_mf, - mem_fetch *m_original_wr_mf) + mem_fetch *m_original_wr_mf):m_access(access) { m_request_uid = sm_next_mf_request_uid++; diff --git a/src/gpgpu-sim/mem_fetch.h b/src/gpgpu-sim/mem_fetch.h index 4eb3a52..1cab9f2 100644 --- a/src/gpgpu-sim/mem_fetch.h +++ b/src/gpgpu-sim/mem_fetch.h @@ -47,6 +47,7 @@ enum mf_type { #undef MF_TUP #undef MF_TUP_END +class memory_config; class mem_fetch { public: mem_fetch( const mem_access_t &access, @@ -55,7 +56,7 @@ public: unsigned wid, unsigned sid, unsigned tpc, - const struct memory_config *config, + const memory_config *config, unsigned long long cycle, mem_fetch *original_mf = NULL, mem_fetch *original_wr_mf = NULL); @@ -149,7 +150,7 @@ private: static unsigned sm_next_mf_request_uid; - const struct memory_config *m_mem_config; + const memory_config *m_mem_config; unsigned icnt_flit_size; mem_fetch* original_mf; //this pointer is set up when a request is divided into sector requests at L2 cache (if the req size > L2 sector size), so the pointer refers to the original request diff --git a/src/gpgpu-sim/mem_latency_stat.cc b/src/gpgpu-sim/mem_latency_stat.cc index d08ba39..4e94991 100644 --- a/src/gpgpu-sim/mem_latency_stat.cc +++ b/src/gpgpu-sim/mem_latency_stat.cc @@ -42,7 +42,7 @@ #include #include -memory_stats_t::memory_stats_t( unsigned n_shader, const shader_core_config *shader_config, const struct memory_config *mem_config, const class gpgpu_sim* gpu ) +memory_stats_t::memory_stats_t( unsigned n_shader, const shader_core_config *shader_config, const memory_config *mem_config, const class gpgpu_sim* gpu ) { assert( mem_config->m_valid ); assert( shader_config->m_valid ); diff --git a/src/gpgpu-sim/mem_latency_stat.h b/src/gpgpu-sim/mem_latency_stat.h index 6ce568d..0c84972 100644 --- a/src/gpgpu-sim/mem_latency_stat.h +++ b/src/gpgpu-sim/mem_latency_stat.h @@ -32,11 +32,12 @@ #include #include +class memory_config; class memory_stats_t { public: memory_stats_t( unsigned n_shader, const class shader_core_config *shader_config, - const struct memory_config *mem_config, + const memory_config *mem_config, const class gpgpu_sim* gpu); unsigned memlatstat_done( class mem_fetch *mf ); @@ -54,7 +55,7 @@ public: unsigned m_n_shader; const shader_core_config *m_shader_config; - const struct memory_config *m_memory_config; + const memory_config *m_memory_config; const class gpgpu_sim* m_gpu; unsigned max_mrq_latency; diff --git a/src/gpgpu-sim/power_stat.cc b/src/gpgpu-sim/power_stat.cc index 007b4c6..2c02082 100644 --- a/src/gpgpu-sim/power_stat.cc +++ b/src/gpgpu-sim/power_stat.cc @@ -42,7 +42,7 @@ -power_mem_stat_t::power_mem_stat_t(const struct memory_config *mem_config, const shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats){ +power_mem_stat_t::power_mem_stat_t(const memory_config *mem_config, const shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats){ assert( mem_config->m_valid ); m_mem_stats = mem_stats; m_config = mem_config; @@ -266,7 +266,7 @@ for(unsigned i=0; inum_shader(); ++i){ } } -power_stat_t::power_stat_t( const shader_core_config *shader_config,float * average_pipeline_duty_cycle,float *active_sms,shader_core_stats * shader_stats, const struct memory_config *mem_config,memory_stats_t * memory_stats) +power_stat_t::power_stat_t( const shader_core_config *shader_config,float * average_pipeline_duty_cycle,float *active_sms,shader_core_stats * shader_stats, const memory_config *mem_config,memory_stats_t * memory_stats) { assert( shader_config->m_valid ); assert( mem_config->m_valid ); diff --git a/src/gpgpu-sim/power_stat.h b/src/gpgpu-sim/power_stat.h index 91fade9..24ade99 100644 --- a/src/gpgpu-sim/power_stat.h +++ b/src/gpgpu-sim/power_stat.h @@ -113,7 +113,7 @@ struct mem_power_stats_pod{ class power_mem_stat_t : public mem_power_stats_pod{ public: - power_mem_stat_t(const struct memory_config *mem_config, const shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats); + power_mem_stat_t(const memory_config *mem_config, const shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats); void visualizer_print( gzFile visualizer_file ); void print (FILE *fout) const; void init(); @@ -128,7 +128,7 @@ private: class power_stat_t { public: - power_stat_t( const shader_core_config *shader_config,float * average_pipeline_duty_cycle,float * active_sms,shader_core_stats * shader_stats, const struct memory_config *mem_config,memory_stats_t * memory_stats); + power_stat_t( const shader_core_config *shader_config,float * average_pipeline_duty_cycle,float * active_sms,shader_core_stats * shader_stats, const memory_config *mem_config,memory_stats_t * memory_stats); void visualizer_print( gzFile visualizer_file ); void print (FILE *fout) const; void save_stats(){ @@ -621,7 +621,7 @@ public: float * m_average_pipeline_duty_cycle; float * m_active_sms; const shader_core_config *m_config; - const struct memory_config *m_mem_config; + const memory_config *m_mem_config; }; diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index f380560..b7ae95d 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -28,7 +28,6 @@ #include #include "shader.h" -#include "gpu-sim.h" #include "addrdec.h" #include "dram.h" #include "stat-tool.h" @@ -53,6 +52,19 @@ #define MIN(a,b) (((a)<(b))?(a):(b)) +mem_fetch *shader_core_mem_fetch_allocator::alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const +{ + mem_access_t access( type, addr, size, wr, m_memory_config->gpgpu_ctx); + mem_fetch *mf = new mem_fetch( access, + NULL, + wr?WRITE_PACKET_SIZE:READ_PACKET_SIZE, + -1, + m_core_id, + m_cluster_id, + m_memory_config, + cycle); + return mf; +} ///////////////////////////////////////////////////////////////////////////// std::list shader_core_ctx::get_regs_written( const inst_t &fvt ) const @@ -71,7 +83,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, unsigned shader_id, unsigned tpc_id, const shader_core_config *config, - const struct memory_config *mem_config, + const memory_config *mem_config, 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 ), @@ -809,7 +821,7 @@ void shader_core_ctx::fetch() // TODO: replace with use of allocator // mem_fetch *mf = m_mem_fetch_allocator->alloc() - mem_access_t acc(INST_ACC_R,ppc,nbytes,false); + mem_access_t acc(INST_ACC_R,ppc,nbytes,false, m_gpu->gpgpu_ctx); mem_fetch *mf = new mem_fetch(acc, NULL/*we don't have an instruction yet*/, READ_PACKET_SIZE, @@ -3787,7 +3799,7 @@ void opndcoll_rfu_t::collector_unit_t::dispatch() simt_core_cluster::simt_core_cluster( class gpgpu_sim *gpu, unsigned cluster_id, const shader_core_config *config, - const struct memory_config *mem_config, + const memory_config *mem_config, shader_core_stats *stats, class memory_stats_t *mstats ) { diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index 2837f1b..b0d7f7f 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1727,6 +1727,7 @@ private: friend class LooseRoundRobbinScheduler; }; +class memory_config; class shader_core_mem_fetch_allocator : public mem_fetch_allocator { public: shader_core_mem_fetch_allocator( unsigned core_id, unsigned cluster_id, const memory_config *config ) @@ -1735,20 +1736,7 @@ public: m_cluster_id = cluster_id; m_memory_config = config; } - mem_fetch *alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const - { - mem_access_t access( type, addr, size, wr ); - mem_fetch *mf = new mem_fetch( access, - NULL, - wr?WRITE_PACKET_SIZE:READ_PACKET_SIZE, - -1, - m_core_id, - m_cluster_id, - m_memory_config, - cycle); - return mf; - } - + mem_fetch *alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const; mem_fetch *alloc( const warp_inst_t &inst, const mem_access_t &access, unsigned long long cycle ) const { warp_inst_t inst_copy = inst; @@ -1777,7 +1765,7 @@ public: unsigned shader_id, unsigned tpc_id, const shader_core_config *config, - const struct memory_config *mem_config, + const memory_config *mem_config, shader_core_stats *stats ); // used by simt_core_cluster: @@ -2072,7 +2060,7 @@ public: simt_core_cluster( class gpgpu_sim *gpu, unsigned cluster_id, const shader_core_config *config, - const struct memory_config *mem_config, + const memory_config *mem_config, shader_core_stats *stats, memory_stats_t *mstats ); -- cgit v1.3 From c4782ba3c78e6e97db9f56780991dc5cf827c755 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Sun, 14 Jul 2019 13:06:46 -0400 Subject: Move warp_inst_sm_next_uid Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 2 ++ src/abstract_hardware_model.cc | 15 +++++++++++++-- src/abstract_hardware_model.h | 18 +++--------------- 3 files changed, 18 insertions(+), 17 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 07473be..0dd6311 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -12,6 +12,7 @@ class gpgpu_context { gpgpu_context() { g_global_allfiles_symbol_table = NULL; sm_next_access_uid=0; + warp_inst_sm_next_uid=0; api = new cuda_runtime_api(this); ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); @@ -23,6 +24,7 @@ class gpgpu_context { symbol_table *g_global_allfiles_symbol_table; const char *g_filename; unsigned sm_next_access_uid; + unsigned warp_inst_sm_next_uid; // objects pointers for each file cuda_runtime_api* api; ptxinfo_data* ptxinfo; diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index fe10daa..1dddecd 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -41,8 +41,6 @@ #include #include "../libcuda/gpgpu_context.h" -unsigned warp_inst_t::sm_next_uid = 0; - void mem_access_t::init(gpgpu_context* ctx) { gpgpu_ctx = ctx; @@ -50,6 +48,19 @@ void mem_access_t::init(gpgpu_context* ctx) m_addr=0; m_req_size=0; } +void warp_inst_t::issue( const active_mask_t &mask, unsigned warp_id, unsigned long long cycle, int dynamic_warp_id, int sch_id ) +{ + m_warp_active_mask = mask; + m_warp_issued_mask = mask; + m_uid = ++(m_config->gpgpu_ctx->warp_inst_sm_next_uid); + m_warp_id = warp_id; + m_dynamic_warp_id = dynamic_warp_id; + issue_cycle = cycle; + cycles = initiation_interval; + m_cache_hit=false; + m_empty=false; + m_scheduler_id=sch_id; +} checkpoint::checkpoint() { diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index a22c8c3..519416d 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -958,19 +958,9 @@ public: { m_empty=true; } - void issue( const active_mask_t &mask, unsigned warp_id, unsigned long long cycle, int dynamic_warp_id, int sch_id ) - { - m_warp_active_mask = mask; - m_warp_issued_mask = mask; - m_uid = ++sm_next_uid; - m_warp_id = warp_id; - m_dynamic_warp_id = dynamic_warp_id; - issue_cycle = cycle; - cycles = initiation_interval; - m_cache_hit=false; - m_empty=false; - m_scheduler_id=sch_id; - } + + void issue( const active_mask_t &mask, unsigned warp_id, unsigned long long cycle, int dynamic_warp_id, int sch_id ); + const active_mask_t & get_active_mask() const { return m_warp_active_mask; @@ -1133,8 +1123,6 @@ protected: bool m_mem_accesses_created; std::list m_accessq; - static unsigned sm_next_uid; - unsigned m_scheduler_id; //the scheduler that issues this inst //Jin: cdp support -- cgit v1.3 From 67950eb12eec495b1976934acc763db481d955f8 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 15 Jul 2019 11:58:10 -0400 Subject: Move operand_info::sm_next_uid Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 2 ++ src/cuda-sim/ptx_ir.cc | 20 ++++++++++------- src/cuda-sim/ptx_ir.h | 54 ++++++++++++++++++++++++---------------------- src/cuda-sim/ptx_parser.cc | 30 +++++++++++++------------- src/cuda-sim/ptx_parser.h | 2 +- src/cuda-sim/ptx_sim.cc | 8 +++---- 6 files changed, 62 insertions(+), 54 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 0dd6311..9c5ae7b 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -13,6 +13,7 @@ class gpgpu_context { g_global_allfiles_symbol_table = NULL; sm_next_access_uid=0; warp_inst_sm_next_uid=0; + operand_info_sm_next_uid = 1; api = new cuda_runtime_api(this); ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); @@ -25,6 +26,7 @@ class gpgpu_context { const char *g_filename; unsigned sm_next_access_uid; unsigned warp_inst_sm_next_uid; + unsigned operand_info_sm_next_uid;//uid for operand_info // objects pointers for each file cuda_runtime_api* api; ptxinfo_data* ptxinfo; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index c537091..0fa29eb 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -39,6 +39,7 @@ typedef void * yyscan_t; #include "assert.h" #include "cuda-sim.h" +#include "../../libcuda/gpgpu_context.h" #define STR_SIZE 1024 @@ -322,11 +323,9 @@ void symbol_table::dump() printf("\n"); } -unsigned operand_info::sm_next_uid=1; - unsigned operand_info::get_uid() { - unsigned result = sm_next_uid++; + unsigned result = (gpgpu_ctx->operand_info_sm_next_uid)++; return result; } @@ -1015,7 +1014,7 @@ void copy_buffer_to_frame(ptx_thread_info * thread, const arg_buffer_t &a) { if( a.is_reg() ) { ptx_reg_t value = a.get_reg(); - operand_info dst_reg = operand_info(a.get_dst()); + operand_info dst_reg = operand_info(a.get_dst(), thread->get_gpu()->gpgpu_ctx); thread->set_reg(dst_reg.get_symbol(),value); } else { const void *buffer = a.get_param_buffer(); @@ -1039,7 +1038,8 @@ void copy_buffer_list_into_frame(ptx_thread_info * thread, arg_buffer_list_t &ar static std::list check_operands( int opcode, const std::list &scalar_type, - const std::list &operands ) + const std::list &operands, + gpgpu_context* ctx) { static int g_warn_literal_operands_two_type_inst; if( (opcode == CVT_OP) || (opcode == SET_OP) || (opcode == SLCT_OP) || (opcode == TEX_OP) || (opcode==MMA_OP) || (opcode == DP4A_OP)) { @@ -1066,7 +1066,7 @@ static std::list check_operands( int opcode, if( (op.get_type() == double_op_t) && (inst_type == F32_TYPE) ) { ptx_reg_t v = op.get_literal_value(); float u = (float)v.f64; - operand_info n(u); + operand_info n(u, ctx); result.push_back(n); } else { result.push_back(op); @@ -1097,7 +1097,7 @@ ptx_instruction::ptx_instruction( int opcode, unsigned line, const char *source, const core_config *config, - gpgpu_context* ctx ) : warp_inst_t(config) + gpgpu_context* ctx ) : warp_inst_t(config), m_return_var(ctx) { gpgpu_ctx = ctx; m_uid = ++g_num_ptx_inst_uid; @@ -1107,7 +1107,7 @@ ptx_instruction::ptx_instruction( int opcode, m_neg_pred = neg_pred; m_pred_mod = pred_mod; m_label = label; - const std::list checked_operands = check_operands(opcode,scalar_type,operands); + const std::list checked_operands = check_operands(opcode,scalar_type,operands, ctx); m_operands.insert(m_operands.begin(), checked_operands.begin(), checked_operands.end() ); m_return_var = return_var; m_options = options; @@ -1371,6 +1371,10 @@ std::string ptx_instruction::to_string() const m_source.c_str() ); return std::string( buf ); } +operand_info ptx_instruction::get_pred() const +{ + return operand_info( m_pred, gpgpu_ctx); +} unsigned function_info::sm_next_uid = 1; diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index babd54b..8ed94d9 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -377,9 +377,9 @@ private: class operand_info { public: - operand_info() + operand_info(gpgpu_context* ctx) { - init(); + init(ctx); m_is_non_arch_reg = false; m_addr_space = undefined_space; m_operand_lohi = 0; @@ -392,9 +392,9 @@ public: m_addr_offset = 0; m_value.m_symbolic=NULL; } - operand_info( const symbol *addr ) + operand_info( const symbol *addr, gpgpu_context* ctx ) { - init(); + init(ctx); m_is_non_arch_reg = false; m_addr_space = undefined_space; m_operand_lohi = 0; @@ -435,9 +435,9 @@ public: m_is_return_var = false; m_immediate_address=false; } - operand_info( const symbol *addr1, const symbol *addr2 ) + operand_info( const symbol *addr1, const symbol *addr2, gpgpu_context* ctx ) { - init(); + init(ctx); m_is_non_arch_reg = false; m_addr_space = undefined_space; m_operand_lohi = 0; @@ -462,9 +462,9 @@ public: m_is_return_var = false; m_immediate_address=false; } - operand_info( int builtin_id, int dim_mod ) + operand_info( int builtin_id, int dim_mod, gpgpu_context* ctx ) { - init(); + init(ctx); m_is_non_arch_reg = false; m_addr_space = undefined_space; m_operand_lohi = 0; @@ -481,9 +481,9 @@ public: m_is_return_var = false; m_immediate_address=false; } - operand_info( const symbol *addr, int offset ) + operand_info( const symbol *addr, int offset, gpgpu_context* ctx ) { - init(); + init(ctx); m_is_non_arch_reg = false; m_addr_space = undefined_space; m_operand_lohi = 0; @@ -500,9 +500,9 @@ public: m_is_return_var = false; m_immediate_address=false; } - operand_info( unsigned x ) + operand_info( unsigned x, gpgpu_context* ctx ) { - init(); + init(ctx); m_is_non_arch_reg = false; m_addr_space = undefined_space; m_operand_lohi = 0; @@ -519,9 +519,9 @@ public: m_is_return_var = false; m_immediate_address=true; } - operand_info( int x ) + operand_info( int x, gpgpu_context* ctx ) { - init(); + init(ctx); m_is_non_arch_reg = false; m_addr_space = undefined_space; m_operand_lohi = 0; @@ -538,9 +538,9 @@ public: m_is_return_var = false; m_immediate_address=false; } - operand_info( float x ) + operand_info( float x, gpgpu_context* ctx ) { - init(); + init(ctx); m_is_non_arch_reg = false; m_addr_space = undefined_space; m_operand_lohi = 0; @@ -557,9 +557,9 @@ public: m_is_return_var = false; m_immediate_address=false; } - operand_info( double x ) + operand_info( double x, gpgpu_context* ctx ) { - init(); + init(ctx); m_is_non_arch_reg = false; m_addr_space = undefined_space; m_operand_lohi = 0; @@ -576,9 +576,9 @@ public: m_is_return_var = false; m_immediate_address=false; } - operand_info( const symbol *s1, const symbol *s2, const symbol *s3, const symbol *s4 ) + operand_info( const symbol *s1, const symbol *s2, const symbol *s3, const symbol *s4, gpgpu_context* ctx ) { - init(); + init(ctx); m_is_non_arch_reg = false; m_addr_space = undefined_space; m_operand_lohi = 0; @@ -603,9 +603,9 @@ public: m_is_return_var = false; m_immediate_address=false; } - operand_info( const symbol *s1, const symbol *s2, const symbol *s3, const symbol *s4 ,const symbol *s5,const symbol *s6,const symbol *s7, const symbol *s8) + operand_info( const symbol *s1, const symbol *s2, const symbol *s3, const symbol *s4 ,const symbol *s5,const symbol *s6,const symbol *s7, const symbol *s8, gpgpu_context* ctx) { - init(); + init(ctx); m_is_non_arch_reg = false; m_addr_space = undefined_space; m_operand_lohi = 0; @@ -631,8 +631,9 @@ public: m_immediate_address=false; } - void init() + void init(gpgpu_context* ctx) { + gpgpu_ctx = ctx; m_uid=(unsigned)-1; m_valid=false; m_vector=false; @@ -842,6 +843,7 @@ public: bool is_non_arch_reg() const { return m_is_non_arch_reg; } private: + gpgpu_context* gpgpu_ctx; unsigned m_uid; bool m_valid; bool m_vector; @@ -957,7 +959,7 @@ public: unsigned source_line() const { return m_source_line;} unsigned get_num_operands() const { return m_operands.size();} bool has_pred() const { return m_pred != NULL;} - operand_info get_pred() const { return operand_info( m_pred );} + operand_info get_pred() const; bool get_pred_neg() const { return m_neg_pred;} int get_pred_mod() const { return m_pred_mod;} const char *get_source() const { return m_source.c_str();} @@ -1448,14 +1450,14 @@ private: class arg_buffer_t { public: - arg_buffer_t() + arg_buffer_t(gpgpu_context* ctx) : m_src_op(ctx) { m_is_reg=false; m_is_param=false; m_param_value=NULL; m_reg_value=ptx_reg_t(); } - arg_buffer_t( const arg_buffer_t &another ) + arg_buffer_t( const arg_buffer_t &another, gpgpu_context* ctx ) : m_src_op(ctx) { make_copy(another); } diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 1d38edf..81b70af 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -102,7 +102,7 @@ void ptx_recognizer::init_instruction_state() g_opcode = -1; g_options.clear(); g_wmma_options.clear(); - g_return_var = operand_info(); + g_return_var = operand_info(gpgpu_ctx); init_directive_state(); } @@ -689,7 +689,7 @@ void ptx_recognizer::add_double_operand( const char *d1, const char *d2 ) const symbol *s1 = g_current_symbol_table->lookup(d1); const symbol *s2 = g_current_symbol_table->lookup(d2); parse_assert( s1 != NULL && s2 != NULL, "component(s) missing declarations."); - g_operands.push_back( operand_info(s1,s2) ); + g_operands.push_back( operand_info(s1,s2,gpgpu_ctx) ); } void ptx_recognizer::add_1vector_operand( const char *d1 ) @@ -698,7 +698,7 @@ void ptx_recognizer::add_1vector_operand( const char *d1 ) PTX_PARSE_DPRINTF("add_1vector_operand"); const symbol *s1 = g_current_symbol_table->lookup(d1); parse_assert( s1 != NULL, "component(s) missing declarations."); - g_operands.push_back( operand_info(s1,NULL,NULL,NULL) ); + g_operands.push_back( operand_info(s1,NULL,NULL,NULL,gpgpu_ctx) ); } void ptx_recognizer::add_2vector_operand( const char *d1, const char *d2 ) @@ -707,7 +707,7 @@ void ptx_recognizer::add_2vector_operand( const char *d1, const char *d2 ) const symbol *s1 = g_current_symbol_table->lookup(d1); const symbol *s2 = g_current_symbol_table->lookup(d2); parse_assert( s1 != NULL && s2 != NULL, "v2 component(s) missing declarations."); - g_operands.push_back( operand_info(s1,s2,NULL,NULL) ); + g_operands.push_back( operand_info(s1,s2,NULL,NULL,gpgpu_ctx) ); } void ptx_recognizer::add_3vector_operand( const char *d1, const char *d2, const char *d3 ) @@ -717,7 +717,7 @@ void ptx_recognizer::add_3vector_operand( const char *d1, const char *d2, const const symbol *s2 = g_current_symbol_table->lookup(d2); const symbol *s3 = g_current_symbol_table->lookup(d3); parse_assert( s1 != NULL && s2 != NULL && s3 != NULL, "v3 component(s) missing declarations."); - g_operands.push_back( operand_info(s1,s2,s3,NULL) ); + g_operands.push_back( operand_info(s1,s2,s3,NULL,gpgpu_ctx) ); } void ptx_recognizer::add_4vector_operand( const char *d1, const char *d2, const char *d3, const char *d4 ) @@ -732,7 +732,7 @@ void ptx_recognizer::add_4vector_operand( const char *d1, const char *d2, const if ( s2 == null_op ) s2 = NULL; if ( s3 == null_op ) s3 = NULL; if ( s4 == null_op ) s4 = NULL; - g_operands.push_back( operand_info(s1,s2,s3,s4) ); + g_operands.push_back( operand_info(s1,s2,s3,s4,gpgpu_ctx) ); } void ptx_recognizer::add_8vector_operand( const char *d1, const char *d2, const char *d3, const char *d4,const char *d5,const char *d6,const char *d7,const char *d8 ) { @@ -754,13 +754,13 @@ void ptx_recognizer::add_8vector_operand( const char *d1, const char *d2, const if ( s6 == null_op ) s6 = NULL; if ( s7 == null_op ) s7 = NULL; if ( s8 == null_op ) s8 = NULL; - g_operands.push_back( operand_info(s1,s2,s3,s4,s5,s6,s7,s8) ); + g_operands.push_back( operand_info(s1,s2,s3,s4,s5,s6,s7,s8,gpgpu_ctx) ); } void ptx_recognizer::add_builtin_operand( int builtin, int dim_modifier ) { PTX_PARSE_DPRINTF("add_builtin_operand"); - g_operands.push_back( operand_info(builtin,dim_modifier) ); + g_operands.push_back( operand_info(builtin,dim_modifier,gpgpu_ctx) ); } void ptx_recognizer::add_memory_operand() @@ -883,19 +883,19 @@ void ptx_recognizer::change_operand_neg( ) void ptx_recognizer::add_literal_int( int value ) { PTX_PARSE_DPRINTF("add_literal_int"); - g_operands.push_back( operand_info(value) ); + g_operands.push_back( operand_info(value,gpgpu_ctx) ); } void ptx_recognizer::add_literal_float( float value ) { PTX_PARSE_DPRINTF("add_literal_float"); - g_operands.push_back( operand_info(value) ); + g_operands.push_back( operand_info(value,gpgpu_ctx) ); } void ptx_recognizer::add_literal_double( double value ) { PTX_PARSE_DPRINTF("add_literal_double"); - g_operands.push_back( operand_info(value) ); + g_operands.push_back( operand_info(value,gpgpu_ctx) ); } void ptx_recognizer::add_scalar_operand( const char *identifier ) @@ -911,7 +911,7 @@ void ptx_recognizer::add_scalar_operand( const char *identifier ) parse_error( msg.c_str() ); } } - g_operands.push_back( operand_info(s) ); + g_operands.push_back( operand_info(s,gpgpu_ctx) ); } void ptx_recognizer::add_neg_pred_operand( const char *identifier ) @@ -921,7 +921,7 @@ void ptx_recognizer::add_neg_pred_operand( const char *identifier ) if ( s == NULL ) { s = g_current_symbol_table->add_variable(identifier,NULL,1,gpgpu_ctx->g_filename,ptx_get_lineno(scanner)); } - operand_info op(s); + operand_info op(s, gpgpu_ctx); op.set_neg_pred(); g_operands.push_back( op ); } @@ -934,13 +934,13 @@ void ptx_recognizer::add_address_operand( const char *identifier, int offset ) std::string msg = std::string("operand \"") + identifier + "\" has no declaration."; parse_error( msg.c_str() ); } - g_operands.push_back( operand_info(s,offset) ); + g_operands.push_back( operand_info(s,offset,gpgpu_ctx) ); } void ptx_recognizer::add_address_operand2( int offset ) { PTX_PARSE_DPRINTF("add_address_operand"); - g_operands.push_back( operand_info((unsigned)offset) ); + g_operands.push_back( operand_info((unsigned)offset,gpgpu_ctx) ); } void ptx_recognizer::add_array_initializer() diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h index 1e559a2..11a3d20 100644 --- a/src/cuda-sim/ptx_parser.h +++ b/src/cuda-sim/ptx_parser.h @@ -35,7 +35,7 @@ class gpgpu_context; typedef void * yyscan_t; class ptx_recognizer { public: - ptx_recognizer( gpgpu_context* ctx ) { + ptx_recognizer( gpgpu_context* ctx ) : g_return_var(ctx) { scanner = NULL; g_size = -1; g_add_identifier_cached__identifier = NULL; diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc index b6cf4bd..62bb529 100644 --- a/src/cuda-sim/ptx_sim.cc +++ b/src/cuda-sim/ptx_sim.cc @@ -422,9 +422,9 @@ bool ptx_thread_info::callstack_pop() assert( !((rv_src != NULL) ^ (rv_dst != NULL)) ); // ensure caller and callee agree on whether there is a return value // read return value from callee frame - arg_buffer_t buffer; + arg_buffer_t buffer(m_gpu->gpgpu_ctx); if( rv_src != NULL ) - buffer = copy_arg_to_buffer(this, operand_info(rv_src), rv_dst ); + buffer = copy_arg_to_buffer(this, operand_info(rv_src, m_gpu->gpgpu_ctx), rv_dst ); m_symbol_table = m_callstack.back().m_symbol_table; m_NPC = m_callstack.back().m_PC; @@ -456,9 +456,9 @@ bool ptx_thread_info::callstack_pop_plus() assert( !((rv_src != NULL) ^ (rv_dst != NULL)) ); // ensure caller and callee agree on whether there is a return value // read return value from callee frame - arg_buffer_t buffer; + arg_buffer_t buffer(m_gpu->gpgpu_ctx); if( rv_src != NULL ) - buffer = copy_arg_to_buffer(this, operand_info(rv_src), rv_dst ); + buffer = copy_arg_to_buffer(this, operand_info(rv_src, m_gpu->gpgpu_ctx), rv_dst ); m_symbol_table = m_callstack.back().m_symbol_table; m_NPC = m_callstack.back().m_PC; -- cgit v1.3 From bd170826a00fe1ea1960ff2e23247ed3f980547b Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 15 Jul 2019 12:16:01 -0400 Subject: Move kernel_info_t::m_next_uid Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 2 ++ src/abstract_hardware_model.cc | 3 +-- src/abstract_hardware_model.h | 1 - 3 files changed, 3 insertions(+), 3 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 9c5ae7b..1e20c62 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -14,6 +14,7 @@ class gpgpu_context { sm_next_access_uid=0; warp_inst_sm_next_uid=0; operand_info_sm_next_uid = 1; + kernel_info_m_next_uid = 1; api = new cuda_runtime_api(this); ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); @@ -27,6 +28,7 @@ class gpgpu_context { unsigned sm_next_access_uid; unsigned warp_inst_sm_next_uid; unsigned operand_info_sm_next_uid;//uid for operand_info + unsigned kernel_info_m_next_uid;//uid for kernel_info_t // objects pointers for each file cuda_runtime_api* api; ptxinfo_data* ptxinfo; diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 1dddecd..733d602 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -709,7 +709,6 @@ void warp_inst_t::completed( unsigned long long cycle ) const ptx_file_line_stats_add_latency(pc, latency * active_count()); } -unsigned kernel_info_t::m_next_uid = 1; /*A snapshot of the texture mappings needs to be stored in the kernel's info as kernels should use the texture bindings seen at the time of launch and textures @@ -724,7 +723,7 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * m_next_cta.z=0; m_next_tid=m_next_cta; m_num_cores_running=0; - m_uid = m_next_uid++; + m_uid = (entry->gpgpu_ctx->kernel_info_m_next_uid)++; m_param_mem = new memory_space_impl<8192>("param",64*1024); //Jin: parent and child kernel management for CDP diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 519416d..d13b8c6 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -299,7 +299,6 @@ private: class function_info *m_kernel_entry; unsigned m_uid; - static unsigned m_next_uid; //These maps contain the snapshot of the texture mappings at kernel launch std::map m_NameToCudaArray; -- cgit v1.3 From 36aaec0a30445533f04aaf6b55c5ef135e20484a Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 15 Jul 2019 13:31:36 -0400 Subject: Move ptx_line_stats_filename and enable_ptx_file_line_stats Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 3 +++ src/cuda-sim/ptx-stats.cc | 9 +++------ src/cuda-sim/ptx-stats.h | 23 ++++++++++++++++------- src/gpgpu-sim/gpu-sim.cc | 4 ++-- 4 files changed, 24 insertions(+), 15 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 1e20c62..337ebb2 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -6,6 +6,7 @@ #include "../src/gpgpusim_entrypoint.h" #include "../src/cuda-sim/cuda-sim.h" #include "../src/cuda-sim/cuda_device_runtime.h" +#include "../src/cuda-sim/ptx-stats.h" class gpgpu_context { public: @@ -21,6 +22,7 @@ class gpgpu_context { the_gpgpusim = new GPGPUsim_ctx(this); func_sim = new cuda_sim(this); device_runtime = new cuda_device_runtime(this); + stats = new ptx_stats(this); } // global list symbol_table *g_global_allfiles_symbol_table; @@ -36,6 +38,7 @@ class gpgpu_context { GPGPUsim_ctx* the_gpgpusim; cuda_sim* func_sim; cuda_device_runtime* device_runtime; + ptx_stats* stats; // 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/src/cuda-sim/ptx-stats.cc b/src/cuda-sim/ptx-stats.cc index 0a9f08d..298729f 100644 --- a/src/cuda-sim/ptx-stats.cc +++ b/src/cuda-sim/ptx-stats.cc @@ -32,12 +32,9 @@ #include #include #include "../tr1_hash_map.h" +#include "../../libcuda/gpgpu_context.h" -// options -bool enable_ptx_file_line_stats; -char * ptx_line_stats_filename = NULL; - -void ptx_file_line_stats_options(option_parser_t opp) +void ptx_stats::ptx_file_line_stats_options(option_parser_t opp) { option_parser_register(opp, "-enable_ptx_file_line_stats", OPT_BOOL, &enable_ptx_file_line_stats, @@ -118,7 +115,7 @@ typedef tr1_hash_map ptx static ptx_file_line_stats_map_t ptx_file_line_stats_tracker; // output statistics to a file -void ptx_file_line_stats_write_file() +void ptx_stats::ptx_file_line_stats_write_file() { // check if stat collection is turned on if (enable_ptx_file_line_stats == 0) return; diff --git a/src/cuda-sim/ptx-stats.h b/src/cuda-sim/ptx-stats.h index 4fc6599..c75fc58 100644 --- a/src/cuda-sim/ptx-stats.h +++ b/src/cuda-sim/ptx-stats.h @@ -29,13 +29,6 @@ #include "../option_parser.h" -extern bool enable_ptx_file_line_stats; - -// set options -void ptx_file_line_stats_options(option_parser_t opp); - -// output stats to a file -void ptx_file_line_stats_write_file(); #ifdef __cplusplus // stat collection interface to cuda-sim @@ -56,3 +49,19 @@ void ptx_file_line_stats_commit_exposed_latency(int sc_id, int exposed_latency); void ptx_file_line_stats_add_warp_divergence(unsigned pc, unsigned n_way_divergence); +class gpgpu_context; +class ptx_stats { + public: + ptx_stats(gpgpu_context* ctx) { + ptx_line_stats_filename = NULL; + gpgpu_ctx = ctx; + } + char * ptx_line_stats_filename; + bool enable_ptx_file_line_stats; + gpgpu_context* gpgpu_ctx; + // set options + void ptx_file_line_stats_options(option_parser_t opp); + + // output stats to a file + void ptx_file_line_stats_write_file(); +}; diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 0644b44..bdf989a 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -545,7 +545,7 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) option_parser_register(opp, "-trace_sampling_memory_partition", OPT_INT32, &Trace::sampling_memory_partition, "The memory partition which is printed using MEMPART_DPRINTF. Default -1 (i.e. all)", "-1"); - ptx_file_line_stats_options(opp); + gpgpu_ctx->stats->ptx_file_line_stats_options(opp); //Jin: kernel launch latency option_parser_register(opp, "-gpgpu_kernel_launch_latency", OPT_INT32, @@ -932,7 +932,7 @@ void gpgpu_sim::update_stats() { void gpgpu_sim::print_stats() { - ptx_file_line_stats_write_file(); + gpgpu_ctx->stats->ptx_file_line_stats_write_file(); gpu_print_stat(); if (g_network_mode) { -- cgit v1.3 From 3b88f9db669d8ec9ffc477352dcfd2a3d423781f Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 15 Jul 2019 13:38:38 -0400 Subject: Move g_num_ptx_inst_uid Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 2 ++ src/cuda-sim/instructions.cc | 1 - src/cuda-sim/ptx_ir.cc | 2 +- src/cuda-sim/ptx_ir.h | 2 -- 4 files changed, 3 insertions(+), 4 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 337ebb2..909f267 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -16,6 +16,7 @@ class gpgpu_context { warp_inst_sm_next_uid=0; operand_info_sm_next_uid = 1; kernel_info_m_next_uid = 1; + g_num_ptx_inst_uid = 0; api = new cuda_runtime_api(this); ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); @@ -31,6 +32,7 @@ class gpgpu_context { unsigned warp_inst_sm_next_uid; unsigned operand_info_sm_next_uid;//uid for operand_info unsigned kernel_info_m_next_uid;//uid for kernel_info_t + unsigned g_num_ptx_inst_uid; //uid for ptx inst inside ptx_instruction // objects pointers for each file cuda_runtime_api* api; ptxinfo_data* ptxinfo; diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 565340c..4d4a80d 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -60,7 +60,6 @@ class ptx_recognizer; using half_float::half; -unsigned ptx_instruction::g_num_ptx_inst_uid=0; bool debug_tensorcore = 0; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 0fa29eb..849cc5d 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1100,7 +1100,7 @@ ptx_instruction::ptx_instruction( int opcode, gpgpu_context* ctx ) : warp_inst_t(config), m_return_var(ctx) { gpgpu_ctx = ctx; - m_uid = ++g_num_ptx_inst_uid; + m_uid = ++(ctx->g_num_ptx_inst_uid); m_PC = 0; m_opcode = opcode; m_pred = pred; diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 8ed94d9..1ffbebc 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -878,7 +878,6 @@ private: }; extern const char *g_opcode_string[]; -extern unsigned g_num_ptx_inst_uid; struct basic_block_t { basic_block_t( unsigned ID, ptx_instruction *begin, ptx_instruction *end, bool entry, bool ex) { @@ -1194,7 +1193,6 @@ private: virtual void pre_decode(); friend class function_info; - static unsigned g_num_ptx_inst_uid; // backward pointer class gpgpu_context* gpgpu_ctx; }; -- cgit v1.3 From 971722f51189e80034a9c80a4846a4ec045f0d4d Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 15 Jul 2019 13:57:48 -0400 Subject: Move g_ptx_cta_info_uid and symbol::sm_next_uid Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 4 ++++ src/cuda-sim/ptx_ir.cc | 8 +++----- src/cuda-sim/ptx_ir.h | 6 +++--- src/cuda-sim/ptx_sim.cc | 3 +-- 4 files changed, 11 insertions(+), 10 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 909f267..e168850 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -17,6 +17,8 @@ class gpgpu_context { operand_info_sm_next_uid = 1; kernel_info_m_next_uid = 1; g_num_ptx_inst_uid = 0; + g_ptx_cta_info_uid = 1; + symbol_sm_next_uid = 1; api = new cuda_runtime_api(this); ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); @@ -33,6 +35,8 @@ class gpgpu_context { unsigned operand_info_sm_next_uid;//uid for operand_info unsigned kernel_info_m_next_uid;//uid for kernel_info_t unsigned g_num_ptx_inst_uid; //uid for ptx inst inside ptx_instruction + unsigned long long g_ptx_cta_info_uid; + unsigned symbol_sm_next_uid; //uid for symbol // objects pointers for each file cuda_runtime_api* api; ptxinfo_data* ptxinfo; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 849cc5d..3719fb2 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -43,11 +43,9 @@ typedef void * yyscan_t; #define STR_SIZE 1024 -unsigned symbol::sm_next_uid = 1; - unsigned symbol::get_uid() { - unsigned result = sm_next_uid++; + unsigned result = (gpgpu_ctx->symbol_sm_next_uid)++; return result; } @@ -152,7 +150,7 @@ symbol *symbol_table::add_variable( const char *identifier, const type_info *typ std::string key(identifier); assert( m_symbols.find(key) == m_symbols.end() ); snprintf(buf,1024,"%s:%u",filename,line); - symbol *s = new symbol(identifier,type,buf,size); + symbol *s = new symbol(identifier,type,buf,size,gpgpu_ctx); m_symbols[ key ] = s; if ( type != NULL && type->get_key().is_global() ) { @@ -173,7 +171,7 @@ void symbol_table::add_function( function_info *func, const char *filename, unsi char buf[1024]; snprintf(buf,1024,"%s:%u",filename,linenumber); type_info *type = add_type( func ); - symbol *s = new symbol(func->get_name().c_str(),type,buf,0); + symbol *s = new symbol(func->get_name().c_str(),type,buf,0,gpgpu_ctx); s->set_function(func); m_symbols[ func->get_name() ] = s; } diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 1ffbebc..2fbda1c 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -152,8 +152,9 @@ class operand_info; class symbol { public: - symbol( const char *name, const type_info *type, const char *location, unsigned size ) + symbol( const char *name, const type_info *type, const char *location, unsigned size, gpgpu_context* ctx ) { + gpgpu_ctx = ctx; m_uid = get_uid(); m_name = name; m_decl_location = location; @@ -273,6 +274,7 @@ public: unsigned uid() const { return m_uid; } private: + gpgpu_context* gpgpu_ctx; unsigned get_uid(); unsigned m_uid; const type_info *m_type; @@ -299,8 +301,6 @@ private: bool m_reg_num_valid; std::list m_initializer; - static unsigned sm_next_uid; - }; class symbol_table { diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc index 62bb529..949ee66 100644 --- a/src/cuda-sim/ptx_sim.cc +++ b/src/cuda-sim/ptx_sim.cc @@ -37,7 +37,6 @@ typedef void * yyscan_t; void feature_not_implemented( const char *f ); -unsigned long long g_ptx_cta_info_uid = 1; ptx_cta_info::ptx_cta_info( unsigned sm_idx, gpgpu_context* ctx ) { @@ -45,7 +44,7 @@ ptx_cta_info::ptx_cta_info( unsigned sm_idx, gpgpu_context* ctx ) ctx->func_sim->g_ptx_cta_info_sm_idx_used.insert(sm_idx); m_sm_idx = sm_idx; - m_uid = g_ptx_cta_info_uid++; + m_uid = (ctx->g_ptx_cta_info_uid)++; m_bar_threads = 0; gpgpu_ctx = ctx; } -- cgit v1.3 From 6ec4563b11fe99e539eb83412acaadc6f67b05ba Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 15 Jul 2019 14:04:43 -0400 Subject: Move function_info::sm_next_id Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 2 ++ src/cuda-sim/ptx_ir.cc | 3 +-- src/cuda-sim/ptx_ir.h | 2 -- 3 files changed, 3 insertions(+), 4 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index e168850..ed4f746 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -19,6 +19,7 @@ class gpgpu_context { g_num_ptx_inst_uid = 0; g_ptx_cta_info_uid = 1; symbol_sm_next_uid = 1; + function_info_sm_next_uid = 1; api = new cuda_runtime_api(this); ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); @@ -37,6 +38,7 @@ class gpgpu_context { unsigned g_num_ptx_inst_uid; //uid for ptx inst inside ptx_instruction unsigned long long g_ptx_cta_info_uid; unsigned symbol_sm_next_uid; //uid for symbol + unsigned function_info_sm_next_uid; // objects pointers for each file cuda_runtime_api* api; ptxinfo_data* ptxinfo; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 3719fb2..5fa9379 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1374,12 +1374,11 @@ operand_info ptx_instruction::get_pred() const return operand_info( m_pred, gpgpu_ctx); } -unsigned function_info::sm_next_uid = 1; function_info::function_info(int entry_point, gpgpu_context* ctx ) { gpgpu_ctx = ctx; - m_uid = sm_next_uid++; + m_uid = (gpgpu_ctx->function_info_sm_next_uid)++; m_entry_point = (entry_point==1)?true:false; m_extern = (entry_point==2)?true:false; num_reconvergence_pairs = 0; diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 2fbda1c..8fc0a06 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -873,7 +873,6 @@ private: bool m_is_return_var; bool m_is_non_arch_reg; - static unsigned sm_next_uid; unsigned get_uid(); }; @@ -1438,7 +1437,6 @@ private: symbol_table *m_symtab; static std::vector s_g_pc_to_insn; // a direct mapping from PC to instruction - static unsigned sm_next_uid; //parameter size for device kernels int m_args_aligned_size; -- cgit v1.3 From 352d2a3336b1c8e5258ca9d92d214973e98837c0 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 15 Jul 2019 16:17:07 -0400 Subject: Move s_g_pc_to_insn Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 5 +++++ src/abstract_hardware_model.cc | 10 +++++----- src/cuda-sim/cuda-sim.cc | 20 +++++++++----------- src/cuda-sim/cuda-sim.h | 2 -- src/cuda-sim/ptx-stats.cc | 28 ++++++++++++++-------------- src/cuda-sim/ptx-stats.h | 15 ++++++++------- src/cuda-sim/ptx_ir.cc | 8 ++++++++ src/cuda-sim/ptx_ir.h | 9 --------- src/gpgpu-sim/gpu-sim.cc | 2 +- src/gpgpu-sim/mem_latency_stat.cc | 3 ++- src/gpgpu-sim/shader.cc | 4 ++-- src/gpgpu-sim/stat-tool.cc | 33 +++++++++++++++++---------------- src/gpgpu-sim/stat-tool.h | 12 +++++++----- 13 files changed, 78 insertions(+), 73 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index ed4f746..346a8a4 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -39,6 +39,8 @@ class gpgpu_context { unsigned long long g_ptx_cta_info_uid; unsigned symbol_sm_next_uid; //uid for symbol unsigned function_info_sm_next_uid; + std::vector s_g_pc_to_insn; // a direct mapping from PC to instruction + // objects pointers for each file cuda_runtime_api* api; ptxinfo_data* ptxinfo; @@ -58,6 +60,9 @@ class gpgpu_context { class gpgpu_sim *gpgpu_ptx_sim_init_perf(); struct _cuda_device_id *GPGPUSim_Init(); void ptx_reg_options(option_parser_t opp); + const ptx_instruction* pc_to_instruction(unsigned pc); + const warp_inst_t *ptx_fetch_inst( address_type pc ); + unsigned translate_pc_to_ptxlineno(unsigned pc); }; gpgpu_context* GPGPU_Context(); diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 733d602..d8d5fbd 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -426,7 +426,7 @@ void warp_inst_t::generate_mem_accesses() } assert( total_accesses > 0 && total_accesses <= m_config->warp_size ); cycles = total_accesses; // shared memory conflicts modeled as larger initiation interval - ptx_file_line_stats_add_smem_bank_conflict( pc, total_accesses ); + m_config->gpgpu_ctx->stats->ptx_file_line_stats_add_smem_bank_conflict( pc, total_accesses ); break; } @@ -471,7 +471,7 @@ void warp_inst_t::generate_mem_accesses() } if ( space.get_type() == global_space ) { - ptx_file_line_stats_add_uncoalesced_gmem( pc, m_accessq.size() - starting_queue_size ); + m_config->gpgpu_ctx->stats->ptx_file_line_stats_add_uncoalesced_gmem( pc, m_accessq.size() - starting_queue_size ); } m_mem_accesses_created=true; } @@ -706,7 +706,7 @@ void warp_inst_t::completed( unsigned long long cycle ) const { unsigned long long latency = cycle - issue_cycle; assert(latency <= cycle); // underflow detection - ptx_file_line_stats_add_latency(pc, latency * active_count()); + m_config->gpgpu_ctx->stats->ptx_file_line_stats_add_latency(pc, latency * active_count()); } @@ -1110,7 +1110,7 @@ void simt_stack::update( simt_mask_t &thread_done, addr_vector_t &next_pc, addre if (warp_diverged) { - ptx_file_line_stats_add_warp_divergence(top_pc, 1); + m_gpu->gpgpu_ctx->stats->ptx_file_line_stats_add_warp_divergence(top_pc, 1); } } @@ -1157,7 +1157,7 @@ warp_inst_t core_t::getExecuteWarp(unsigned warpId) { unsigned pc,rpc; m_simt_stack[warpId]->get_pdom_stack_top_info(&pc,&rpc); - warp_inst_t wi= *ptx_fetch_inst(pc); + warp_inst_t wi= *(m_gpu->gpgpu_ctx->ptx_fetch_inst(pc)); wi.set_active(m_simt_stack[warpId]->get_active_mask()); return wi; } diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index c06f093..b9e6552 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -215,8 +215,6 @@ void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) m_NameToTextureInfo.erase(texname); } -std::vector function_info::s_g_pc_to_insn; - #define MAX_INST_SIZE 8 /*bytes*/ void function_info::ptx_assemble() @@ -237,14 +235,14 @@ void function_info::ptx_assemble() addr_t PC = gpgpu_ctx->func_sim->g_assemble_code_next_pc; // globally unique address (across functions) // start function on an aligned address for( unsigned i=0; i < (PC%MAX_INST_SIZE); i++ ) - s_g_pc_to_insn.push_back((ptx_instruction*)NULL); + gpgpu_ctx->s_g_pc_to_insn.push_back((ptx_instruction*)NULL); PC += PC%MAX_INST_SIZE; m_start_PC = PC; addr_t n=0; // offset in m_instr_mem //Why s_g_pc_to_insn.size() is needed to reserve additional memory for insts? reserve is cumulative. //s_g_pc_to_insn.reserve(s_g_pc_to_insn.size() + MAX_INST_SIZE*m_instructions.size()); - s_g_pc_to_insn.reserve(MAX_INST_SIZE*m_instructions.size()); + gpgpu_ctx->s_g_pc_to_insn.reserve(MAX_INST_SIZE*m_instructions.size()); for ( i=m_instructions.begin(); i != m_instructions.end(); i++ ) { ptx_instruction *pI = *i; if ( pI->is_label() ) { @@ -253,13 +251,13 @@ void function_info::ptx_assemble() } else { gpgpu_ctx->func_sim->g_pc_to_finfo[PC] = this; m_instr_mem[n] = pI; - s_g_pc_to_insn.push_back(pI); - assert(pI == s_g_pc_to_insn[PC]); + gpgpu_ctx->s_g_pc_to_insn.push_back(pI); + assert(pI == gpgpu_ctx->s_g_pc_to_insn[PC]); pI->set_m_instr_mem_index(n); pI->set_PC(PC); assert( pI->inst_size() <= MAX_INST_SIZE ); for( unsigned i=1; i < pI->inst_size(); i++ ) { - s_g_pc_to_insn.push_back((ptx_instruction*)NULL); + gpgpu_ctx->s_g_pc_to_insn.push_back((ptx_instruction*)NULL); m_instr_mem[n+i]=NULL; } n += pI->inst_size(); @@ -1738,9 +1736,9 @@ const struct gpgpu_ptx_sim_info* ptx_sim_kernel_info(const function_info *kernel return kernel->get_kernel_info(); } -const warp_inst_t *ptx_fetch_inst( address_type pc ) +const warp_inst_t *gpgpu_context::ptx_fetch_inst( address_type pc ) { - return function_info::pc_to_instruction(pc); + return pc_to_instruction(pc); } unsigned ptx_sim_init_thread( kernel_info_t &kernel, @@ -2366,11 +2364,11 @@ void functionalCoreSim::executeWarp(unsigned i, bool &allAtBarrier, bool & someO if(!m_warpAtBarrier[i]&& m_liveThreadCount[i]>0) allAtBarrier = false; } -unsigned translate_pc_to_ptxlineno(unsigned pc) +unsigned gpgpu_context::translate_pc_to_ptxlineno(unsigned pc) { // this function assumes that the kernel fits inside a single PTX file // function_info *pFunc = g_func_info; // assume that the current kernel is the one in query - const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); + const ptx_instruction *pInsn = pc_to_instruction(pc); unsigned ptx_line_number = pInsn->source_line(); return ptx_line_number; diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 5bd4cb2..1be3d19 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -58,10 +58,8 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, unsigned hw_warp_id, gpgpu_t *gpu, bool functionalSimulationMode = false); -const warp_inst_t *ptx_fetch_inst( address_type pc ); const struct gpgpu_ptx_sim_info* ptx_sim_kernel_info(const class function_info *kernel); - /*! * This class functionally executes a kernel. It uses the basic data structures and procedures in core_t */ diff --git a/src/cuda-sim/ptx-stats.cc b/src/cuda-sim/ptx-stats.cc index 298729f..22517df 100644 --- a/src/cuda-sim/ptx-stats.cc +++ b/src/cuda-sim/ptx-stats.cc @@ -151,27 +151,27 @@ void ptx_file_line_stats_add_exec_count(const ptx_instruction *pInsn) // attribute pipeline latency to this ptx instruction (specified by the pc) // pipeline latency is the number of cycles a warp with this instruction spent in the pipeline -void ptx_file_line_stats_add_latency(unsigned pc, unsigned latency) +void ptx_stats::ptx_file_line_stats_add_latency(unsigned pc, unsigned latency) { - const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); + const ptx_instruction *pInsn = gpgpu_ctx->pc_to_instruction(pc); ptx_file_line_stats_tracker[ptx_file_line(pInsn->source_file(), pInsn->source_line())].latency += latency; } // attribute dram traffic to this ptx instruction (specified by the pc) // dram traffic is counted in number of requests -void ptx_file_line_stats_add_dram_traffic(unsigned pc, unsigned dram_traffic) +void ptx_stats::ptx_file_line_stats_add_dram_traffic(unsigned pc, unsigned dram_traffic) { - const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); + const ptx_instruction *pInsn = gpgpu_ctx->pc_to_instruction(pc); ptx_file_line_stats_tracker[ptx_file_line(pInsn->source_file(), pInsn->source_line())].dram_traffic += dram_traffic; } // attribute the number of shared memory access cycles to a ptx instruction // counts both the number of warps doing shared memory access and the number of cycles involved -void ptx_file_line_stats_add_smem_bank_conflict(unsigned pc, unsigned n_way_bkconflict) +void ptx_stats::ptx_file_line_stats_add_smem_bank_conflict(unsigned pc, unsigned n_way_bkconflict) { - const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); + const ptx_instruction *pInsn = gpgpu_ctx->pc_to_instruction(pc); ptx_file_line_stats& line_stats = ptx_file_line_stats_tracker[ptx_file_line(pInsn->source_file(), pInsn->source_line())]; line_stats.smem_n_way_bank_conflict_total += n_way_bkconflict; @@ -180,9 +180,9 @@ void ptx_file_line_stats_add_smem_bank_conflict(unsigned pc, unsigned n_way_bkco // attribute a non-coalesced mem access to a ptx instruction // counts both the number of warps causing this and the number of memory requests generated -void ptx_file_line_stats_add_uncoalesced_gmem(unsigned pc, unsigned n_access) +void ptx_stats::ptx_file_line_stats_add_uncoalesced_gmem(unsigned pc, unsigned n_access) { - const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); + const ptx_instruction *pInsn = gpgpu_ctx->pc_to_instruction(pc); ptx_file_line_stats& line_stats = ptx_file_line_stats_tracker[ptx_file_line(pInsn->source_file(), pInsn->source_line())]; line_stats.gmem_n_access_total += n_access; @@ -239,17 +239,17 @@ void ptx_file_line_stats_create_exposed_latency_tracker(int n_shader_cores) } // add an inflight memory instruction -void ptx_file_line_stats_add_inflight_memory_insn(int sc_id, unsigned pc) +void ptx_stats::ptx_file_line_stats_add_inflight_memory_insn(int sc_id, unsigned pc) { - const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); + const ptx_instruction *pInsn = gpgpu_ctx->pc_to_instruction(pc); inflight_mem_tracker[sc_id].add_count(pInsn); } // remove an inflight memory instruction -void ptx_file_line_stats_sub_inflight_memory_insn(int sc_id, unsigned pc) +void ptx_stats::ptx_file_line_stats_sub_inflight_memory_insn(int sc_id, unsigned pc) { - const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); + const ptx_instruction *pInsn = gpgpu_ctx->pc_to_instruction(pc); inflight_mem_tracker[sc_id].sub_count(pInsn); } @@ -262,9 +262,9 @@ void ptx_file_line_stats_commit_exposed_latency(int sc_id, int exposed_latency) } // attribute the number of warp divergence to a ptx instruction -void ptx_file_line_stats_add_warp_divergence(unsigned pc, unsigned n_way_divergence) +void ptx_stats::ptx_file_line_stats_add_warp_divergence(unsigned pc, unsigned n_way_divergence) { - const ptx_instruction *pInsn = function_info::pc_to_instruction(pc); + const ptx_instruction *pInsn = gpgpu_ctx->pc_to_instruction(pc); ptx_file_line_stats& line_stats = ptx_file_line_stats_tracker[ptx_file_line(pInsn->source_file(), pInsn->source_line())]; line_stats.warp_divergence += n_way_divergence; diff --git a/src/cuda-sim/ptx-stats.h b/src/cuda-sim/ptx-stats.h index c75fc58..246b4ce 100644 --- a/src/cuda-sim/ptx-stats.h +++ b/src/cuda-sim/ptx-stats.h @@ -37,17 +37,10 @@ void ptx_file_line_stats_add_exec_count(const ptx_instruction *pInsn); #endif // stat collection interface to gpgpu-sim -void ptx_file_line_stats_add_latency(unsigned pc, unsigned latency); -void ptx_file_line_stats_add_dram_traffic(unsigned pc, unsigned dram_traffic); -void ptx_file_line_stats_add_smem_bank_conflict(unsigned pc, unsigned n_way_bkconflict); -void ptx_file_line_stats_add_uncoalesced_gmem(unsigned pc, unsigned n_access); void ptx_file_line_stats_create_exposed_latency_tracker(int n_shader_cores); -void ptx_file_line_stats_add_inflight_memory_insn(int sc_id, unsigned pc); -void ptx_file_line_stats_sub_inflight_memory_insn(int sc_id, unsigned pc); void ptx_file_line_stats_commit_exposed_latency(int sc_id, int exposed_latency); -void ptx_file_line_stats_add_warp_divergence(unsigned pc, unsigned n_way_divergence); class gpgpu_context; class ptx_stats { @@ -64,4 +57,12 @@ class ptx_stats { // output stats to a file void ptx_file_line_stats_write_file(); + // stat collection interface to gpgpu-sim + void ptx_file_line_stats_add_latency(unsigned pc, unsigned latency); + void ptx_file_line_stats_add_dram_traffic(unsigned pc, unsigned dram_traffic); + void ptx_file_line_stats_add_smem_bank_conflict(unsigned pc, unsigned n_way_bkconflict); + void ptx_file_line_stats_add_uncoalesced_gmem(unsigned pc, unsigned n_access); + void ptx_file_line_stats_add_inflight_memory_insn(int sc_id, unsigned pc); + void ptx_file_line_stats_sub_inflight_memory_insn(int sc_id, unsigned pc); + void ptx_file_line_stats_add_warp_divergence(unsigned pc, unsigned n_way_divergence); }; diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 5fa9379..3384d49 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -43,6 +43,14 @@ typedef void * yyscan_t; #define STR_SIZE 1024 +const ptx_instruction* gpgpu_context::pc_to_instruction(unsigned pc) +{ + if( pc < s_g_pc_to_insn.size() ) + return s_g_pc_to_insn[pc]; + else + return NULL; +} + unsigned symbol::get_uid() { unsigned result = (gpgpu_ctx->symbol_sm_next_uid)++; diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 8fc0a06..f4c5c37 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1372,13 +1372,6 @@ public: return m_symtab; } - static const ptx_instruction* pc_to_instruction(unsigned pc) - { - if( pc < s_g_pc_to_insn.size() ) - return s_g_pc_to_insn[pc]; - else - return NULL; - } unsigned local_mem_framesize() const { return m_local_mem_framesize; @@ -1436,8 +1429,6 @@ private: symbol_table *m_symtab; - static std::vector s_g_pc_to_insn; // a direct mapping from PC to instruction - //parameter size for device kernels int m_args_aligned_size; diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index bdf989a..e4ae04f 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -886,7 +886,7 @@ void gpgpu_sim::init() m_shader_stats->new_grid(); // initialize the control-flow, memory access, memory latency logger if (m_config.g_visualizer_enabled) { - create_thread_CFlogger( m_config.num_shader(), m_shader_config->n_thread_per_shader, 0, m_config.gpgpu_cflog_interval ); + create_thread_CFlogger( gpgpu_ctx, m_config.num_shader(), m_shader_config->n_thread_per_shader, 0, m_config.gpgpu_cflog_interval ); } shader_CTA_count_create( m_config.num_shader(), m_config.gpgpu_cflog_interval); if (m_config.gpgpu_cflog_interval != 0) { diff --git a/src/gpgpu-sim/mem_latency_stat.cc b/src/gpgpu-sim/mem_latency_stat.cc index 4e94991..a1b43a8 100644 --- a/src/gpgpu-sim/mem_latency_stat.cc +++ b/src/gpgpu-sim/mem_latency_stat.cc @@ -41,6 +41,7 @@ #include #include #include +#include "../../libcuda/gpgpu_context.h" memory_stats_t::memory_stats_t( unsigned n_shader, const shader_core_config *shader_config, const memory_config *mem_config, const class gpgpu_sim* gpu ) { @@ -195,7 +196,7 @@ void memory_stats_t::memlatstat_dram_access(mem_fetch *mf) mem_access_type_stats[mf->get_access_type()][dram_id][bank]++; } if (mf->get_pc() != (unsigned)-1) - ptx_file_line_stats_add_dram_traffic(mf->get_pc(), mf->get_data_size()); + m_gpu->gpgpu_ctx->stats->ptx_file_line_stats_add_dram_traffic(mf->get_pc(), mf->get_data_size()); } void memory_stats_t::memlatstat_icnt2mem_pop(mem_fetch *mf) diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index b7ae95d..c697450 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -744,7 +744,7 @@ void shader_core_ctx::decode() if( m_inst_fetch_buffer.m_valid ) { // decode 1 or 2 instructions and place them into ibuffer address_type pc = m_inst_fetch_buffer.m_pc; - const warp_inst_t* pI1 = ptx_fetch_inst(pc); + const warp_inst_t* pI1 = m_gpu->gpgpu_ctx->ptx_fetch_inst(pc); m_warp[m_inst_fetch_buffer.m_warp_id].ibuffer_fill(0,pI1); m_warp[m_inst_fetch_buffer.m_warp_id].inc_inst_in_pipeline(); if( pI1 ) { @@ -754,7 +754,7 @@ void shader_core_ctx::decode() }else if(pI1->oprnd_type==FP_OP) { m_stats->m_num_FPdecoded_insn[m_sid]++; } - const warp_inst_t* pI2 = ptx_fetch_inst(pc+pI1->isize); + const warp_inst_t* pI2 = m_gpu->gpgpu_ctx->ptx_fetch_inst(pc+pI1->isize); if( pI2 ) { m_warp[m_inst_fetch_buffer.m_warp_id].ibuffer_fill(1,pI2); m_warp[m_inst_fetch_buffer.m_warp_id].inc_inst_in_pipeline(); diff --git a/src/gpgpu-sim/stat-tool.cc b/src/gpgpu-sim/stat-tool.cc index 6a4c75b..35a4cc3 100644 --- a/src/gpgpu-sim/stat-tool.cc +++ b/src/gpgpu-sim/stat-tool.cc @@ -37,6 +37,7 @@ #include #include #include +#include "../../libcuda/gpgpu_context.h" //////////////////////////////////////////////////////////////////////////////// @@ -110,12 +111,10 @@ void spill_log_to_file (FILE *fout, int final, unsigned long long current_cycle //////////////////////////////////////////////////////////////////////////////// -unsigned translate_pc_to_ptxlineno(unsigned pc); - static int n_thread_CFloggers = 0; static thread_CFlocality** thread_CFlogger = NULL; -void create_thread_CFlogger( int n_loggers, int n_threads, address_type start_pc, unsigned long long logging_interval) +void create_thread_CFlogger(gpgpu_context* ctx, int n_loggers, int n_threads, address_type start_pc, unsigned long long logging_interval) { destroy_thread_CFlogger(); @@ -126,7 +125,7 @@ void create_thread_CFlogger( int n_loggers, int n_threads, address_type start_pc char buffer[32]; for (int i = 0; i < n_thread_CFloggers; i++) { snprintf(buffer, 32, "%02d", i); - thread_CFlogger[i] = new thread_CFlocality( name_tpl + buffer, logging_interval, n_threads, start_pc); + thread_CFlogger[i] = new thread_CFlocality( ctx, name_tpl + buffer, logging_interval, n_threads, start_pc); if (logging_interval != 0) { add_snap_shot_trigger(thread_CFlogger[i]); add_spill_log(thread_CFlogger[i]); @@ -368,10 +367,10 @@ static int s_cache_access_logger_n_types = 0; static std::vector s_cache_access_logger; enum cache_access_logger_types { - NORMAL, TEXTURE, CONSTANT, INSTRUCTION + NORMALS, TEXTURE, CONSTANT, INSTRUCTION }; -int get_shader_normal_cache_id() { return NORMAL; } +int get_shader_normal_cache_id() { return NORMALS; } int get_shader_texture_cache_id() { return TEXTURE; } int get_shader_constant_cache_id() { return CONSTANT; } int get_shader_instruction_cache_id() { return INSTRUCTION; } @@ -394,7 +393,7 @@ void shader_cache_access_log( int logger_id, int type, int miss) { if (s_cache_access_logger_n_types == 0) return; if (logger_id < 0) return; - assert(type == NORMAL || type == TEXTURE || type == CONSTANT || type == INSTRUCTION); + assert(type == NORMALS || type == TEXTURE || type == CONSTANT || type == INSTRUCTION); assert(miss == 0 || miss == 1); s_cache_access_logger[logger_id].log(2 * type + miss); @@ -404,7 +403,7 @@ void shader_cache_access_unlog( int logger_id, int type, int miss) { if (s_cache_access_logger_n_types == 0) return; if (logger_id < 0) return; - assert(type == NORMAL || type == TEXTURE || type == CONSTANT || type == INSTRUCTION); + assert(type == NORMALS || type == TEXTURE || type == CONSTANT || type == INSTRUCTION); assert(miss == 0 || miss == 1); s_cache_access_logger[logger_id].unlog(2 * type + miss); @@ -477,22 +476,24 @@ void shader_CTA_count_visualizer_gzprint( gzFile fout ) //////////////////////////////////////////////////////////////////////////////// //////////////////////////////////////////////////////////////////////////////// -thread_insn_span::thread_insn_span(unsigned long long cycle) +thread_insn_span::thread_insn_span(unsigned long long cycle, gpgpu_context* ctx) : m_cycle(cycle), #if (tr1_hash_map_ismap == 1) m_insn_span_count() #else m_insn_span_count(32*1024) #endif -{ +{ + gpgpu_ctx = ctx; } thread_insn_span::~thread_insn_span() { } -thread_insn_span::thread_insn_span(const thread_insn_span& other) +thread_insn_span::thread_insn_span(const thread_insn_span& other, gpgpu_context* ctx) : m_cycle(other.m_cycle), - m_insn_span_count(other.m_insn_span_count) + m_insn_span_count(other.m_insn_span_count) { + gpgpu_ctx = ctx; } thread_insn_span& thread_insn_span::operator=(const thread_insn_span& other) @@ -551,7 +552,7 @@ void thread_insn_span::print_sparse_histo(FILE *fout) const int n_printed_entries = 0; span_count_map::const_iterator i_sc = m_insn_span_count.begin(); for (; i_sc != m_insn_span_count.end(); ++i_sc) { - unsigned ptx_lineno = translate_pc_to_ptxlineno(i_sc->first); + unsigned ptx_lineno = gpgpu_ctx->translate_pc_to_ptxlineno(i_sc->first); fprintf(fout, "%u %d ", ptx_lineno, i_sc->second); n_printed_entries++; } @@ -566,7 +567,7 @@ void thread_insn_span::print_sparse_histo(gzFile fout) const int n_printed_entries = 0; span_count_map::const_iterator i_sc = m_insn_span_count.begin(); for (; i_sc != m_insn_span_count.end(); ++i_sc) { - unsigned ptx_lineno = translate_pc_to_ptxlineno(i_sc->first); + unsigned ptx_lineno = gpgpu_ctx->translate_pc_to_ptxlineno(i_sc->first); gzprintf(fout, "%u %d ", ptx_lineno, i_sc->second); n_printed_entries++; } @@ -578,14 +579,14 @@ void thread_insn_span::print_sparse_histo(gzFile fout) const //////////////////////////////////////////////////////////////////////////////// -thread_CFlocality::thread_CFlocality(std::string name, +thread_CFlocality::thread_CFlocality( gpgpu_context* ctx, std::string name, unsigned long long snap_shot_interval, int nthreads, address_type start_pc, unsigned long long start_cycle) : snap_shot_trigger(snap_shot_interval), m_name(name), m_nthreads(nthreads), m_thread_pc(nthreads, start_pc), m_cycle(start_cycle), - m_thd_span(start_cycle) + m_thd_span(start_cycle, ctx) { std::fill(m_thread_pc.begin(), m_thread_pc.end(), -1); // so that hw thread with no work assigned will not clobber results } diff --git a/src/gpgpu-sim/stat-tool.h b/src/gpgpu-sim/stat-tool.h index 5646f01..67b3923 100644 --- a/src/gpgpu-sim/stat-tool.h +++ b/src/gpgpu-sim/stat-tool.h @@ -35,6 +35,7 @@ #include #include +class gpgpu_context; ///////////////////////////////////////////////////////////////////////////////////// // logger snapshot trigger: // - automate the snap_shot part of loggers to avoid modifying simulation loop everytime @@ -80,8 +81,8 @@ public: class thread_insn_span { public: - thread_insn_span(unsigned long long cycle); - thread_insn_span(const thread_insn_span& other); + thread_insn_span(unsigned long long cycle, gpgpu_context* ctx); + thread_insn_span(const thread_insn_span& other, gpgpu_context* ctx); ~thread_insn_span(); thread_insn_span& operator=(const thread_insn_span& other); @@ -94,7 +95,8 @@ public: void print_sparse_histo(FILE *fout) const; void print_sparse_histo(gzFile fout) const; -private: +private: + gpgpu_context* gpgpu_ctx; typedef tr1_hash_map span_count_map; unsigned long long m_cycle; span_count_map m_insn_span_count; @@ -102,7 +104,7 @@ private: class thread_CFlocality : public snap_shot_trigger, public spill_log_interface { public: - thread_CFlocality(std::string name, unsigned long long snap_shot_interval, + thread_CFlocality(gpgpu_context* ctx, std::string name, unsigned long long snap_shot_interval, int nthreads, address_type start_pc, unsigned long long start_cycle = 0); ~thread_CFlocality(); @@ -270,7 +272,7 @@ void try_snap_shot (unsigned long long current_cycle); void set_spill_interval (unsigned long long interval); void spill_log_to_file (FILE *fout, int final, unsigned long long current_cycle); -void create_thread_CFlogger( int n_loggers, int n_threads, address_type start_pc, unsigned long long logging_interval); +void create_thread_CFlogger(gpgpu_context* ctx, int n_loggers, int n_threads, address_type start_pc, unsigned long long logging_interval); void destroy_thread_CFlogger( ); void cflog_update_thread_pc( int logger_id, int thread_id, address_type pc ); void cflog_snapshot( int logger_id, unsigned long long cycle ); -- cgit v1.3 From e1fa1a3cc7c509417064a8e4cdab71e3f7feb881 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 15 Jul 2019 22:16:55 -0400 Subject: Move debug_tensorcore Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 2 ++ src/cuda-sim/instructions.cc | 43 +++++++++++++++++++++---------------------- 2 files changed, 23 insertions(+), 22 deletions(-) (limited to 'libcuda/gpgpu_context.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 346a8a4..45c5cdd 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -20,6 +20,7 @@ class gpgpu_context { g_ptx_cta_info_uid = 1; symbol_sm_next_uid = 1; function_info_sm_next_uid = 1; + debug_tensorcore = 0; api = new cuda_runtime_api(this); ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); @@ -40,6 +41,7 @@ class gpgpu_context { unsigned symbol_sm_next_uid; //uid for symbol unsigned function_info_sm_next_uid; std::vector s_g_pc_to_insn; // a direct mapping from PC to instruction + bool debug_tensorcore; // objects pointers for each file cuda_runtime_api* api; diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 4d4a80d..58a077e 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -60,7 +60,6 @@ class ptx_recognizer; using half_float::half; -bool debug_tensorcore = 0; const char *g_opcode_string[NUM_OPCODES] = { @@ -1840,14 +1839,14 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) for (thrd=0; thrd < core->get_warp_size(); thrd++){ thread = core->get_thread_info()[tid+thrd]; - if(debug_tensorcore) + if(core->get_gpu()->gpgpu_ctx->debug_tensorcore) printf("THREAD=%d\n:",thrd); for(int operand_num=1;operand_num<=3;operand_num++){ const operand_info &src_a= pI->operand_lookup(operand_num); unsigned nelem = src_a.get_vect_nelem(); ptx_reg_t v[8]; thread->get_vector_operand_values( src_a, v, nelem ); - if(debug_tensorcore){ + if(core->get_gpu()->gpgpu_ctx->debug_tensorcore){ printf("Thread%d_Iteration=%d\n:",thrd,operand_num); for(k=0;kget_gpu()->gpgpu_ctx->debug_tensorcore) printf("%.2f ",temp); } - if(debug_tensorcore) + if(core->get_gpu()->gpgpu_ctx->debug_tensorcore) printf("\n"); } else{ - if(debug_tensorcore){ + if(core->get_gpu()->gpgpu_ctx->debug_tensorcore){ for(k=0;k<8;k++){ printf("%.2f ",v[k].f32); } @@ -1887,7 +1886,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) case 1 ://operand 1 for(k=0;k<8;k++){ mapping(thrd,LOAD_A,a_layout,F16_TYPE,k,16,row,col,offset); - if(debug_tensorcore) + if(core->get_gpu()->gpgpu_ctx->debug_tensorcore) printf("A:thread=%d,row=%d,col=%d,offset=%d\n",thrd,row,col,offset); matrix_a[row][col]=nw_v[offset]; } @@ -1895,7 +1894,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) case 2 ://operand 2 for(k=0;k<8;k++){ mapping(thrd,LOAD_B,b_layout,F16_TYPE,k,16,row,col,offset); - if(debug_tensorcore) + if(core->get_gpu()->gpgpu_ctx->debug_tensorcore) printf("B:thread=%d,row=%d,col=%d,offset=%d\n",thrd,row,col,offset); matrix_b[row][col]=nw_v[offset]; } @@ -1903,7 +1902,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) case 3 ://operand 3 for(k=0;k<8;k++){ mapping(thrd,LOAD_C,ROW,type2,k,16,row,col,offset); - if(debug_tensorcore) + if(core->get_gpu()->gpgpu_ctx->debug_tensorcore) printf("C:thread=%d,row=%d,col=%d,offset=%d\n",thrd,row,col,offset); if(type2!=F16_TYPE){ matrix_c[row][col]=v[offset]; @@ -1917,10 +1916,10 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) printf("Invalid Operand Index\n" ); } } - if(debug_tensorcore) + if(core->get_gpu()->gpgpu_ctx->debug_tensorcore) printf("\n"); } - if(debug_tensorcore){ + if(core->get_gpu()->gpgpu_ctx->debug_tensorcore){ printf("MATRIX_A\n"); for (i=0;i<16;i++){ for(j=0;j<16;j++){ @@ -1980,7 +1979,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) } } } - if(debug_tensorcore){ + if(core->get_gpu()->gpgpu_ctx->debug_tensorcore){ printf("MATRIX_D\n"); for (i=0;i<16;i++){ for(j=0;j<16;j++){ @@ -1999,7 +1998,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) int col_t[8]; for(k=0;k<8;k++){ mapping(thrd,LOAD_C,ROW,type,k,16,row_t[k],col_t[k],offset); - if(debug_tensorcore) + if(core->get_gpu()->gpgpu_ctx->debug_tensorcore) printf("mma:store:row:%d,col%d\n",row_t[k],col_t[k]); } thread = core->get_thread_info()[tid+thrd]; @@ -2008,7 +2007,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) if(type==F32_TYPE){ thread->set_wmma_vector_operand_values(dst,matrix_d[row_t[0]][col_t[0]],matrix_d[row_t[1]][col_t[1]],matrix_d[row_t[2]][col_t[2]],matrix_d[row_t[3]][col_t[3]],matrix_d[row_t[4]][col_t[4]],matrix_d[row_t[5]][col_t[5]],matrix_d[row_t[6]][col_t[6]],matrix_d[row_t[7]][col_t[7]]); - if(debug_tensorcore) + if(core->get_gpu()->gpgpu_ctx->debug_tensorcore) { printf("thread%d:",thrd); for(k=0;k<8;k++){ @@ -2018,7 +2017,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) } } else if(type==F16_TYPE){ - if(debug_tensorcore){ + if(core->get_gpu()->gpgpu_ctx->debug_tensorcore){ printf("thread%d:",thrd); for(k=0;k<8;k++){ temp=matrix_d[row_t[k]][col_t[k]].f16; @@ -2038,7 +2037,7 @@ void mma_impl( const ptx_instruction *pI, core_t *core, warp_inst_t inst ) nw_data3.s64=((matrix_d[row_t[4]][col_t[4]].s64 & 0xffff))|((matrix_d[row_t[5]][col_t[5]].s64&0xffff)<<16); 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(debug_tensorcore) + 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); } @@ -3132,7 +3131,7 @@ 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); - if(debug_tensorcore) + 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); addr_t new_addr = addr+thread_group_offset(thrd,wmma_type,wmma_layout,type,stride)*size/8; addr_t push_addr; @@ -3152,7 +3151,7 @@ void mma_st_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) mem->write(push_addr,size/8,&v[k].s64,thread,pI); mem_txn_addr[num_mem_txn++]=push_addr; - if(debug_tensorcore){ + 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); float temp; int l; @@ -3179,7 +3178,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(debug_tensorcore) + 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); } } @@ -3242,7 +3241,7 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) type_info_key::type_decode(type,size,t); ptx_reg_t data[16]; - if(debug_tensorcore) + 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); addr_t new_addr = addr+thread_group_offset(thrd,wmma_type,wmma_layout,type,stride)*size/8; @@ -3338,7 +3337,7 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) inst.data_size = 4; // 4 byte transaction assert( inst.memory_op == insn_memory_op ); - if(debug_tensorcore){ + if(core->get_gpu()->gpgpu_ctx->debug_tensorcore){ if(type==F16_TYPE){ printf("\nmma_ld:thread%d= ",thrd); for(i=0;i<16;i++){ @@ -3388,7 +3387,7 @@ void mma_ld_impl( const ptx_instruction *pI, core_t *core, warp_inst_t &inst ) thread->set_vector_operand_values(dst,nw_data[0],nw_data[1],nw_data[2],nw_data[3]); 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(debug_tensorcore){ + 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); -- cgit v1.3