From 0c16df3c8b108d8720846bb44b9abcc60ddf42f9 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Wed, 15 May 2019 20:16:54 -0400 Subject: make gpu_tot_cycle local variable not global variable --- src/abstract_hardware_model.h | 15 ++++++++++++--- 1 file changed, 12 insertions(+), 3 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 22ef509..1735c2f 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -399,7 +399,7 @@ typedef std::vector addr_vector_t; class simt_stack { public: - simt_stack( unsigned wid, unsigned warpSize); + simt_stack( unsigned wid, unsigned warpSize, class gpgpu_sim * gpu); void reset(); void launch( address_type start_pc, const simt_mask_t &active_mask ); @@ -416,6 +416,7 @@ protected: unsigned m_warp_id; unsigned m_warp_size; + enum stack_entry_type { STACK_ENTRY_TYPE_NORMAL = 0, STACK_ENTRY_TYPE_CALL @@ -433,6 +434,8 @@ protected: }; std::deque m_stack; + + class gpgpu_sim * m_gpu; }; #define GLOBAL_HEAP_START 0xC0000000 @@ -571,6 +574,12 @@ public: int resume_CTA; int checkpoint_CTA_t; int checkpoint_insn_Y; + + //Move some cycle core stats here instead of being global + unsigned long long gpu_sim_cycle; + unsigned long long gpu_tot_sim_cycle; + + void* gpu_malloc( size_t size ); void* gpu_mallocarray( size_t count ); void gpu_memset( size_t dst_start_addr, int c, size_t count ); @@ -835,8 +844,8 @@ public: class mem_fetch_allocator { public: - virtual mem_fetch *alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr ) const = 0; - virtual mem_fetch *alloc( const class warp_inst_t &inst, const mem_access_t &access ) const = 0; + virtual mem_fetch *alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const = 0; + virtual mem_fetch *alloc( const class warp_inst_t &inst, const mem_access_t &access, unsigned long long cycle ) const = 0; }; // the maximum number of destination, source, or address uarch operands in a instruction -- cgit v1.3 From ed9f0e6b2a99840e9649551825a40a04e236dcd9 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Thu, 30 May 2019 18:28:15 -0400 Subject: adding new values to gpu context --- libcuda/cuda_runtime_api.cc | 42 +++++++++++++++++++------------------ src/abstract_hardware_model.cc | 9 ++++---- src/abstract_hardware_model.h | 2 +- src/cuda-sim/cuda-sim.cc | 36 +++++++++++++++---------------- src/cuda-sim/cuda_device_runtime.cc | 5 +++-- src/gpgpusim_entrypoint.cc | 12 +++++++++-- src/gpgpusim_entrypoint.h | 24 ++++++++++++++++----- 7 files changed, 78 insertions(+), 52 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index fcd5b07..4990edf 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -185,7 +185,7 @@ struct cudaArray cudaError_t g_last_cudaError = cudaSuccess; -extern stream_manager *g_stream_manager; +//extern stream_manager *g_stream_manager(); void register_ptx_function( const char *name, function_info *impl ) { @@ -375,7 +375,8 @@ private: struct _cuda_device_id *GPGPUSim_Init() { - static _cuda_device_id *the_device = NULL; + //static _cuda_device_id *the_device = NULL; + _cuda_device_id *the_device = GPGPUsim_ctx_ptr()->the_cude_device; if( !the_device ) { gpgpu_sim *the_gpu = gpgpu_ptx_sim_init_perf(); @@ -428,7 +429,8 @@ struct _cuda_device_id *GPGPUSim_Init() static CUctx_st* GPGPUSim_Context() { - static CUctx_st *the_context = NULL; + //static CUctx_st *the_context = NULL; + CUctx_st *the_context = GPGPUsim_ctx_ptr()->the_context; if( the_context == NULL ) { _cuda_device_id *the_gpu = GPGPUSim_Init(); the_context = new CUctx_st(the_gpu); @@ -699,21 +701,21 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t cou if(g_debug_execution >= 3) printf("GPGPU-Sim PTX: cudaMemcpy(): devPtr = %p\n", dst); if( kind == cudaMemcpyHostToDevice ) - g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) ); + g_stream_manager()->push( stream_operation(src,(size_t)dst,count,0) ); else if( kind == cudaMemcpyDeviceToHost ) - g_stream_manager->push( stream_operation((size_t)src,dst,count,0) ); + g_stream_manager()->push( stream_operation((size_t)src,dst,count,0) ); else if( kind == cudaMemcpyDeviceToDevice ) - g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) ); + g_stream_manager()->push( stream_operation((size_t)src,(size_t)dst,count,0) ); else if ( kind == cudaMemcpyDefault ) { if ((size_t)src >= GLOBAL_HEAP_START) { if ((size_t)dst >= GLOBAL_HEAP_START) - g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) ); // device to device + g_stream_manager()->push( stream_operation((size_t)src,(size_t)dst,count,0) ); // device to device else - g_stream_manager->push( stream_operation((size_t)src,dst,count,0) ); // device to host + g_stream_manager()->push( stream_operation((size_t)src,dst,count,0) ); // device to host } else { if ((size_t)dst >= GLOBAL_HEAP_START) - g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) ); + g_stream_manager()->push( stream_operation(src,(size_t)dst,count,0) ); else { printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported transfer: host to host\n"); abort(); @@ -855,7 +857,7 @@ __host__ cudaError_t CUDARTAPI cudaMemcpyToSymbol(const char *symbol, const void assert(kind == cudaMemcpyHostToDevice); printf("GPGPU-Sim PTX: cudaMemcpyToSymbol: symbol = %p\n", symbol); //stream_operation( const char *symbol, const void *src, size_t count, size_t offset ) - g_stream_manager->push( stream_operation(src,symbol,count,offset,0) ); + g_stream_manager()->push( stream_operation(src,symbol,count,offset,0) ); //gpgpu_ptx_sim_memcpy_symbol(symbol,src,count,offset,1,context->get_device()->get_gpgpu()); return g_last_cudaError = cudaSuccess; } @@ -869,7 +871,7 @@ __host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbol(void *dst, const char *symbo //CUctx_st *context = GPGPUSim_Context(); assert(kind == cudaMemcpyDeviceToHost); printf("GPGPU-Sim PTX: cudaMemcpyFromSymbol: symbol = %p\n", symbol); - g_stream_manager->push( stream_operation(symbol,dst,count,offset,0) ); + g_stream_manager()->push( stream_operation(symbol,dst,count,offset,0) ); //gpgpu_ptx_sim_memcpy_symbol(symbol,dst,count,offset,0,context->get_device()->get_gpgpu()); return g_last_cudaError = cudaSuccess; } @@ -898,9 +900,9 @@ __host__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_ } struct CUstream_st *s = (struct CUstream_st *)stream; switch( kind ) { - case cudaMemcpyHostToDevice: g_stream_manager->push( stream_operation(src,(size_t)dst,count,s) ); break; - case cudaMemcpyDeviceToHost: g_stream_manager->push( stream_operation((size_t)src,dst,count,s) ); break; - case cudaMemcpyDeviceToDevice: g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,s) ); break; + case cudaMemcpyHostToDevice: g_stream_manager()->push( stream_operation(src,(size_t)dst,count,s) ); break; + case cudaMemcpyDeviceToHost: g_stream_manager()->push( stream_operation((size_t)src,dst,count,s) ); break; + case cudaMemcpyDeviceToDevice: g_stream_manager()->push( stream_operation((size_t)src,(size_t)dst,count,s) ); break; default: abort(); } @@ -1611,7 +1613,7 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) 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); - g_stream_manager->push(op); + g_stream_manager()->push(op); context->g_cuda_launch_stack.pop_back(); return g_last_cudaError = cudaSuccess; } @@ -1650,7 +1652,7 @@ __host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream) printf("GPGPU-Sim PTX: cudaStreamCreate\n"); #if (CUDART_VERSION >= 3000) *stream = new struct CUstream_st(); - g_stream_manager->add_stream(*stream); + g_stream_manager()->add_stream(*stream); #else *stream = 0; printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); @@ -1689,7 +1691,7 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) //per-stream synchronization required for application using external libraries without explicit synchronization in the code to //avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. stream->synchronize(); - g_stream_manager->destroy_stream(stream); + g_stream_manager()->destroy_stream(stream); #endif return g_last_cudaError = cudaSuccess; } @@ -1769,7 +1771,7 @@ __host__ cudaError_t CUDARTAPI cudaEventRecord(cudaEvent_t event, cudaStream_t s if( !e ) return g_last_cudaError = cudaErrorUnknown; struct CUstream_st *s = (struct CUstream_st *)stream; stream_operation op(e,s); - g_stream_manager->push(op); + g_stream_manager()->push(op); return g_last_cudaError = cudaSuccess; } @@ -1785,11 +1787,11 @@ __host__ cudaError_t CUDARTAPI cudaStreamWaitEvent(cudaStream_t stream, cudaEven return g_last_cudaError = cudaSuccess; } if (!stream){ - g_stream_manager->pushCudaStreamWaitEventToAllStreams(e, flags); + g_stream_manager()->pushCudaStreamWaitEventToAllStreams(e, flags); } else { struct CUstream_st *s = (struct CUstream_st *)stream; stream_operation op(s,e,flags); - g_stream_manager->push(op); + g_stream_manager()->push(op); } return g_last_cudaError = cudaSuccess; } diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 63b139e..7755477 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -34,6 +34,7 @@ #include "cuda-sim/cuda-sim.h" #include "gpgpu-sim/gpu-sim.h" #include "option_parser.h" +#include "gpgpusim_entrypoint.h" #include #include #include @@ -771,14 +772,14 @@ void kernel_info_t::notify_parent_finished() { extern unsigned long long g_total_param_size; g_total_param_size -= ((m_kernel_entry->get_args_aligned_size() + 255)/256*256); m_parent_kernel->remove_child(this); - g_stream_manager->register_finished_kernel(m_parent_kernel->get_uid()); + g_stream_manager()->register_finished_kernel(m_parent_kernel->get_uid()); } } CUstream_st * kernel_info_t::create_stream_cta(dim3 ctaid) { assert(get_default_stream_cta(ctaid)); CUstream_st * stream = new CUstream_st(); - g_stream_manager->add_stream(stream); + g_stream_manager()->add_stream(stream); assert(m_cta_streams.find(ctaid) != m_cta_streams.end()); assert(m_cta_streams[ctaid].size() >= 1); //must have default stream m_cta_streams[ctaid].push_back(stream); @@ -794,7 +795,7 @@ CUstream_st * kernel_info_t::get_default_stream_cta(dim3 ctaid) { else { m_cta_streams[ctaid] = std::list(); CUstream_st * stream = new CUstream_st(); - g_stream_manager->add_stream(stream); + g_stream_manager()->add_stream(stream); m_cta_streams[ctaid].push_back(stream); return stream; } @@ -826,7 +827,7 @@ void kernel_info_t::destroy_cta_streams() { for(auto s = m_cta_streams.begin(); s != m_cta_streams.end(); s++) { stream_size += s->second.size(); for(auto ss = s->second.begin(); ss != s->second.end(); ss++) - g_stream_manager->destroy_stream(*ss); + g_stream_manager()->destroy_stream(*ss); s->second.clear(); } printf("size %lu\n", stream_size); diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 1735c2f..77d5f58 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -197,7 +197,7 @@ void increment_x_then_y_then_z( dim3 &i, const dim3 &bound); #include "stream_manager.h" class stream_manager; struct CUstream_st; -extern stream_manager * g_stream_manager; +//extern stream_manager * g_stream_manager; //support for pinned memories added extern std::map pinned_memory; extern std::map pinned_memory_size; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index e733b7f..a456978 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -448,8 +448,8 @@ void gpgpu_t::memcpy_to_gpu( size_t dst_start_addr, const void *src, size_t coun m_global_mem->write(dst_start_addr+n,1, src_data+n,NULL,NULL); // Copy into the performance model. - extern gpgpu_sim* g_the_gpu; - g_the_gpu->perf_memcpy_to_gpu(dst_start_addr, count); + //extern gpgpu_sim* g_the_gpu; + g_the_gpu()->perf_memcpy_to_gpu(dst_start_addr, count); if(g_debug_execution >= 3) { printf( " done.\n"); fflush(stdout); @@ -467,8 +467,8 @@ void gpgpu_t::memcpy_from_gpu( void *dst, size_t src_start_addr, size_t count ) m_global_mem->read(src_start_addr+n,1,dst_data+n); // Copy into the performance model. - extern gpgpu_sim* g_the_gpu; - g_the_gpu->perf_memcpy_to_gpu(src_start_addr, count); + //extern gpgpu_sim* g_the_gpu; + g_the_gpu()->perf_memcpy_to_gpu(src_start_addr, count); if(g_debug_execution >= 3) { printf( " done.\n"); fflush(stdout); @@ -1270,8 +1270,8 @@ void function_info::finalize( memory_space *param_mem ) void function_info::param_to_shared( memory_space *shared_mem, symbol_table *symtab ) { // TODO: call this only for PTXPlus with GT200 models - extern gpgpu_sim* g_the_gpu; - if (not g_the_gpu->get_config().convert_to_ptxplus()) return; + //extern gpgpu_sim* g_the_gpu; + if (not g_the_gpu()->get_config().convert_to_ptxplus()) return; // copies parameters into simulated shared memory for( std::map::iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) { @@ -2150,7 +2150,7 @@ void 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()); //using a shader core object for book keeping, it is not needed but as most function built for performance simulation need it we use it here - extern gpgpu_sim *g_the_gpu; + //extern gpgpu_sim *g_the_gpu; //before we execute, we should do PDOM analysis for functional simulation scenario. function_info *kernel_func_info = kernel.entry(); const struct gpgpu_ptx_sim_info *kernel_info = ptx_sim_kernel_info(kernel_func_info); @@ -2165,7 +2165,7 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) kernel_func_info->set_pdom(); } - unsigned max_cta_tot = max_cta(kernel_info,kernel.threads_per_cta(), g_the_gpu->getShaderCoreConfig()->warp_size, g_the_gpu->getShaderCoreConfig()->n_thread_per_shader, g_the_gpu->getShaderCoreConfig()->gpgpu_shmem_size, g_the_gpu->getShaderCoreConfig()->gpgpu_shader_registers, g_the_gpu->getShaderCoreConfig()->max_cta_per_core); + unsigned max_cta_tot = max_cta(kernel_info,kernel.threads_per_cta(), g_the_gpu()->getShaderCoreConfig()->warp_size, g_the_gpu()->getShaderCoreConfig()->n_thread_per_shader, g_the_gpu()->getShaderCoreConfig()->gpgpu_shmem_size, g_the_gpu()->getShaderCoreConfig()->gpgpu_shader_registers, g_the_gpu()->getShaderCoreConfig()->max_cta_per_core); printf("Max CTA : %d\n",max_cta_tot); @@ -2173,11 +2173,11 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) int inst_count=50; - int cp_op= g_the_gpu->checkpoint_option; - int cp_CTA = g_the_gpu->checkpoint_CTA; - int cp_kernel= g_the_gpu->checkpoint_kernel; - cp_count= g_the_gpu->checkpoint_insn_Y; - cp_cta_resume= g_the_gpu->checkpoint_CTA_t; + int cp_op= g_the_gpu()->checkpoint_option; + int cp_CTA = g_the_gpu()->checkpoint_CTA; + int cp_kernel= g_the_gpu()->checkpoint_kernel; + cp_count= g_the_gpu()->checkpoint_insn_Y; + cp_cta_resume= g_the_gpu()->checkpoint_CTA_t; int cta_launched =0; //we excute the kernel one CTA (Block) at the time, as synchronization functions work block wise @@ -2189,8 +2189,8 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) { functionalCoreSim cta( &kernel, - g_the_gpu, - g_the_gpu->getShaderCoreConfig()->warp_size + g_the_gpu(), + g_the_gpu()->getShaderCoreConfig()->warp_size ); cta.execute(cp_count,temp); @@ -2211,7 +2211,7 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) { char f1name[2048]; snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", kernel.get_uid() ); - g_checkpoint->store_global_mem(g_the_gpu->get_global_memory(), f1name , "%08x"); + g_checkpoint->store_global_mem(g_the_gpu()->get_global_memory(), f1name , "%08x"); } @@ -2221,8 +2221,8 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) //openCL kernel simulation calls don't register the kernel so we don't register its exit if(!openCL) { - extern stream_manager *g_stream_manager; - g_stream_manager->register_finished_kernel(kernel.get_uid()); + //extern stream_manager *g_stream_manager; + g_stream_manager()->register_finished_kernel(kernel.get_uid()); } //******PRINTING******* diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 86e8147..be8369f 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -18,6 +18,7 @@ unsigned long long g_max_total_param_size = 0; #include "cuda-sim.h" #include "ptx_ir.h" #include "../stream_manager.h" +#include "../gpgpusim_entrypoint.h" #include "cuda_device_runtime.h" #define DEV_RUNTIME_REPORT(a) \ @@ -64,7 +65,7 @@ public: std::map g_cuda_device_launch_param_map; std::list g_cuda_device_launch_op; -extern stream_manager *g_stream_manager; +//extern stream_manager *g_stream_manager(); //Handling device runtime api: //void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize) @@ -322,7 +323,7 @@ void launch_one_device_kernel() { 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); - g_stream_manager->push(stream_op); + g_stream_manager()->push(stream_op); g_cuda_device_launch_op.pop_front(); } } diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index c8770e2..de937b0 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -41,8 +41,6 @@ struct GPGPUsim_ctx* the_gpgpusim = NULL; -static void print_simulation_time(); - struct GPGPUsim_ctx* GPGPUsim_ctx_ptr(){ if(the_gpgpusim == NULL) the_gpgpusim = new GPGPUsim_ctx(); @@ -50,6 +48,16 @@ struct GPGPUsim_ctx* GPGPUsim_ctx_ptr(){ return the_gpgpusim; } +class gpgpu_sim* g_the_gpu() { + return GPGPUsim_ctx_ptr()->g_the_gpu; +} + +class stream_manager* g_stream_manager() { + return GPGPUsim_ctx_ptr()->g_stream_manager; +} + +static void print_simulation_time(); + void *gpgpu_sim_thread_sequential(void*) { // at most one kernel running at a time diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h index e29159b..2ad0fdf 100644 --- a/src/gpgpusim_entrypoint.h +++ b/src/gpgpusim_entrypoint.h @@ -46,9 +46,15 @@ struct GPGPUsim_ctx { sg_argc = 3; sg_argv = {"", "-config","gpgpusim.config"}; + + g_the_gpu_config=NULL; + g_the_gpu=NULL; + g_stream_manager=NULL; + the_cude_device=NULL; + the_context=NULL; } - struct gpgpu_ptx_sim_arg *grid_params; + //struct gpgpu_ptx_sim_arg *grid_params; sem_t g_sim_signal_start; sem_t g_sim_signal_finish; @@ -60,19 +66,27 @@ struct GPGPUsim_ctx { class gpgpu_sim *g_the_gpu; class stream_manager *g_stream_manager; + struct _cuda_device_id *the_cude_device; + struct CUctx_st* the_context; + + int sg_argc; const char *sg_argv[3]; - pthread_mutex_t g_sim_lock; - bool g_sim_active; - bool g_sim_done; - bool break_limit; + pthread_mutex_t g_sim_lock; + bool g_sim_active; + bool g_sim_done; + bool break_limit; }; class gpgpu_sim *gpgpu_ptx_sim_init_perf(); void start_sim_thread(int api); +class gpgpu_sim* g_the_gpu(); +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 ); -- cgit v1.3 From c68377b8eb2ce322e195a23062f5ba635dba546a Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 10 Jun 2019 19:43:36 -0400 Subject: Remove those duplicated defined Signed-off-by: Mengchi Zhang --- libcuda/cuda_api_object.h | 2 ++ src/Makefile | 2 ++ src/abstract_hardware_model.h | 75 ++++++++++++++++++++++--------------------- src/cuda-sim/cuda-math.h | 16 ++++----- src/cuda-sim/instructions.cc | 10 +++--- src/gpgpu-sim/Makefile | 2 ++ src/intersim2/Makefile | 1 + 7 files changed, 59 insertions(+), 49 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/libcuda/cuda_api_object.h b/libcuda/cuda_api_object.h index 73c077e..d931fd5 100644 --- a/libcuda/cuda_api_object.h +++ b/libcuda/cuda_api_object.h @@ -6,6 +6,8 @@ #include #include +#include "builtin_types.h" + #include "../src/gpgpu-sim/gpu-sim.h" #include "../src/cuda-sim/ptx_ir.h" #include "../src/abstract_hardware_model.h" diff --git a/src/Makefile b/src/Makefile index 6001669..3ad511e 100644 --- a/src/Makefile +++ b/src/Makefile @@ -51,6 +51,8 @@ else CXXFLAGS += endif +CXXFLAGS += -I$(CUDA_INSTALL_PATH)/include + OPTFLAGS += -g3 -fPIC CPP = g++ $(SNOW) diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 77d5f58..7455f25 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -173,9 +173,10 @@ enum _memory_op_t { #include #if !defined(__VECTOR_TYPES_H__) -struct dim3 { - unsigned int x, y, z; -}; +#include "vector_types.h" +//struct dim3 { +// unsigned int x, y, z; +//}; #endif struct dim3comp { bool operator() (const dim3 & a, const dim3 & b) const @@ -454,19 +455,21 @@ protected: #if !defined(__CUDA_RUNTIME_API_H__) -enum cudaChannelFormatKind { - cudaChannelFormatKindSigned, - cudaChannelFormatKindUnsigned, - cudaChannelFormatKindFloat -}; +#include "builtin_types.h" -struct cudaChannelFormatDesc { - int x; - int y; - int z; - int w; - enum cudaChannelFormatKind f; -}; +//enum cudaChannelFormatKind { +// cudaChannelFormatKindSigned, +// cudaChannelFormatKindUnsigned, +// cudaChannelFormatKindFloat +//}; + +//struct cudaChannelFormatDesc { +// int x; +// int y; +// int z; +// int w; +// enum cudaChannelFormatKind f; +//}; struct cudaArray { void *devPtr; @@ -478,27 +481,27 @@ struct cudaArray { unsigned dimensions; }; -enum cudaTextureAddressMode { - cudaAddressModeWrap, - cudaAddressModeClamp -}; - -enum cudaTextureFilterMode { - cudaFilterModePoint, - cudaFilterModeLinear -}; - -enum cudaTextureReadMode { - cudaReadModeElementType, - cudaReadModeNormalizedFloat -}; - -struct textureReference { - int normalized; - enum cudaTextureFilterMode filterMode; - enum cudaTextureAddressMode addressMode[3]; - struct cudaChannelFormatDesc channelDesc; -}; +//enum cudaTextureAddressMode { +// cudaAddressModeWrap, +// cudaAddressModeClamp +//}; + +//enum cudaTextureFilterMode { +// cudaFilterModePoint, +// cudaFilterModeLinear +//}; + +//enum cudaTextureReadMode { +// cudaReadModeElementType, +// cudaReadModeNormalizedFloat +//}; + +//struct textureReference { +// int normalized; +// enum cudaTextureFilterMode filterMode; +// enum cudaTextureAddressMode addressMode[3]; +// struct cudaChannelFormatDesc channelDesc; +//}; #endif diff --git a/src/cuda-sim/cuda-math.h b/src/cuda-sim/cuda-math.h index a5db337..9a5468c 100644 --- a/src/cuda-sim/cuda-math.h +++ b/src/cuda-sim/cuda-math.h @@ -277,10 +277,10 @@ int float2int(float a, enum cudaRoundMode mode) { int tmp; switch (mode) { - case cuda_math::cudaRoundZero: tmp = truncf(a); break; - case cuda_math::cudaRoundNearest: tmp = nearbyintf(a); break; - case cuda_math::cudaRoundMinInf: tmp = floorf(a); break; - case cuda_math::cudaRoundPosInf: tmp = ceilf(a); break; + case cudaRoundZero: tmp = truncf(a); break; + case cudaRoundNearest: tmp = nearbyintf(a); break; + case cudaRoundMinInf: tmp = floorf(a); break; + case cudaRoundPosInf: tmp = ceilf(a); break; default: abort(); } return tmp; @@ -296,10 +296,10 @@ unsigned int float2uint(float a, enum cudaRoundMode mode) { unsigned int tmp; switch (mode) { - case cuda_math::cudaRoundZero: tmp = truncf(a); break; - case cuda_math::cudaRoundNearest: tmp = nearbyintf(a); break; - case cuda_math::cudaRoundMinInf: tmp = floorf(a); break; - case cuda_math::cudaRoundPosInf: tmp = ceilf(a); break; + case cudaRoundZero: tmp = truncf(a); break; + case cudaRoundNearest: tmp = nearbyintf(a); break; + case cudaRoundMinInf: tmp = floorf(a); break; + case cudaRoundPosInf: tmp = ceilf(a); break; default: abort(); } return tmp; diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 69a97b6..186bb1e 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -2302,12 +2302,12 @@ ptx_reg_t f2x( ptx_reg_t x, unsigned from_width, unsigned to_width, int to_sign, half_float::half tmp_h; //assert( from_width == 32); - enum cuda_math::cudaRoundMode mode = cuda_math::cudaRoundZero; + enum cudaRoundMode mode = cudaRoundZero; switch (rounding_mode) { - case RZI_OPTION: mode = cuda_math::cudaRoundZero; break; - case RNI_OPTION: mode = cuda_math::cudaRoundNearest; break; - case RMI_OPTION: mode = cuda_math::cudaRoundMinInf; break; - case RPI_OPTION: mode = cuda_math::cudaRoundPosInf; break; + case RZI_OPTION: mode = cudaRoundZero; break; + case RNI_OPTION: mode = cudaRoundNearest; break; + case RMI_OPTION: mode = cudaRoundMinInf; break; + case RPI_OPTION: mode = cudaRoundPosInf; break; default: break; } diff --git a/src/gpgpu-sim/Makefile b/src/gpgpu-sim/Makefile index f10a8a4..4994577 100644 --- a/src/gpgpu-sim/Makefile +++ b/src/gpgpu-sim/Makefile @@ -53,6 +53,8 @@ else CXXFLAGS += endif +CXXFLAGS += -I$(CUDA_INSTALL_PATH)/include + POWER_FLAGS= ifneq ($(GPGPUSIM_POWER_MODEL),) POWER_FLAGS = -I$(GPGPUSIM_POWER_MODEL) -DGPGPUSIM_POWER_MODEL diff --git a/src/intersim2/Makefile b/src/intersim2/Makefile index 7d10b3f..3eeeb70 100644 --- a/src/intersim2/Makefile +++ b/src/intersim2/Makefile @@ -48,6 +48,7 @@ CPPFLAGS += -O3 endif CPPFLAGS += -g CPPFLAGS += -fPIC +CPPFLAGS += -I$(CUDA_INSTALL_PATH)/include LFLAGS += -- cgit v1.3 From a34d4295ca856297ad76d5a6408cdaa163554809 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 10 Jun 2019 20:58:22 -0400 Subject: Remove definitions Signed-off-by: Mengchi Zhang --- src/abstract_hardware_model.h | 39 --------------------------------------- 1 file changed, 39 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 7455f25..68cb693 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -174,9 +174,6 @@ enum _memory_op_t { #if !defined(__VECTOR_TYPES_H__) #include "vector_types.h" -//struct dim3 { -// unsigned int x, y, z; -//}; #endif struct dim3comp { bool operator() (const dim3 & a, const dim3 & b) const @@ -457,20 +454,6 @@ protected: #include "builtin_types.h" -//enum cudaChannelFormatKind { -// cudaChannelFormatKindSigned, -// cudaChannelFormatKindUnsigned, -// cudaChannelFormatKindFloat -//}; - -//struct cudaChannelFormatDesc { -// int x; -// int y; -// int z; -// int w; -// enum cudaChannelFormatKind f; -//}; - struct cudaArray { void *devPtr; int devPtr32; @@ -481,28 +464,6 @@ struct cudaArray { unsigned dimensions; }; -//enum cudaTextureAddressMode { -// cudaAddressModeWrap, -// cudaAddressModeClamp -//}; - -//enum cudaTextureFilterMode { -// cudaFilterModePoint, -// cudaFilterModeLinear -//}; - -//enum cudaTextureReadMode { -// cudaReadModeElementType, -// cudaReadModeNormalizedFloat -//}; - -//struct textureReference { -// int normalized; -// enum cudaTextureFilterMode filterMode; -// enum cudaTextureAddressMode addressMode[3]; -// struct cudaChannelFormatDesc channelDesc; -//}; - #endif // Struct that record other attributes in the textureReference declaration -- cgit v1.3 From 8f6668941cf2728dba9700e45f11f61401a1fcf4 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 8 Jul 2019 15:45:05 -0400 Subject: Move g_ptx_sim_num_insn Signed-off-by: Mengchi Zhang --- src/abstract_hardware_model.cc | 3 ++- src/abstract_hardware_model.h | 6 +++++- src/cuda-sim/cuda-sim.cc | 9 ++++----- src/cuda-sim/cuda-sim.h | 4 ++++ src/gpgpu-sim/gpu-sim.cc | 2 +- src/gpgpu-sim/gpu-sim.h | 2 ++ 6 files changed, 18 insertions(+), 8 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 7755477..248e7a5 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -173,9 +173,10 @@ void gpgpu_functional_sim_config::ptx_set_tex_cache_linesize(unsigned linesize) m_texcache_linesize = linesize; } -gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config ) +gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config, gpgpu_context* ctx ) : m_function_model_config(config) { + gpgpu_ctx = ctx; m_global_mem = new memory_space_impl<8192>("global",64*1024); m_tex_mem = new memory_space_impl<8192>("tex",64*1024); diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 68cb693..da29a11 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -31,6 +31,8 @@ // Forward declarations class gpgpu_sim; class kernel_info_t; +class gpgpu_context; + //Set a hard limit of 32 CTAs per shader [cuda only has 8] #define MAX_CTA_PER_SHADER 32 @@ -529,7 +531,9 @@ private: class gpgpu_t { public: - gpgpu_t( const gpgpu_functional_sim_config &config ); + gpgpu_t( const gpgpu_functional_sim_config &config, gpgpu_context* ctx ); + // backward pointer + class gpgpu_context* gpgpu_ctx; int checkpoint_option; int checkpoint_kernel; int checkpoint_CTA; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index a143aa5..939358b 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -62,7 +62,6 @@ int g_debug_thread_uid = 0; addr_t g_debug_pc = 0xBEEF1518; // Output debug information to file options -unsigned g_ptx_sim_num_insn = 0; unsigned gpgpu_param_num_shaders = 0; char *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor; @@ -1629,7 +1628,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) dim3 ctaid = get_ctaid(); dim3 tid = get_tid(); printf("%u [thd=%u][i=%u] : ctaid=(%u,%u,%u) tid=(%u,%u,%u) icount=%u [pc=%u] (%s:%u - %s) [0x%llx]\n", - g_ptx_sim_num_insn, + m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn, get_uid(), pI->uid(), ctaid.x,ctaid.y,ctaid.z,tid.x,tid.y,tid.z, get_icount(), @@ -1687,7 +1686,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) dump_regs(stdout); } update_pc(); - g_ptx_sim_num_insn++; + m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn++; //not using it with functional simulation mode if(!(this->m_functionalSimulationMode)) @@ -1714,11 +1713,11 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) 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() ); } - if ( (g_ptx_sim_num_insn % 100000) == 0 ) { + if ( (m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn % 100000) == 0 ) { dim3 ctaid = get_ctaid(); dim3 tid = get_tid(); DPRINTF(LIVENESS, "GPGPU-Sim PTX: %u instructions simulated : ctaid=(%u,%u,%u) tid=(%u,%u,%u)\n", - g_ptx_sim_num_insn, ctaid.x,ctaid.y,ctaid.z,tid.x,tid.y,tid.z ); + m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn, ctaid.x,ctaid.y,ctaid.z,tid.x,tid.y,tid.z ); fflush(stdout); } diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 76450dc..4566dc2 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -134,11 +134,15 @@ struct gpgpu_ptx_sim_info get_ptxinfo(); class cuda_sim { public: + cuda_sim() { + g_ptx_sim_num_insn = 0; + } //global variables char *opcode_latency_int; int cp_count; int cp_cta_resume; int g_ptxinfo_error_detected; + unsigned g_ptx_sim_num_insn; //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 ); diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 39acdd9..4f9ccbf 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -699,7 +699,7 @@ void gpgpu_sim::stop_all_running_kernels(){ } gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config, gpgpu_context* ctx ) - : gpgpu_t(config), m_config(config) + : gpgpu_t(config, ctx), m_config(config) { gpgpu_ctx = ctx; m_shader_config = &m_config.m_shader_config; diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index b47ab16..119b934 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -62,6 +62,8 @@ #define SAMPLELOG 222 #define DUMPLOG 333 +class gpgpu_context; + extern tr1_hash_map address_random_interleaving; enum dram_ctrl_t { -- cgit v1.3 From 87726d32ada00fcd93f2cf24ccae4ba593c4f9ec Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Wed, 10 Jul 2019 15:03:16 -0400 Subject: Move g_pc_to_finfo Signed-off-by: Mengchi Zhang --- src/abstract_hardware_model.cc | 3 ++- src/abstract_hardware_model.h | 13 ++++++++----- src/cuda-sim/cuda-sim.cc | 7 +++---- src/cuda-sim/cuda-sim.h | 5 +++-- src/cuda-sim/ptx_ir.cc | 12 +++++++----- src/cuda-sim/ptx_ir.h | 9 +++++++-- src/cuda-sim/ptx_parser.cc | 2 +- src/gpgpu-sim/shader.cc | 4 ++-- src/gpgpu-sim/shader.h | 2 +- 9 files changed, 34 insertions(+), 23 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 248e7a5..ef09051 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -39,6 +39,7 @@ #include #include #include +#include "../libcuda/gpgpu_context.h" unsigned mem_access_t::sm_next_access_uid = 0; unsigned warp_inst_t::sm_next_uid = 0; @@ -945,7 +946,7 @@ void simt_stack::print (FILE *fout) const } else { fprintf(fout," " ); } - ptx_print_insn( stack_entry.m_pc, fout ); + m_gpu->gpgpu_ctx->func_sim->ptx_print_insn( stack_entry.m_pc, fout ); fprintf(fout,"\n"); } diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index da29a11..8ef8376 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -347,9 +347,11 @@ public: mutable bool volta_cache_config_set; }; -struct core_config { - core_config() - { +class core_config { + public: + core_config(gpgpu_context* ctx) + { + gpgpu_ctx = ctx; m_valid = false; num_shmem_bank=16; shmem_limited_broadcast = false; @@ -361,6 +363,8 @@ struct core_config { bool m_valid; unsigned warp_size; + // backward pointer + class gpgpu_context* gpgpu_ctx; // off-chip memory request architecture parameters int gpgpu_coalesce_arch; @@ -934,7 +938,7 @@ public: m_empty=true; m_config=NULL; } - warp_inst_t( const core_config *config ) + warp_inst_t( const core_config *config ) { m_uid=0; assert(config->warp_size<=MAX_WARP_SIZE); @@ -1105,7 +1109,6 @@ public: unsigned get_uid() const { return m_uid; } unsigned get_schd_id() const { return m_scheduler_id; } - protected: unsigned m_uid; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 3d1da62..fb9bc9e 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -220,7 +220,6 @@ void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) } unsigned g_assemble_code_next_pc=0; -std::map g_pc_to_finfo; std::vector function_info::s_g_pc_to_insn; #define MAX_INST_SIZE 8 /*bytes*/ @@ -257,7 +256,7 @@ void function_info::ptx_assemble() const symbol *l = pI->get_label(); labels[l->name()] = n; } else { - g_pc_to_finfo[PC] = this; + gpgpu_ctx->func_sim->g_pc_to_finfo[PC] = this; m_instr_mem[n] = pI; s_g_pc_to_insn.push_back(pI); assert(pI == s_g_pc_to_insn[PC]); @@ -497,7 +496,7 @@ void gpgpu_t::gpu_memset( size_t dst_start_addr, int c, size_t count ) } } -void ptx_print_insn( address_type pc, FILE *fp ) +void cuda_sim::ptx_print_insn( address_type pc, FILE *fp ) { std::map::iterator f = g_pc_to_finfo.find(pc); if( f == g_pc_to_finfo.end() ) { @@ -509,7 +508,7 @@ void ptx_print_insn( address_type pc, FILE *fp ) finfo->print_insn(pc,fp); } -std::string ptx_get_insn_str( address_type pc ) +std::string cuda_sim::ptx_get_insn_str( address_type pc ) { std::map::iterator f = g_pc_to_finfo.find(pc); if( f == g_pc_to_finfo.end() ) { diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index 5c95100..c578524 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -61,8 +61,6 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, bool functionalSimulationMode = false); const warp_inst_t *ptx_fetch_inst( address_type pc ); const struct gpgpu_ptx_sim_info* ptx_sim_kernel_info(const class function_info *kernel); -void ptx_print_insn( address_type pc, FILE *fp ); -std::string ptx_get_insn_str( address_type pc ); /*! @@ -162,6 +160,7 @@ class cuda_sim { void ** g_inst_op_classification_stat; std::set g_globals; std::set g_constants; + std::map g_pc_to_finfo; // backward pointer class gpgpu_context* gpgpu_ctx; //global functions @@ -181,6 +180,8 @@ class cuda_sim { 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 ); + void ptx_print_insn( address_type pc, FILE *fp ); + std::string ptx_get_insn_str( address_type pc ); }; #endif diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 1bd409e..c537091 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -83,8 +83,9 @@ symbol_table::symbol_table() assert(0); } -symbol_table::symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent ) +symbol_table::symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent, gpgpu_context* ctx ) { + gpgpu_ctx = ctx; m_scope_name = std::string(scope_name); m_reg_allocator=0; m_shared_next = 0; @@ -183,7 +184,7 @@ symbol_table* symbol_table::start_inst_group() { //previous added assert(m_inst_group_symtab.find(std::string(inst_group_name)) == m_inst_group_symtab.end()); - symbol_table *sym_table = new symbol_table(inst_group_name, 3/*inst group*/, this ); + symbol_table *sym_table = new symbol_table(inst_group_name, 3/*inst group*/, this, gpgpu_ctx ); sym_table->m_global_next = m_global_next; sym_table->m_shared_next = m_shared_next; @@ -221,7 +222,7 @@ bool symbol_table::add_function_decl( const char *name, int entry_point, functio *func_info = m_function_info_lookup[key]; prior_decl = true; } else { - *func_info = new function_info(entry_point); + *func_info = new function_info(entry_point, gpgpu_ctx); (*func_info)->set_name(name); (*func_info)->set_maxnt_id(0); m_function_info_lookup[key] = *func_info; @@ -232,7 +233,7 @@ 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, this ); + *sym_table = new symbol_table( "", entry_point, this, gpgpu_ctx ); // Initial setup code to support a register represented as "_". // This register is used when an instruction operand is @@ -1373,8 +1374,9 @@ std::string ptx_instruction::to_string() const unsigned function_info::sm_next_uid = 1; -function_info::function_info(int entry_point ) +function_info::function_info(int entry_point, gpgpu_context* ctx ) { + gpgpu_ctx = ctx; m_uid = sm_next_uid++; m_entry_point = (entry_point==1)?true:false; m_extern = (entry_point==2)?true:false; diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index 1604551..babd54b 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -306,7 +306,7 @@ private: class symbol_table { public: symbol_table(); - symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent ); + symbol_table( const char *scope_name, unsigned entry_point, symbol_table *parent, gpgpu_context* ctx); void set_name( const char *name ); const ptx_version &get_ptx_version() const; unsigned get_sm_target() const; @@ -348,6 +348,9 @@ public: symbol_table* start_inst_group(); symbol_table* end_inst_group(); + // backward pointer + class gpgpu_context* gpgpu_ctx; + private: unsigned m_reg_allocator; unsigned m_shared_next; @@ -1233,7 +1236,7 @@ private: class function_info { public: - function_info(int entry_point ); + function_info(int entry_point, gpgpu_context* ctx ); const ptx_version &get_ptx_version() const { return m_symtab->get_ptx_version(); } unsigned get_sm_target() const { return m_symtab->get_sm_target(); } bool is_extern() const { return m_extern; } @@ -1403,6 +1406,8 @@ public: void set_maxnt_id(unsigned maxthreads) { maxnt_id = maxthreads;} unsigned get_maxnt_id() { return maxnt_id;} + // backward pointer + class gpgpu_context* gpgpu_ctx; private: unsigned maxnt_id; diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc index 269ec4d..5a94679 100644 --- a/src/cuda-sim/ptx_parser.cc +++ b/src/cuda-sim/ptx_parser.cc @@ -112,7 +112,7 @@ symbol_table * gpgpu_context::init_parser( const char *ptx_filename ) { g_filename = strdup(ptx_filename); if (g_global_allfiles_symbol_table == NULL) { - g_global_allfiles_symbol_table = new symbol_table("global_allfiles", 0, NULL); + g_global_allfiles_symbol_table = new symbol_table("global_allfiles", 0, NULL, this); ptx_parser->g_global_symbol_table = ptx_parser->g_current_symbol_table = g_global_allfiles_symbol_table; } /*else { diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 69790fc..6cd6d8f 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -1025,7 +1025,7 @@ void scheduler_unit::cycle() m_simt_stack[warp_id]->get_pdom_stack_top_info(&pc,&rpc); SCHED_DPRINTF( "Warp (warp_id %u, dynamic_warp_id %u) has valid instruction (%s)\n", (*iter)->get_warp_id(), (*iter)->get_dynamic_warp_id(), - ptx_get_insn_str( pc).c_str() ); + m_shader->m_config->gpgpu_ctx->func_sim->ptx_get_insn_str( pc).c_str() ); if( pI ) { assert(valid); if( pc != pI->pc ) { @@ -2690,7 +2690,7 @@ void warp_inst_t::print( FILE *fout ) const for (unsigned j=0; jwarp_size; j++) fprintf(fout, "%c", (active(j)?'1':'0') ); fprintf(fout, "]: "); - ptx_print_insn( pc, fout ); + m_config->gpgpu_ctx->func_sim->ptx_print_insn( pc, fout ); fprintf(fout, "\n"); } void shader_core_ctx::incexecstat(warp_inst_t *&inst) diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index e0cefac..2837f1b 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1367,7 +1367,7 @@ const char* const pipeline_stage_name_decode[] = { class shader_core_config : public core_config { public: - shader_core_config(gpgpu_context* ctx){ + shader_core_config(gpgpu_context* ctx):core_config(ctx){ pipeline_widths_string = NULL; gpgpu_ctx = ctx; } -- cgit v1.3 From 23c0bb224295dde9651fd915d854e4f7eafdf88f Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Sun, 14 Jul 2019 11:25:42 -0400 Subject: Move sm_next_access_uid Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 2 ++ src/abstract_hardware_model.cc | 13 ++++++++++--- src/abstract_hardware_model.h | 22 +++++++++------------- src/gpgpu-sim/dram.cc | 2 +- src/gpgpu-sim/dram.h | 5 +++-- src/gpgpu-sim/gpu-cache.cc | 6 ++++-- src/gpgpu-sim/gpu-sim.cc | 2 +- src/gpgpu-sim/gpu-sim.h | 15 +++++++++------ src/gpgpu-sim/l2cache.cc | 11 ++++++----- src/gpgpu-sim/l2cache.h | 12 ++++++------ src/gpgpu-sim/mem_fetch.cc | 4 ++-- src/gpgpu-sim/mem_fetch.h | 5 +++-- src/gpgpu-sim/mem_latency_stat.cc | 2 +- src/gpgpu-sim/mem_latency_stat.h | 5 +++-- src/gpgpu-sim/power_stat.cc | 4 ++-- src/gpgpu-sim/power_stat.h | 6 +++--- src/gpgpu-sim/shader.cc | 20 ++++++++++++++++---- src/gpgpu-sim/shader.h | 20 ++++---------------- 18 files changed, 85 insertions(+), 71 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index c0b250a..07473be 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -11,6 +11,7 @@ class gpgpu_context { public: gpgpu_context() { g_global_allfiles_symbol_table = NULL; + sm_next_access_uid=0; api = new cuda_runtime_api(this); ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); @@ -21,6 +22,7 @@ class gpgpu_context { // global list symbol_table *g_global_allfiles_symbol_table; const char *g_filename; + unsigned sm_next_access_uid; // objects pointers for each file cuda_runtime_api* api; ptxinfo_data* ptxinfo; diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index fde7874..fe10daa 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -41,9 +41,16 @@ #include #include "../libcuda/gpgpu_context.h" -unsigned mem_access_t::sm_next_access_uid = 0; unsigned warp_inst_t::sm_next_uid = 0; +void mem_access_t::init(gpgpu_context* ctx) +{ + gpgpu_ctx = ctx; + m_uid=++(gpgpu_ctx->sm_next_access_uid); + m_addr=0; + m_req_size=0; +} + checkpoint::checkpoint() { @@ -449,7 +456,7 @@ void warp_inst_t::generate_mem_accesses() byte_mask.set(idx+i); } for( a=accesses.begin(); a != accesses.end(); ++a ) - m_accessq.push_back( mem_access_t(access_type,a->first,cache_block_size,is_write,a->second, byte_mask, mem_access_sector_mask_t())); + m_accessq.push_back( mem_access_t(access_type,a->first,cache_block_size,is_write,a->second, byte_mask, mem_access_sector_mask_t(), m_config->gpgpu_ctx)); } if ( space.get_type() == global_space ) { @@ -681,7 +688,7 @@ void warp_inst_t::memory_coalescing_arch_reduce_and_send( bool is_write, mem_acc assert(lower_half_used && upper_half_used); } } - m_accessq.push_back( mem_access_t(access_type,addr,size,is_write,info.active,info.bytes, info.chunks) ); + m_accessq.push_back( mem_access_t(access_type,addr,size,is_write,info.active,info.bytes, info.chunks,m_config->gpgpu_ctx) ); } void warp_inst_t::completed( unsigned long long cycle ) const diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 8ef8376..a22c8c3 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -732,13 +732,14 @@ enum cache_operator_type { class mem_access_t { public: - mem_access_t() { init(); } + mem_access_t(gpgpu_context* ctx) { init(ctx); } mem_access_t( mem_access_type type, new_addr_type address, unsigned size, - bool wr ) + bool wr, + gpgpu_context* ctx) { - init(); + init(ctx); m_type = type; m_addr = address; m_req_size = size; @@ -750,10 +751,11 @@ public: bool wr, const active_mask_t &active_mask, const mem_access_byte_mask_t &byte_mask, - const mem_access_sector_mask_t §or_mask) + const mem_access_sector_mask_t §or_mask, + gpgpu_context* ctx) : m_warp_mask(active_mask), m_byte_mask(byte_mask), m_sector_mask(sector_mask) { - init(); + init(ctx); m_type = type; m_addr = address; m_req_size = size; @@ -786,13 +788,9 @@ public: } } + gpgpu_context* gpgpu_ctx; private: - void init() - { - m_uid=++sm_next_access_uid; - m_addr=0; - m_req_size=0; - } + void init(gpgpu_context* ctx); unsigned m_uid; new_addr_type m_addr; // request address @@ -802,8 +800,6 @@ private: active_mask_t m_warp_mask; mem_access_byte_mask_t m_byte_mask; mem_access_sector_mask_t m_sector_mask; - - static unsigned sm_next_access_uid; }; class mem_fetch; diff --git a/src/gpgpu-sim/dram.cc b/src/gpgpu-sim/dram.cc index 5e36d4b..d443d79 100644 --- a/src/gpgpu-sim/dram.cc +++ b/src/gpgpu-sim/dram.cc @@ -41,7 +41,7 @@ int PRINT_CYCLE = 0; template class fifo_pipeline; template class fifo_pipeline; -dram_t::dram_t( unsigned int partition_id, const struct memory_config *config, memory_stats_t *stats, +dram_t::dram_t( unsigned int partition_id, const memory_config *config, memory_stats_t *stats, memory_partition_unit *mp, gpgpu_sim* gpu ) { id = partition_id; diff --git a/src/gpgpu-sim/dram.h b/src/gpgpu-sim/dram.h index 7a3a2da..0bd9725 100644 --- a/src/gpgpu-sim/dram.h +++ b/src/gpgpu-sim/dram.h @@ -106,11 +106,12 @@ enum bank_grp_bits_position{ }; class mem_fetch; +class memory_config; class dram_t { public: - dram_t( unsigned int parition_id, const struct memory_config *config, class memory_stats_t *stats, + dram_t( unsigned int parition_id, const memory_config *config, class memory_stats_t *stats, class memory_partition_unit *mp, class gpgpu_sim* gpu ); bool full(bool is_write) const; @@ -145,7 +146,7 @@ public: - const struct memory_config *m_config; + const memory_config *m_config; private: bankgrp_t **bkgrp; diff --git a/src/gpgpu-sim/gpu-cache.cc b/src/gpgpu-sim/gpu-cache.cc index f1f6e19..1705821 100644 --- a/src/gpgpu-sim/gpu-cache.cc +++ b/src/gpgpu-sim/gpu-cache.cc @@ -1251,7 +1251,8 @@ data_cache::wr_miss_wa_naive( new_addr_type addr, false, // Now performing a read mf->get_access_warp_mask(), mf->get_access_byte_mask(), - mf->get_access_sector_mask()); + mf->get_access_sector_mask(), + m_gpu->gpgpu_ctx); mem_fetch *n_mf = new mem_fetch( *ma, NULL, @@ -1365,7 +1366,8 @@ data_cache::wr_miss_wa_fetch_on_write( new_addr_type addr, false, // Now performing a read mf->get_access_warp_mask(), mf->get_access_byte_mask(), - mf->get_access_sector_mask()); + mf->get_access_sector_mask(), + m_gpu->gpgpu_ctx); mem_fetch *n_mf = new mem_fetch( *ma, NULL, diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 0481259..0644b44 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1828,7 +1828,7 @@ const shader_core_config * gpgpu_sim::getShaderCoreConfig() return m_shader_config; } -const struct memory_config * gpgpu_sim::getMemoryConfig() +const memory_config * gpgpu_sim::getMemoryConfig() { return m_memory_config; } diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 78f0505..19e1eb3 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -33,6 +33,7 @@ #include "../trace.h" #include "addrdec.h" #include "shader.h" +#include "gpu-cache.h" #include #include #include @@ -143,13 +144,14 @@ struct power_config { }; - -struct memory_config { - memory_config() +class memory_config { + public: + memory_config(gpgpu_context* ctx) { m_valid = false; gpgpu_dram_timing_opt=NULL; gpgpu_L2_queue_config=NULL; + gpgpu_ctx = ctx; } void init() { @@ -291,13 +293,14 @@ struct memory_config { unsigned write_high_watermark; unsigned write_low_watermark; bool m_perf_sim_memcpy; + gpgpu_context* gpgpu_ctx; }; extern bool g_interactive_debugger_enabled; class gpgpu_sim_config : public power_config, public gpgpu_functional_sim_config { public: - gpgpu_sim_config(gpgpu_context* ctx): m_shader_config(ctx) { + gpgpu_sim_config(gpgpu_context* ctx): m_shader_config(ctx), m_memory_config(ctx) { m_valid = false; gpgpu_ctx = ctx; } @@ -507,7 +510,7 @@ public: /*! * Returning the memory configuration of the shader core, used by the functional simulation only so far */ - const struct memory_config * getMemoryConfig(); + const memory_config * getMemoryConfig(); //! Get shader core SIMT cluster @@ -567,7 +570,7 @@ private: const struct cudaDeviceProp *m_cuda_properties; const shader_core_config *m_shader_config; - const struct memory_config *m_memory_config; + const memory_config *m_memory_config; // stats class shader_core_stats *m_shader_stats; diff --git a/src/gpgpu-sim/l2cache.cc b/src/gpgpu-sim/l2cache.cc index 62e70a7..6540b52 100644 --- a/src/gpgpu-sim/l2cache.cc +++ b/src/gpgpu-sim/l2cache.cc @@ -49,7 +49,7 @@ mem_fetch * partition_mf_allocator::alloc(new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const { assert( wr ); - mem_access_t access( type, addr, size, wr ); + mem_access_t access( type, addr, size, wr, m_memory_config->gpgpu_ctx ); mem_fetch *mf = new mem_fetch( access, NULL, WRITE_PACKET_SIZE, @@ -62,7 +62,7 @@ mem_fetch * partition_mf_allocator::alloc(new_addr_type addr, mem_access_type ty } memory_partition_unit::memory_partition_unit( unsigned partition_id, - const struct memory_config *config, + const memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu) : m_id(partition_id), m_config(config), m_stats(stats), m_arbitration_metadata(config), m_gpu(gpu) @@ -95,7 +95,7 @@ memory_partition_unit::~memory_partition_unit() delete[] m_sub_partition; } -memory_partition_unit::arbitration_metadata::arbitration_metadata(const struct memory_config *config) +memory_partition_unit::arbitration_metadata::arbitration_metadata(const memory_config *config) : m_last_borrower(config->m_n_sub_partition_per_memory_channel - 1), m_private_credit(config->m_n_sub_partition_per_memory_channel, 0), m_shared_credit(0) @@ -312,7 +312,7 @@ void memory_partition_unit::print( FILE *fp ) const } memory_sub_partition::memory_sub_partition( unsigned sub_partition_id, - const struct memory_config *config, + const memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu) { @@ -640,7 +640,8 @@ std::vector memory_sub_partition::breakdown_request_to_sector_reques mf->is_write(), mf->get_access_warp_mask(), mf->get_access_byte_mask() & byte_sector_mask, - std::bitset().set(j)); + std::bitset().set(j), + m_gpu->gpgpu_ctx); mem_fetch *n_mf = new mem_fetch( *ma, NULL, diff --git a/src/gpgpu-sim/l2cache.h b/src/gpgpu-sim/l2cache.h index 8ff2666..1f74c47 100644 --- a/src/gpgpu-sim/l2cache.h +++ b/src/gpgpu-sim/l2cache.h @@ -58,7 +58,7 @@ private: class memory_partition_unit { public: - memory_partition_unit( unsigned partition_id, const struct memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu ); + memory_partition_unit( unsigned partition_id, const memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu ); ~memory_partition_unit(); bool busy() const; @@ -98,7 +98,7 @@ public: private: unsigned m_id; - const struct memory_config *m_config; + const memory_config *m_config; class memory_stats_t *m_stats; class memory_sub_partition **m_sub_partition; class dram_t *m_dram; @@ -106,7 +106,7 @@ private: class arbitration_metadata { public: - arbitration_metadata(const struct memory_config *config); + arbitration_metadata(const memory_config *config); // check if a subpartition still has credit bool has_credits(int inner_sub_partition_id) const; @@ -130,7 +130,7 @@ private: std::vector m_private_credit; int m_shared_credit; }; - arbitration_metadata m_arbitration_metadata; + arbitration_metadata m_arbitration_metadata; // determine wheither a given subpartition can issue to DRAM bool can_issue_to_dram(int inner_sub_partition_id); @@ -149,7 +149,7 @@ private: class memory_sub_partition { public: - memory_sub_partition( unsigned sub_partition_id, const struct memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu ); + memory_sub_partition( unsigned sub_partition_id, const memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu ); ~memory_sub_partition(); unsigned get_id() const { return m_id; } @@ -197,7 +197,7 @@ public: private: // data unsigned m_id; //< the global sub partition ID - const struct memory_config *m_config; + const memory_config *m_config; class l2_cache *m_L2cache; class L2interface *m_L2interface; class gpgpu_sim* m_gpu; diff --git a/src/gpgpu-sim/mem_fetch.cc b/src/gpgpu-sim/mem_fetch.cc index c9b0484..6a00889 100644 --- a/src/gpgpu-sim/mem_fetch.cc +++ b/src/gpgpu-sim/mem_fetch.cc @@ -39,10 +39,10 @@ mem_fetch::mem_fetch( const mem_access_t &access, unsigned wid, unsigned sid, unsigned tpc, - const struct memory_config *config, + const memory_config *config, unsigned long long cycle, mem_fetch *m_original_mf, - mem_fetch *m_original_wr_mf) + mem_fetch *m_original_wr_mf):m_access(access) { m_request_uid = sm_next_mf_request_uid++; diff --git a/src/gpgpu-sim/mem_fetch.h b/src/gpgpu-sim/mem_fetch.h index 4eb3a52..1cab9f2 100644 --- a/src/gpgpu-sim/mem_fetch.h +++ b/src/gpgpu-sim/mem_fetch.h @@ -47,6 +47,7 @@ enum mf_type { #undef MF_TUP #undef MF_TUP_END +class memory_config; class mem_fetch { public: mem_fetch( const mem_access_t &access, @@ -55,7 +56,7 @@ public: unsigned wid, unsigned sid, unsigned tpc, - const struct memory_config *config, + const memory_config *config, unsigned long long cycle, mem_fetch *original_mf = NULL, mem_fetch *original_wr_mf = NULL); @@ -149,7 +150,7 @@ private: static unsigned sm_next_mf_request_uid; - const struct memory_config *m_mem_config; + const memory_config *m_mem_config; unsigned icnt_flit_size; mem_fetch* original_mf; //this pointer is set up when a request is divided into sector requests at L2 cache (if the req size > L2 sector size), so the pointer refers to the original request diff --git a/src/gpgpu-sim/mem_latency_stat.cc b/src/gpgpu-sim/mem_latency_stat.cc index d08ba39..4e94991 100644 --- a/src/gpgpu-sim/mem_latency_stat.cc +++ b/src/gpgpu-sim/mem_latency_stat.cc @@ -42,7 +42,7 @@ #include #include -memory_stats_t::memory_stats_t( unsigned n_shader, const shader_core_config *shader_config, const struct memory_config *mem_config, const class gpgpu_sim* gpu ) +memory_stats_t::memory_stats_t( unsigned n_shader, const shader_core_config *shader_config, const memory_config *mem_config, const class gpgpu_sim* gpu ) { assert( mem_config->m_valid ); assert( shader_config->m_valid ); diff --git a/src/gpgpu-sim/mem_latency_stat.h b/src/gpgpu-sim/mem_latency_stat.h index 6ce568d..0c84972 100644 --- a/src/gpgpu-sim/mem_latency_stat.h +++ b/src/gpgpu-sim/mem_latency_stat.h @@ -32,11 +32,12 @@ #include #include +class memory_config; class memory_stats_t { public: memory_stats_t( unsigned n_shader, const class shader_core_config *shader_config, - const struct memory_config *mem_config, + const memory_config *mem_config, const class gpgpu_sim* gpu); unsigned memlatstat_done( class mem_fetch *mf ); @@ -54,7 +55,7 @@ public: unsigned m_n_shader; const shader_core_config *m_shader_config; - const struct memory_config *m_memory_config; + const memory_config *m_memory_config; const class gpgpu_sim* m_gpu; unsigned max_mrq_latency; diff --git a/src/gpgpu-sim/power_stat.cc b/src/gpgpu-sim/power_stat.cc index 007b4c6..2c02082 100644 --- a/src/gpgpu-sim/power_stat.cc +++ b/src/gpgpu-sim/power_stat.cc @@ -42,7 +42,7 @@ -power_mem_stat_t::power_mem_stat_t(const struct memory_config *mem_config, const shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats){ +power_mem_stat_t::power_mem_stat_t(const memory_config *mem_config, const shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats){ assert( mem_config->m_valid ); m_mem_stats = mem_stats; m_config = mem_config; @@ -266,7 +266,7 @@ for(unsigned i=0; inum_shader(); ++i){ } } -power_stat_t::power_stat_t( const shader_core_config *shader_config,float * average_pipeline_duty_cycle,float *active_sms,shader_core_stats * shader_stats, const struct memory_config *mem_config,memory_stats_t * memory_stats) +power_stat_t::power_stat_t( const shader_core_config *shader_config,float * average_pipeline_duty_cycle,float *active_sms,shader_core_stats * shader_stats, const memory_config *mem_config,memory_stats_t * memory_stats) { assert( shader_config->m_valid ); assert( mem_config->m_valid ); diff --git a/src/gpgpu-sim/power_stat.h b/src/gpgpu-sim/power_stat.h index 91fade9..24ade99 100644 --- a/src/gpgpu-sim/power_stat.h +++ b/src/gpgpu-sim/power_stat.h @@ -113,7 +113,7 @@ struct mem_power_stats_pod{ class power_mem_stat_t : public mem_power_stats_pod{ public: - power_mem_stat_t(const struct memory_config *mem_config, const shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats); + power_mem_stat_t(const memory_config *mem_config, const shader_core_config *shdr_config, memory_stats_t *mem_stats, shader_core_stats *shdr_stats); void visualizer_print( gzFile visualizer_file ); void print (FILE *fout) const; void init(); @@ -128,7 +128,7 @@ private: class power_stat_t { public: - power_stat_t( const shader_core_config *shader_config,float * average_pipeline_duty_cycle,float * active_sms,shader_core_stats * shader_stats, const struct memory_config *mem_config,memory_stats_t * memory_stats); + power_stat_t( const shader_core_config *shader_config,float * average_pipeline_duty_cycle,float * active_sms,shader_core_stats * shader_stats, const memory_config *mem_config,memory_stats_t * memory_stats); void visualizer_print( gzFile visualizer_file ); void print (FILE *fout) const; void save_stats(){ @@ -621,7 +621,7 @@ public: float * m_average_pipeline_duty_cycle; float * m_active_sms; const shader_core_config *m_config; - const struct memory_config *m_mem_config; + const memory_config *m_mem_config; }; diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index f380560..b7ae95d 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -28,7 +28,6 @@ #include #include "shader.h" -#include "gpu-sim.h" #include "addrdec.h" #include "dram.h" #include "stat-tool.h" @@ -53,6 +52,19 @@ #define MIN(a,b) (((a)<(b))?(a):(b)) +mem_fetch *shader_core_mem_fetch_allocator::alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const +{ + mem_access_t access( type, addr, size, wr, m_memory_config->gpgpu_ctx); + mem_fetch *mf = new mem_fetch( access, + NULL, + wr?WRITE_PACKET_SIZE:READ_PACKET_SIZE, + -1, + m_core_id, + m_cluster_id, + m_memory_config, + cycle); + return mf; +} ///////////////////////////////////////////////////////////////////////////// std::list shader_core_ctx::get_regs_written( const inst_t &fvt ) const @@ -71,7 +83,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, unsigned shader_id, unsigned tpc_id, const shader_core_config *config, - const struct memory_config *mem_config, + const memory_config *mem_config, shader_core_stats *stats ) : core_t( gpu, NULL, config->warp_size, config->n_thread_per_shader ), m_barriers( this, config->max_warps_per_shader, config->max_cta_per_core, config->max_barriers_per_cta, config->warp_size ), @@ -809,7 +821,7 @@ void shader_core_ctx::fetch() // TODO: replace with use of allocator // mem_fetch *mf = m_mem_fetch_allocator->alloc() - mem_access_t acc(INST_ACC_R,ppc,nbytes,false); + mem_access_t acc(INST_ACC_R,ppc,nbytes,false, m_gpu->gpgpu_ctx); mem_fetch *mf = new mem_fetch(acc, NULL/*we don't have an instruction yet*/, READ_PACKET_SIZE, @@ -3787,7 +3799,7 @@ void opndcoll_rfu_t::collector_unit_t::dispatch() simt_core_cluster::simt_core_cluster( class gpgpu_sim *gpu, unsigned cluster_id, const shader_core_config *config, - const struct memory_config *mem_config, + const memory_config *mem_config, shader_core_stats *stats, class memory_stats_t *mstats ) { diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index 2837f1b..b0d7f7f 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1727,6 +1727,7 @@ private: friend class LooseRoundRobbinScheduler; }; +class memory_config; class shader_core_mem_fetch_allocator : public mem_fetch_allocator { public: shader_core_mem_fetch_allocator( unsigned core_id, unsigned cluster_id, const memory_config *config ) @@ -1735,20 +1736,7 @@ public: m_cluster_id = cluster_id; m_memory_config = config; } - mem_fetch *alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const - { - mem_access_t access( type, addr, size, wr ); - mem_fetch *mf = new mem_fetch( access, - NULL, - wr?WRITE_PACKET_SIZE:READ_PACKET_SIZE, - -1, - m_core_id, - m_cluster_id, - m_memory_config, - cycle); - return mf; - } - + mem_fetch *alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const; mem_fetch *alloc( const warp_inst_t &inst, const mem_access_t &access, unsigned long long cycle ) const { warp_inst_t inst_copy = inst; @@ -1777,7 +1765,7 @@ public: unsigned shader_id, unsigned tpc_id, const shader_core_config *config, - const struct memory_config *mem_config, + const memory_config *mem_config, shader_core_stats *stats ); // used by simt_core_cluster: @@ -2072,7 +2060,7 @@ public: simt_core_cluster( class gpgpu_sim *gpu, unsigned cluster_id, const shader_core_config *config, - const struct memory_config *mem_config, + const memory_config *mem_config, shader_core_stats *stats, memory_stats_t *mstats ); -- cgit v1.3 From c4782ba3c78e6e97db9f56780991dc5cf827c755 Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Sun, 14 Jul 2019 13:06:46 -0400 Subject: Move warp_inst_sm_next_uid Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 2 ++ src/abstract_hardware_model.cc | 15 +++++++++++++-- src/abstract_hardware_model.h | 18 +++--------------- 3 files changed, 18 insertions(+), 17 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 07473be..0dd6311 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -12,6 +12,7 @@ class gpgpu_context { gpgpu_context() { g_global_allfiles_symbol_table = NULL; sm_next_access_uid=0; + warp_inst_sm_next_uid=0; api = new cuda_runtime_api(this); ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); @@ -23,6 +24,7 @@ class gpgpu_context { symbol_table *g_global_allfiles_symbol_table; const char *g_filename; unsigned sm_next_access_uid; + unsigned warp_inst_sm_next_uid; // objects pointers for each file cuda_runtime_api* api; ptxinfo_data* ptxinfo; diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index fe10daa..1dddecd 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -41,8 +41,6 @@ #include #include "../libcuda/gpgpu_context.h" -unsigned warp_inst_t::sm_next_uid = 0; - void mem_access_t::init(gpgpu_context* ctx) { gpgpu_ctx = ctx; @@ -50,6 +48,19 @@ void mem_access_t::init(gpgpu_context* ctx) m_addr=0; m_req_size=0; } +void warp_inst_t::issue( const active_mask_t &mask, unsigned warp_id, unsigned long long cycle, int dynamic_warp_id, int sch_id ) +{ + m_warp_active_mask = mask; + m_warp_issued_mask = mask; + m_uid = ++(m_config->gpgpu_ctx->warp_inst_sm_next_uid); + m_warp_id = warp_id; + m_dynamic_warp_id = dynamic_warp_id; + issue_cycle = cycle; + cycles = initiation_interval; + m_cache_hit=false; + m_empty=false; + m_scheduler_id=sch_id; +} checkpoint::checkpoint() { diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index a22c8c3..519416d 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -958,19 +958,9 @@ public: { m_empty=true; } - void issue( const active_mask_t &mask, unsigned warp_id, unsigned long long cycle, int dynamic_warp_id, int sch_id ) - { - m_warp_active_mask = mask; - m_warp_issued_mask = mask; - m_uid = ++sm_next_uid; - m_warp_id = warp_id; - m_dynamic_warp_id = dynamic_warp_id; - issue_cycle = cycle; - cycles = initiation_interval; - m_cache_hit=false; - m_empty=false; - m_scheduler_id=sch_id; - } + + void issue( const active_mask_t &mask, unsigned warp_id, unsigned long long cycle, int dynamic_warp_id, int sch_id ); + const active_mask_t & get_active_mask() const { return m_warp_active_mask; @@ -1133,8 +1123,6 @@ protected: bool m_mem_accesses_created; std::list m_accessq; - static unsigned sm_next_uid; - unsigned m_scheduler_id; //the scheduler that issues this inst //Jin: cdp support -- cgit v1.3 From bd170826a00fe1ea1960ff2e23247ed3f980547b Mon Sep 17 00:00:00 2001 From: Mengchi Zhang Date: Mon, 15 Jul 2019 12:16:01 -0400 Subject: Move kernel_info_t::m_next_uid Signed-off-by: Mengchi Zhang --- libcuda/gpgpu_context.h | 2 ++ src/abstract_hardware_model.cc | 3 +-- src/abstract_hardware_model.h | 1 - 3 files changed, 3 insertions(+), 3 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index 9c5ae7b..1e20c62 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -14,6 +14,7 @@ class gpgpu_context { sm_next_access_uid=0; warp_inst_sm_next_uid=0; operand_info_sm_next_uid = 1; + kernel_info_m_next_uid = 1; api = new cuda_runtime_api(this); ptxinfo = new ptxinfo_data(this); ptx_parser = new ptx_recognizer(this); @@ -27,6 +28,7 @@ class gpgpu_context { unsigned sm_next_access_uid; unsigned warp_inst_sm_next_uid; unsigned operand_info_sm_next_uid;//uid for operand_info + unsigned kernel_info_m_next_uid;//uid for kernel_info_t // objects pointers for each file cuda_runtime_api* api; ptxinfo_data* ptxinfo; diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 1dddecd..733d602 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -709,7 +709,6 @@ void warp_inst_t::completed( unsigned long long cycle ) const ptx_file_line_stats_add_latency(pc, latency * active_count()); } -unsigned kernel_info_t::m_next_uid = 1; /*A snapshot of the texture mappings needs to be stored in the kernel's info as kernels should use the texture bindings seen at the time of launch and textures @@ -724,7 +723,7 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * m_next_cta.z=0; m_next_tid=m_next_cta; m_num_cores_running=0; - m_uid = m_next_uid++; + m_uid = (entry->gpgpu_ctx->kernel_info_m_next_uid)++; m_param_mem = new memory_space_impl<8192>("param",64*1024); //Jin: parent and child kernel management for CDP diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 519416d..d13b8c6 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -299,7 +299,6 @@ private: class function_info *m_kernel_entry; unsigned m_uid; - static unsigned m_next_uid; //These maps contain the snapshot of the texture mappings at kernel launch std::map m_NameToCudaArray; -- cgit v1.3 From 93597743b24b7d7726096d337e6c9c4516cb2273 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Fri, 30 Aug 2019 18:43:21 -0400 Subject: update the adaptive cache behaviour and make L1 fully assoc in Volta --- configs/tested-cfgs/SM7_QV100/gpgpusim.config | 2 +- configs/tested-cfgs/SM7_TITANV/gpgpusim.config | 2 +- src/abstract_hardware_model.cc | 4 +- src/abstract_hardware_model.h | 9 ++++- src/gpgpu-sim/addrdec.cc | 4 +- src/gpgpu-sim/gpu-cache.cc | 7 +++- src/gpgpu-sim/gpu-cache.h | 5 +++ src/gpgpu-sim/gpu-sim.cc | 2 +- src/gpgpu-sim/shader.cc | 54 ++++++++++++++++---------- 9 files changed, 59 insertions(+), 30 deletions(-) (limited to 'src/abstract_hardware_model.h') diff --git a/configs/tested-cfgs/SM7_QV100/gpgpusim.config b/configs/tested-cfgs/SM7_QV100/gpgpusim.config index f807e11..c0d22ee 100644 --- a/configs/tested-cfgs/SM7_QV100/gpgpusim.config +++ b/configs/tested-cfgs/SM7_QV100/gpgpusim.config @@ -91,7 +91,7 @@ # Volta unified cache has four banks -l1_banks 4 #-mem_unit_ports 4 --gpgpu_cache:dl1 S:4:128:64,L:L:s:N:L,A:256:8,16:0,32 +-gpgpu_cache:dl1 S:1:128:256,L:L:s:N:L,A:256:8,16:0,32 -gpgpu_shmem_size 98304 -gpgpu_shmem_sizeDefault 98304 -gpgpu_shmem_per_block 65536 diff --git a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config index 888ce71..0339b0d 100644 --- a/configs/tested-cfgs/SM7_TITANV/gpgpusim.config +++ b/configs/tested-cfgs/SM7_TITANV/gpgpusim.config @@ -91,7 +91,7 @@ # Volta unified cache has four banks -l1_banks 4 #-mem_unit_ports 4 --gpgpu_cache:dl1 S:4:128:64,L:L:s:N:L,A:256:8,16:0,32 +-gpgpu_cache:dl1 S:1:128:256,L:L:s:N:L,A:256:8,16:0,32 -gpgpu_shmem_size 98304 -gpgpu_shmem_sizeDefault 98304 -gpgpu_shmem_per_block 65536 diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 023f51b..35a3984 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -710,7 +710,7 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * //Jin: launch latency management m_launch_latency = g_kernel_launch_latency; - volta_cache_config_set=false; + cache_config_set=false; } /*A snapshot of the texture mappings needs to be stored in the kernel's info as @@ -735,7 +735,7 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * //Jin: launch latency management m_launch_latency = g_kernel_launch_latency; - volta_cache_config_set=false; + cache_config_set=false; m_NameToCudaArray = nameToCudaArray; m_NameToTextureInfo = nameToTextureInfo; } diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 27a1ba6..231b6a2 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -65,6 +65,11 @@ enum FuncCache FuncCachePreferL1 = 2 }; +enum AdaptiveCache +{ + FIXED = 0, + VOLTA = 1 +}; #ifdef __cplusplus @@ -345,7 +350,7 @@ public: unsigned long long end_cycle; unsigned m_launch_latency; - mutable bool volta_cache_config_set; + mutable bool cache_config_set; }; struct core_config { @@ -388,7 +393,7 @@ struct core_config { unsigned gpgpu_max_insn_issue_per_warp; bool gmem_skip_L1D; // on = global memory access always skip the L1 cache - bool adaptive_volta_cache_config; + unsigned adaptive_cache_config; }; // bounded stack that implements simt reconvergence using pdom mechanism from MICRO'07 paper diff --git a/src/gpgpu-sim/addrdec.cc b/src/gpgpu-sim/addrdec.cc index b0db034..3262456 100644 --- a/src/gpgpu-sim/addrdec.cc +++ b/src/gpgpu-sim/addrdec.cc @@ -124,8 +124,8 @@ void linear_to_raw_address_translation::addrdec_tlx(new_addr_type addr, addrdec_ * Rau, B. R et al. * ISCA 1991 * - * equations are adopted from: - * "Sacat: streaming-aware conflict-avoiding thrashing-resistant gpgpu cache management scheme." + * equations are corresponding to IPOLY(37) and are adopted from: + * "SACAT: streaming-aware conflict-avoiding thrashing-resistant gpgpu cache management scheme." * Khairy et al. * IEEE TPDS 2017. */ diff --git a/src/gpgpu-sim/gpu-cache.cc b/src/gpgpu-sim/gpu-cache.cc index db9701d..8d00ea9 100644 --- a/src/gpgpu-sim/gpu-cache.cc +++ b/src/gpgpu-sim/gpu-cache.cc @@ -65,8 +65,13 @@ const char * cache_fail_status_str(enum cache_reservation_fail_reason status) unsigned l1d_cache_config::set_bank(new_addr_type addr) const{ - if(m_cache_type == SECTOR) + //For sector cache, we select one sector per bank (sector interleaving) + //This is what was found in Volta (one sector per bank, sector interleaving) + //otherwise, line interleaving + if(m_cache_type == SECTOR) { + //assert(l1_banks == SECTOR_CHUNCK_SIZE); return (addr >> m_sector_sz_log2) & (l1_banks-1); + } else return (addr >> m_line_sz_log2) & (l1_banks-1); } diff --git a/src/gpgpu-sim/gpu-cache.h b/src/gpgpu-sim/gpu-cache.h index 90adbb5..dd22886 100644 --- a/src/gpgpu-sim/gpu-cache.h +++ b/src/gpgpu-sim/gpu-cache.h @@ -656,6 +656,11 @@ public: assert( m_valid ); return MAX_DEFAULT_CACHE_SIZE_MULTIBLIER * m_nset * original_m_assoc; } + unsigned get_max_assoc() const + { + assert( m_valid ); + return MAX_DEFAULT_CACHE_SIZE_MULTIBLIER * original_m_assoc; + } void print( FILE *fp ) const { fprintf( fp, "Size = %d B (%d Set x %d-way x %d byte line)\n", diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 4f071c7..ed94865 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -323,7 +323,7 @@ void shader_core_config::reg_options(class OptionParser * opp) option_parser_register(opp, "-gpgpu_shmem_size", OPT_UINT32, &gpgpu_shmem_size, "Size of shared memory per shader core (default 16kB)", "16384"); - option_parser_register(opp, "-adaptive_cache_config", OPT_BOOL, &adaptive_volta_cache_config, + option_parser_register(opp, "-adaptive_cache_config", OPT_UINT32, &adaptive_cache_config, "adaptive_cache_config", "0"); option_parser_register(opp, "-gpgpu_shmem_sizeDefault", OPT_UINT32, &gpgpu_shmem_sizeDefault, diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index ffd3035..ff2bf3f 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -2980,33 +2980,47 @@ unsigned int shader_core_config::max_cta( const kernel_info_t &k ) const abort(); } - if(adaptive_volta_cache_config && !k.volta_cache_config_set) { - //For Volta, we assign the remaining shared memory to L1 cache - //For more info, see https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory-7-x + if(adaptive_cache_config && !k.cache_config_set) { + //For more info about adaptive cache, see https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory-7-x unsigned total_shmed = kernel_info->smem * result; assert(total_shmed >=0 && total_shmed <= gpgpu_shmem_size); - assert(gpgpu_shmem_size == 98304); //Volta has 96 KB shared - assert(m_L1D_config.get_nset() == 4); //Volta L1 has four sets + //assert(gpgpu_shmem_size == 98304); //Volta has 96 KB shared + //assert(m_L1D_config.get_nset() == 4); //Volta L1 has four sets if(total_shmed < gpgpu_shmem_size){ - if(total_shmed == 0) - m_L1D_config.set_assoc(256); //L1 is 128KB ans shd=0 - else if(total_shmed > 0 && total_shmed <= 8192) - m_L1D_config.set_assoc(240); //L1 is 120KB ans shd=8KB - else if(total_shmed > 8192 && total_shmed <= 16384) - m_L1D_config.set_assoc(224); //L1 is 112KB ans shd=16KB - else if(total_shmed > 16384 && total_shmed <= 32768) - m_L1D_config.set_assoc(192); //L1 is 96KB ans shd=32KB - else if(total_shmed > 32768 && total_shmed <= 65536) - m_L1D_config.set_assoc(128); //L1 is 64KB ans shd=64KB - else if(total_shmed > 65536 && total_shmed <= gpgpu_shmem_size) - m_L1D_config.set_assoc(64); //L1 is 32KB and shd=96KB - else + switch (adaptive_cache_config) { + case FIXED: + break; + case VOLTA: { + //For Volta, we assign the remaining shared memory to L1 cache + //For more info about adaptive cache, see https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory-7-x + assert(gpgpu_shmem_size == 98304); //Volta has 96 KB shared + + //To Do: make it flexible and not tuned to 9KB share memory + unsigned max_assoc = m_L1D_config.get_max_assoc(); + if(total_shmed == 0) + m_L1D_config.set_assoc(max_assoc); //L1 is 128KB and shd=0 + else if(total_shmed > 0 && total_shmed <= 8192) + m_L1D_config.set_assoc(0.9375 * max_assoc); //L1 is 120KB and shd=8KB + else if(total_shmed > 8192 && total_shmed <= 16384) + m_L1D_config.set_assoc(0.875 * max_assoc); //L1 is 112KB and shd=16KB + else if(total_shmed > 16384 && total_shmed <= 32768) + m_L1D_config.set_assoc(0.75 * max_assoc); //L1 is 96KB and shd=32KB + else if(total_shmed > 32768 && total_shmed <= 65536) + m_L1D_config.set_assoc(0.5 * max_assoc); //L1 is 64KB and shd=64KB + else if(total_shmed > 65536 && total_shmed <= gpgpu_shmem_size) + m_L1D_config.set_assoc(0.25 * max_assoc); //L1 is 32KB and shd=96KB + else + assert(0); + break; + } + default: assert(0); + } - printf ("GPGPU-Sim: Reconfigure L1 cache in Volta Archi to %uKB\n", m_L1D_config.get_total_size_inKB()); + printf ("GPGPU-Sim: Reconfigure L1 cache to %uKB\n", m_L1D_config.get_total_size_inKB()); } - k.volta_cache_config_set = true; + k.cache_config_set = true; } return result; -- cgit v1.3