From b400820185db543ff98fd823a15958444e12fee6 Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Sun, 8 Aug 2010 21:15:00 -0800 Subject: ** fully decouple kernel compilation from kernel launch for opencl changes required to do this (this enables oclReduction SDK 3.1 to pass and is a prerequisite to Fermi style concurrent kernel launching) - separate kernel launch for cuda and opencl (and init_grid functions) - add pgm_info structure for holding kernel function_info pointer in _cl_program object after ptx compilation ** added support for clEnqueueNDRangeKernel with local_work_size==NULL which is part of OpenCL spec and used in oclSortingNetworks (however, this is a braindead implementation that only handles the case where global_work_size[0] is smaller than the max number of threads per shader; moreover oclSortingNetworks is still not working but for what looks like another reason) ** refactoring / cleanup - g_global_symbol_table made static to ptx_parser.cc - remove g_kernel_name_to_symtab_lookup (not really being used) - moving various function prototypes into headers (e.g., if foo() defined in bar.cc, then place prototype in bar.h) - adding icnt_reg_options() [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7168] --- src/cuda-sim/cuda-sim.cc | 147 +++++++++++++++++++++++++++++++----------- src/cuda-sim/cuda-sim.h | 11 ++-- src/cuda-sim/ptx_ir.cc | 137 ++------------------------------------- src/cuda-sim/ptx_ir.h | 10 +-- src/cuda-sim/ptx_loader.cc | 139 ++++++++++++++++++++++++++++++++++++--- src/cuda-sim/ptx_loader.h | 1 + src/cuda-sim/ptx_parser.cc | 7 +- src/cuda-sim/ptx_parser.h | 7 +- src/gpgpu-sim/gpu-sim.cc | 18 ++++-- src/gpgpu-sim/gpu-sim.h | 13 +--- src/gpgpu-sim/icnt_wrapper.cc | 12 +++- src/gpgpu-sim/icnt_wrapper.h | 10 +-- src/gpgpu-sim/visualizer.cc | 1 + src/gpgpusim_entrypoint.cc | 72 ++++++++++++--------- src/gpgpusim_entrypoint.h | 18 +++++- src/option_parser.h | 2 +- 16 files changed, 354 insertions(+), 251 deletions(-) (limited to 'src') diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 6a0ee36..737b94a 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1114,11 +1114,10 @@ void init_inst_classification_stat() { std::map *g_kernel_name_to_function_lookup=NULL; -std::map g_kernel_name_to_symtab_lookup; std::map *g_host_to_kernel_entrypoint_name_lookup=NULL; extern unsigned g_ptx_thread_info_uid_next; -void gpgpu_ptx_sim_init_grid( const char *kernel_key, struct gpgpu_ptx_sim_arg* args, +void gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, struct gpgpu_ptx_sim_arg* args, struct dim3 gridDim, struct dim3 blockDim ) { g_gx=0; @@ -1152,7 +1151,39 @@ void gpgpu_ptx_sim_init_grid( const char *kernel_key, struct gpgpu_ptx_sim_arg* abort(); } g_entrypoint_func_info = g_func_info = (*g_kernel_name_to_function_lookup)[kname]; - g_entrypoint_symbol_table = g_kernel_name_to_symtab_lookup[kname]; + + unsigned argcount=0; + struct gpgpu_ptx_sim_arg *tmparg = args; + while (tmparg) { + tmparg = tmparg->m_next; + argcount++; + } + + unsigned argn=1; + while (args) { + g_func_info->add_param_data(argcount-argn,args); + args = args->m_next; + argn++; + } + g_func_info->finalize(g_param_mem); + g_ptx_kernel_count++; + if ( gpgpu_ptx_instruction_classification ) { + init_inst_classification_stat(); + } + fflush(stdout); +} + +void gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,struct gpgpu_ptx_sim_arg *args, struct dim3 gridDim, struct dim3 blockDim ) +{ + g_gx=0; + g_gy=0; + g_gz=0; + g_cudaGridDim = gridDim; + g_cudaBlockDim = blockDim; + g_sm_idx_offset_next.clear(); + g_sm_next_index = 0; + + g_entrypoint_func_info = g_func_info = entry; unsigned argcount=0; struct gpgpu_ptx_sim_arg *tmparg = args; @@ -1209,20 +1240,6 @@ void gpgpu_ptx_sim_register_kernel(const char *hostFun, const char *deviceFun) printf("GPGPU-Sim PTX: __cudaRegisterFunction %s : 0x%Lx\n", deviceFun, (unsigned long long)hostFun); } -void register_ptx_function( const char *name, function_info *impl, symbol_table *symtab ) -{ - printf("GPGPU-Sim PTX: parsing function %s\n", name ); - if( g_kernel_name_to_function_lookup == NULL ) - g_kernel_name_to_function_lookup = new std::map; - - std::map::iterator i_kernel = g_kernel_name_to_function_lookup->find(name); - if (i_kernel != g_kernel_name_to_function_lookup->end() && i_kernel->second != NULL) { - printf("GPGPU-Sim PTX: WARNING: Function already parsed once. Overwriting.\n"); - } - (*g_kernel_name_to_function_lookup)[name] = impl; - g_kernel_name_to_symtab_lookup[name] = symtab; -} - std::map g_const_name_lookup; // indexed by hostVar std::map g_global_name_lookup; // indexed by hostVar std::set g_globals; @@ -1314,9 +1331,7 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co fflush(stdout); } -int g_ptx_sim_mode=0; -// used by libcuda.a if non-zer cudaLaunch() will call gpgpu_ptx_sim_main_func() -// if zero it calls gpgpu_ptx_sim_main_perf() +int g_ptx_sim_mode=0; // if non-zero run functional simulation only (i.e., no notion of a clock cycle) extern "C" int ptx_debug; @@ -1378,13 +1393,13 @@ void read_sim_environment_variables() -extern time_t simulation_starttime; +extern time_t g_simulation_starttime; ptx_cta_info *g_func_cta_info = NULL; #define MAX(a,b) (((a)>(b))?(a):(b)) -void gpgpu_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, struct gpgpu_ptx_sim_arg *args) +void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, struct gpgpu_ptx_sim_arg *args) { printf("GPGPU-Sim: Performing Functional Simulation...\n"); @@ -1396,7 +1411,7 @@ void gpgpu_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDi int i1, i2, i3, i4, o1, o2, o3, o4; int vectorin, vectorout; - gpgpu_ptx_sim_init_grid(kernel_key, args,gridDim,blockDim); + gpgpu_cuda_ptx_sim_init_grid(kernel_key, args,gridDim,blockDim); memory_space *shared_mem = new memory_space_impl<16*1024>("shared",4); @@ -1499,7 +1514,7 @@ void gpgpu_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDi StatDisp ( g_inst_op_classification_stat[g_ptx_kernel_count]); } end_time = time((time_t *)NULL); - elapsed_time = MAX(end_time - simulation_starttime, 1); + elapsed_time = MAX(end_time - g_simulation_starttime, 1); days = elapsed_time/(3600*24); hrs = elapsed_time/3600 - 24*days; @@ -1609,21 +1624,8 @@ extern "C" void ptxinfo_cmem( unsigned nbytes, unsigned bank ) g_ptxinfo_kinfo.cmem+=nbytes; } -extern "C" void ptxinfo_addinfo() +void clear_ptxinfo() { - if ( g_kernel_name_to_function_lookup ) { - std::map::iterator i=g_kernel_name_to_function_lookup->find(g_ptxinfo_kname); - if ( (g_kernel_name_to_function_lookup == NULL) || (i == g_kernel_name_to_function_lookup->end()) ) { - printf ("GPGPU-Sim PTX: Kernel '%s' in %s not found. Ignoring.\n", g_ptxinfo_kname, g_filename); - } else { - printf ("GPGPU-Sim PTX: Kernel %s\n", g_ptxinfo_kname); - function_info *fi = i->second; - fi->set_kernel_info(&g_ptxinfo_kinfo); - } - } else { - printf ("GPGPU-Sim PTX: Kernel '%s' in %s not found (no kernels registered).\n", g_ptxinfo_kname, g_filename); - } - free(g_ptxinfo_kname); g_ptxinfo_kname=NULL; g_ptxinfo_kinfo.regs=0; @@ -1632,6 +1634,62 @@ extern "C" void ptxinfo_addinfo() g_ptxinfo_kinfo.cmem=0; } +void ptxinfo_opencl_addinfo( std::map &kernels ) +{ + if( !strcmp("__cuda_dummy_entry__",g_ptxinfo_kname) ) { + // this string produced by ptxas for empty ptx files (e.g., bandwidth test) + clear_ptxinfo(); + return; + } + std::map::iterator k=kernels.find(g_ptxinfo_kname); + if( k==kernels.end() ) { + printf ("GPGPU-Sim PTX: ERROR ** implementation for '%s' not found.\n", g_ptxinfo_kname ); + abort(); + } else { + printf ("GPGPU-Sim PTX: Kernel \'%s\' : regs=%u, lmem=%u, smem=%u, cmem=%u\n", + g_ptxinfo_kname, + g_ptxinfo_kinfo.regs, + g_ptxinfo_kinfo.lmem, + g_ptxinfo_kinfo.smem, + g_ptxinfo_kinfo.cmem ); + function_info *finfo = k->second; + assert(finfo!=NULL); + finfo->set_kernel_info( g_ptxinfo_kinfo ); + } + clear_ptxinfo(); +} + +void ptxinfo_cuda_addinfo() +{ + if( !strcmp("__cuda_dummy_entry__",g_ptxinfo_kname) ) { + // this string produced by ptxas for empty ptx files (e.g., bandwidth test) + clear_ptxinfo(); + return; + } + if ( g_kernel_name_to_function_lookup ) { + std::map::iterator i=g_kernel_name_to_function_lookup->find(g_ptxinfo_kname); + if ( (g_kernel_name_to_function_lookup == NULL) || (i == g_kernel_name_to_function_lookup->end()) ) { + printf ("GPGPU-Sim PTX: ERROR ** implementation for '%s' not found.\n", g_ptxinfo_kname ); + abort(); + } else { + printf ("GPGPU-Sim PTX: Kernel \'%s\' : regs=%u, lmem=%u, smem=%u, cmem=%u\n", + g_ptxinfo_kname, + g_ptxinfo_kinfo.regs, + g_ptxinfo_kinfo.lmem, + g_ptxinfo_kinfo.smem, + g_ptxinfo_kinfo.cmem ); + function_info *fi = i->second; + assert(fi!=NULL); + fi->set_kernel_info(g_ptxinfo_kinfo); + } + } else { + printf ("GPGPU-Sim PTX: ERROR ** Kernel '%s' not found (no kernels registered).\n", g_ptxinfo_kname ); + abort(); + } + clear_ptxinfo(); +} + + void dwf_insert_reconv_pt(address_type pc); struct rec_pts { @@ -1704,6 +1762,19 @@ void dwf_process_reconv_pts() } } +void register_function_implementation( const char *name, function_info *impl ) +{ + printf("GPGPU-Sim PTX: parsing function %s\n", name ); + if( g_kernel_name_to_function_lookup == NULL ) + g_kernel_name_to_function_lookup = new std::map; + + std::map::iterator i_kernel = g_kernel_name_to_function_lookup->find(name); + if (i_kernel != g_kernel_name_to_function_lookup->end() && i_kernel->second != NULL) { + printf("GPGPU-Sim PTX: WARNING: Function already parsed once. Overwriting.\n"); + } + (*g_kernel_name_to_function_lookup)[name] = impl; +} + void *gpgpusim_opencl_getkernel_Object( const char *kernel_name ) { std::map::iterator i=g_kernel_name_to_function_lookup->find(kernel_name); diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 6662b4c..2c5e0a7 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -16,10 +16,11 @@ extern memory_space *g_global_mem; extern int g_debug_execution; extern int g_debug_thread_uid; extern std::map *g_kernel_name_to_function_lookup; -extern std::map g_kernel_name_to_symtab_lookup; -extern void gpgpu_ptx_sim_add_ptxstring( const char *ptx, const char *source_fname ); -extern void gpgpu_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, struct gpgpu_ptx_sim_arg *); +extern void gpgpu_cuda_ptx_sim_init_grid(const char *kernel_key,struct gpgpu_ptx_sim_arg *args, struct dim3 gridDim, struct dim3 blockDim ); +extern void gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,struct gpgpu_ptx_sim_arg *args, struct dim3 gridDim, struct dim3 blockDim ); +extern void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, struct gpgpu_ptx_sim_arg *); +extern void print_splash(); extern void* gpgpu_ptx_sim_malloc( size_t count ); extern void* gpgpu_ptx_sim_mallocarray( size_t count ); extern void gpgpu_ptx_sim_memcpy_to_gpu( size_t dst_start_addr, const void *src, size_t count ); @@ -39,6 +40,8 @@ extern int gpgpu_ptx_sim_sizeofTexture(const char* name); extern const char* gpgpu_ptx_sim_findNamefromTexture(const struct textureReference* texref); extern const struct textureReference* gpgpu_ptx_sim_accessTextureofName(const char* name); extern void read_sim_environment_variables(); -extern void register_ptx_function( const char *name, function_info *impl, symbol_table *symtab ); +extern void register_function_implementation( const char *name, function_info *impl ); +extern void ptxinfo_cuda_addinfo(); +extern void ptxinfo_opencl_addinfo( std::map &kernels ); #endif diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index cd70daf..5945b72 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -191,6 +191,8 @@ void symbol_table::add_function( function_info *func, const char *filename, unsi m_symbols[ func->get_name() ] = s; } +void register_ptx_function( const char *name, function_info *impl ); // either libcuda or libopencl + bool symbol_table::add_function_decl( const char *name, int entry_point, function_info **func_info, symbol_table **sym_table ) { std::string key = std::string(name); @@ -209,13 +211,14 @@ bool symbol_table::add_function_decl( const char *name, int entry_point, functio *sym_table = m_function_symtab_lookup[key]; } else { assert( !prior_decl ); - *sym_table = new symbol_table( "", entry_point, g_global_symbol_table ); + *sym_table = new symbol_table( "", entry_point, this ); symbol *null_reg = (*sym_table)->add_variable("_",NULL,0,"",0); null_reg->set_regno(0, 0); (*sym_table)->set_name(name); (*func_info)->set_symtab(*sym_table); m_function_symtab_lookup[key] = *sym_table; - register_ptx_function(name,*func_info,*sym_table); + assert( (*func_info)->get_symtab() == *sym_table ); + register_ptx_function(name,*func_info); } return prior_decl; } @@ -1010,8 +1013,6 @@ void function_info::print_insn( unsigned pc, FILE * fp ) const } } -static int g_save_embedded_ptx = 0; - extern "C" int ptx_parse(); extern "C" int ptx__scan_string(const char*); extern "C" FILE *ptx_in; @@ -1021,128 +1022,6 @@ extern "C" int ptxinfo_parse(); extern "C" int ptxinfo_debug; extern "C" FILE *ptxinfo_in; -static void print_ptx_file( const char *p, unsigned source_num, const char *filename ); - -void 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"); -} - -void gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ) -{ - char buf[1024]; - snprintf(buf,1024,"_%u.ptx", source_num ); - if( g_save_embedded_ptx ) { - FILE *fp = fopen(buf,"w"); - fprintf(fp,"%s",p); - fclose(fp); - } - init_parser(buf); - ptx__scan_string(p); - int errors = ptx_parse (); - if ( errors ) { - char fname[1024]; - snprintf(fname,1024,"_ptx_XXXXXX"); - int fd=mkstemp(fname); - close(fd); - printf("GPGPU-Sim PTX: parser error detected, exiting... but first extracting .ptx to \"%s\"\n", fname); - FILE *ptxfile = fopen(fname,"w"); - fprintf(ptxfile,"%s", p ); - fclose(ptxfile); - abort(); - exit(40); - } - - if ( g_debug_execution >= 100 ) - print_ptx_file(p,source_num,buf); - - printf("GPGPU-Sim PTX: finished parsing EMBEDDED .ptx file %s\n",buf); - - char fname[1024]; - snprintf(fname,1024,"_ptx_XXXXXX"); - int fd=mkstemp(fname); - close(fd); - - printf("GPGPU-Sim PTX: extracting embedded .ptx to temporary file \"%s\"\n", fname); - FILE *ptxfile = fopen(fname,"w"); - fprintf(ptxfile,"%s",p); - fclose(ptxfile); - - char fname2[1024]; - snprintf(fname2,1024,"_ptx2_XXXXXX"); - fd=mkstemp(fname2); - close(fd); - char commandline2[4096]; - snprintf(commandline2,4096,"cat %s | sed 's/.version 1.5/.version 1.4/' | sed 's/, texmode_independent//' | sed 's/\\(\\.extern \\.const\\[1\\] .b8 \\w\\+\\)\\[\\]/\\1\\[1\\]/' | sed 's/const\\[.\\]/const\\[0\\]/g' > %s", fname, fname2); - int result = system(commandline2); - if( result != 0 ) { - printf("GPGPU-Sim PTX: ERROR ** while loading PTX (a) %d\n", result); - printf(" Ensure you have write access to simulation directory\n"); - printf(" and have \'cat\' and \'sed\' in your path.\n"); - exit(1); - } - - char tempfile_ptxinfo[1024]; - snprintf(tempfile_ptxinfo,1024,"%sinfo",fname); - char commandline[1024]; - char extra_flags[1024]; - extra_flags[0]=0; -#if CUDART_VERSION >= 3000 - snprintf(extra_flags,1024,"--gpu-name=sm_20"); -#endif - snprintf(commandline,1024,"ptxas %s -v %s --output-file /dev/null 2> %s", - extra_flags, fname2, tempfile_ptxinfo); - printf("GPGPU-Sim PTX: generating ptxinfo using \"%s\"\n", commandline); - result = system(commandline); - if( result != 0 ) { - printf("GPGPU-Sim PTX: ERROR ** while loading PTX (b) %d\n", result); - printf(" Ensure ptxas is in your path.\n"); - exit(1); - } - - ptxinfo_in = fopen(tempfile_ptxinfo,"r"); - g_ptxinfo_filename = tempfile_ptxinfo; - ptxinfo_parse(); - snprintf(commandline,1024,"rm -f %s %s %s", fname, fname2, tempfile_ptxinfo); - printf("GPGPU-Sim PTX: removing ptxinfo using \"%s\"\n", commandline); - result = system(commandline); - if( result != 0 ) { - printf("GPGPU-Sim PTX: ERROR ** while loading PTX (c) %d\n", result); - exit(1); - } -} - - -const ptx_instruction *ptx_instruction_lookup( const char *filename, unsigned linenumber ); - -void print_ptx_file( const char *p, unsigned source_num, const char *filename ) -{ - printf("\nGPGPU-Sim PTX: file _%u.ptx contents:\n\n", source_num ); - char *s = strdup(p); - char *t = s; - unsigned n=1; - while ( *t != '\0' ) { - char *u = t; - while ( (*u != '\n') && (*u != '\0') ) u++; - unsigned last = (*u == '\0'); - *u = '\0'; - const ptx_instruction *pI = ptx_instruction_lookup(filename,n); - char pc[64]; - if( pI && pI->get_PC() ) - snprintf(pc,64,"%4u", pI->get_PC() ); - else - snprintf(pc,64," "); - printf(" _%u.ptx %4u (pc=%s): %s\n", source_num, n, pc, t ); - if ( last ) break; - t = u+1; - n++; - } - free(s); - fflush(stdout); -} - void gpgpu_ptx_assemble( std::string kname, void *kinfo ) { function_info *func_info = (function_info *)kinfo; @@ -1150,12 +1029,6 @@ void gpgpu_ptx_assemble( std::string kname, void *kinfo ) printf("GPGPU-Sim PTX: skipping assembly for extern declared function \'%s\'\n", func_info->get_name().c_str() ); return; } - if( g_kernel_name_to_symtab_lookup[ kname ] == NULL ) { - printf("\nGPGPU-Sim PTX: ERROR no information for kernel \'%s\'\n" - " this can happen for kernels contained in CUDA\n" - " libraries (such as CUBLAS, CUFFT, or CUDPP)\n", kname.c_str() ); - exit(1); - } g_func_info = func_info; func_info->ptx_assemble(); } diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index fc7f192..1b1d864 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -1015,8 +1015,8 @@ public: return &m_kernel_info; } - const void set_kernel_info (const struct gpgpu_ptx_sim_kernel_info *info) { - m_kernel_info = *info; + const void set_kernel_info (const struct gpgpu_ptx_sim_kernel_info &info) { + m_kernel_info = info; } symbol_table *get_symtab() { @@ -1254,14 +1254,11 @@ struct textureInfo { extern function_info *g_func_info; -extern int g_error_detected; -extern symbol_table *g_entrypoint_symbol_table; extern function_info *g_entrypoint_func_info; -extern symbol_table *g_global_symbol_table; extern std::map g_sym_name_to_symbol_table; + #define GLOBAL_HEAP_START 0x10000000 // start allocating from this address (lower values used for allocating globals in .ptx file) - #define SHARED_MEM_SIZE_MAX (64*1024) #define LOCAL_MEM_SIZE_MAX 1024 #define MAX_STREAMING_MULTIPROCESSORS 64 @@ -1274,7 +1271,6 @@ extern std::map g_sym_name_to_symbol_table; #define STATIC_ALLOC_LIMIT (GLOBAL_HEAP_START - (TOTAL_LOCAL_MEM+TOTAL_SHARED_MEM)) -void gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); void gpgpu_ptx_assemble( std::string kname, void *kinfo ); #include "../option_parser.h" void ptx_reg_options(option_parser_t opp); diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index fb4e8dd..713bdb3 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -212,20 +212,20 @@ void gpgpu_ptx_sim_load_gpu_kernels() ptx_in = fopen( g_filename, "r" ); gpgpu_ptx_sim_init_memory(); if (ptx_in) { - init_parser(g_filename); + symbol_table *symtab=init_parser(g_filename); ptx_parse(); ptxinfo_in = open_ptxinfo(g_filename); ptxinfo_parse(); - load_static_globals(g_global_symbol_table,STATIC_ALLOC_LIMIT,0xFFFFFFFF); - load_constants(g_global_symbol_table,STATIC_ALLOC_LIMIT); + load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF); + load_constants(symtab,STATIC_ALLOC_LIMIT); } else { if (!g_override_embedded_ptx) { printf("GPGPU-Sim PTX: USING EMBEDDED .ptx files...\n"); ptx_info_t *s; for ( s=g_ptx_source_array; s!=NULL; s=s->next ) { - gpgpu_ptx_sim_load_ptx_from_string(s->str, ++source_num); - load_static_globals(g_global_symbol_table,STATIC_ALLOC_LIMIT,0xFFFFFFFF); - load_constants(g_global_symbol_table,STATIC_ALLOC_LIMIT); + symbol_table *symtab=gpgpu_ptx_sim_load_ptx_from_string(s->str, ++source_num); + load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF); + load_constants(symtab,STATIC_ALLOC_LIMIT); } } else { const char *filename = NULL; @@ -243,12 +243,12 @@ void gpgpu_ptx_sim_load_gpu_kernels() printf("Parsing %s..\n", filename); ptx_in = fopen( filename, "r" ); free(namelist[n]); - init_parser(filename); + symbol_table *symtab=init_parser(filename); ptx_parse (); ptxinfo_in = open_ptxinfo(filename); ptxinfo_parse(); - load_static_globals(g_global_symbol_table,STATIC_ALLOC_LIMIT,0xFFFFFFFF); - load_constants(g_global_symbol_table,STATIC_ALLOC_LIMIT); + load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF); + load_constants(symtab,STATIC_ALLOC_LIMIT); } free(namelist); } @@ -298,3 +298,124 @@ void gpgpu_ptx_sim_add_ptxstring( const char *ptx_string, const char *sourcefnam l_ptx_source->next = t; } } + +static int g_save_embedded_ptx = 0; + +void 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"); +} + +void print_ptx_file( const char *p, unsigned source_num, const char *filename ) +{ + printf("\nGPGPU-Sim PTX: file _%u.ptx contents:\n\n", source_num ); + char *s = strdup(p); + char *t = s; + unsigned n=1; + while ( *t != '\0' ) { + char *u = t; + while ( (*u != '\n') && (*u != '\0') ) u++; + unsigned last = (*u == '\0'); + *u = '\0'; + const ptx_instruction *pI = ptx_instruction_lookup(filename,n); + char pc[64]; + if( pI && pI->get_PC() ) + snprintf(pc,64,"%4u", pI->get_PC() ); + else + snprintf(pc,64," "); + printf(" _%u.ptx %4u (pc=%s): %s\n", source_num, n, pc, t ); + if ( last ) break; + t = u+1; + n++; + } + free(s); + fflush(stdout); +} + +symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ) +{ + char buf[1024]; + snprintf(buf,1024,"_%u.ptx", source_num ); + if( g_save_embedded_ptx ) { + FILE *fp = fopen(buf,"w"); + fprintf(fp,"%s",p); + fclose(fp); + } + symbol_table *symtab=init_parser(buf); + ptx__scan_string(p); + int errors = ptx_parse (); + if ( errors ) { + char fname[1024]; + snprintf(fname,1024,"_ptx_XXXXXX"); + int fd=mkstemp(fname); + close(fd); + printf("GPGPU-Sim PTX: parser error detected, exiting... but first extracting .ptx to \"%s\"\n", fname); + FILE *ptxfile = fopen(fname,"w"); + fprintf(ptxfile,"%s", p ); + fclose(ptxfile); + abort(); + exit(40); + } + + if ( g_debug_execution >= 100 ) + print_ptx_file(p,source_num,buf); + + printf("GPGPU-Sim PTX: finished parsing EMBEDDED .ptx file %s\n",buf); + + char fname[1024]; + snprintf(fname,1024,"_ptx_XXXXXX"); + int fd=mkstemp(fname); + close(fd); + + printf("GPGPU-Sim PTX: extracting embedded .ptx to temporary file \"%s\"\n", fname); + FILE *ptxfile = fopen(fname,"w"); + fprintf(ptxfile,"%s",p); + fclose(ptxfile); + + char fname2[1024]; + snprintf(fname2,1024,"_ptx2_XXXXXX"); + fd=mkstemp(fname2); + close(fd); + char commandline2[4096]; + snprintf(commandline2,4096,"cat %s | sed 's/.version 1.5/.version 1.4/' | sed 's/, texmode_independent//' | sed 's/\\(\\.extern \\.const\\[1\\] .b8 \\w\\+\\)\\[\\]/\\1\\[1\\]/' | sed 's/const\\[.\\]/const\\[0\\]/g' > %s", fname, fname2); + int result = system(commandline2); + if( result != 0 ) { + printf("GPGPU-Sim PTX: ERROR ** while loading PTX (a) %d\n", result); + printf(" Ensure you have write access to simulation directory\n"); + printf(" and have \'cat\' and \'sed\' in your path.\n"); + exit(1); + } + + char tempfile_ptxinfo[1024]; + snprintf(tempfile_ptxinfo,1024,"%sinfo",fname); + char commandline[1024]; + char extra_flags[1024]; + extra_flags[0]=0; +#if CUDART_VERSION >= 3000 + snprintf(extra_flags,1024,"--gpu-name=sm_20"); +#endif + snprintf(commandline,1024,"ptxas %s -v %s --output-file /dev/null 2> %s", + extra_flags, fname2, tempfile_ptxinfo); + printf("GPGPU-Sim PTX: generating ptxinfo using \"%s\"\n", commandline); + result = system(commandline); + if( result != 0 ) { + printf("GPGPU-Sim PTX: ERROR ** while loading PTX (b) %d\n", result); + printf(" Ensure ptxas is in your path.\n"); + exit(1); + } + + ptxinfo_in = fopen(tempfile_ptxinfo,"r"); + g_ptxinfo_filename = tempfile_ptxinfo; + ptxinfo_parse(); + snprintf(commandline,1024,"rm -f %s %s %s", fname, fname2, tempfile_ptxinfo); + printf("GPGPU-Sim PTX: removing ptxinfo using \"%s\"\n", commandline); + result = system(commandline); + if( result != 0 ) { + printf("GPGPU-Sim PTX: ERROR ** while loading PTX (c) %d\n", result); + exit(1); + } + return symtab; +} + diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index 33c6dc9..95130f1 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -76,5 +76,6 @@ extern bool g_override_embedded_ptx; void gpgpu_ptx_sim_load_gpu_kernels(); void gpgpu_ptx_sim_add_ptxstring( const char *ptx_string, const char *sourcefname ); +class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num ); #endif diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 17e4ca2..05ce397 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -73,7 +73,7 @@ const char *g_filename; unsigned g_max_regs_per_thread = 0; // the program intermediate representation... -symbol_table *g_global_symbol_table = NULL; +static symbol_table *g_global_symbol_table = NULL; std::map g_sym_name_to_symbol_table; static symbol_table *g_current_symbol_table = NULL; static std::list g_instructions; @@ -133,7 +133,7 @@ void read_parser_environment_variables() } } -void init_parser( const char *ptx_filename ) +symbol_table *init_parser( const char *ptx_filename ) { g_filename = strdup(ptx_filename); g_global_symbol_table = g_current_symbol_table = new symbol_table("global",0,NULL); @@ -142,6 +142,8 @@ void init_parser( const char *ptx_filename ) #define DEF(X,Y) g_ptx_token_decode[X] = Y; #include "ptx_parser_decode.def" #undef DEF + + return g_global_symbol_table; } void init_directive_state() @@ -192,7 +194,6 @@ void add_function_name( const char *name ) bool prior_decl = g_global_symbol_table->add_function_decl( name, g_entry_point, &g_func_info, &g_current_symbol_table ); if( g_entry_point ) { g_entrypoint_func_info = g_func_info; - g_entrypoint_symbol_table = g_current_symbol_table; } if( g_add_identifier_cached__identifier ) { add_identifier( g_add_identifier_cached__identifier, diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h index 076fa11..8ffe593 100644 --- a/src/cuda-sim/ptx_parser.h +++ b/src/cuda-sim/ptx_parser.h @@ -66,14 +66,17 @@ #define ptx_parser_INCLUDED #include "../abstract_hardware_model.h" - extern const char *g_filename; +extern int g_error_detected; + #ifdef __cplusplus +class symbol_table* init_parser(const char*); +const class ptx_instruction *ptx_instruction_lookup( const char *filename, unsigned linenumber ); extern "C" { #endif + const char *decode_token( int type ); void read_parser_environment_variables(); -void init_parser(const char*); void start_function( int entry_point ); void add_function_name( const char *fname ); void init_directive_state(); diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index cd4851b..bb228a5 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -66,6 +66,16 @@ #include "gpu-sim.h" +#include +#include +#include +#include "zlib.h" + +#include "../option_parser.h" +#include "shader.h" +#include "dram.h" +#include "mem_fetch.h" + #include #include "gpu-cache.h" #include "gpu-misc.h" @@ -83,6 +93,7 @@ #include "../intersim/statwraper.h" #include "../abstract_hardware_model.h" #include "../debug.h" +#include "../gpgpusim_entrypoint.h" #include #include @@ -299,10 +310,8 @@ unsigned int more_thread = 1; extern unsigned int n_regconflict_stall; unsigned int warp_conflict_at_writeback = 0; unsigned int gpgpu_commit_pc_beyond_two = 0; -extern int g_network_mode; int gpgpu_cache_wt_through = 0; - //memory access classification int gpgpu_n_mem_read_local = 0; int gpgpu_n_mem_write_local = 0; @@ -595,6 +604,7 @@ inline int mem2device(int memid) { /* Allocate memory for uArch structures */ void init_gpu () { + // initialize the GPU microarchitecture model int i; gpu_max_cycle = gpu_max_cycle_opt; @@ -724,6 +734,7 @@ void init_once(void ) { // return the number of cycle required to run all the trace on the gpu unsigned int run_gpu_sim(int grid_num) { + // run a CUDA grid on the GPU microarchitecture simulator int not_completed; int mem_busy; @@ -1432,7 +1443,6 @@ inline int next_clock_domain(void) return mask; } -extern time_t simulation_starttime; void gpu_sim_loop( int grid_num ) { int clock_mask = next_clock_domain(); @@ -1598,7 +1608,7 @@ void gpu_sim_loop( int grid_num ) time_t days, hrs, minutes, sec; time_t curr_time; time(&curr_time); - unsigned long long elapsed_time = MAX(curr_time - simulation_starttime, 1); + unsigned long long elapsed_time = MAX(curr_time - g_simulation_starttime, 1); days = elapsed_time/(3600*24); hrs = elapsed_time/3600 - 24*days; minutes = elapsed_time/60 - 60*(hrs + 24*days); diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 27d83ad..2e9794f 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -66,16 +66,6 @@ * Vancouver, BC V6T 1Z4 */ -#include -#include -#include -#include "zlib.h" - -#include "../option_parser.h" -#include "shader.h" -#include "dram.h" -#include "mem_fetch.h" - #ifndef GPU_SIM_H #define GPU_SIM_H @@ -113,5 +103,8 @@ int mem_ctrl_full( int mc_id ); void dramqueue_latency_log_dump(); void dump_pipeline_impl( int mask, int s, int m ); +unsigned int run_gpu_sim(int grid_num); +extern void gpu_reg_options(class OptionParser * opp); +extern void init_gpu(); #endif diff --git a/src/gpgpu-sim/icnt_wrapper.cc b/src/gpgpu-sim/icnt_wrapper.cc index 3e38269..aa55110 100644 --- a/src/gpgpu-sim/icnt_wrapper.cc +++ b/src/gpgpu-sim/icnt_wrapper.cc @@ -74,8 +74,16 @@ icnt_pop_p icnt_pop; icnt_transfer_p icnt_transfer; icnt_busy_p icnt_busy; -extern int g_network_mode; -extern char* g_network_config_filename; +int g_network_mode; +char* g_network_config_filename; + +#include "../option_parser.h" + +void icnt_reg_options( class OptionParser * opp ) +{ + option_parser_register(opp, "-network_mode", OPT_INT32, &g_network_mode, "Interconnection network mode", "1"); + option_parser_register(opp, "-inter_config_file", OPT_CSTR, &g_network_config_filename, "Interconnection network config file", "mesh"); +} void icnt_init( unsigned int n_shader, unsigned int n_mem ) { diff --git a/src/gpgpu-sim/icnt_wrapper.h b/src/gpgpu-sim/icnt_wrapper.h index 998a4f3..0961fc4 100644 --- a/src/gpgpu-sim/icnt_wrapper.h +++ b/src/gpgpu-sim/icnt_wrapper.h @@ -75,26 +75,20 @@ typedef void (*icnt_transfer_p)( ); typedef unsigned (*icnt_busy_p)( ); typedef void (*icnt_drain_p)( ); - extern icnt_has_buffer_p icnt_has_buffer; extern icnt_push_p icnt_push; extern icnt_pop_p icnt_pop; extern icnt_transfer_p icnt_transfer; extern icnt_busy_p icnt_busy; extern icnt_drain_p icnt_drain; +extern int g_network_mode; -/* -// definition of valid gpu network mode. -extern enum { - INTERSIM = 1, - N_NETWORK_MODE -} gpu_network_mode; -*/ enum network_mode { INTERSIM = 1, N_NETWORK_MODE }; void icnt_init( unsigned int n_shader, unsigned int n_mem ); +void icnt_reg_options( class OptionParser * opp ); #endif diff --git a/src/gpgpu-sim/visualizer.cc b/src/gpgpu-sim/visualizer.cc index d0f406c..0a2d920 100644 --- a/src/gpgpu-sim/visualizer.cc +++ b/src/gpgpu-sim/visualizer.cc @@ -60,6 +60,7 @@ */ #include "gpu-sim.h" +#include "shader.h" #include "../option_parser.h" #include #include diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 6ce564b..29d9787 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -64,69 +64,52 @@ #include "gpgpusim_entrypoint.h" #include -#include #include "option_parser.h" #include "cuda-sim/cuda-sim.h" #include "cuda-sim/ptx_ir.h" #include "cuda-sim/ptx_parser.h" +#include "gpgpu-sim/gpu-sim.h" +#include "gpgpu-sim/icnt_wrapper.h" #define MAX(a,b) (((a)>(b))?(a):(b)) struct gpgpu_ptx_sim_arg *grid_params; -int g_grid_num=0; -int g_argc = 3; -const char *g_argv[] = {"", "-config","gpgpusim.config"}; +static int sg_grid_num=0; +static int sg_argc = 3; +static const char *sg_argv[] = {"", "-config","gpgpusim.config"}; -unsigned int run_gpu_sim(int grid_num); -void gpgpu_ptx_sim_init_grid(const char *kernel_key,struct gpgpu_ptx_sim_arg *args, struct dim3 gridDim, struct dim3 blockDim ); - -int g_network_mode = 0; -char* g_network_config_filename; -option_parser_t opp; -extern void print_splash(); - -extern void gpu_reg_options(option_parser_t opp); -extern void init_gpu(); - -time_t simulation_starttime; +time_t g_simulation_starttime; void gpgpu_ptx_sim_init_perf() { print_splash(); read_sim_environment_variables(); read_parser_environment_variables(); - opp = option_parser_create(); - option_parser_register(opp, "-network_mode", OPT_INT32, &g_network_mode, "Interconnection network mode", "1"); - option_parser_register(opp, "-inter_config_file", OPT_CSTR, &g_network_config_filename, "Interconnection network config file", "mesh"); + option_parser_t opp = option_parser_create(); + icnt_reg_options(opp); gpu_reg_options(opp); // register GPU microrachitecture options ptx_reg_options(opp); - option_parser_cmdline(opp, g_argc, g_argv); // parse configuration options + option_parser_cmdline(opp, sg_argc, sg_argv); // parse configuration options srand(1); fprintf(stdout, "GPGPU-Sim: Configuration options:\n\n"); option_parser_print(opp, stdout); - init_gpu(); // initialize the GPU microarchitecture model + init_gpu(); fprintf(stdout, "GPU performance model initialization complete.\n"); - simulation_starttime = time((time_t *)NULL); + g_simulation_starttime = time((time_t *)NULL); } extern unsigned long long gpu_tot_sim_insn; extern unsigned long long gpu_tot_sim_cycle; -int gpgpu_ptx_sim_main_perf( const char *kernel_key, struct dim3 gridDim, struct dim3 blockDim, struct gpgpu_ptx_sim_arg *grid_params ) +static void print_simulation_time() { time_t current_time, difference, d, h, m, s; - gpgpu_ptx_sim_init_grid(kernel_key,grid_params,gridDim,blockDim); - - run_gpu_sim(g_grid_num); // run a CUDA grid on the GPU microarchitecture simulator - - g_grid_num++; - current_time = time((time_t *)NULL); - difference = MAX(current_time - simulation_starttime, 1); + difference = MAX(current_time - g_simulation_starttime, 1); d = difference/(3600*24); h = difference/3600 - 24*d; @@ -139,6 +122,35 @@ int gpgpu_ptx_sim_main_perf( const char *kernel_key, struct dim3 gridDim, struct printf("gpgpu_simulation_rate = %u (inst/sec)\n", (unsigned)(gpu_tot_sim_insn / difference) ); printf("gpgpu_simulation_rate = %u (cycle/sec)\n", (unsigned)(gpu_tot_sim_cycle / difference) ); fflush(stdout); +} +int gpgpu_cuda_ptx_sim_main_perf( const char *kernel_key, + struct dim3 gridDim, + struct dim3 blockDim, + struct gpgpu_ptx_sim_arg *grid_params ) +{ + gpgpu_cuda_ptx_sim_init_grid(kernel_key,grid_params,gridDim,blockDim); + run_gpu_sim(sg_grid_num++); + print_simulation_time(); return 0; } + +int gpgpu_opencl_ptx_sim_main_perf( class function_info *entry, + struct dim3 gridDim, + struct dim3 blockDim, + struct gpgpu_ptx_sim_arg *grid_params ) +{ + gpgpu_opencl_ptx_sim_init_grid(entry,grid_params,gridDim,blockDim); + run_gpu_sim(sg_grid_num++); + print_simulation_time(); + return 0; +} + +int gpgpu_opencl_ptx_sim_main_func( class function_info *entry, + struct dim3 gridDim, + struct dim3 blockDim, + struct gpgpu_ptx_sim_arg *grid_params ) +{ + printf("GPGPU-Sim PTX API: OpenCL functional-only simulation not yet implemented (use performance simulation)\n"); + exit(1); +} diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h index 8de05c7..2563c5e 100644 --- a/src/gpgpusim_entrypoint.h +++ b/src/gpgpusim_entrypoint.h @@ -67,7 +67,23 @@ #include "abstract_hardware_model.h" +#include +extern time_t g_simulation_starttime; + void gpgpu_ptx_sim_init_perf(); -int gpgpu_ptx_sim_main_perf( const char *kernel_key, struct dim3 gridDim, struct dim3 blockDim, struct gpgpu_ptx_sim_arg *grid_params ); +int gpgpu_cuda_ptx_sim_main_perf( const char *kernel_key, + struct dim3 gridDim, + struct dim3 blockDim, + struct gpgpu_ptx_sim_arg *grid_params ); + +int gpgpu_opencl_ptx_sim_main_perf( class function_info *entry, + struct dim3 gridDim, + struct dim3 blockDim, + struct gpgpu_ptx_sim_arg *grid_params ); + +int gpgpu_opencl_ptx_sim_main_func( class function_info *entry, + struct dim3 gridDim, + struct dim3 blockDim, + struct gpgpu_ptx_sim_arg *grid_params ); #endif diff --git a/src/option_parser.h b/src/option_parser.h index e14c08e..5e27ef1 100644 --- a/src/option_parser.h +++ b/src/option_parser.h @@ -67,7 +67,7 @@ #include // pointer to C++ class -typedef void* option_parser_t; +typedef class OptionParser *option_parser_t; // data type of the option enum option_dtype { -- cgit v1.3