From 112a3a2e3c9d17af420d2389f68a35407c692744 Mon Sep 17 00:00:00 2001 From: tgrogers Date: Fri, 29 Jun 2018 17:34:22 -0400 Subject: Allowing modern GPU configrations to properly run PTXPLUS. There are several problems with this: 1) To get proper occupancy, based on register file usage, you must querry ptxas with the sm version that corresponds to the register usage in the config file. To enable this, a new config option has been added that determines what sm version you pass to ptxas to compute reg-usage. This configuration option is always required in the gpgpusim.config file 2) If you are running PTXPLUS with a modern card (i.e. volta/pascal), you need ptxas from CUDA 9.1. However since PTXPLUS only supports sm_13 - you need a version of CUDA where cudaobjectdump supports sm_1x. This ended at CUDA 5 - and PTXPLUS requires CUDA 4.2. Therefore, to run PTXPLUS on a modern card, you need CUDA 4.2 + modern CUDA installed. To fascilitate this, a new envronment varaible is added and the setup envrionment script prints an appropraite warning if you are using a newer CUDA. We have tried to make this as fail-proof as possible - and die appropraitely when something is wrong. --- libopencl/opencl_runtime_api.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'libopencl/opencl_runtime_api.cc') diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 63f7c81..97a54d8 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -577,7 +577,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 ); - gpgpu_ptxinfo_load_from_string( tmp, source_num); + gpgpu_ptxinfo_load_from_string( tmp, source_num ); free(tmp); } printf("GPGPU-Sim OpenCL API: finished compiling OpenCL kernels.\n"); -- 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 'libopencl/opencl_runtime_api.cc') 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 'libopencl/opencl_runtime_api.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 14e3329..0a6eabb 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -640,7 +640,7 @@ void** cudaRegisterFatBinaryInternal( void *fatCubin, gpgpu_context* gpgpu_ctx = } else { symtab=ctx->gpgpu_ptx_sim_load_ptx_from_string(ptx,source_num); context->add_binary(symtab,fat_cubin_handle); - gpgpu_ptxinfo_load_from_string( ptx, source_num, max_capability, context->no_of_ptx ); + ctx->gpgpu_ptxinfo_load_from_string( ptx, source_num, max_capability, context->no_of_ptx ); } source_num++; ctx->api->load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu()); diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 16626eb..d3db1ad 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -1,19 +1,24 @@ #ifndef __gpgpu_context_h__ #define __gpgpu_context_h__ #include "cuda_api_object.h" +#include "../src/cuda-sim/ptx_loader.h" class gpgpu_context { public: gpgpu_context() { api = new cuda_runtime_api(); + ptxinfo = new ptxinfo_data(); } // global list // objects pointers for each file cuda_runtime_api* api; + ptxinfo_data* ptxinfo; // member function list void cuobjdumpParseBinary(unsigned int handle); class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); class symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename ); + void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version); + void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version=20, int no_of_ptx=0 ); }; gpgpu_context* GPGPU_Context(); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 0ec635e..50a02fa 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -580,7 +580,7 @@ void _cl_program::Build(const char *options) } info.m_asm = tmp; info.m_symtab = ctx->gpgpu_ptx_sim_load_ptx_from_string( tmp, source_num ); - gpgpu_ptxinfo_load_from_string( tmp, source_num ); + ctx->gpgpu_ptxinfo_load_from_string( tmp, source_num ); free(tmp); } printf("GPGPU-Sim OpenCL API: finished compiling OpenCL kernels.\n"); diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 9857741..2a6d930 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -332,7 +332,7 @@ char* get_app_binary_name(){ return self_exe_path; } -void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version) +void gpgpu_context::gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version) { std::string ptxas_filename(std::string(filename) + "as"); char buff[1024], extra_flags[1024]; @@ -352,17 +352,16 @@ void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_versio } FILE *ptxinfo_in; - ptxinfo_data ptxinfo; - ptxinfo.g_ptxinfo_filename = strdup(ptxas_filename.c_str()); - ptxinfo_in = fopen(ptxinfo.g_ptxinfo_filename,"r"); - ptxinfo_lex_init(&(ptxinfo.scanner)); - ptxinfo_set_in(ptxinfo_in, ptxinfo.scanner); - ptxinfo_parse(ptxinfo.scanner, &ptxinfo); - ptxinfo_lex_destroy(ptxinfo.scanner); + ptxinfo->g_ptxinfo_filename = strdup(ptxas_filename.c_str()); + ptxinfo_in = fopen(ptxinfo->g_ptxinfo_filename,"r"); + ptxinfo_lex_init(&(ptxinfo->scanner)); + ptxinfo_set_in(ptxinfo_in, ptxinfo->scanner); + ptxinfo_parse(ptxinfo->scanner, ptxinfo); + ptxinfo_lex_destroy(ptxinfo->scanner); fclose(ptxinfo_in); } -void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version, int no_of_ptx ) +void gpgpu_context::gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version, int no_of_ptx ) { //do ptxas for individual files instead of one big embedded ptx. This prevents the duplicate defs and declarations. char ptx_file[1000]; @@ -420,14 +419,13 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num if( result != 0 ) { // 65280 = duplicate errors if (result == 65280) { - ptxinfo_data ptxinfo; FILE *ptxinfo_in; ptxinfo_in = fopen(tempfile_ptxinfo,"r"); - ptxinfo.g_ptxinfo_filename = tempfile_ptxinfo; - ptxinfo_lex_init(&(ptxinfo.scanner)); - ptxinfo_set_in(ptxinfo_in, ptxinfo.scanner); - ptxinfo_parse(ptxinfo.scanner, &ptxinfo); - ptxinfo_lex_destroy(ptxinfo.scanner); + ptxinfo->g_ptxinfo_filename = tempfile_ptxinfo; + ptxinfo_lex_init(&(ptxinfo->scanner)); + ptxinfo_set_in(ptxinfo_in, ptxinfo->scanner); + ptxinfo_parse(ptxinfo->scanner, ptxinfo); + ptxinfo_lex_destroy(ptxinfo->scanner); fclose(ptxinfo_in); fix_duplicate_errors(fname2); @@ -507,18 +505,17 @@ void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num } } - ptxinfo_data ptxinfo; if(no_of_ptx>0) - ptxinfo.g_ptxinfo_filename = final_tempfile_ptxinfo; + ptxinfo->g_ptxinfo_filename = final_tempfile_ptxinfo; else - ptxinfo.g_ptxinfo_filename = tempfile_ptxinfo; + ptxinfo->g_ptxinfo_filename = tempfile_ptxinfo; FILE *ptxinfo_in; - ptxinfo_in = fopen(ptxinfo.g_ptxinfo_filename,"r"); + ptxinfo_in = fopen(ptxinfo->g_ptxinfo_filename,"r"); - ptxinfo_lex_init(&(ptxinfo.scanner)); - ptxinfo_set_in(ptxinfo_in, ptxinfo.scanner); - ptxinfo_parse(ptxinfo.scanner, &ptxinfo); - ptxinfo_lex_destroy(ptxinfo.scanner); + ptxinfo_lex_init(&(ptxinfo->scanner)); + ptxinfo_set_in(ptxinfo_in, ptxinfo->scanner); + ptxinfo_parse(ptxinfo->scanner, ptxinfo); + ptxinfo_lex_destroy(ptxinfo->scanner); fclose(ptxinfo_in); snprintf(commandline,1024,"rm -f *info"); diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index c3ce888..3637bff 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -44,8 +44,6 @@ class ptxinfo_data{ extern bool g_override_embedded_ptx; extern int no_of_ptx; //counter to track number of ptx files to be extracted in an application. -void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version=20, int no_of_ptx=0 ); -void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version ); char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str); bool keep_intermediate_files(); -- cgit v1.3 From cb678c3670de4a435a3260ed80dc476da3860082 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 1 Jul 2019 16:23:33 -0400 Subject: Move g_debug_ir_generation and GPGPUSim_Init Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 674 ++++++++++++++++++++++------------------ libcuda/gpgpu_context.h | 2 + libopencl/opencl_runtime_api.cc | 4 +- src/cuda-sim/ptx_parser.cc | 3 +- src/cuda-sim/ptx_parser.h | 3 + src/gpgpusim_entrypoint.cc | 5 +- src/gpgpusim_entrypoint.h | 1 - 7 files changed, 375 insertions(+), 317 deletions(-) (limited to 'libopencl/opencl_runtime_api.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 61f8415..a34727f 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -197,7 +197,7 @@ void register_ptx_function( const char *name, function_info *impl ) # endif #endif -struct _cuda_device_id *GPGPUSim_Init() +struct _cuda_device_id *gpgpu_context::GPGPUSim_Init() { //static _cuda_device_id *the_device = NULL; _cuda_device_id *the_device = GPGPUsim_ctx_ptr()->the_cude_device; @@ -255,9 +255,10 @@ struct _cuda_device_id *GPGPUSim_Init() static CUctx_st* GPGPUSim_Context() { //static CUctx_st *the_context = NULL; + gpgpu_context *cur_ctx = GPGPU_Context(); CUctx_st *the_context = GPGPUsim_ctx_ptr()->the_context; if( the_context == NULL ) { - _cuda_device_id *the_gpu = GPGPUSim_Init(); + _cuda_device_id *the_gpu = cur_ctx->GPGPUSim_Init(); GPGPUsim_ctx_ptr()->the_context = new CUctx_st(the_gpu); the_context = GPGPUsim_ctx_ptr()->the_context; } @@ -496,7 +497,7 @@ cudaError_t cudaSetDeviceInternal(int device, gpgpu_context* gpgpu_ctx = NULL) announce_call(__my_func__); } //set the active device to run cuda - if ( device <= GPGPUSim_Init()->num_devices() ) { + if ( device <= ctx->GPGPUSim_Init()->num_devices() ) { ctx->api->g_active_device = device; return g_last_cudaError = cudaSuccess; } else { @@ -519,6 +520,55 @@ cudaError_t cudaGetDeviceInternal(int *device, gpgpu_context* gpgpu_ctx = NULL) return g_last_cudaError = cudaSuccess; } +__host__ cudaError_t CUDARTAPI cudaDeviceGetLimitInternal( size_t* pValue, cudaLimit limit, gpgpu_context* gpgpu_ctx = NULL ) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + _cuda_device_id *dev = ctx->GPGPUSim_Init(); + const struct cudaDeviceProp *prop = dev->get_prop(); + const gpgpu_sim_config& config=dev->get_gpgpu()->get_config(); + switch(limit) { + case 0: // cudaLimitStackSize + *pValue=config.stack_limit(); + break; + case 2: // cudaLimitMallocHeapSize + *pValue=config.heap_limit(); + break; +#if (CUDART_VERSION > 5050) + case 3: // cudaLimitDevRuntimeSyncDepth + if(prop->major > 2){ + *pValue=config.sync_depth_limit(); + break; + } + else{ + printf("ERROR:Limit %s is not supported on this architecture \n",limit); + abort(); + } + case 4: // cudaLimitDevRuntimePendingLaunchCount + if(prop->major > 2){ + *pValue=config.pending_launch_count_limit(); + break; + } + else{ + printf("ERROR:Limit %s is not supported on this architecture \n",limit); + abort(); + } +#endif + default: + printf("ERROR:Limit %s unimplemented \n",limit); + abort(); + } + return g_last_cudaError = cudaSuccess; + +} + void** cudaRegisterFatBinaryInternal( void *fatCubin, gpgpu_context* gpgpu_ctx = NULL) { @@ -739,6 +789,260 @@ cudaError_t cudaConfigureCallInternal(dim3 gridDim, dim3 blockDim, size_t shared return g_last_cudaError = cudaSuccess; } +__host__ cudaError_t CUDARTAPI cudaGetDeviceCountInternal(int *count, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + _cuda_device_id *dev = ctx->GPGPUSim_Init(); + *count = dev->num_devices(); + return g_last_cudaError = cudaSuccess; +} + +__host__ cudaError_t CUDARTAPI cudaGetDevicePropertiesInternal(struct cudaDeviceProp *prop, int device, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + _cuda_device_id *dev = ctx->GPGPUSim_Init(); + if (device <= dev->num_devices() ) { + *prop= *dev->get_prop(); + return g_last_cudaError = cudaSuccess; + } else { + return g_last_cudaError = cudaErrorInvalidDevice; + } +} + +#if (CUDART_VERSION > 5000) +__host__ cudaError_t CUDARTAPI cudaDeviceGetAttributeInternal(int *value, enum cudaDeviceAttr attr, int device, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + const struct cudaDeviceProp *prop; + _cuda_device_id *dev = ctx->GPGPUSim_Init(); + if (device <= dev->num_devices() ) { + prop = dev->get_prop(); + switch (attr) { + case 1: + *value= prop->maxThreadsDim[0] * prop->maxThreadsDim[1] * prop->maxThreadsDim[2] * prop->maxGridSize[0] * prop->maxGridSize[1] * prop->maxGridSize[2]; + break; + case 2: + *value= prop->maxThreadsDim[0]; + break; + case 3: + *value= prop->maxThreadsDim[1]; + break; + case 4: + *value= prop->maxThreadsDim[2]; + break; + case 5: + *value= prop->maxGridSize[0]; + break; + case 6: + *value= prop->maxGridSize[1]; + break; + case 7: + *value= prop->maxGridSize[2]; + break; + case 8: + *value= prop->sharedMemPerBlock; + break; + case 9: + *value= prop->totalConstMem; + break; + case 10: + *value= prop->warpSize; + break; + case 11: + *value= 16;//dummy value + break; + case 12: + *value= prop->regsPerBlock; + break; + case 13: + *value= 1480000;//for 1080ti + break; + case 14: + *value= prop->textureAlignment ; + break; + case 15: + *value = 0; + break; + case 16: + *value= prop->multiProcessorCount ; + break; + case 17: + case 18: + case 19: + *value = 0; + break; + case 21: + case 22: + case 23: + case 24: + case 25: + case 26: + case 27: + case 28: + case 42: + case 45: + case 46: + case 47: + case 48: + case 49: + case 52: + case 53: + case 55: + case 56: + case 57: + case 58: + case 59: + case 60: + case 61: + case 62: + case 63: + case 64: + case 66: + case 67: + case 69: + case 70: + case 71: + case 73: + case 74: + case 77: + *value = 1000;//dummy value + break; + case 29: + case 43: + case 54: + case 65: + case 68: + case 72: + *value = 10;//dummy value + break; + case 30: + case 51: + *value = 128;//dummy value + break; + case 31: + *value = 1; + break; + case 32: + *value = 0; + break; + case 33: + case 50: + *value = 0;//dummy value + break; + case 34: + *value= 0; + break; + case 35: + *value = 0; + break; + case 36: + *value = 1250000;//CK value for 1080ti + break; + case 37: + *value = 352;//value for 1080ti + break; + case 38: + *value = 3000000;//value for 1080ti + break; + case 39: + *value= dev->get_gpgpu()->threads_per_core(); + break; + case 40: + *value= 0; + break; + case 41: + *value= 0; + break; + case 75://cudaDevAttrComputeCapabilityMajor + *value= prop->major ; + break; + case 76://cudaDevAttrComputeCapabilityMinor + *value= prop->minor ; + break; + case 78: + *value= 0 ; //TODO: as of now, we dont support stream priorities. + break; + case 79: + *value= 0; + break; + case 80: + *value= 0; + break; + #if (CUDART_VERSION > 5050) + case 81: + *value= prop->sharedMemPerMultiprocessor; + break; + case 82: + *value= prop->regsPerMultiprocessor; + break; + #endif + case 83: + case 84: + case 85: + case 86: + *value= 0; + break; + case 87: + *value= 4;//dummy value + break; + case 88: + case 89: + case 90: + case 91: + case 95: + *value= 0; + break; + default: + printf("ERROR: Attribute number %d unimplemented \n",attr); + abort(); + } + return g_last_cudaError = cudaSuccess; + } else { + return g_last_cudaError = cudaErrorInvalidDevice; + } +} +#endif + +__host__ cudaError_t CUDARTAPI cudaChooseDeviceInternal(int *device, const struct cudaDeviceProp *prop, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + _cuda_device_id *dev = ctx->GPGPUSim_Init(); + *device = dev->get_id(); + return g_last_cudaError = cudaSuccess; +} + cudaError_t cudaSetupArgumentInternal(const void *arg, size_t size, size_t offset, gpgpu_context* gpgpu_ctx = NULL) { gpgpu_context *ctx; @@ -1058,6 +1362,52 @@ cudaError_t cudaHostAllocInternal(void **pHost, size_t bytes, unsigned int flag #endif +size_t getMaxThreadsPerBlock(struct cudaFuncAttributes *attr, gpgpu_context *ctx) { + _cuda_device_id *dev = ctx->GPGPUSim_Init(); + struct cudaDeviceProp prop; + + prop = *dev->get_prop(); + + size_t max = prop.maxThreadsPerBlock; + + if ((prop.regsPerBlock / attr->numRegs) < max) + max = prop.regsPerBlock / attr->numRegs; + + return max; +} + +cudaError_t CUDARTAPI cudaFuncGetAttributesInternal(struct cudaFuncAttributes *attr, const char *hostFun, gpgpu_context* gpgpu_ctx = NULL ) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + CUctx_st *context = GPGPUSim_Context(); + function_info *entry = context->get_kernel(hostFun); + if( entry ) { + const struct gpgpu_ptx_sim_info *kinfo = entry->get_kernel_info(); + attr->sharedSizeBytes = kinfo->smem; + attr->constSizeBytes = kinfo->cmem; + attr->localSizeBytes = kinfo->lmem; + attr->numRegs = kinfo->regs; + if(kinfo->maxthreads > 0) + attr->maxThreadsPerBlock = kinfo->maxthreads; + else + attr->maxThreadsPerBlock = getMaxThreadsPerBlock(attr, ctx); +#if CUDART_VERSION >= 3000 + attr->ptxVersion = kinfo->ptx_version; + attr->binaryVersion = kinfo->sm_target; +#endif + } + return g_last_cudaError = cudaSuccess; +} + + /******************************************************************************* * * * * @@ -1503,235 +1853,24 @@ __host__ cudaError_t CUDARTAPI cudaGetSymbolSize(size_t *size, const char *symbo *******************************************************************************/ __host__ cudaError_t CUDARTAPI cudaGetDeviceCount(int *count) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - _cuda_device_id *dev = GPGPUSim_Init(); - *count = dev->num_devices(); - return g_last_cudaError = cudaSuccess; + return cudaGetDeviceCountInternal(count); } - -__host__ cudaError_t CUDARTAPI cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device) -{ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - _cuda_device_id *dev = GPGPUSim_Init(); - if (device <= dev->num_devices() ) { - *prop= *dev->get_prop(); - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorInvalidDevice; - } + +__host__ cudaError_t CUDARTAPI cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device) +{ + return cudaGetDevicePropertiesInternal(prop, device); } #if (CUDART_VERSION > 5000) __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDeviceAttr attr, int device) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - const struct cudaDeviceProp *prop; - _cuda_device_id *dev = GPGPUSim_Init(); - if (device <= dev->num_devices() ) { - prop = dev->get_prop(); - switch (attr) { - case 1: - *value= prop->maxThreadsDim[0] * prop->maxThreadsDim[1] * prop->maxThreadsDim[2] * prop->maxGridSize[0] * prop->maxGridSize[1] * prop->maxGridSize[2]; - break; - case 2: - *value= prop->maxThreadsDim[0]; - break; - case 3: - *value= prop->maxThreadsDim[1]; - break; - case 4: - *value= prop->maxThreadsDim[2]; - break; - case 5: - *value= prop->maxGridSize[0]; - break; - case 6: - *value= prop->maxGridSize[1]; - break; - case 7: - *value= prop->maxGridSize[2]; - break; - case 8: - *value= prop->sharedMemPerBlock; - break; - case 9: - *value= prop->totalConstMem; - break; - case 10: - *value= prop->warpSize; - break; - case 11: - *value= 16;//dummy value - break; - case 12: - *value= prop->regsPerBlock; - break; - case 13: - *value= 1480000;//for 1080ti - break; - case 14: - *value= prop->textureAlignment ; - break; - case 15: - *value = 0; - break; - case 16: - *value= prop->multiProcessorCount ; - break; - case 17: - case 18: - case 19: - *value = 0; - break; - case 21: - case 22: - case 23: - case 24: - case 25: - case 26: - case 27: - case 28: - case 42: - case 45: - case 46: - case 47: - case 48: - case 49: - case 52: - case 53: - case 55: - case 56: - case 57: - case 58: - case 59: - case 60: - case 61: - case 62: - case 63: - case 64: - case 66: - case 67: - case 69: - case 70: - case 71: - case 73: - case 74: - case 77: - *value = 1000;//dummy value - break; - case 29: - case 43: - case 54: - case 65: - case 68: - case 72: - *value = 10;//dummy value - break; - case 30: - case 51: - *value = 128;//dummy value - break; - case 31: - *value = 1; - break; - case 32: - *value = 0; - break; - case 33: - case 50: - *value = 0;//dummy value - break; - case 34: - *value= 0; - break; - case 35: - *value = 0; - break; - case 36: - *value = 1250000;//CK value for 1080ti - break; - case 37: - *value = 352;//value for 1080ti - break; - case 38: - *value = 3000000;//value for 1080ti - break; - case 39: - *value= dev->get_gpgpu()->threads_per_core(); - break; - case 40: - *value= 0; - break; - case 41: - *value= 0; - break; - case 75://cudaDevAttrComputeCapabilityMajor - *value= prop->major ; - break; - case 76://cudaDevAttrComputeCapabilityMinor - *value= prop->minor ; - break; - case 78: - *value= 0 ; //TODO: as of now, we dont support stream priorities. - break; - case 79: - *value= 0; - break; - case 80: - *value= 0; - break; - #if (CUDART_VERSION > 5050) - case 81: - *value= prop->sharedMemPerMultiprocessor; - break; - case 82: - *value= prop->regsPerMultiprocessor; - break; - #endif - case 83: - case 84: - case 85: - case 86: - *value= 0; - break; - case 87: - *value= 4;//dummy value - break; - case 88: - case 89: - case 90: - case 91: - case 92: - case 93: - case 94: - case 95: - *value= 0; - break; - default: - printf("ERROR: Attribute number %d unimplemented \n",attr); - abort(); - } - return g_last_cudaError = cudaSuccess; - } else { - return g_last_cudaError = cudaErrorInvalidDevice; - } + return cudaDeviceGetAttributeInternal(value, attr, device); } #endif __host__ cudaError_t CUDARTAPI cudaChooseDevice(int *device, const struct cudaDeviceProp *prop) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - _cuda_device_id *dev = GPGPUSim_Init(); - *device = dev->get_id(); - return g_last_cudaError = cudaSuccess; + return cudaChooseDeviceInternal(device, prop); } __host__ cudaError_t CUDARTAPI cudaSetDevice(int device) @@ -1744,47 +1883,9 @@ __host__ cudaError_t CUDARTAPI cudaGetDevice(int *device) return cudaGetDeviceInternal(device); } -__host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit limit ) +__host__ cudaError_t CUDARTAPI cudaDeviceGetLimit( size_t* pValue, cudaLimit limit ) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - _cuda_device_id *dev = GPGPUSim_Init(); - const struct cudaDeviceProp *prop = dev->get_prop(); - const gpgpu_sim_config& config=dev->get_gpgpu()->get_config(); - switch(limit) { - case 0: // cudaLimitStackSize - *pValue=config.stack_limit(); - break; - case 2: // cudaLimitMallocHeapSize - *pValue=config.heap_limit(); - break; -#if (CUDART_VERSION > 5050) - case 3: // cudaLimitDevRuntimeSyncDepth - if(prop->major > 2){ - *pValue=config.sync_depth_limit(); - break; - } - else{ - printf("ERROR:Limit %s is not supported on this architecture \n",limit); - abort(); - } - case 4: // cudaLimitDevRuntimePendingLaunchCount - if(prop->major > 2){ - *pValue=config.pending_launch_count_limit(); - break; - } - else{ - printf("ERROR:Limit %s is not supported on this architecture \n",limit); - abort(); - } -#endif - default: - printf("ERROR:Limit %s unimplemented \n",limit); - abort(); - } - return g_last_cudaError = cudaSuccess; - + return cudaDeviceGetLimitInternal( pValue, limit ); } __host__ cudaError_t CUDARTAPI cudaStreamGetPriority ( cudaStream_t hStream, int* priority ) @@ -3167,58 +3268,9 @@ cudaError_t CUDARTAPI cudaSetDeviceFlags( int flags ) } } -size_t getMaxThreadsPerBlock(struct cudaFuncAttributes *attr) { - _cuda_device_id *dev = GPGPUSim_Init(); - struct cudaDeviceProp prop; - - prop = *dev->get_prop(); - - size_t max = prop.maxThreadsPerBlock; - - if ((prop.regsPerBlock / attr->numRegs) < max) - max = prop.regsPerBlock / attr->numRegs; - - return max; -} - cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *attr, const char *hostFun ) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - CUctx_st *context = GPGPUSim_Context(); - function_info *entry = context->get_kernel(hostFun); - if( entry ) { - const struct gpgpu_ptx_sim_info *kinfo = entry->get_kernel_info(); - attr->sharedSizeBytes = kinfo->smem; - attr->constSizeBytes = kinfo->cmem; - attr->localSizeBytes = kinfo->lmem; - attr->numRegs = kinfo->regs; - if(kinfo->maxthreads > 0) - attr->maxThreadsPerBlock = kinfo->maxthreads; - else - attr->maxThreadsPerBlock = getMaxThreadsPerBlock(attr); -#if CUDART_VERSION >= 3000 - attr->ptxVersion = kinfo->ptx_version; - attr->binaryVersion = kinfo->sm_target; -#endif - } - return g_last_cudaError = cudaSuccess; -} - -cudaError_t CUDARTAPI cudaEventCreateWithFlags(cudaEvent_t *event, int flags) -{ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - CUevent_st *e = new CUevent_st(flags==cudaEventBlockingSync); - g_timer_events[e->get_uid()] = e; -#if CUDART_VERSION >= 3000 - *event = e; -#else - *event = e->get_uid(); -#endif - return g_last_cudaError = cudaSuccess; + return cudaFuncGetAttributesInternal(attr, hostFun ); } cudaError_t CUDARTAPI cudaDriverGetVersion(int *driverVersion) diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index fa5d02b..dd8b51d 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -26,6 +26,8 @@ class gpgpu_context { void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version=20, int no_of_ptx=0 ); void print_ptx_file( const char *p, unsigned source_num, const char *filename ); class symbol_table* init_parser(const char*); + class gpgpu_sim *gpgpu_ptx_sim_init_perf(); + struct _cuda_device_id *GPGPUSim_Init(); }; gpgpu_context* GPGPU_Context(); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 50a02fa..e91e2e0 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -648,7 +648,9 @@ class _cl_device_id *GPGPUSim_Init() { static _cl_device_id *the_device = NULL; if( !the_device ) { - gpgpu_sim *the_gpu = gpgpu_ptx_sim_init_perf(); + gpgpu_context *ctx; + ctx = GPGPU_Context(); + gpgpu_sim *the_gpu = ctx->gpgpu_ptx_sim_init_perf(); the_device = new _cl_device_id(the_gpu); } start_sim_thread(2); diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 2872b84..9094ec3 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -48,7 +48,6 @@ void set_ptx_warp_size(const struct core_config * warp_size) g_shader_core_config=warp_size; } -static bool g_debug_ir_generation=false; const char *g_filename; // the program intermediate representation... @@ -70,7 +69,7 @@ const char *decode_token( int type ) return g_ptx_token_decode[type].c_str(); } -void read_parser_environment_variables() +void ptx_recognizer::read_parser_environment_variables() { g_filename = getenv("PTX_SIM_KERNELFILE"); char *dbg_level = getenv("PTX_SIM_DEBUG"); diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h index 25c01fe..bc9a872 100644 --- a/src/cuda-sim/ptx_parser.h +++ b/src/cuda-sim/ptx_parser.h @@ -62,6 +62,7 @@ class ptx_recognizer { g_error_detected = 0; g_entry_func_param_index=0; g_func_info = NULL; + g_debug_ir_generation=false; } // global list yyscan_t scanner; @@ -103,6 +104,7 @@ class ptx_recognizer { unsigned g_entry_func_param_index; function_info *g_func_info; operand_info g_return_var; + bool g_debug_ir_generation; // member function list void init_directive_state(); @@ -169,6 +171,7 @@ class ptx_recognizer { void start_inst_group(); void end_inst_group(); bool check_for_duplicates( const char *identifier ); + void read_parser_environment_variables(); }; diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index b09674a..476a9d4 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -35,6 +35,7 @@ #include "gpgpu-sim/gpu-sim.h" #include "gpgpu-sim/icnt_wrapper.h" #include "stream_manager.h" +#include "../libcuda/gpgpu_context.h" #define MAX(a,b) (((a)>(b))?(a):(b)) @@ -208,12 +209,12 @@ void exit_simulation() extern bool g_cuda_launch_blocking; -gpgpu_sim *gpgpu_ptx_sim_init_perf() +gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf() { srand(1); print_splash(); read_sim_environment_variables(); - read_parser_environment_variables(); + ptx_parser->read_parser_environment_variables(); option_parser_t opp = option_parser_create(); ptx_reg_options(opp); diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h index 9ce7fef..a443151 100644 --- a/src/gpgpusim_entrypoint.h +++ b/src/gpgpusim_entrypoint.h @@ -74,7 +74,6 @@ struct GPGPUsim_ctx { }; -class gpgpu_sim *gpgpu_ptx_sim_init_perf(); void start_sim_thread(int api); class gpgpu_sim* g_the_gpu(); -- cgit v1.3 From 7e9591a24ffbe645cfaa7a84b33a82edb9fca0bc Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Sun, 7 Jul 2019 23:56:21 -0400 Subject: g_keep_intermediate_files Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 1 + libopencl/opencl_runtime_api.cc | 2 +- src/cuda-sim/ptx_ir.h | 3 --- src/cuda-sim/ptx_loader.cc | 7 +++---- src/cuda-sim/ptx_loader.h | 3 ++- 5 files changed, 7 insertions(+), 9 deletions(-) (limited to 'libopencl/opencl_runtime_api.cc') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index d819559..a8b60f4 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -32,6 +32,7 @@ class gpgpu_context { class symbol_table* init_parser(const char*); class gpgpu_sim *gpgpu_ptx_sim_init_perf(); struct _cuda_device_id *GPGPUSim_Init(); + void ptx_reg_options(option_parser_t opp); }; gpgpu_context* GPGPU_Context(); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index e91e2e0..03ec80c 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -537,7 +537,7 @@ void _cl_program::Build(const char *options) exit(1); } } - if( !g_keep_intermediate_files ) { + if( !ctx->ptxinfo->g_keep_intermediate_files ) { // clean up files... snprintf(commandline,1024,"rm -f %s", cl_fname ); int result = system(commandline); diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 1af85de..fd869c6 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1575,11 +1575,8 @@ struct textureInfo { extern std::map g_sym_name_to_symbol_table; -extern bool g_keep_intermediate_files; - void gpgpu_ptx_assemble( std::string kname, void *kinfo ); #include "../option_parser.h" -void ptx_reg_options(option_parser_t opp); unsigned ptx_kernel_shmem_size( void *kernel_impl ); unsigned ptx_kernel_nregs( void *kernel_impl ); diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index 921d1e6..e0d9a11 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -58,17 +58,16 @@ extern int ptxinfo_lex_destroy(yyscan_t scanner); static bool g_save_embedded_ptx; static int g_occupancy_sm_number; -bool g_keep_intermediate_files; bool m_ptx_save_converted_ptxplus; -bool keep_intermediate_files() {return g_keep_intermediate_files;} +bool ptxinfo_data::keep_intermediate_files() {return g_keep_intermediate_files;} -void ptx_reg_options(option_parser_t opp) +void gpgpu_context::ptx_reg_options(option_parser_t opp) { option_parser_register(opp, "-save_embedded_ptx", OPT_BOOL, &g_save_embedded_ptx, "saves ptx files embedded in binary as .ptx", "0"); - option_parser_register(opp, "-keep", OPT_BOOL, &g_keep_intermediate_files, + option_parser_register(opp, "-keep", OPT_BOOL, &(ptxinfo->g_keep_intermediate_files), "keep intermediate files created by GPGPU-Sim when interfacing with external programs", "0"); option_parser_register(opp, "-gpgpu_ptx_save_converted_ptxplus", OPT_BOOL, diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index c214b95..aa3e2f3 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -42,7 +42,9 @@ class ptxinfo_data{ unsigned col; const char *g_ptxinfo_filename; class gpgpu_context* gpgpu_ctx; + bool g_keep_intermediate_files; void ptxinfo_addinfo(); + bool keep_intermediate_files(); }; @@ -50,6 +52,5 @@ extern bool g_override_embedded_ptx; extern int no_of_ptx; //counter to track number of ptx files to be extracted in an application. char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str); -bool keep_intermediate_files(); #endif -- cgit v1.3 From 98bb2046a70715c8599561ca78425a7ff99ce53b Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 8 Jul 2019 14:21:08 -0400 Subject: Move cp_count Signed-off-by: Mengchi Zhang --- libopencl/opencl_runtime_api.cc | 4 +++- src/cuda-sim/cuda-sim.cc | 5 ++--- src/cuda-sim/cuda-sim.h | 4 +++- src/gpgpu-sim/gpu-sim.h | 4 ++-- src/gpgpusim_entrypoint.cc | 4 ++-- src/gpgpusim_entrypoint.h | 1 - 6 files changed, 12 insertions(+), 10 deletions(-) (limited to 'libopencl/opencl_runtime_api.cc') diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 03ec80c..d302ff8 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -877,6 +877,8 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, const cl_event * event_wait_list, cl_event * event) CL_API_SUFFIX__VERSION_1_0 { + gpgpu_context *ctx; + ctx = GPGPU_Context(); int _global_size[3]; int zeros[3] = { 0, 0, 0}; printf("\n\n\n"); @@ -956,7 +958,7 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, } kernel_info_t *grid = gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu); if ( g_ptx_sim_mode ) - gpgpu_opencl_ptx_sim_main_func( grid ); + ctx->func_sim->gpgpu_opencl_ptx_sim_main_func( grid ); else gpgpu_opencl_ptx_sim_main_perf( grid ); return CL_SUCCESS; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index df0bbd7..12f38af 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -61,7 +61,6 @@ int g_debug_execution = 0; int g_debug_thread_uid = 0; addr_t g_debug_pc = 0xBEEF1518; // Output debug information to file options -int cp_count; int cp_cta_resume; unsigned g_ptx_sim_num_insn = 0; @@ -2146,7 +2145,7 @@ unsigned max_cta (const struct gpgpu_ptx_sim_info *kernel_info, unsigned threads This function simulates the CUDA code functionally, it takes a kernel_info_t parameter which holds the data for the CUDA kernel to be executed !*/ -void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) +void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) { printf("GPGPU-Sim: Performing Functional Simulation, executing kernel %s...\n",kernel.name().c_str()); @@ -2312,7 +2311,7 @@ void functionalCoreSim::createWarp(unsigned warpId) void functionalCoreSim::execute(int inst_count, unsigned ctaid_cp) { - cp_count= m_gpu->checkpoint_insn_Y; + m_gpu->gpgpu_ctx->func_sim->cp_count= m_gpu->checkpoint_insn_Y; cp_cta_resume= m_gpu->checkpoint_CTA_t; initializeCTA(ctaid_cp); diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 96d34f5..4c8d953 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -55,7 +55,6 @@ extern class kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info * struct dim3 gridDim, struct dim3 blockDim, class gpgpu_t *gpu ); -extern void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL = false ); extern void print_splash(); extern void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size ); extern void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ); @@ -137,8 +136,11 @@ class cuda_sim { public: //global variables char *opcode_latency_int; + int cp_count; //global functions void ptx_opcocde_latency_options (option_parser_t opp); + void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL = false ); + int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid ); }; #endif diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 7eeb7dd..b47ab16 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -494,6 +494,8 @@ public: */ simt_core_cluster * getSIMTCluster(); + // backward pointer + class gpgpu_context* gpgpu_ctx; private: // clocks @@ -511,8 +513,6 @@ private: void gpgpu_debug(); ///// data ///// -// backward pointer - class gpgpu_context* gpgpu_ctx; class simt_core_cluster **m_cluster; class memory_partition_unit **m_memory_partition_unit; diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 5018305..d9d1023 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -134,7 +134,7 @@ void *gpgpu_sim_thread_concurrent(void*) if( GPGPUsim_ctx_ptr()->g_the_gpu->is_functional_sim()) { kernel_info_t * kernel = GPGPUsim_ctx_ptr()->g_the_gpu->get_functional_kernel(); assert(kernel); - gpgpu_cuda_ptx_sim_main_func(*kernel); + GPGPUsim_ctx_ptr()->gpgpu_ctx->func_sim->gpgpu_cuda_ptx_sim_main_func(*kernel); GPGPUsim_ctx_ptr()->g_the_gpu->finish_functional_sim(kernel); } @@ -287,7 +287,7 @@ int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ) /*! * This function call the CUDA PTX functional simulator */ -int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid ) +int cuda_sim::gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid ) { //calling the CUDA PTX simulator, sending the kernel by reference and a flag set to true, //the flag used by the function to distinguish OpenCL calls from the CUDA simulation calls which diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h index dfb82d0..887b3c8 100644 --- a/src/gpgpusim_entrypoint.h +++ b/src/gpgpusim_entrypoint.h @@ -83,6 +83,5 @@ struct GPGPUsim_ctx* GPGPUsim_ctx_ptr(); class stream_manager* g_stream_manager(); int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ); -int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid ); #endif -- cgit v1.3 From cda7a145b9e28eff0f3e9ac8197c2b6215755fc8 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 9 Jul 2019 14:32:24 -0400 Subject: Move g_ptx_kernel_count Signed-off-by: Mengchi Zhang --- libcuda/cuda_api_object.h | 5 ++++- libcuda/cuda_runtime_api.cc | 2 +- libcuda/gpgpu_context.h | 2 +- libopencl/opencl_runtime_api.cc | 2 +- src/cuda-sim/cuda-sim.cc | 13 ++++++------- src/cuda-sim/cuda-sim.h | 14 ++++++++------ src/gpgpu-sim/gpu-sim.cc | 4 ++-- 7 files changed, 23 insertions(+), 19 deletions(-) (limited to 'libopencl/opencl_runtime_api.cc') diff --git a/libcuda/cuda_api_object.h b/libcuda/cuda_api_object.h index 0054697..db5e6a4 100644 --- a/libcuda/cuda_api_object.h +++ b/libcuda/cuda_api_object.h @@ -169,9 +169,10 @@ private: class cuda_runtime_api { public: - cuda_runtime_api() { + cuda_runtime_api( gpgpu_context* ctx ) { g_glbmap = NULL; g_active_device = 0; //active gpu that runs the code + gpgpu_ctx = ctx; } // global list std::list cuobjdumpSectionList; @@ -187,6 +188,8 @@ class cuda_runtime_api { std::map pinned_memory_size; glbmap_entry_t* g_glbmap; int g_active_device; //active gpu that runs the code + // backward pointer + class gpgpu_context* gpgpu_ctx; // member function list void cuobjdumpInit(); void extract_code_using_cuobjdump(); diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 43c5e1f..43c8bae 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -3545,7 +3545,7 @@ kernel_info_t * cuda_runtime_api::gpgpu_cuda_ptx_sim_init_grid( const char *host } entry->finalize(result->get_param_memory()); - g_ptx_kernel_count++; + gpgpu_ctx->func_sim->g_ptx_kernel_count++; fflush(stdout); if(g_debug_execution >= 4){ diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 2e21009..a2ae7b6 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -10,7 +10,7 @@ class gpgpu_context { public: gpgpu_context() { g_global_allfiles_symbol_table = NULL; - api = new cuda_runtime_api(); + api = new cuda_runtime_api(this); ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); the_gpgpusim = new GPGPUsim_ctx(this); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index d302ff8..0a6eb3e 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -956,7 +956,7 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); } - kernel_info_t *grid = gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu); + kernel_info_t *grid = ctx->func_sim->gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu); if ( g_ptx_sim_mode ) ctx->func_sim->gpgpu_opencl_ptx_sim_main_func( grid ); else diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index e86395d..0ed125a 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -56,7 +56,6 @@ typedef void * yyscan_t; int gpgpu_ptx_instruction_classification; void ** g_inst_classification_stat = NULL; void ** g_inst_op_classification_stat= NULL; -int g_ptx_kernel_count = -1; // used for classification stat collection purposes int g_debug_execution = 0; int g_debug_thread_uid = 0; addr_t g_debug_pc = 0xBEEF1518; @@ -1480,7 +1479,7 @@ bool ptx_debug_exec_dump_cond(int thd_uid, addr_t pc) return false; } -void init_inst_classification_stat() +void cuda_sim::init_inst_classification_stat() { static std::set init; if( init.find(g_ptx_kernel_count) != init.end() ) @@ -1690,7 +1689,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) ptx_file_line_stats_add_exec_count(pI); if ( gpgpu_ptx_instruction_classification ) { - init_inst_classification_stat(); + m_gpu->gpgpu_ctx->func_sim->init_inst_classification_stat(); unsigned space_type=0; switch ( pI->get_space().get_type() ) { case global_space: space_type = 10; break; @@ -1706,9 +1705,9 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) space_type = 0 ; break; } - StatAddSample( g_inst_classification_stat[g_ptx_kernel_count], op_classification); - if (space_type) StatAddSample( g_inst_classification_stat[g_ptx_kernel_count], ( int )space_type); - StatAddSample( g_inst_op_classification_stat[g_ptx_kernel_count], (int) pI->get_opcode() ); + StatAddSample( g_inst_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], op_classification); + if (space_type) StatAddSample( g_inst_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], ( int )space_type); + StatAddSample( g_inst_op_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], (int) pI->get_opcode() ); } if ( (m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn % 100000) == 0 ) { dim3 ctaid = get_ctaid(); @@ -1917,7 +1916,7 @@ size_t get_kernel_code_size( class function_info *entry ) } -kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, +kernel_info_t *cuda_sim::gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, gpgpu_ptx_sim_arg_list_t args, struct dim3 gridDim, struct dim3 blockDim, diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 16ee46e..e259f1f 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -46,13 +46,7 @@ extern int g_debug_execution; extern int g_debug_thread_uid; extern void ** g_inst_classification_stat; extern void ** g_inst_op_classification_stat; -extern int g_ptx_kernel_count; // used for classification stat collection purposes -extern class kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, - gpgpu_ptx_sim_arg_list_t args, - struct dim3 gridDim, - struct dim3 blockDim, - class gpgpu_t *gpu ); extern void print_splash(); extern void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size ); extern void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ); @@ -134,6 +128,7 @@ class cuda_sim { public: cuda_sim() { g_ptx_sim_num_insn = 0; + g_ptx_kernel_count = -1; // used for classification stat collection purposes } //global variables char *opcode_latency_int; @@ -151,10 +146,17 @@ class cuda_sim { int g_ptxinfo_error_detected; unsigned g_ptx_sim_num_insn; char *cdp_latency_str; + int g_ptx_kernel_count; // used for classification stat collection purposes //global functions void ptx_opcocde_latency_options (option_parser_t opp); void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL = false ); int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid ); + void init_inst_classification_stat(); + kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry, + gpgpu_ptx_sim_arg_list_t args, + struct dim3 gridDim, + struct dim3 blockDim, + gpgpu_t *gpu ); }; #endif diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 4f9ccbf..9f47067 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1201,8 +1201,8 @@ void gpgpu_sim::gpu_print_stat() insn_warp_occ_print(stdout); } if ( gpgpu_ptx_instruction_classification ) { - StatDisp( g_inst_classification_stat[g_ptx_kernel_count]); - StatDisp( g_inst_op_classification_stat[g_ptx_kernel_count]); + StatDisp( g_inst_classification_stat[gpgpu_ctx->func_sim->g_ptx_kernel_count]); + StatDisp( g_inst_op_classification_stat[gpgpu_ctx->func_sim->g_ptx_kernel_count]); } #ifdef GPGPUSIM_POWER_MODEL -- cgit v1.3 From 4d4d5938d715d2b79a617c32583184426b4a642d Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Tue, 9 Jul 2019 23:16:17 -0400 Subject: Move g_ptx_sim_mode Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 6 +++--- libcuda/gpgpu_context.h | 9 ++++++++- libopencl/opencl_runtime_api.cc | 6 +++--- src/cuda-sim/cuda-sim.cc | 6 ++---- src/cuda-sim/cuda-sim.h | 10 +++++++--- src/cuda-sim/cuda_device_runtime.cc | 7 ++++--- src/cuda-sim/cuda_device_runtime.h | 18 ++++++++++++++++-- src/gpgpu-sim/gpu-sim.cc | 4 ++-- src/gpgpusim_entrypoint.cc | 2 +- 9 files changed, 46 insertions(+), 22 deletions(-) (limited to 'libopencl/opencl_runtime_api.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index bbbaf23..59d2a60 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1076,7 +1076,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx = CUctx_st* context = GPGPUSim_Context(); char *mode = getenv("PTX_SIM_MODE_FUNC"); if( mode ) - sscanf(mode,"%u", &g_ptx_sim_mode); + sscanf(mode,"%u", &(ctx->func_sim->g_ptx_sim_mode)); gpgpusim_ptx_assert( !ctx->api->g_cuda_launch_stack.empty(), "empty launch stack" ); kernel_config config = ctx->api->g_cuda_launch_stack.back(); { @@ -1092,7 +1092,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx = } struct CUstream_st *stream = config.get_stream(); printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun, - g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 ); + (ctx->func_sim->g_ptx_sim_mode)?"functional simulation":"performance simulation", stream?stream->get_uid():0 ); kernel_info_t *grid = ctx->api->gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context); //do dynamic PDOM analysis for performance simulation scenario std::string kname = grid->name(); @@ -1143,7 +1143,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx = } printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n", kname.c_str(), stream?stream->get_uid():0, gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z ); - stream_operation op(grid,g_ptx_sim_mode,stream); + stream_operation op(grid,ctx->func_sim->g_ptx_sim_mode,stream); g_stream_manager()->push(op); ctx->api->g_cuda_launch_stack.pop_back(); return g_last_cudaError = cudaSuccess; diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index a2ae7b6..3c9f87c 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -5,6 +5,7 @@ #include "../src/cuda-sim/ptx_parser.h" #include "../src/gpgpusim_entrypoint.h" #include "../src/cuda-sim/cuda-sim.h" +#include "../src/cuda-sim/cuda_device_runtime.h" class gpgpu_context { public: @@ -14,7 +15,10 @@ class gpgpu_context { ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); the_gpgpusim = new GPGPUsim_ctx(this); - func_sim = new cuda_sim(); + func_sim = new cuda_sim(this); +#if (CUDART_VERSION >= 5000) + device_runtime = new cuda_device_runtime(this); +#endif } // global list symbol_table *g_global_allfiles_symbol_table; @@ -25,6 +29,9 @@ class gpgpu_context { ptx_recognizer* ptx_parser; GPGPUsim_ctx* the_gpgpusim; cuda_sim* func_sim; +#if (CUDART_VERSION >= 5000) + cuda_device_runtime* device_runtime; +#endif // member function list void cuobjdumpParseBinary(unsigned int handle); class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 0a6eb3e..aaaec4f 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -884,9 +884,9 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, printf("\n\n\n"); char *mode = getenv("PTX_SIM_MODE_FUNC"); if ( mode ) - sscanf(mode,"%u", &g_ptx_sim_mode); + sscanf(mode,"%u", &(ctx->func_sim->g_ptx_sim_mode)); printf("GPGPU-Sim OpenCL API: clEnqueueNDRangeKernel '%s' (mode=%s)\n", kernel->name().c_str(), - g_ptx_sim_mode?"functional simulation":"performance simulation"); + (ctx->func_sim->g_ptx_sim_mode)?"functional simulation":"performance simulation"); if ( !work_dim || work_dim > 3 ) return CL_INVALID_WORK_DIMENSION; size_t _local_size[3]; if( local_work_size != NULL ) { @@ -957,7 +957,7 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); } kernel_info_t *grid = ctx->func_sim->gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu); - if ( g_ptx_sim_mode ) + if ( ctx->func_sim->g_ptx_sim_mode ) ctx->func_sim->gpgpu_opencl_ptx_sim_main_func( grid ); else gpgpu_opencl_ptx_sim_main_perf( grid ); diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index b3e2965..7a7d205 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -2033,13 +2033,11 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co fflush(stdout); } -int g_ptx_sim_mode; // if non-zero run functional simulation only (i.e., no notion of a clock cycle) - extern int ptx_debug; bool g_cuda_launch_blocking = false; -void read_sim_environment_variables() +void cuda_sim::read_sim_environment_variables() { ptx_debug = 0; g_debug_execution = 0; @@ -2185,7 +2183,7 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL cta.execute(cp_count,temp); #if (CUDART_VERSION >= 5000) - launch_all_device_kernels(); + gpgpu_ctx->device_runtime->launch_all_device_kernels(); #endif } else diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 25ebf7b..3c4336d 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -36,12 +36,12 @@ #include #include"ptx_sim.h" +class gpgpu_context; class memory_space; class function_info; class symbol_table; extern const char *g_gpgpusim_version_string; -extern int g_ptx_sim_mode; extern int g_debug_execution; extern int g_debug_thread_uid; extern void ** g_inst_classification_stat; @@ -50,7 +50,6 @@ extern void ** g_inst_op_classification_stat; extern void print_splash(); extern void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu ); -extern void read_sim_environment_variables(); extern void ptxinfo_opencl_addinfo( std::map &kernels ); unsigned ptx_sim_init_thread( kernel_info_t &kernel, class ptx_thread_info** thread_info, @@ -124,9 +123,10 @@ struct gpgpu_ptx_sim_info get_ptxinfo(); class cuda_sim { public: - cuda_sim() { + cuda_sim( gpgpu_context* ctx ) { g_ptx_sim_num_insn = 0; g_ptx_kernel_count = -1; // used for classification stat collection purposes + gpgpu_ctx = ctx; } //global variables char *opcode_latency_int; @@ -147,6 +147,9 @@ class cuda_sim { int g_ptx_kernel_count; // used for classification stat collection purposes std::map g_global_name_lookup; // indexed by hostVar std::map g_const_name_lookup; // indexed by hostVar + int g_ptx_sim_mode; // if non-zero run functional simulation only (i.e., no notion of a clock cycle) + // backward pointer + class gpgpu_context* gpgpu_ctx; //global functions void ptx_opcocde_latency_options (option_parser_t opp); void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL = false ); @@ -159,6 +162,7 @@ class cuda_sim { gpgpu_t *gpu ); void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ); void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size ); + void read_sim_environment_variables(); }; #endif diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index be8369f..354fa79 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -20,6 +20,7 @@ unsigned long long g_max_total_param_size = 0; #include "../stream_manager.h" #include "../gpgpusim_entrypoint.h" #include "cuda_device_runtime.h" +#include "../../libcuda/gpgpu_context.h" #define DEV_RUNTIME_REPORT(a) \ if( g_debug_execution ) { \ @@ -318,17 +319,17 @@ void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_ } -void launch_one_device_kernel() { +void cuda_device_runtime::launch_one_device_kernel() { if(!g_cuda_device_launch_op.empty()) { device_launch_operation_t &op = g_cuda_device_launch_op.front(); - stream_operation stream_op = stream_operation(op.grid, g_ptx_sim_mode, op.stream); + stream_operation stream_op = stream_operation(op.grid, gpgpu_ctx->func_sim->g_ptx_sim_mode, op.stream); g_stream_manager()->push(stream_op); g_cuda_device_launch_op.pop_front(); } } -void launch_all_device_kernels() { +void cuda_device_runtime::launch_all_device_kernels() { while(!g_cuda_device_launch_op.empty()) { launch_one_device_kernel(); } diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h index 6dbcd71..851fed2 100644 --- a/src/cuda-sim/cuda_device_runtime.h +++ b/src/cuda-sim/cuda_device_runtime.h @@ -6,6 +6,20 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); -void launch_all_device_kernels(); -void launch_one_device_kernel(); +#endif +#if (CUDART_VERSION >= 5000) + +class gpgpu_context; + +class cuda_device_runtime { + public: + cuda_device_runtime( gpgpu_context* ctx ) { + gpgpu_ctx = ctx; + } + // backward pointer + class gpgpu_context* gpgpu_ctx; + void launch_all_device_kernels(); + void launch_one_device_kernel(); +}; + #endif diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 9f47067..bbcc078 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -503,7 +503,7 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) &gpgpu_ptx_instruction_classification, "if enabled will classify ptx instruction types per kernel (Max 255 kernels now)", "0"); - option_parser_register(opp, "-gpgpu_ptx_sim_mode", OPT_INT32, &g_ptx_sim_mode, + option_parser_register(opp, "-gpgpu_ptx_sim_mode", OPT_INT32, &(gpgpu_ctx->func_sim->g_ptx_sim_mode), "Select between Performance (default) or Functional simulation (1)", "0"); option_parser_register(opp, "-gpgpu_clock_domains", OPT_CSTR, &gpgpu_clock_domains, @@ -1753,7 +1753,7 @@ void gpgpu_sim::cycle() #if (CUDART_VERSION >= 5000) //launch device kernel - launch_one_device_kernel(); + gpgpu_ctx->device_runtime->launch_one_device_kernel(); #endif } } diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index d9d1023..683a695 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -213,7 +213,7 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf() { srand(1); print_splash(); - read_sim_environment_variables(); + func_sim->read_sim_environment_variables(); ptx_parser->read_parser_environment_variables(); option_parser_t opp = option_parser_create(); -- cgit v1.3 From c58d52ccde1eea6c45d0f8d2d5411bcf1ebab4f9 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 10 Jul 2019 00:45:39 -0400 Subject: Move g_globals and g_constants Signed-off-by: Mengchi Zhang --- libopencl/opencl_runtime_api.cc | 10 +++++----- src/cuda-sim/cuda-sim.cc | 5 +---- src/cuda-sim/cuda-sim.h | 4 +++- src/cuda-sim/ptx_parser.cc | 7 ++----- src/stream_manager.cc | 5 +++-- 5 files changed, 14 insertions(+), 17 deletions(-) (limited to 'libopencl/opencl_runtime_api.cc') diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index aaaec4f..7f029c7 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -950,11 +950,11 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, gpgpu_t *gpu = command_queue->get_device()->the_device(); if (kernel->get_implementation()->get_ptx_version().ver() <3.0){ - gpgpu_ptx_sim_memcpy_symbol( "%_global_size", _global_size, 3 * sizeof(int), 0, 1, gpu ); - gpgpu_ptx_sim_memcpy_symbol( "%_work_dim", &work_dim, 1 * sizeof(int), 0, 1, gpu ); - gpgpu_ptx_sim_memcpy_symbol( "%_global_num_groups", &GridDim, 3 * sizeof(int), 0, 1, gpu ); - gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); - gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); + ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_size", _global_size, 3 * sizeof(int), 0, 1, gpu ); + ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_work_dim", &work_dim, 1 * sizeof(int), 0, 1, gpu ); + ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_num_groups", &GridDim, 3 * sizeof(int), 0, 1, gpu ); + ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); + ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu ); } kernel_info_t *grid = ctx->func_sim->gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu); if ( ctx->func_sim->g_ptx_sim_mode ) diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 7e8aab9..3d1da62 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1942,9 +1942,6 @@ void print_splash() } } -std::set g_globals; -std::set g_constants; - void cuda_sim::gpgpu_ptx_sim_register_const_variable(void *hostVar, const char *deviceName, size_t size ) { printf("GPGPU-Sim PTX registering constant %s (%zu bytes) to name mapping\n", deviceName, size ); @@ -1957,7 +1954,7 @@ void cuda_sim::gpgpu_ptx_sim_register_global_variable(void *hostVar, const char g_global_name_lookup[hostVar] = deviceName; } -void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu ) +void cuda_sim::gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu ) { printf("GPGPU-Sim PTX: starting gpgpu_ptx_sim_memcpy_symbol with hostVar 0x%p\n", hostVar); bool found_sym = false; diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 628e434..5c95100 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -46,7 +46,6 @@ extern int g_debug_execution; extern int g_debug_thread_uid; extern void print_splash(); -extern void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu ); extern void ptxinfo_opencl_addinfo( std::map &kernels ); unsigned ptx_sim_init_thread( kernel_info_t &kernel, @@ -161,6 +160,8 @@ class cuda_sim { bool g_cuda_launch_blocking; void ** g_inst_classification_stat; void ** g_inst_op_classification_stat; + std::set g_globals; + std::set g_constants; // backward pointer class gpgpu_context* gpgpu_ctx; //global functions @@ -179,6 +180,7 @@ class cuda_sim { void set_param_gpgpu_num_shaders(int num_shaders); struct rec_pts find_reconvergence_points( function_info *finfo ); address_type get_converge_point( address_type pc ); + void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu ); }; #endif diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 0139534..269ec4d 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -322,9 +322,6 @@ bool ptx_recognizer::check_for_duplicates( const char *identifier ) return ( s != NULL ); } -extern std::set g_globals; -extern std::set g_constants; - // Returns padding that needs to be inserted ahead of address to make it aligned to min(size, maxalign) /* * @param address the address in bytes @@ -453,7 +450,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi g_current_symbol_table->alloc_global( num_bits/8 + addr_pad ); } if( g_current_symbol_table == g_global_symbol_table ) { - g_constants.insert( identifier ); + gpgpu_ctx->func_sim->g_constants.insert( identifier ); } assert( g_current_symbol_table != NULL ); g_sym_name_to_symbol_table[ identifier ] = g_current_symbol_table; @@ -471,7 +468,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi fflush(stdout); g_last_symbol->set_address( addr+addr_pad ); g_current_symbol_table->alloc_global( num_bits/8 + addr_pad ); - g_globals.insert( identifier ); + gpgpu_ctx->func_sim->g_globals.insert( identifier ); assert( g_current_symbol_table != NULL ); g_sym_name_to_symbol_table[ identifier ] = g_current_symbol_table; break; diff --git a/src/stream_manager.cc b/src/stream_manager.cc index e07f4e4..be3dd71 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -29,6 +29,7 @@ #include "gpgpusim_entrypoint.h" #include "cuda-sim/cuda-sim.h" #include "gpgpu-sim/gpu-sim.h" +#include "../libcuda/gpgpu_context.h" unsigned CUstream_st::sm_next_stream_uid = 0; @@ -150,13 +151,13 @@ bool stream_operation::do_operation( gpgpu_sim *gpu ) case stream_memcpy_to_symbol: if(g_debug_execution >= 3) printf("memcpy to symbol\n"); - gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_src,m_cnt,m_offset,1,gpu); + gpu->gpgpu_ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_src,m_cnt,m_offset,1,gpu); m_stream->record_next_done(); break; case stream_memcpy_from_symbol: if(g_debug_execution >= 3) printf("memcpy from symbol\n"); - gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_dst,m_cnt,m_offset,0,gpu); + gpu->gpgpu_ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_dst,m_cnt,m_offset,0,gpu); m_stream->record_next_done(); break; case stream_kernel_launch: -- cgit v1.3 From 48887cfe0261abde6de23fc5d1d76694426b7e8e Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Thu, 12 Sep 2019 02:01:01 -0400 Subject: Remove GPGPUsim_ctx_ptr() Signed-off-by: Mengchi Zhang --- libcuda/cuda_runtime_api.cc | 116 ++++++++++++++++++++--------- libcuda/gpgpu_context.h | 5 ++ libopencl/opencl_runtime_api.cc | 8 +- src/cuda-sim/cuda-sim.cc | 2 +- src/gpgpu-sim/gpu-sim.cc | 4 +- src/gpgpusim_entrypoint.cc | 161 +++++++++++++++++++--------------------- src/gpgpusim_entrypoint.h | 6 -- 7 files changed, 167 insertions(+), 135 deletions(-) (limited to 'libopencl/opencl_runtime_api.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 9459b64..5cefe60 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -150,9 +150,6 @@ #endif -extern void synchronize(); -extern void exit_simulation(); - /*DEVICE_BUILTIN*/ struct cudaArray { @@ -242,8 +239,8 @@ struct _cuda_device_id *gpgpu_context::GPGPUSim_Init() prop->maxThreadsPerMultiProcessor = the_gpu->threads_per_core(); #endif the_gpu->set_prop(prop); - GPGPUsim_ctx_ptr()->the_cude_device = new _cuda_device_id(the_gpu); - the_device = GPGPUsim_ctx_ptr()->the_cude_device; + the_gpgpusim->the_cude_device = new _cuda_device_id(the_gpu); + the_device = the_gpgpusim->the_cude_device; } start_sim_thread(1); return the_device; @@ -1876,6 +1873,28 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroyInternal(cudaStream_t stream, gp return g_last_cudaError = cudaSuccess; } +__host__ cudaError_t CUDARTAPI cudaStreamSynchronizeInternal(cudaStream_t stream, gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } +#if (CUDART_VERSION >= 3000) + if( stream == NULL ) + ctx->synchronize(); + return g_last_cudaError = cudaSuccess; + stream->synchronize(); +#else + printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); +#endif + return g_last_cudaError = cudaSuccess; +} + void __cudaRegisterTextureInternal( void **fatCubinHandle, const struct textureReference *hostVar, @@ -2065,6 +2084,53 @@ __host__ cudaError_t CUDARTAPI cudaStreamWaitEventInternal(cudaStream_t stream, return g_last_cudaError = cudaSuccess; } +__host__ cudaError_t CUDARTAPI cudaThreadExitInternal(gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + ctx->exit_simulation(); + return g_last_cudaError = cudaSuccess; +} + +__host__ cudaError_t CUDARTAPI cudaThreadSynchronizeInternal(gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + //Called on host side + ctx->synchronize(); + return g_last_cudaError = cudaSuccess; +} + +cudaError_t CUDARTAPI cudaDeviceSynchronizeInternal(gpgpu_context* gpgpu_ctx = NULL) +{ + gpgpu_context *ctx; + if (gpgpu_ctx){ + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if(g_debug_execution >= 3){ + announce_call(__my_func__); + } + //Blocks until the device has completed all preceding requested tasks + ctx->synchronize(); + return g_last_cudaError = cudaSuccess; +} + /******************************************************************************* * * * * @@ -2592,18 +2658,7 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) __host__ cudaError_t CUDARTAPI cudaStreamSynchronize(cudaStream_t stream) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } -#if (CUDART_VERSION >= 3000) - if( stream == NULL ) - synchronize(); - return g_last_cudaError = cudaSuccess; - stream->synchronize(); -#else - printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); -#endif - return g_last_cudaError = cudaSuccess; + return cudaStreamSynchronizeInternal(stream); } __host__ cudaError_t CUDARTAPI cudaStreamQuery(cudaStream_t stream) @@ -2722,22 +2777,13 @@ __host__ cudaError_t CUDARTAPI cudaEventElapsedTime(float *ms, cudaEvent_t start __host__ cudaError_t CUDARTAPI cudaThreadExit(void) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - exit_simulation(); - return g_last_cudaError = cudaSuccess; + return cudaThreadExitInternal(); } __host__ cudaError_t CUDARTAPI cudaThreadSynchronize(void) { - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - //Called on host side - synchronize(); - return g_last_cudaError = cudaSuccess; -}; + return cudaThreadSynchronizeInternal(); +} int CUDARTAPI __cudaSynchronizeThreads(void**, void*) { @@ -3441,15 +3487,11 @@ cudaError_t cudaDeviceReset ( void ) { } return g_last_cudaError = cudaSuccess; } -cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ - if(g_debug_execution >= 3){ - announce_call(__my_func__); - } - //Blocks until the device has completed all preceding requested tasks - synchronize(); - return g_last_cudaError = cudaSuccess; -} +cudaError_t CUDARTAPI cudaDeviceSynchronize(void) +{ + return cudaDeviceSynchronizeInternal(); +} void __cudaRegisterShared( void **fatCubinHandle, diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 45c5cdd..61d7507 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -52,6 +52,10 @@ class gpgpu_context { cuda_device_runtime* device_runtime; ptx_stats* stats; // member function list + void synchronize(); + void exit_simulation(); + void print_simulation_time(); + int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ); void cuobjdumpParseBinary(unsigned int handle); class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); class symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename ); @@ -60,6 +64,7 @@ class gpgpu_context { void print_ptx_file( const char *p, unsigned source_num, const char *filename ); class symbol_table* init_parser(const char*); class gpgpu_sim *gpgpu_ptx_sim_init_perf(); + void start_sim_thread(int api); struct _cuda_device_id *GPGPUSim_Init(); void ptx_reg_options(option_parser_t opp); const ptx_instruction* pc_to_instruction(unsigned pc); diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 7f029c7..b032c05 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -647,13 +647,13 @@ unsigned _cl_kernel::sm_context_uid = 0; class _cl_device_id *GPGPUSim_Init() { static _cl_device_id *the_device = NULL; + gpgpu_context *ctx; + ctx = GPGPU_Context(); if( !the_device ) { - gpgpu_context *ctx; - ctx = GPGPU_Context(); gpgpu_sim *the_gpu = ctx->gpgpu_ptx_sim_init_perf(); the_device = new _cl_device_id(the_gpu); } - start_sim_thread(2); + ctx->start_sim_thread(2); return the_device; } @@ -960,7 +960,7 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, if ( ctx->func_sim->g_ptx_sim_mode ) ctx->func_sim->gpgpu_opencl_ptx_sim_main_func( grid ); else - gpgpu_opencl_ptx_sim_main_perf( grid ); + ctx->gpgpu_opencl_ptx_sim_main_perf( grid ); return CL_SUCCESS; } diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index fc1ea8e..7a130ea 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -2203,7 +2203,7 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL //g_simulation_starttime is initilized by gpgpu_ptx_sim_init_perf() in gpgpusim_entrypoint.cc upon starting gpgpu-sim time_t end_time, elapsed_time, days, hrs, minutes, sec; end_time = time((time_t *)NULL); - elapsed_time = MAX(end_time - GPGPUsim_ctx_ptr()->g_simulation_starttime, 1); + elapsed_time = MAX(end_time - gpgpu_ctx->the_gpgpusim->g_simulation_starttime, 1); //calculating and printing simulation time in terms of days, hours, minutes and seconds diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index d236c74..56ea8c4 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1144,7 +1144,7 @@ void gpgpu_sim::gpu_print_stat() time_t curr_time; time(&curr_time); - unsigned long long elapsed_time = MAX( curr_time - GPGPUsim_ctx_ptr()->g_simulation_starttime, 1 ); + unsigned long long elapsed_time = MAX( curr_time - gpgpu_ctx->the_gpgpusim->g_simulation_starttime, 1 ); printf( "gpu_total_sim_rate=%u\n", (unsigned)( ( gpu_tot_sim_insn + gpu_sim_insn ) / elapsed_time ) ); //shader_print_l1_miss_stat( stdout ); @@ -1717,7 +1717,7 @@ void gpgpu_sim::cycle() time_t days, hrs, minutes, sec; time_t curr_time; time(&curr_time); - unsigned long long elapsed_time = MAX(curr_time - GPGPUsim_ctx_ptr()->g_simulation_starttime, 1); + unsigned long long elapsed_time = MAX(curr_time - gpgpu_ctx->the_gpgpusim->g_simulation_starttime, 1); if ( (elapsed_time - last_liveness_message_time) >= m_config.liveness_message_freq && DTRACE(LIVENESS) ) { days = elapsed_time/(3600*24); hrs = elapsed_time/3600 - 24*days; diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 5ce40c6..846773d 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -43,38 +43,28 @@ static int sg_argc = 3; static const char *sg_argv[] = {"", "-config","gpgpusim.config"}; -GPGPUsim_ctx* the_gpgpusim = NULL; - -GPGPUsim_ctx* GPGPUsim_ctx_ptr(){ - if(the_gpgpusim == NULL) - the_gpgpusim = GPGPU_Context()->the_gpgpusim; - - return the_gpgpusim; -} - -static void print_simulation_time(); - -void *gpgpu_sim_thread_sequential(void*) +void * gpgpu_sim_thread_sequential(void * ctx_ptr) { + gpgpu_context * ctx = (gpgpu_context *)ctx_ptr; // at most one kernel running at a time bool done; do { - sem_wait(&(GPGPUsim_ctx_ptr()->g_sim_signal_start)); + sem_wait(&(ctx->the_gpgpusim->g_sim_signal_start)); done = true; - if( GPGPUsim_ctx_ptr()->g_the_gpu->get_more_cta_left() ) { + if( ctx->the_gpgpusim->g_the_gpu->get_more_cta_left() ) { done = false; - GPGPUsim_ctx_ptr()->g_the_gpu->init(); - while( GPGPUsim_ctx_ptr()->g_the_gpu->active() ) { - GPGPUsim_ctx_ptr()->g_the_gpu->cycle(); - GPGPUsim_ctx_ptr()->g_the_gpu->deadlock_check(); + ctx->the_gpgpusim->g_the_gpu->init(); + while( ctx->the_gpgpusim->g_the_gpu->active() ) { + ctx->the_gpgpusim->g_the_gpu->cycle(); + ctx->the_gpgpusim->g_the_gpu->deadlock_check(); } - GPGPUsim_ctx_ptr()->g_the_gpu->print_stats(); - GPGPUsim_ctx_ptr()->g_the_gpu->update_stats(); - print_simulation_time(); + ctx->the_gpgpusim->g_the_gpu->print_stats(); + ctx->the_gpgpusim->g_the_gpu->update_stats(); + ctx->print_simulation_time(); } - sem_post(&(GPGPUsim_ctx_ptr()->g_sim_signal_finish)); + sem_post(&(ctx->the_gpgpusim->g_sim_signal_finish)); } while(!done); - sem_post(&(GPGPUsim_ctx_ptr()->g_sim_signal_exit)); + sem_post(&(ctx->the_gpgpusim->g_sim_signal_exit)); return NULL; } @@ -86,8 +76,9 @@ static void termination_callback() fflush(stdout); } -void *gpgpu_sim_thread_concurrent(void*) +void *gpgpu_sim_thread_concurrent(void * ctx_ptr) { + gpgpu_context * ctx = (gpgpu_context *)ctx_ptr; atexit(termination_callback); // concurrent kernel execution simulation thread do { @@ -95,19 +86,19 @@ void *gpgpu_sim_thread_concurrent(void*) printf("GPGPU-Sim: *** simulation thread starting and spinning waiting for work ***\n"); fflush(stdout); } - while( GPGPUsim_ctx_ptr()->g_stream_manager->empty_protected() && !GPGPUsim_ctx_ptr()->g_sim_done ) + while( ctx->the_gpgpusim->g_stream_manager->empty_protected() && !ctx->the_gpgpusim->g_sim_done ) ; if(g_debug_execution >= 3) { printf("GPGPU-Sim: ** START simulation thread (detected work) **\n"); - GPGPUsim_ctx_ptr()->g_stream_manager->print(stdout); + ctx->the_gpgpusim->g_stream_manager->print(stdout); fflush(stdout); } - pthread_mutex_lock(&(GPGPUsim_ctx_ptr()->g_sim_lock)); - GPGPUsim_ctx_ptr()->g_sim_active = true; - pthread_mutex_unlock(&(GPGPUsim_ctx_ptr()->g_sim_lock)); + pthread_mutex_lock(&(ctx->the_gpgpusim->g_sim_lock)); + ctx->the_gpgpusim->g_sim_active = true; + pthread_mutex_unlock(&(ctx->the_gpgpusim->g_sim_lock)); bool active = false; bool sim_cycles = false; - GPGPUsim_ctx_ptr()->g_the_gpu->init(); + ctx->the_gpgpusim->g_the_gpu->init(); do { // check if a kernel has completed // launch operation on device if one is pending and can be run @@ -119,82 +110,82 @@ void *gpgpu_sim_thread_concurrent(void*) // another kernel, the gpu is not re-initialized and the inter-kernel // behaviour may be incorrect. Check that a kernel has finished and // no other kernel is currently running. - if(GPGPUsim_ctx_ptr()->g_stream_manager->operation(&sim_cycles) && !GPGPUsim_ctx_ptr()->g_the_gpu->active()) + if(ctx->the_gpgpusim->g_stream_manager->operation(&sim_cycles) && !ctx->the_gpgpusim->g_the_gpu->active()) break; //functional simulation - if( GPGPUsim_ctx_ptr()->g_the_gpu->is_functional_sim()) { - kernel_info_t * kernel = GPGPUsim_ctx_ptr()->g_the_gpu->get_functional_kernel(); + if( ctx->the_gpgpusim->g_the_gpu->is_functional_sim()) { + kernel_info_t * kernel = ctx->the_gpgpusim->g_the_gpu->get_functional_kernel(); assert(kernel); - GPGPUsim_ctx_ptr()->gpgpu_ctx->func_sim->gpgpu_cuda_ptx_sim_main_func(*kernel); - GPGPUsim_ctx_ptr()->g_the_gpu->finish_functional_sim(kernel); + ctx->the_gpgpusim->gpgpu_ctx->func_sim->gpgpu_cuda_ptx_sim_main_func(*kernel); + ctx->the_gpgpusim->g_the_gpu->finish_functional_sim(kernel); } //performance simulation - if( GPGPUsim_ctx_ptr()->g_the_gpu->active() ) { - GPGPUsim_ctx_ptr()->g_the_gpu->cycle(); + if( ctx->the_gpgpusim->g_the_gpu->active() ) { + ctx->the_gpgpusim->g_the_gpu->cycle(); sim_cycles = true; - GPGPUsim_ctx_ptr()->g_the_gpu->deadlock_check(); + ctx->the_gpgpusim->g_the_gpu->deadlock_check(); }else { - if(GPGPUsim_ctx_ptr()->g_the_gpu->cycle_insn_cta_max_hit()){ - GPGPUsim_ctx_ptr()->g_stream_manager->stop_all_running_kernels(); - GPGPUsim_ctx_ptr()->g_sim_done = true; - GPGPUsim_ctx_ptr()->break_limit = true; + if(ctx->the_gpgpusim->g_the_gpu->cycle_insn_cta_max_hit()){ + ctx->the_gpgpusim->g_stream_manager->stop_all_running_kernels(); + ctx->the_gpgpusim->g_sim_done = true; + ctx->the_gpgpusim->break_limit = true; } } - active=GPGPUsim_ctx_ptr()->g_the_gpu->active() || !(GPGPUsim_ctx_ptr()->g_stream_manager->empty_protected()); + active=ctx->the_gpgpusim->g_the_gpu->active() || !(ctx->the_gpgpusim->g_stream_manager->empty_protected()); - } while( active && !GPGPUsim_ctx_ptr()->g_sim_done); + } while( active && !ctx->the_gpgpusim->g_sim_done); if(g_debug_execution >= 3) { printf("GPGPU-Sim: ** STOP simulation thread (no work) **\n"); fflush(stdout); } if(sim_cycles) { - GPGPUsim_ctx_ptr()->g_the_gpu->print_stats(); - GPGPUsim_ctx_ptr()->g_the_gpu->update_stats(); - print_simulation_time(); + ctx->the_gpgpusim->g_the_gpu->print_stats(); + ctx->the_gpgpusim->g_the_gpu->update_stats(); + ctx->print_simulation_time(); } - pthread_mutex_lock(&(GPGPUsim_ctx_ptr()->g_sim_lock)); - GPGPUsim_ctx_ptr()->g_sim_active = false; - pthread_mutex_unlock(&(GPGPUsim_ctx_ptr()->g_sim_lock)); - } while( !GPGPUsim_ctx_ptr()->g_sim_done ); + pthread_mutex_lock(&(ctx->the_gpgpusim->g_sim_lock)); + ctx->the_gpgpusim->g_sim_active = false; + pthread_mutex_unlock(&(ctx->the_gpgpusim->g_sim_lock)); + } while( !ctx->the_gpgpusim->g_sim_done ); printf("GPGPU-Sim: *** simulation thread exiting ***\n"); fflush(stdout); - if(GPGPUsim_ctx_ptr()->break_limit) { + if(ctx->the_gpgpusim->break_limit) { printf("GPGPU-Sim: ** break due to reaching the maximum cycles (or instructions) **\n"); exit(1); } - sem_post(&(GPGPUsim_ctx_ptr()->g_sim_signal_exit)); + sem_post(&(ctx->the_gpgpusim->g_sim_signal_exit)); return NULL; } -void synchronize() +void gpgpu_context::synchronize() { printf("GPGPU-Sim: synchronize waiting for inactive GPU simulation\n"); - GPGPUsim_ctx_ptr()->g_stream_manager->print(stdout); + the_gpgpusim->g_stream_manager->print(stdout); fflush(stdout); // sem_wait(&g_sim_signal_finish); bool done = false; do { - pthread_mutex_lock(&(GPGPUsim_ctx_ptr()->g_sim_lock)); - done = ( GPGPUsim_ctx_ptr()->g_stream_manager->empty() && !GPGPUsim_ctx_ptr()->g_sim_active ) || GPGPUsim_ctx_ptr()->g_sim_done; - pthread_mutex_unlock(&(GPGPUsim_ctx_ptr()->g_sim_lock)); + pthread_mutex_lock(&(the_gpgpusim->g_sim_lock)); + done = ( the_gpgpusim->g_stream_manager->empty() && !the_gpgpusim->g_sim_active ) || the_gpgpusim->g_sim_done; + pthread_mutex_unlock(&(the_gpgpusim->g_sim_lock)); } while (!done); printf("GPGPU-Sim: detected inactive GPU simulation thread\n"); fflush(stdout); // sem_post(&g_sim_signal_start); } -void exit_simulation() +void gpgpu_context::exit_simulation() { - GPGPUsim_ctx_ptr()->g_sim_done=true; + the_gpgpusim->g_sim_done=true; printf("GPGPU-Sim: exit_simulation called\n"); fflush(stdout); - sem_wait(&(GPGPUsim_ctx_ptr()->g_sim_signal_exit)); + sem_wait(&(the_gpgpusim->g_sim_signal_exit)); printf("GPGPU-Sim: simulation thread signaled exit\n"); fflush(stdout); } @@ -211,8 +202,8 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf() func_sim->ptx_opcocde_latency_options(opp); icnt_reg_options(opp); - GPGPUsim_ctx_ptr()->g_the_gpu_config = new gpgpu_sim_config(this); - GPGPUsim_ctx_ptr()->g_the_gpu_config->reg_options(opp); // register GPU microrachitecture options + the_gpgpusim->g_the_gpu_config = new gpgpu_sim_config(this); + the_gpgpusim->g_the_gpu_config->reg_options(opp); // register GPU microrachitecture options option_parser_cmdline(opp, sg_argc, sg_argv); // parse configuration options fprintf(stdout, "GPGPU-Sim: Configuration options:\n\n"); @@ -220,37 +211,37 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf() // Set the Numeric locale to a standard locale where a decimal point is a "dot" not a "comma" // so it does the parsing correctly independent of the system environment variables assert(setlocale(LC_NUMERIC,"C")); - GPGPUsim_ctx_ptr()->g_the_gpu_config->init(); + the_gpgpusim->g_the_gpu_config->init(); - GPGPUsim_ctx_ptr()->g_the_gpu = new gpgpu_sim(*(GPGPUsim_ctx_ptr()->g_the_gpu_config), this); - GPGPUsim_ctx_ptr()->g_stream_manager = new stream_manager((GPGPUsim_ctx_ptr()->g_the_gpu), func_sim->g_cuda_launch_blocking); + the_gpgpusim->g_the_gpu = new gpgpu_sim(*(the_gpgpusim->g_the_gpu_config), this); + the_gpgpusim->g_stream_manager = new stream_manager((the_gpgpusim->g_the_gpu), func_sim->g_cuda_launch_blocking); - GPGPUsim_ctx_ptr()->g_simulation_starttime = time((time_t *)NULL); + the_gpgpusim->g_simulation_starttime = time((time_t *)NULL); - sem_init(&(GPGPUsim_ctx_ptr()->g_sim_signal_start),0,0); - sem_init(&(GPGPUsim_ctx_ptr()->g_sim_signal_finish),0,0); - sem_init(&(GPGPUsim_ctx_ptr()->g_sim_signal_exit),0,0); + sem_init(&(the_gpgpusim->g_sim_signal_start),0,0); + sem_init(&(the_gpgpusim->g_sim_signal_finish),0,0); + sem_init(&(the_gpgpusim->g_sim_signal_exit),0,0); - return GPGPUsim_ctx_ptr()->g_the_gpu; + return the_gpgpusim->g_the_gpu; } -void start_sim_thread(int api) +void gpgpu_context::start_sim_thread(int api) { - if( GPGPUsim_ctx_ptr()->g_sim_done ) { - GPGPUsim_ctx_ptr()->g_sim_done = false; + if( the_gpgpusim->g_sim_done ) { + the_gpgpusim->g_sim_done = false; if( api == 1 ) { - pthread_create(&(GPGPUsim_ctx_ptr()->g_simulation_thread),NULL,gpgpu_sim_thread_concurrent,NULL); + pthread_create(&(the_gpgpusim->g_simulation_thread),NULL,gpgpu_sim_thread_concurrent,(void *)this); } else { - pthread_create(&(GPGPUsim_ctx_ptr()->g_simulation_thread),NULL,gpgpu_sim_thread_sequential,NULL); + pthread_create(&(the_gpgpusim->g_simulation_thread),NULL,gpgpu_sim_thread_sequential,(void *)this); } } } -void print_simulation_time() +void gpgpu_context::print_simulation_time() { time_t current_time, difference, d, h, m, s; current_time = time((time_t *)NULL); - difference = MAX(current_time - GPGPUsim_ctx_ptr()->g_simulation_starttime, 1); + difference = MAX(current_time - the_gpgpusim->g_simulation_starttime, 1); d = difference/(3600*24); h = difference/3600 - 24*d; @@ -260,18 +251,18 @@ void print_simulation_time() fflush(stderr); printf("\n\ngpgpu_simulation_time = %u days, %u hrs, %u min, %u sec (%u sec)\n", (unsigned)d, (unsigned)h, (unsigned)m, (unsigned)s, (unsigned)difference ); - printf("gpgpu_simulation_rate = %u (inst/sec)\n", (unsigned)(GPGPUsim_ctx_ptr()->g_the_gpu->gpu_tot_sim_insn / difference) ); - const unsigned cycles_per_sec = (unsigned)(GPGPUsim_ctx_ptr()->g_the_gpu->gpu_tot_sim_cycle / difference); + printf("gpgpu_simulation_rate = %u (inst/sec)\n", (unsigned)(the_gpgpusim->g_the_gpu->gpu_tot_sim_insn / difference) ); + const unsigned cycles_per_sec = (unsigned)(the_gpgpusim->g_the_gpu->gpu_tot_sim_cycle / difference); printf("gpgpu_simulation_rate = %u (cycle/sec)\n", cycles_per_sec ); - printf("gpgpu_silicon_slowdown = %ux\n", GPGPUsim_ctx_ptr()->g_the_gpu->shader_clock() * 1000 / cycles_per_sec); + printf("gpgpu_silicon_slowdown = %ux\n", the_gpgpusim->g_the_gpu->shader_clock() * 1000 / cycles_per_sec); fflush(stdout); } -int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ) +int gpgpu_context::gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ) { - GPGPUsim_ctx_ptr()->g_the_gpu->launch(grid); - sem_post(&(GPGPUsim_ctx_ptr()->g_sim_signal_start)); - sem_wait(&(GPGPUsim_ctx_ptr()->g_sim_signal_finish)); + the_gpgpusim->g_the_gpu->launch(grid); + sem_post(&(the_gpgpusim->g_sim_signal_start)); + sem_wait(&(the_gpgpusim->g_sim_signal_finish)); return 0; } diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h index b2e22e6..9f408df 100644 --- a/src/gpgpusim_entrypoint.h +++ b/src/gpgpusim_entrypoint.h @@ -76,10 +76,4 @@ class GPGPUsim_ctx { }; -void start_sim_thread(int api); - -struct GPGPUsim_ctx* GPGPUsim_ctx_ptr(); - -int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid ); - #endif -- cgit v1.3 From 690774e8f70624ef759f0f76f88d804f89db8f46 Mon Sep 17 00:00:00 2001 From: Roy Spliet Date: Thu, 25 Jul 2019 16:43:06 +0100 Subject: libopencl: Stub clGetProgramBuildInfo() Assumes everything is ok, even if it isn't. Allows for some Rodinia benchmarks to fail slightly later. v2: Use available case macros. Signed-off-by: Roy Spliet --- libopencl/opencl_runtime_api.cc | 29 +++++++++++++++++++++++++++++ 1 file changed, 29 insertions(+) (limited to 'libopencl/opencl_runtime_api.cc') diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 97a54d8..a21e11d 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -1255,6 +1255,35 @@ clGetProgramInfo(cl_program program, return CL_SUCCESS; } +extern CL_API_ENTRY cl_int CL_API_CALL +clGetProgramBuildInfo (cl_program program, + cl_device_id device, + cl_program_build_info param_name, + size_t param_value_size, + void * param_value, + size_t * param_value_size_ret) CL_API_SUFFIX__VERSION_1_0 +{ + char *buf = (char*)param_value; + + switch( param_name ) { + case CL_PROGRAM_BUILD_STATUS: + CL_CASE( cl_build_status, CL_BUILD_SUCCESS ); + break; + case CL_PROGRAM_BUILD_OPTIONS: + case CL_PROGRAM_BUILD_LOG: + CL_STRING_CASE( "" ); + break; + case CL_PROGRAM_BINARY_TYPE: + CL_CASE( cl_program_binary_type, CL_PROGRAM_BINARY_TYPE_EXECUTABLE ); + break; + default: + return CL_INVALID_VALUE; + break; + } + + return CL_SUCCESS; +} + extern CL_API_ENTRY cl_int CL_API_CALL clEnqueueCopyBuffer(cl_command_queue command_queue, cl_mem src_buffer, -- cgit v1.3 From 6cdb7b18f70f8a1da200020f1b72f56fc33864d8 Mon Sep 17 00:00:00 2001 From: Roy Spliet Date: Thu, 25 Jul 2019 18:20:56 +0100 Subject: libopencl: Perform PDOM analysis Signed-off-by: Roy Spliet --- libopencl/opencl_runtime_api.cc | 11 +++++++++++ 1 file changed, 11 insertions(+) (limited to 'libopencl/opencl_runtime_api.cc') diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index a21e11d..23e52f6 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -950,6 +950,17 @@ 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 = gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu); + + //do dynamic PDOM analysis for performance simulation scenario + std::string kname = grid->name(); + function_info *kernel_func_info = grid->entry(); + if (kernel_func_info->is_pdom_set()) { + printf("GPGPU-Sim PTX: PDOM analysis already done for %s \n", kname.c_str() ); + } else { + printf("GPGPU-Sim PTX: finding reconvergence points for \'%s\'...\n", kname.c_str() ); + kernel_func_info->do_pdom(); + kernel_func_info->set_pdom(); + } if ( g_ptx_sim_mode ) gpgpu_opencl_ptx_sim_main_func( grid ); else -- cgit v1.3 From 55a05c11b0af3159d81bce823c1fda4fb3cd04e3 Mon Sep 17 00:00:00 2001 From: Roy Spliet Date: Thu, 25 Jul 2019 18:34:44 +0100 Subject: libopencl: Calculate a valid offset in bind_args() v2: Adhere to alignment requirements. v3: Small style fix. Signed-off-by: Roy Spliet --- libopencl/opencl_runtime_api.cc | 14 +++++++++++++- 1 file changed, 13 insertions(+), 1 deletion(-) (limited to 'libopencl/opencl_runtime_api.cc') diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index 23e52f6..752bfdf 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -263,15 +263,27 @@ void _cl_kernel::SetKernelArg( cl_int _cl_kernel::bind_args( gpgpu_ptx_sim_arg_list_t &arg_list ) { + size_t offset = 0; + assert( arg_list.empty() ); unsigned k=0; std::map::iterator i; for( i = m_args.begin(); i!=m_args.end(); i++ ) { if( i->first != k ) return CL_INVALID_KERNEL_ARGS; + arg_info arg = i->second; - gpgpu_ptx_sim_arg param( arg.m_arg_value, arg.m_arg_size, 0); + const symbol *sym = m_kernel_impl->get_arg(i->first); + const type_info_key &t = sym->type()->get_key(); + + int align = (t.get_alignment_spec() == -1) ? arg.m_arg_size : t.get_alignment_spec(); + if( offset % align ) + offset += (align - (offset % align)); + + gpgpu_ptx_sim_arg param( arg.m_arg_value, arg.m_arg_size, offset ); arg_list.push_front( param ); + + offset += arg.m_arg_size; k++; } return CL_SUCCESS; -- cgit v1.3