diff options
| -rw-r--r-- | libcuda/Makefile | 14 | ||||
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 447 | ||||
| -rw-r--r-- | src/abstract_hardware_model.h | 71 | ||||
| -rw-r--r-- | src/cuda-sim/cuda-sim.cc | 232 | ||||
| -rw-r--r-- | src/cuda-sim/cuda-sim.h | 15 | ||||
| -rw-r--r-- | src/cuda-sim/ptx_ir.cc | 9 | ||||
| -rw-r--r-- | src/cuda-sim/ptx_loader.cc | 268 | ||||
| -rw-r--r-- | src/cuda-sim/ptx_loader.h | 1 | ||||
| -rw-r--r-- | src/gpgpu-sim/gpu-sim.h | 70 | ||||
| -rw-r--r-- | src/gpgpusim_entrypoint.cc | 3 | ||||
| -rw-r--r-- | src/gpgpusim_entrypoint.h | 2 |
11 files changed, 560 insertions, 572 deletions
diff --git a/libcuda/Makefile b/libcuda/Makefile index e82a0e3..bdb48f4 100644 --- a/libcuda/Makefile +++ b/libcuda/Makefile @@ -66,19 +66,25 @@ ifeq ($(OPENGL_SUPPORT),1) GL = -DOPENGL_SUPPORT endif +GNUC_CPP0X := $(shell gcc --version | perl -ne 'if (/gcc\s+\(.*\)\s+([0-9.]+)/){ if($$1 >= 4.3) {$$n=1} else {$$n=0;} } END { print $$n; }') +ifeq ($(GNUC_CPP0X), 1) + CXXFLAGS = -std=c++0x +endif + + CPP = g++ $(SNOW) CC = gcc &(SNOW) CREATELIBRARY = 1 DEBUG ?= 0 -CCFLAGS = -O3 -g -Wall -fPIC $(GL) +CXXFLAGS += -O3 -g -Wall -fPIC $(GL) ifeq ($(DEBUG),1) - CCFLAGS = -Wall -g -fPIC $(GL) + CXXFLAGS += -Wall -g -fPIC $(GL) endif PROG =cuda CXX_SRCS = cuda_runtime_api.cc -CCFLAGS += -DCUDART_VERSION=$(CUDART_VERSION) +CXXFLAGS += -DCUDART_VERSION=$(CUDART_VERSION) .PHONY: clean @@ -90,7 +96,7 @@ lib$(PROG).a: $(OBJS) ar rcs lib$(PROG).a $(OBJS) %.o: %.cc - $(CPP) $(CCFLAGS) -I./ -I$(CUDAHOME)/include -c $< -o $@ + $(CPP) $(CXXFLAGS) -I./ -I$(CUDAHOME)/include -c $< -o $@ clean: rm -f $(PROG) diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index d8092e8..3a857d7 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -100,6 +100,7 @@ * include, in the user documentation and internal comments to the code, * the above Disclaimer and U.S. Government End Users Notice. */ + #include <stdlib.h> #include <stdio.h> #include <string.h> @@ -114,21 +115,33 @@ #include <GL/gl.h> #endif #endif +#include <dirent.h> #define __CUDA_RUNTIME_API_H__ -/******************************************************************************* -* * -* * -* * -*******************************************************************************/ - - #include "host_defines.h" #include "builtin_types.h" #include "driver_types.h" #include "__cudaFatFormat.h" #include "../src/gpgpu-sim/gpu-sim.h" +#include "../src/cuda-sim/ptx_loader.h" +#include "../src/cuda-sim/cuda-sim.h" +#include "../src/cuda-sim/ptx_ir.h" +#include "../src/cuda-sim/ptx_parser.h" +#include "../src/gpgpusim_entrypoint.h" + +static void gpgpu_ptx_sim_load_gpu_kernels(); +static const struct gpgpu_ptx_sim_kernel_info * get_kernel_info(const char *kernel_key); +static void register_function_implementation( const char *name, class function_info *impl ); +static void ptxinfo_cuda_addinfo(); +static void gpgpu_ptx_sim_add_ptxstring( unsigned fat_cubin_handle, + const char *ptx_string, + const char *cubin_string, + const char *sourcefname, + unsigned capability ); +static kernel_info_t gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, + gpgpu_ptx_sim_arg_list_t args, + struct dim3 gridDim, struct dim3 blockDim ); /*DEVICE_BUILTIN*/ struct cudaArray @@ -142,7 +155,6 @@ struct cudaArray unsigned dimensions; }; - #if !defined(__dv) #if defined(__cplusplus) #define __dv(v) \ @@ -152,18 +164,8 @@ struct cudaArray #endif /* __cplusplus */ #endif /* !__dv */ -/******************************************************************************* -* * -* * -* * -*******************************************************************************/ - cudaError_t g_last_cudaError = cudaSuccess; -#include "../src/cuda-sim/ptx_loader.h" -#include "../src/cuda-sim/cuda-sim.h" -#include "../src/gpgpusim_entrypoint.h" - void register_ptx_function( const char *name, function_info *impl ) { register_function_implementation( name, impl ); @@ -873,10 +875,11 @@ __host__ cudaError_t CUDARTAPI cudaLaunch(const char *symbol ) g_ptx_sim_mode?"functional simulation":"performance simulation"); gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); kernel_config &config = g_cuda_launch_stack.back(); + kernel_info_t grid = gpgpu_cuda_ptx_sim_init_grid(symbol,config.get_args(),config.grid_dim(),config.block_dim()); if( g_ptx_sim_mode ) - gpgpu_cuda_ptx_sim_main_func( symbol, config.grid_dim(), config.block_dim(), config.get_args() ); + gpgpu_cuda_ptx_sim_main_func( grid, config.grid_dim(), config.block_dim(), config.get_args() ); else - gpgpu_cuda_ptx_sim_main_perf( symbol, config.grid_dim(), config.block_dim(), config.get_args() ); + gpgpu_cuda_ptx_sim_main_perf( grid, config.grid_dim(), config.block_dim(), config.get_args() ); return g_last_cudaError = cudaSuccess; } @@ -1358,3 +1361,407 @@ int CUDARTAPI __cudaSynchronizeThreads(void**, void*) } } + +//////// + + + +struct ptx_info_t { + unsigned fat_cubin_handle; + char *str; + char *cubin_str; + char *fname; + ptx_info_t *next; + unsigned capability; +}; + +static ptx_info_t *g_ptx_source_array = NULL; +static unsigned g_ptx_max_capability = 0; + +extern "C" int ptx_parse(); +extern "C" int ptx__scan_string(const char*); +extern "C" FILE *ptx_in; + +extern "C" const char *g_ptxinfo_filename; +extern "C" int ptxinfo_parse(); +extern "C" int ptxinfo_debug; +extern "C" FILE *ptxinfo_in; + +/// static functions + + + +std::map<std::string,function_info*> *g_kernel_name_to_function_lookup=NULL; +std::map<const void*,std::string> *g_host_to_kernel_entrypoint_name_lookup=NULL; + +function_info *get_kernel(const char *kernel_key, std::string &kernel_func_name_mangled ) +{ + if ( g_host_to_kernel_entrypoint_name_lookup->find(kernel_key) == + g_host_to_kernel_entrypoint_name_lookup->end() ) { + printf("GPGPU-Sim PTX: ERROR ** cannot locate __global__ function from hostPtr\n" ); + printf("GPGPU-Sim PTX: registered PTX kernels: \n"); + std::map<const void*,std::string>::iterator i_eptr = g_host_to_kernel_entrypoint_name_lookup->begin(); + for (; i_eptr != g_host_to_kernel_entrypoint_name_lookup->end(); ++i_eptr) { + printf("GPGPU-Sim PTX: (%p,%s)\n", i_eptr->first, i_eptr->second.c_str()); + } + printf("\n"); + abort(); + } + kernel_func_name_mangled = (*g_host_to_kernel_entrypoint_name_lookup)[kernel_key]; + if ( g_kernel_name_to_function_lookup->find(kernel_func_name_mangled) == + g_kernel_name_to_function_lookup->end() ) { + printf("GPGPU-Sim PTX: ERROR ** function \'%s\' not found in ptx file\n", + kernel_func_name_mangled.c_str() ); + abort(); + } + return (*g_kernel_name_to_function_lookup)[kernel_func_name_mangled]; +} + +void gpgpu_ptx_sim_register_kernel(void **fatCubinHandle,const char *hostFun, const char *deviceFun) +{ + const void* key=hostFun; + print_splash(); + if ( g_host_to_kernel_entrypoint_name_lookup == NULL ) + g_host_to_kernel_entrypoint_name_lookup = new std::map<const void*,std::string>; + if( g_kernel_name_to_function_lookup == NULL ) + g_kernel_name_to_function_lookup = new std::map<std::string,function_info*>; + if ( g_host_to_kernel_entrypoint_name_lookup->find(key) != + g_host_to_kernel_entrypoint_name_lookup->end() ) { + printf("GPGPU-Sim Loader error: Don't know how to identify PTX kernels during cudaLaunch\n" + " for this application.\n"); + abort(); + } + (*g_host_to_kernel_entrypoint_name_lookup)[key] = deviceFun; + if( g_kernel_name_to_function_lookup->find(deviceFun) == + g_kernel_name_to_function_lookup->end() ) { + (*g_kernel_name_to_function_lookup)[deviceFun] = NULL; // we set this later, set keys now for error checking + } + + printf("GPGPU-Sim PTX: __cudaRegisterFunction %s : 0x%Lx\n", deviceFun, (unsigned long long)hostFun); +} + +void ptxinfo_cuda_addinfo() +{ + if( !strcmp("__cuda_dummy_entry__",get_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<std::string,function_info*>::iterator i=g_kernel_name_to_function_lookup->find(get_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", get_ptxinfo_kname() ); + abort(); + } else { + print_ptxinfo(); + function_info *fi = i->second; + assert(fi!=NULL); + fi->set_kernel_info(get_ptxinfo_kinfo()); + } + } else { + printf ("GPGPU-Sim PTX: ERROR ** Kernel '%s' not found (no kernels registered).\n", get_ptxinfo_kname() ); + abort(); + } + clear_ptxinfo(); +} + +static 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::string,function_info*>; + + std::map<std::string,function_info*>::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; +} + + +static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr) +{ + printf( "GPGPU-Sim PTX: loading globals with explicit initializers... \n" ); + fflush(stdout); + int ng_bytes=0; + symbol_table::iterator g=symtab->global_iterator_begin(); + + for ( ; g!=symtab->global_iterator_end(); g++) { + symbol *global = *g; + if ( global->has_initializer() ) { + printf( "GPGPU-Sim PTX: initializing '%s' ... ", global->name().c_str() ); + unsigned addr=global->get_address(); + const type_info *type = global->type(); + type_info_key ti=type->get_key(); + size_t size; + int t; + ti.type_decode(size,t); + int nbytes = size/8; + int offset=0; + std::list<operand_info> init_list = global->get_initializer(); + for ( std::list<operand_info>::iterator i=init_list.begin(); i!=init_list.end(); i++ ) { + operand_info op = *i; + ptx_reg_t value = op.get_literal_value(); + assert( (addr+offset+nbytes) < min_gaddr ); // min_gaddr is start of "heap" for cudaMalloc + g_global_mem->write(addr+offset,nbytes,&value,NULL,NULL); // assuming little endian here + offset+=nbytes; + ng_bytes+=nbytes; + } + printf(" wrote %u bytes\n", offset ); + } + } + printf( "GPGPU-Sim PTX: finished loading globals (%u bytes total).\n", ng_bytes ); + fflush(stdout); + return ng_bytes; +} + +static int load_constants( symbol_table *symtab, addr_t min_gaddr ) +{ + printf( "GPGPU-Sim PTX: loading constants with explicit initializers... " ); + fflush(stdout); + int nc_bytes = 0; + symbol_table::iterator g=symtab->const_iterator_begin(); + + for ( ; g!=symtab->const_iterator_end(); g++) { + symbol *constant = *g; + if ( constant->is_const() && constant->has_initializer() ) { + + // get the constant element data size + int basic_type; + size_t num_bits; + constant->type()->get_key().type_decode(num_bits,basic_type); + + std::list<operand_info> init_list = constant->get_initializer(); + int nbytes_written = 0; + for ( std::list<operand_info>::iterator i=init_list.begin(); i!=init_list.end(); i++ ) { + operand_info op = *i; + ptx_reg_t value = op.get_literal_value(); + int nbytes = num_bits/8; + switch ( op.get_type() ) { + case int_t: assert(nbytes >= 1); break; + case float_op_t: assert(nbytes == 4); break; + case double_op_t: assert(nbytes >= 4); break; // account for double DEMOTING + default: + abort(); + } + unsigned addr=constant->get_address() + nbytes_written; + assert( addr+nbytes < min_gaddr ); + + g_global_mem->write(addr,nbytes,&value,NULL,NULL); // assume little endian (so u8 is the first byte in u32) + nc_bytes+=nbytes; + nbytes_written += nbytes; + } + } + } + printf( " done.\n"); + fflush(stdout); + return nc_bytes; +} + +static FILE *open_ptxinfo (const char* ptx_filename) +{ + const int ptx_fnamelen = strlen(ptx_filename); + char *ptxi_fname = new char[ptx_fnamelen+5]; + strcpy (ptxi_fname, ptx_filename); + strcpy (ptxi_fname+ptx_fnamelen, "info"); + + //ptxinfo_debug=1; + g_ptxinfo_filename = ptxi_fname; + FILE *f = fopen (ptxi_fname, "rt"); + return f; +} + + +static int ptx_file_filter( +#if !defined(__APPLE__) + const +#endif + struct dirent *de ) +{ + const char *tmp = strstr(de->d_name,".ptx"); + if ( tmp != NULL && tmp[4] == 0 ) { + return 1; + } + return 0; +} + +extern unsigned g_ptx_force_max_capability; +extern int g_ptx_convert_to_ptxplus; + +void gpgpu_ptx_sim_add_ptxstring( unsigned fat_cubin_handle, const char *ptx_string, const char *cubin_string, const char *sourcefname, unsigned capability ) +{ + ptx_info_t *t = new ptx_info_t; + t->next = NULL; + t->str = strdup(ptx_string); + t->fat_cubin_handle = fat_cubin_handle; + if (cubin_string != NULL) { + t->cubin_str = strdup(cubin_string); + } else { + if(g_ptx_convert_to_ptxplus != 0) { + printf("GPGPU-Sim PTX: ERROR cubin information, required for ptxplus, missing from fat bin object\n"); + printf("GPGPU-Sim PTX: ptxplus currently requires CUDA 2.3 or earlier... disable ptxplus by removing\n"); + printf("GPGPU-Sim PTX: \'-gpgpu_ptx_convert_to_ptxplus 1\' in \'gpgpusim.config\'\n"); + abort(); + } + t->cubin_str = NULL; + } + t->fname = strdup(sourcefname); + t->capability = capability; + + // Set overall max capability + if(g_ptx_max_capability < capability) + g_ptx_max_capability = capability; + + // put ptx source into a fifo + if (g_ptx_source_array == NULL) { + // first ptx source + g_ptx_source_array = t; + } else { + // insert subsequent ptx source at the end of queue + ptx_info_t *l_ptx_source = g_ptx_source_array; + while (l_ptx_source->next != NULL) { + l_ptx_source = l_ptx_source->next; + } + l_ptx_source->next = t; + } +} + +void gpgpu_ptx_sim_load_gpu_kernels() +{ + static unsigned source_num = 0; + ptx_in = NULL; + if ( g_filename ) + ptx_in = fopen( g_filename, "r" ); + gpgpu_ptx_sim_init_memory(); + if (ptx_in) { + symbol_table *symtab=init_parser(g_filename); + ptx_parse(); + ptxinfo_in = open_ptxinfo(g_filename); + ptxinfo_parse(); + 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"); + unsigned selected_capability = 0; + + if(g_ptx_force_max_capability == 0) { + // No forced max capability, selected the highest capability + //assert(g_ptx_max_capability > 0); + selected_capability = g_ptx_max_capability; + printf("GPGPU-Sim PTX: Loading PTX, selected capability: compute_%u\n", selected_capability); + } else { + // Forced max capability, select the highest capability less than or equal to forced capability + ptx_info_t *pti; + for( pti=g_ptx_source_array; pti!=NULL; pti=pti->next ){ + if(selected_capability < pti->capability && pti->capability <= g_ptx_force_max_capability) + selected_capability = pti->capability; + } + //assert(selected_capability > 0); + printf("GPGPU-Sim PTX: Loading PTX, max forced capability: compute_%u, selected capability: compute_%u\n", + g_ptx_force_max_capability, selected_capability); + } + + ptx_info_t *s; + for ( s=g_ptx_source_array; s!=NULL; s=s->next ) { + if(s->capability != selected_capability) + continue; + + printf("GPGPU-Sim PTX: Loading PTX for %s, capability = compute_%u\n", s->fname, s->capability); + + symbol_table *symtab; + source_num++; + if(g_ptx_convert_to_ptxplus) { + char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_to_ptxplus(s->str, s->cubin_str, source_num); + symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxplus_str, s->str, source_num); + delete[] ptxplus_str; + } else { + symtab=gpgpu_ptx_sim_load_ptx_from_string(s->str, s->str, source_num); + } + load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF); + load_constants(symtab,STATIC_ALLOC_LIMIT); + } + } else { + if(g_ptx_convert_to_ptxplus) { + perror("GPGPU-Sim PTX: convert_to_ptxplus option enabled. Cannot use this option with external ptx files.\n"); + assert(0); + } + const char *filename = NULL; + struct dirent **namelist; + int n = scandir(".", &namelist, ptx_file_filter, alphasort); + if (n < 0) + perror("GPGPU-Sim PTX: no PTX files returned by scandir"); + else { + while (n--) { + filename = namelist[n]->d_name; + printf("Parsing %s..\n", filename); + ptx_in = fopen( filename, "r" ); + symbol_table *symtab=init_parser(filename); + ptx_parse (); + ptxinfo_in = open_ptxinfo(filename); + ptxinfo_parse(); + load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF); + load_constants(symtab,STATIC_ALLOC_LIMIT); + + free(namelist[n]); + } + free(namelist); + } + } + } + + if ( ptx_in == NULL && g_override_embedded_ptx ) { + printf("GPGPU-Sim PTX Simulator error: Could find/open .ptx file for reading\n"); + printf(" This means there are no .ptx files in the current directory.\n"); + printf(" Either place a .ptx file in the current directory, or ensure\n" ); + printf(" the PTX_SIM_KERNELFILE environment variable points to .ptx file.\n"); + printf(" PTX_SIM_KERNELFILE=\"%s\"\n", g_filename ); + exit(1); + } + + if ( g_error_detected ) { + printf( "GPGPU-Sim PTX: PTX parsing errors detected -- exiting.\n" ); + exit(1); + } + printf( "GPGPU-Sim PTX: Program parsing completed\n" ); + + if ( g_kernel_name_to_function_lookup ) { + for ( std::map<std::string,function_info*>::iterator f=g_kernel_name_to_function_lookup->begin(); + f != g_kernel_name_to_function_lookup->end(); f++ ) { + gpgpu_ptx_assemble(f->first,f->second); + } + } +} + +const struct gpgpu_ptx_sim_kernel_info * get_kernel_info(const char *kernel_key) +{ + std::string kname; + function_info *finfo = get_kernel(kernel_key,kname); + return finfo->get_kernel_info(); +} + +kernel_info_t gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, gpgpu_ptx_sim_arg_list_t args, + struct dim3 gridDim, struct dim3 blockDim ) +{ + std::string kname; + function_info *entry = get_kernel(kernel_key,kname); + + printf("GPGPU-Sim PTX: Launching kernel \'%s\' gridDim= (%u,%u,%u) blockDim = (%u,%u,%u); ntuid=%u\n", + kname.c_str(), gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z, + g_ptx_thread_info_uid_next ); + + + unsigned argcount=args.size(); + unsigned argn=1; + for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) { + entry->add_param_data(argcount-argn,&(*a)); + argn++; + } + + entry->finalize(g_param_mem); + g_ptx_kernel_count++; + fflush(stdout); + + return kernel_info_t(gridDim,blockDim,entry); +} + diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index c2f5b1a..5d73da7 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -55,6 +55,76 @@ struct dim3 { }; #endif +void increment_x_then_y_then_z( dim3 &i, const dim3 &bound); + +class kernel_info_t { +public: + kernel_info_t() + { + m_valid=false; + m_kernel_entry=NULL; + } + kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry ) + { + m_valid=true; + m_kernel_entry=entry; + m_grid_dim=gridDim; + m_block_dim=blockDim; + m_next_cta.x=0; + m_next_cta.y=0; + m_next_cta.z=0; + m_next_tid=m_next_cta; + } + + class function_info *entry() { return m_kernel_entry; } + + size_t num_blocks() const + { + return m_grid_dim.x * m_grid_dim.y * m_grid_dim.z; + } + + size_t threads_per_cta() const + { + return m_block_dim.x * m_block_dim.y * m_block_dim.z; + } + + dim3 get_grid_dim() const { return m_grid_dim; } + dim3 get_cta_dim() const { return m_block_dim; } + + void increment_cta_id() + { + increment_x_then_y_then_z(m_next_cta,m_grid_dim); + m_next_tid.x=0; + m_next_tid.y=0; + m_next_tid.z=0; + } + dim3 get_next_cta_id() const { return m_next_cta; } + bool no_more_ctas_to_run() const + { + return (m_next_cta.x >= m_grid_dim.x || m_next_cta.y >= m_grid_dim.y || m_next_cta.z >= m_grid_dim.z ); + } + + void increment_thread_id() { increment_x_then_y_then_z(m_next_tid,m_block_dim); } + dim3 get_next_thread_id_3d() const { return m_next_tid; } + unsigned get_next_thread_id() const + { + return m_next_tid.x + m_block_dim.x*m_next_tid.y + m_block_dim.x*m_block_dim.y*m_next_tid.z; + } + bool more_threads_in_cta() const + { + return m_next_tid.z < m_block_dim.z && m_next_tid.y < m_block_dim.y && m_next_tid.z < m_block_dim.x; + } + +private: + bool m_valid; + class function_info *m_kernel_entry; + + dim3 m_grid_dim; + dim3 m_block_dim; + dim3 m_next_cta; + dim3 m_next_tid; +}; + class core_t { public: virtual ~core_t() {} @@ -190,7 +260,6 @@ protected: }; -const struct gpgpu_ptx_sim_kernel_info * get_kernel_info(const char *kernel_key); size_t get_kernel_code_size( class function_info *entry ); #endif // #ifdef __cplusplus diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index afadbe4..80aa60b 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -902,6 +902,22 @@ unsigned datatype2size( unsigned data_type ) return data_size; } +void init_inst_classification_stat() +{ + static bool init=false; + if( init ) return; + init=true; + char kernelname[256] =""; +#define MAX_CLASS_KER 1024 + if (!g_inst_classification_stat) g_inst_classification_stat = (void**)calloc(MAX_CLASS_KER, sizeof(void*)); + snprintf(kernelname, MAX_CLASS_KER, "Kernel %d Classification\n",g_ptx_kernel_count ); + assert( g_ptx_kernel_count < MAX_CLASS_KER ) ; // a static limit on number of kernels increase it if it fails! + g_inst_classification_stat[g_ptx_kernel_count] = StatCreate(kernelname,1,20); + if (!g_inst_op_classification_stat) g_inst_op_classification_stat = (void**)calloc(MAX_CLASS_KER, sizeof(void*)); + snprintf(kernelname, MAX_CLASS_KER, "Kernel %d OP Classification\n",g_ptx_kernel_count ); + g_inst_op_classification_stat[g_ptx_kernel_count] = StatCreate(kernelname,1,100); +} + unsigned g_warp_active_mask; void ptx_thread_info::ptx_exec_inst( inst_t &inst ) @@ -1058,6 +1074,7 @@ void ptx_thread_info::ptx_exec_inst( inst_t &inst ) g_ptx_sim_num_insn++; ptx_file_line_stats_add_exec_count(pI); if ( gpgpu_ptx_instruction_classification ) { + init_inst_classification_stat(); unsigned space_type=0; switch ( pI->get_space().get_type() ) { case global_space: space_type = 10; break; @@ -1104,10 +1121,7 @@ void ptx_thread_info::ptx_exec_inst( inst_t &inst ) } } -unsigned g_cta_launch_sid; std::list<ptx_thread_info *> g_active_threads; -std::map<unsigned,unsigned> g_sm_idx_offset_next; -unsigned g_sm_next_index; std::map<unsigned,memory_space*> g_shared_memory_lookup; std::map<unsigned,ptx_cta_info*> g_ptx_cta_lookup; std::map<unsigned,std::map<unsigned,memory_space*> > g_local_memory_lookup; @@ -1155,9 +1169,6 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, if ( !g_active_threads.empty() ) { //if g_active_threads not empty... assert( g_active_threads.size() <= threads_left ); - if ( g_cta_launch_sid == (unsigned)-1 ) - g_cta_launch_sid = sid; - assert( g_cta_launch_sid == (unsigned)sid ); ptx_thread_info *thd = g_active_threads.front(); g_active_threads.pop_front(); *thread_info = thd; @@ -1187,7 +1198,6 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, memory_space *shared_mem = NULL; unsigned cta_size = kernel.threads_per_cta(); - unsigned sm_offset = g_sm_idx_offset_next[sid]; unsigned max_cta_per_sm = num_threads/cta_size; // e.g., 256 / 48 = 5 assert( max_cta_per_sm > 0 ); @@ -1195,8 +1205,8 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, if ( g_shared_memory_lookup.find(sm_idx) == g_shared_memory_lookup.end() ) { if ( g_debug_execution >= 1 ) { - printf(" <CTA alloc> : sm_idx=%u sid=%u sm_offset=%u max_cta_per_sm=%u\n", - sm_idx, sid, sm_offset, max_cta_per_sm ); + printf(" <CTA alloc> : sm_idx=%u sid=%u max_cta_per_sm=%u\n", + sm_idx, sid, max_cta_per_sm ); } char buf[512]; snprintf(buf,512,"shared_%u", sid); @@ -1206,8 +1216,8 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, g_ptx_cta_lookup[sm_idx] = cta_info; } else { if ( g_debug_execution >= 1 ) { - printf(" <CTA realloc> : sm_idx=%u sid=%u sm_offset=%u max_cta_per_sm=%u\n", - sm_idx, sid, sm_offset, max_cta_per_sm ); + printf(" <CTA realloc> : sm_idx=%u sid=%u max_cta_per_sm=%u\n", + sm_idx, sid, max_cta_per_sm ); } shared_mem = g_shared_memory_lookup[sm_idx]; cta_info = g_ptx_cta_lookup[sm_idx]; @@ -1267,11 +1277,8 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, kernel.increment_cta_id(); - g_cta_launch_sid = -1; - assert( g_active_threads.size() <= threads_left ); - g_cta_launch_sid = sid; *thread_info = g_active_threads.front(); (*thread_info)->set_hw_tid(tid); (*thread_info)->set_hw_wid(hw_warp_id); @@ -1283,92 +1290,14 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, return 1; } -void init_inst_classification_stat() { - char kernelname[256] =""; -#define MAX_CLASS_KER 1024 - if (!g_inst_classification_stat) g_inst_classification_stat = (void**)calloc(MAX_CLASS_KER, sizeof(void*)); - snprintf(kernelname, MAX_CLASS_KER, "Kernel %d Classification\n",g_ptx_kernel_count ); - assert( g_ptx_kernel_count < MAX_CLASS_KER ) ; // a static limit on number of kernels increase it if it fails! - g_inst_classification_stat[g_ptx_kernel_count] = StatCreate(kernelname,1,20); - if (!g_inst_op_classification_stat) g_inst_op_classification_stat = (void**)calloc(MAX_CLASS_KER, sizeof(void*)); - snprintf(kernelname, MAX_CLASS_KER, "Kernel %d OP Classification\n",g_ptx_kernel_count ); - g_inst_op_classification_stat[g_ptx_kernel_count] = StatCreate(kernelname,1,100); -} - - -std::map<std::string,function_info*> *g_kernel_name_to_function_lookup=NULL; -std::map<const void*,std::string> *g_host_to_kernel_entrypoint_name_lookup=NULL; - -function_info *get_kernel(const char *kernel_key, std::string &kernel_func_name_mangled ) -{ - if ( g_host_to_kernel_entrypoint_name_lookup->find(kernel_key) == - g_host_to_kernel_entrypoint_name_lookup->end() ) { - printf("GPGPU-Sim PTX: ERROR ** cannot locate __global__ function from hostPtr\n" ); - printf("GPGPU-Sim PTX: registered PTX kernels: \n"); - std::map<const void*,std::string>::iterator i_eptr = g_host_to_kernel_entrypoint_name_lookup->begin(); - for (; i_eptr != g_host_to_kernel_entrypoint_name_lookup->end(); ++i_eptr) { - printf("GPGPU-Sim PTX: (%p,%s)\n", i_eptr->first, i_eptr->second.c_str()); - } - printf("\n"); - abort(); - } - kernel_func_name_mangled = (*g_host_to_kernel_entrypoint_name_lookup)[kernel_key]; - if ( g_kernel_name_to_function_lookup->find(kernel_func_name_mangled) == - g_kernel_name_to_function_lookup->end() ) { - printf("GPGPU-Sim PTX: ERROR ** function \'%s\' not found in ptx file\n", - kernel_func_name_mangled.c_str() ); - abort(); - } - return (*g_kernel_name_to_function_lookup)[kernel_func_name_mangled]; -} - -const struct gpgpu_ptx_sim_kernel_info * get_kernel_info(const char *kernel_key) -{ - std::string kname; - function_info *finfo = get_kernel(kernel_key,kname); - return finfo->get_kernel_info(); -} - size_t get_kernel_code_size( class function_info *entry ) { return entry->get_function_size(); } -kernel_info_t gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, gpgpu_ptx_sim_arg_list_t args, - struct dim3 gridDim, struct dim3 blockDim ) -{ - g_sm_idx_offset_next.clear(); - g_sm_next_index = 0; - std::string kname; - function_info *entry = get_kernel(kernel_key,kname); - - printf("GPGPU-Sim PTX: Launching kernel \'%s\' gridDim= (%u,%u,%u) blockDim = (%u,%u,%u); ntuid=%u\n", - kname.c_str(), gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z, - g_ptx_thread_info_uid_next ); - - - unsigned argcount=args.size(); - unsigned argn=1; - for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) { - entry->add_param_data(argcount-argn,&(*a)); - argn++; - } - - entry->finalize(g_param_mem); - g_ptx_kernel_count++; - if ( gpgpu_ptx_instruction_classification ) { - init_inst_classification_stat(); - } - fflush(stdout); - - return kernel_info_t(gridDim,blockDim,entry); -} 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 ) { - g_sm_idx_offset_next.clear(); - g_sm_next_index = 0; - unsigned argcount=args.size(); unsigned argn=1; for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) { @@ -1377,9 +1306,6 @@ kernel_info_t gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,gpgpu_pt } entry->finalize(g_param_mem); g_ptx_kernel_count++; - if ( gpgpu_ptx_instruction_classification ) { - init_inst_classification_stat(); - } fflush(stdout); return kernel_info_t(gridDim,blockDim,entry); @@ -1396,29 +1322,6 @@ void print_splash() } } -void gpgpu_ptx_sim_register_kernel(void **fatCubinHandle,const char *hostFun, const char *deviceFun) -{ - const void* key=hostFun; - print_splash(); - if ( g_host_to_kernel_entrypoint_name_lookup == NULL ) - g_host_to_kernel_entrypoint_name_lookup = new std::map<const void*,std::string>; - if( g_kernel_name_to_function_lookup == NULL ) - g_kernel_name_to_function_lookup = new std::map<std::string,function_info*>; - if ( g_host_to_kernel_entrypoint_name_lookup->find(key) != - g_host_to_kernel_entrypoint_name_lookup->end() ) { - printf("GPGPU-Sim Loader error: Don't know how to identify PTX kernels during cudaLaunch\n" - " for this application.\n"); - abort(); - } - (*g_host_to_kernel_entrypoint_name_lookup)[key] = deviceFun; - if( g_kernel_name_to_function_lookup->find(deviceFun) == - g_kernel_name_to_function_lookup->end() ) { - (*g_kernel_name_to_function_lookup)[deviceFun] = NULL; // we set this later, set keys now for error checking - } - - printf("GPGPU-Sim PTX: __cudaRegisterFunction %s : 0x%Lx\n", deviceFun, (unsigned long long)hostFun); -} - std::map<const void*,std::string> g_const_name_lookup; // indexed by hostVar std::map<const void*,std::string> g_global_name_lookup; // indexed by hostVar std::set<std::string> g_globals; @@ -1567,21 +1470,12 @@ ptx_cta_info *g_func_cta_info = NULL; #define MAX(a,b) (((a)>(b))?(a):(b)) -void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, gpgpu_ptx_sim_arg_list_t args) +void gpgpu_cuda_ptx_sim_main_func( kernel_info_t kernel, dim3 gridDim, dim3 blockDim, gpgpu_ptx_sim_arg_list_t args) { printf("GPGPU-Sim: Performing Functional Simulation...\n"); - printf("ERROR: Need to derived core_t for functional simulation, functional simulation no longer operational\n"); - // also: need PDOM stack, etc... for functional simulation - exit(1); - time_t end_time, elapsed_time, days, hrs, minutes, sec; - - kernel_info_t kernel = gpgpu_cuda_ptx_sim_init_grid(kernel_key, args,gridDim,blockDim); - - std::string kname; - function_info *finfo = get_kernel(kernel_key,kname); - + function_info *finfo = kernel.entry(); memory_space *shared_mem = new memory_space_impl<16*1024>("shared",4); std::map<unsigned,memory_space*> lm_lookup; @@ -1698,13 +1592,33 @@ unsigned translate_pc_to_ptxlineno(unsigned pc) return ptx_line_number; } -int g_ptxinfo_error_detected; +// ptxinfo parser +int g_ptxinfo_error_detected; static char *g_ptxinfo_kname = NULL; static struct gpgpu_ptx_sim_kernel_info g_ptxinfo_kinfo; -static void clear_ptxinfo(); +const char *get_ptxinfo_kname() +{ + return g_ptxinfo_kname; +} + +void print_ptxinfo() +{ + printf ("GPGPU-Sim PTX: Kernel \'%s\' : regs=%u, lmem=%u, smem=%u, cmem=%u\n", + get_ptxinfo_kname(), + g_ptxinfo_kinfo.regs, + g_ptxinfo_kinfo.lmem, + g_ptxinfo_kinfo.smem, + g_ptxinfo_kinfo.cmem ); +} + + +struct gpgpu_ptx_sim_kernel_info get_ptxinfo_kinfo() +{ + return g_ptxinfo_kinfo; +} extern "C" void ptxinfo_function(const char *fname ) { @@ -1744,6 +1658,7 @@ void clear_ptxinfo() g_ptxinfo_kinfo.sm_target=0; } + void ptxinfo_opencl_addinfo( std::map<std::string,function_info*> &kernels ) { if( !strcmp("__cuda_dummy_entry__",g_ptxinfo_kname) ) { @@ -1769,37 +1684,6 @@ void ptxinfo_opencl_addinfo( std::map<std::string,function_info*> &kernels ) 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<std::string,function_info*>::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 { @@ -1867,27 +1751,3 @@ void dwf_process_reconv_pts(function_info *entry) } } -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::string,function_info*>; - - std::map<std::string,function_info*>::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<std::string,function_info*>::iterator i=g_kernel_name_to_function_lookup->find(kernel_name); - - if( i == g_kernel_name_to_function_lookup->end() ) { - abort(); - return NULL; - } - return i->second; -} - diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index fa42341..4a0185a 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -16,22 +16,17 @@ extern int g_ptx_sim_mode; extern memory_space *g_global_mem; extern int g_debug_execution; extern int g_debug_thread_uid; -extern std::map<std::string,function_info*> *g_kernel_name_to_function_lookup; 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 FILE* ptx_inst_debug_file; -extern class kernel_info_t gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, - gpgpu_ptx_sim_arg_list_t args, - struct dim3 gridDim, - struct dim3 blockDim ); 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 ); -extern void gpgpu_cuda_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDim, gpgpu_ptx_sim_arg_list_t ); +extern void gpgpu_cuda_ptx_sim_main_func( kernel_info_t kernel, dim3 gridDim, dim3 blockDim, gpgpu_ptx_sim_arg_list_t args); extern void print_splash(); extern void* gpgpu_ptx_sim_malloc( size_t count ); extern void* gpgpu_ptx_sim_mallocarray( size_t count ); @@ -40,7 +35,6 @@ extern void gpgpu_ptx_sim_memcpy_from_gpu( void *dst, size_t src_start_addr, s extern void gpgpu_ptx_sim_memcpy_gpu_to_gpu( size_t dst, size_t src, size_t count ); extern void gpgpu_ptx_sim_memset( size_t dst_start_addr, int c, size_t count ); extern void gpgpu_ptx_sim_init_memory(); -extern void gpgpu_ptx_sim_load_gpu_kernels(); extern void gpgpu_ptx_sim_register_kernel(void **fatCubinHandle,const char *hostFun, const char *deviceFun); 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 ); @@ -52,8 +46,6 @@ 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_function_implementation( const char *name, function_info *impl ); -extern void ptxinfo_cuda_addinfo(); extern void ptxinfo_opencl_addinfo( std::map<std::string,function_info*> &kernels ); unsigned ptx_sim_init_thread( kernel_info_t &kernel, class ptx_thread_info** thread_info, @@ -76,9 +68,12 @@ void ptx_thread_release_barrier( void *thd ); void ptx_print_insn( address_type pc, FILE *fp ); unsigned int ptx_set_tex_cache_linesize( unsigned linesize); -function_info *get_kernel(const char *kernel_key, std::string &kernel_func_name_mangled ); void dwf_process_reconv_pts(function_info *entry); void set_param_gpgpu_num_shaders(int num_shaders); unsigned int get_converge_point(unsigned int pc, void *thd); +const char *get_ptxinfo_kname(); +void print_ptxinfo(); +void clear_ptxinfo(); +struct gpgpu_ptx_sim_kernel_info get_ptxinfo_kinfo(); #endif diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index fd51de5..2d63227 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -1206,15 +1206,6 @@ void function_info::print_insn( unsigned pc, FILE * fp ) const } } -extern "C" int ptx_parse(); -extern "C" int ptx__scan_string(const char*); -extern "C" FILE *ptx_in; - -extern "C" const char *g_ptxinfo_filename; -extern "C" int ptxinfo_parse(); -extern "C" int ptxinfo_debug; -extern "C" FILE *ptxinfo_in; - void gpgpu_ptx_assemble( std::string kname, void *kinfo ) { function_info *func_info = (function_info *)kinfo; diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index c836fac..d4f3822 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -77,286 +77,18 @@ memory_space *g_surf_mem; memory_space *g_param_mem; bool g_override_embedded_ptx = false; -struct ptx_info_t { - unsigned fat_cubin_handle; - char *str; - char *cubin_str; - char *fname; - ptx_info_t *next; - unsigned capability; -}; - /// extern prototypes extern "C" int ptx_parse(); extern "C" int ptx__scan_string(const char*); -extern "C" FILE *ptx_in; extern "C" const char *g_ptxinfo_filename = NULL; extern "C" int ptxinfo_parse(); extern "C" int ptxinfo_debug; extern "C" FILE *ptxinfo_in; -extern int g_ptx_convert_to_ptxplus; extern int g_ptx_save_converted_ptxplus; -extern unsigned g_ptx_force_max_capability; - -/// static functions - -static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr) -{ - printf( "GPGPU-Sim PTX: loading globals with explicit initializers... \n" ); - fflush(stdout); - int ng_bytes=0; - symbol_table::iterator g=symtab->global_iterator_begin(); - - for ( ; g!=symtab->global_iterator_end(); g++) { - symbol *global = *g; - if ( global->has_initializer() ) { - printf( "GPGPU-Sim PTX: initializing '%s' ... ", global->name().c_str() ); - unsigned addr=global->get_address(); - const type_info *type = global->type(); - type_info_key ti=type->get_key(); - size_t size; - int t; - ti.type_decode(size,t); - int nbytes = size/8; - int offset=0; - std::list<operand_info> init_list = global->get_initializer(); - for ( std::list<operand_info>::iterator i=init_list.begin(); i!=init_list.end(); i++ ) { - operand_info op = *i; - ptx_reg_t value = op.get_literal_value(); - assert( (addr+offset+nbytes) < min_gaddr ); // min_gaddr is start of "heap" for cudaMalloc - g_global_mem->write(addr+offset,nbytes,&value,NULL,NULL); // assuming little endian here - offset+=nbytes; - ng_bytes+=nbytes; - } - printf(" wrote %u bytes\n", offset ); - } - } - printf( "GPGPU-Sim PTX: finished loading globals (%u bytes total).\n", ng_bytes ); - fflush(stdout); - return ng_bytes; -} - -static int load_constants( symbol_table *symtab, addr_t min_gaddr ) -{ - printf( "GPGPU-Sim PTX: loading constants with explicit initializers... " ); - fflush(stdout); - int nc_bytes = 0; - symbol_table::iterator g=symtab->const_iterator_begin(); - - for ( ; g!=symtab->const_iterator_end(); g++) { - symbol *constant = *g; - if ( constant->is_const() && constant->has_initializer() ) { - // get the constant element data size - int basic_type; - size_t num_bits; - constant->type()->get_key().type_decode(num_bits,basic_type); - - std::list<operand_info> init_list = constant->get_initializer(); - int nbytes_written = 0; - for ( std::list<operand_info>::iterator i=init_list.begin(); i!=init_list.end(); i++ ) { - operand_info op = *i; - ptx_reg_t value = op.get_literal_value(); - int nbytes = num_bits/8; - switch ( op.get_type() ) { - case int_t: assert(nbytes >= 1); break; - case float_op_t: assert(nbytes == 4); break; - case double_op_t: assert(nbytes >= 4); break; // account for double DEMOTING - default: - abort(); - } - unsigned addr=constant->get_address() + nbytes_written; - assert( addr+nbytes < min_gaddr ); - - g_global_mem->write(addr,nbytes,&value,NULL,NULL); // assume little endian (so u8 is the first byte in u32) - nc_bytes+=nbytes; - nbytes_written += nbytes; - } - } - } - printf( " done.\n"); - fflush(stdout); - return nc_bytes; -} - -static FILE *open_ptxinfo (const char* ptx_filename) -{ - const int ptx_fnamelen = strlen(ptx_filename); - char *ptxi_fname = new char[ptx_fnamelen+5]; - strcpy (ptxi_fname, ptx_filename); - strcpy (ptxi_fname+ptx_fnamelen, "info"); - - //ptxinfo_debug=1; - g_ptxinfo_filename = ptxi_fname; - FILE *f = fopen (ptxi_fname, "rt"); - return f; -} - -static int ptx_file_filter( -#if !defined(__APPLE__) - const -#endif - struct dirent *de ) -{ - const char *tmp = strstr(de->d_name,".ptx"); - if ( tmp != NULL && tmp[4] == 0 ) { - return 1; - } - return 0; -} - -// global functions - -static ptx_info_t *g_ptx_source_array = NULL; -static unsigned g_ptx_max_capability = 0; - -void gpgpu_ptx_sim_load_gpu_kernels() -{ - static unsigned source_num = 0; - ptx_in = NULL; - if ( g_filename ) - ptx_in = fopen( g_filename, "r" ); - gpgpu_ptx_sim_init_memory(); - if (ptx_in) { - symbol_table *symtab=init_parser(g_filename); - ptx_parse(); - ptxinfo_in = open_ptxinfo(g_filename); - ptxinfo_parse(); - 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"); - unsigned selected_capability = 0; - - if(g_ptx_force_max_capability == 0) { - // No forced max capability, selected the highest capability - //assert(g_ptx_max_capability > 0); - selected_capability = g_ptx_max_capability; - printf("GPGPU-Sim PTX: Loading PTX, selected capability: compute_%u\n", selected_capability); - } else { - // Forced max capability, select the highest capability less than or equal to forced capability - ptx_info_t *pti; - for( pti=g_ptx_source_array; pti!=NULL; pti=pti->next ){ - if(selected_capability < pti->capability && pti->capability <= g_ptx_force_max_capability) - selected_capability = pti->capability; - } - //assert(selected_capability > 0); - printf("GPGPU-Sim PTX: Loading PTX, max forced capability: compute_%u, selected capability: compute_%u\n", - g_ptx_force_max_capability, selected_capability); - } - - ptx_info_t *s; - for ( s=g_ptx_source_array; s!=NULL; s=s->next ) { - if(s->capability != selected_capability) - continue; - - printf("GPGPU-Sim PTX: Loading PTX for %s, capability = compute_%u\n", s->fname, s->capability); - - symbol_table *symtab; - source_num++; - if(g_ptx_convert_to_ptxplus) { - char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_to_ptxplus(s->str, s->cubin_str, source_num); - symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxplus_str, s->str, source_num); - delete[] ptxplus_str; - } else { - symtab=gpgpu_ptx_sim_load_ptx_from_string(s->str, s->str, source_num); - } - load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF); - load_constants(symtab,STATIC_ALLOC_LIMIT); - } - } else { - if(g_ptx_convert_to_ptxplus) { - perror("GPGPU-Sim PTX: convert_to_ptxplus option enabled. Cannot use this option with external ptx files.\n"); - assert(0); - } - const char *filename = NULL; - struct dirent **namelist; - int n = scandir(".", &namelist, ptx_file_filter, alphasort); - if (n < 0) - perror("GPGPU-Sim PTX: no PTX files returned by scandir"); - else { - while (n--) { - filename = namelist[n]->d_name; - printf("Parsing %s..\n", filename); - ptx_in = fopen( filename, "r" ); - symbol_table *symtab=init_parser(filename); - ptx_parse (); - ptxinfo_in = open_ptxinfo(filename); - ptxinfo_parse(); - load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF); - load_constants(symtab,STATIC_ALLOC_LIMIT); - - free(namelist[n]); - } - free(namelist); - } - } - } - - if ( ptx_in == NULL && g_override_embedded_ptx ) { - printf("GPGPU-Sim PTX Simulator error: Could find/open .ptx file for reading\n"); - printf(" This means there are no .ptx files in the current directory.\n"); - printf(" Either place a .ptx file in the current directory, or ensure\n" ); - printf(" the PTX_SIM_KERNELFILE environment variable points to .ptx file.\n"); - printf(" PTX_SIM_KERNELFILE=\"%s\"\n", g_filename ); - exit(1); - } - - if ( g_error_detected ) { - printf( "GPGPU-Sim PTX: PTX parsing errors detected -- exiting.\n" ); - exit(1); - } - printf( "GPGPU-Sim PTX: Program parsing completed\n" ); - - if ( g_kernel_name_to_function_lookup ) { - for ( std::map<std::string,function_info*>::iterator f=g_kernel_name_to_function_lookup->begin(); - f != g_kernel_name_to_function_lookup->end(); f++ ) { - gpgpu_ptx_assemble(f->first,f->second); - } - } -} - -void gpgpu_ptx_sim_add_ptxstring( unsigned fat_cubin_handle, const char *ptx_string, const char *cubin_string, const char *sourcefname, unsigned capability ) -{ - ptx_info_t *t = new ptx_info_t; - t->next = NULL; - t->str = strdup(ptx_string); - t->fat_cubin_handle = fat_cubin_handle; - if (cubin_string != NULL) { - t->cubin_str = strdup(cubin_string); - } else { - if(g_ptx_convert_to_ptxplus != 0) { - printf("GPGPU-Sim PTX: ERROR cubin information, required for ptxplus, missing from fat bin object\n"); - printf("GPGPU-Sim PTX: ptxplus currently requires CUDA 2.3 or earlier... disable ptxplus by removing\n"); - printf("GPGPU-Sim PTX: \'-gpgpu_ptx_convert_to_ptxplus 1\' in \'gpgpusim.config\'\n"); - abort(); - } - t->cubin_str = NULL; - } - t->fname = strdup(sourcefname); - t->capability = capability; - - // Set overall max capability - if(g_ptx_max_capability < capability) - g_ptx_max_capability = capability; - - // put ptx source into a fifo - if (g_ptx_source_array == NULL) { - // first ptx source - g_ptx_source_array = t; - } else { - // insert subsequent ptx source at the end of queue - ptx_info_t *l_ptx_source = g_ptx_source_array; - while (l_ptx_source->next != NULL) { - l_ptx_source = l_ptx_source->next; - } - l_ptx_source->next = t; - } -} static bool g_save_embedded_ptx; bool g_keep_intermediate_files; diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h index 6b9d2f8..e9efd9b 100644 --- a/src/cuda-sim/ptx_loader.h +++ b/src/cuda-sim/ptx_loader.h @@ -74,7 +74,6 @@ extern memory_space *g_surf_mem; extern memory_space *g_param_mem; extern bool g_override_embedded_ptx; -void gpgpu_ptx_sim_add_ptxstring( unsigned fat_cubin_handle, const char *ptx_string, const char *cubin_string, const char *sourcefname, unsigned capability ); class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, const char *p_for_info, unsigned source_num ); char* gpgpu_ptx_sim_convert_ptx_to_ptxplus(const char *ptx_str, const char *cubin_str, unsigned source_num); diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 48a4c52..081cacf 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -96,76 +96,6 @@ #define SAMPLELOG 222 #define DUMPLOG 333 -void increment_x_then_y_then_z( dim3 &i, const dim3 &bound); - -class kernel_info_t { -public: - kernel_info_t() - { - m_valid=false; - m_kernel_entry=NULL; - } - kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry ) - { - m_valid=true; - m_kernel_entry=entry; - m_grid_dim=gridDim; - m_block_dim=blockDim; - m_next_cta.x=0; - m_next_cta.y=0; - m_next_cta.z=0; - m_next_tid=m_next_cta; - } - - class function_info *entry() { return m_kernel_entry; } - - size_t num_blocks() const - { - return m_grid_dim.x * m_grid_dim.y * m_grid_dim.z; - } - - size_t threads_per_cta() const - { - return m_block_dim.x * m_block_dim.y * m_block_dim.z; - } - - dim3 get_grid_dim() const { return m_grid_dim; } - dim3 get_cta_dim() const { return m_block_dim; } - - void increment_cta_id() - { - increment_x_then_y_then_z(m_next_cta,m_grid_dim); - m_next_tid.x=0; - m_next_tid.y=0; - m_next_tid.z=0; - } - dim3 get_next_cta_id() const { return m_next_cta; } - bool no_more_ctas_to_run() const - { - return (m_next_cta.x >= m_grid_dim.x || m_next_cta.y >= m_grid_dim.y || m_next_cta.z >= m_grid_dim.z ); - } - - void increment_thread_id() { increment_x_then_y_then_z(m_next_tid,m_block_dim); } - dim3 get_next_thread_id_3d() const { return m_next_tid; } - unsigned get_next_thread_id() const - { - return m_next_tid.x + m_block_dim.x*m_next_tid.y + m_block_dim.x*m_block_dim.y*m_next_tid.z; - } - bool more_threads_in_cta() const - { - return m_next_tid.z < m_block_dim.z && m_next_tid.y < m_block_dim.y && m_next_tid.z < m_block_dim.x; - } - -private: - bool m_valid; - class function_info *m_kernel_entry; - - dim3 m_grid_dim; - dim3 m_block_dim; - dim3 m_next_cta; - dim3 m_next_tid; -}; - enum divergence_support_t { POST_DOMINATOR = 1, MIMD = 2, diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index fe5d112..4f5cffe 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -156,12 +156,11 @@ void print_simulation_time() fflush(stdout); } -int gpgpu_cuda_ptx_sim_main_perf( const char *kernel_key, +int gpgpu_cuda_ptx_sim_main_perf( kernel_info_t grid, struct dim3 gridDim, struct dim3 blockDim, gpgpu_ptx_sim_arg_list_t grid_params ) { - kernel_info_t grid = gpgpu_cuda_ptx_sim_init_grid(kernel_key,grid_params,gridDim,blockDim); g_the_gpu.launch(grid); sem_post(&g_sim_signal_start); sem_wait(&g_sim_signal_finish); diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h index 22b82f2..f77a906 100644 --- a/src/gpgpusim_entrypoint.h +++ b/src/gpgpusim_entrypoint.h @@ -72,7 +72,7 @@ extern time_t g_simulation_starttime; class gpgpu_sim *gpgpu_ptx_sim_init_perf(); -int gpgpu_cuda_ptx_sim_main_perf( const char *kernel_key, +int gpgpu_cuda_ptx_sim_main_perf( kernel_info_t grid, struct dim3 gridDim, struct dim3 blockDim, gpgpu_ptx_sim_arg_list_t grid_params ); |