From 642818ae5ff61c1544bcce9e7ba2dd0aea47ea6a Mon Sep 17 00:00:00 2001 From: Deval Shah Date: Fri, 9 Nov 2018 21:29:18 -0800 Subject: Adding checkpoint support --- libcuda/cuda_runtime_api.cc | 35 +++++++++++++++- src/abstract_hardware_model.h | 30 ++++++++++++-- src/cuda-sim/cuda-sim.cc | 94 +++++++++++++++++++++++++++++++++++++++---- src/cuda-sim/instructions.cc | 56 ++++++++++++++++++++++++++ src/gpgpu-sim/gpu-sim.cc | 25 ++++++++++-- 5 files changed, 224 insertions(+), 16 deletions(-) diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index f00fe52..4b50e34 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1425,7 +1425,40 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) dim3 gridDim = config.grid_dim(); dim3 blockDim = config.block_dim(); - + gpgpu_t *gpu = context->get_device()->get_gpgpu(); + checkpoint *g_checkpoint; + g_checkpoint = new checkpoint(); + class memory_space *global_mem; + global_mem = gpu->get_global_memory(); + + if(gpu->resume_option ==1 && (grid->get_uid()==gpu->resume_kernel)) + { + + char f1name[2048]; + snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", grid->get_uid()); + + g_checkpoint->load_global_mem(global_mem, f1name); + for (int i=0;iresume_CTA;i++) + grid->increment_cta_id(); + } + if(gpu->resume_option==1 && (grid->get_uid()resume_kernel)) + { + char f1name[2048]; + snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", grid->get_uid()); + + g_checkpoint->load_global_mem(global_mem, f1name); + printf("Skipping kernel %d as resuming from kernel %d\n",grid->get_uid(),gpu->resume_kernel ); + g_cuda_launch_stack.pop_back(); + return g_last_cudaError = cudaSuccess; + + } + if(gpu->checkpoint_option==1 && (grid->get_uid()>gpu->checkpoint_kernel)) + { + printf("Skipping kernel %d as checkpoint from kernel %d\n",grid->get_uid(),gpu->checkpoint_kernel ); + g_cuda_launch_stack.pop_back(); + return g_last_cudaError = cudaSuccess; + + } 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); diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 45fba76..2350db4 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -381,7 +381,8 @@ public: void get_pdom_stack_top_info( unsigned *pc, unsigned *rpc ) const; unsigned get_rp() const; void print(FILE *fp) const; - + void resume(char * fname) ; + void print_checkpoint (FILE *fout) const; protected: unsigned m_warp_id; @@ -501,14 +502,28 @@ public: const char* get_ptx_inst_debug_file() const { return g_ptx_inst_debug_file; } int get_ptx_inst_debug_thread_uid() const { return g_ptx_inst_debug_thread_uid; } unsigned get_texcache_linesize() const { return m_texcache_linesize; } - + int get_checkpoint_option() const {return checkpoint_option; } + int get_checkpoint_kernel() const {return checkpoint_kernel; } + int get_checkpoint_CTA() const {return checkpoint_CTA; } + int get_resume_option() const {return resume_option; } + int get_resume_kernel() const {return resume_kernel; } + int get_resume_CTA() const {return resume_CTA; } + int get_checkpoint_CTA_t() const {return checkpoint_CTA_t; } + int get_checkpoint_insn_Y() const {return checkpoint_insn_Y; } private: // PTX options int m_ptx_convert_to_ptxplus; int m_ptx_use_cuobjdump; int m_experimental_lib_support; unsigned m_ptx_force_max_capability; - + int checkpoint_option; + int checkpoint_kernel; + int checkpoint_CTA; + int resume_option; + int resume_kernel; + int resume_CTA; + int checkpoint_CTA_t; + int checkpoint_insn_Y; int g_ptx_inst_debug_to_file; char* g_ptx_inst_debug_file; int g_ptx_inst_debug_thread_uid; @@ -520,7 +535,14 @@ private: class gpgpu_t { public: gpgpu_t( const gpgpu_functional_sim_config &config ); - + int checkpoint_option; + int checkpoint_kernel; + int checkpoint_CTA; + int resume_option; + int resume_kernel; + int resume_CTA; + int checkpoint_CTA_t; + int checkpoint_insn_Y; 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 ); diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 6a6b307..642e301 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -2125,7 +2125,8 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) //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); - + checkpoint *g_checkpoint; + g_checkpoint = new checkpoint(); if (kernel_func_info->is_pdom_set()) { printf("GPGPU-Sim PTX: PDOM analysis already done for %s \n", kernel.name().c_str() ); @@ -2142,12 +2143,21 @@ 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 cta_launched =0; //we excute the kernel one CTA (Block) at the time, as synchronization functions work block wise while(!kernel.no_more_ctas_to_run()){ unsigned temp=kernel.get_next_cta_id_single(); + if(cp_op==0 || (cp_op==1 && cta_launched= 5000) launch_all_device_kernels(); #endif - + } + else + { + kernel.increment_cta_id(); + } + cta_launched++; } - + if(cp_op==1) + { + 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"); + } @@ -2207,6 +2227,7 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) void functionalCoreSim::initializeCTA(unsigned ctaid_cp) { int ctaLiveThreads=0; + symbol_table * symtab= m_kernel->entry()->get_symtab(); for(int i=0; i< m_warp_count; i++){ m_warpAtBarrier[i]=false; @@ -2219,7 +2240,10 @@ void functionalCoreSim::initializeCTA(unsigned ctaid_cp) for(unsigned i=0; ithreads_per_cta();i++) { ptx_sim_init_thread(*m_kernel,&m_thread[i],0,i,m_kernel->threads_per_cta()-i,m_kernel->threads_per_cta(),this,0,i/m_warp_size,(gpgpu_t*)m_gpu, true); assert(m_thread[i]!=NULL && !m_thread[i]->is_done()); - + char fname[2048]; + snprintf(fname,2048,"checkpoint_files/thread_%d_0_reg.txt",i ); + if(cp_cta_resume==1) + m_thread[i]->resume_reg_thread(fname,symtab); ctaLiveThreads++; } @@ -2242,25 +2266,40 @@ void functionalCoreSim::createWarp(unsigned warpId) char fname[2048]; snprintf(fname,2048,"checkpoint_files/warp_%d_0_simt.txt",warpId ); + if(cp_cta_resume==1) + { + unsigned pc,rpc; + m_simt_stack[warpId]->resume(fname); + m_simt_stack[warpId]->get_pdom_stack_top_info(&pc,&rpc); + for(int i=warpId*m_warp_size; iset_npc(pc); + m_thread[i]->update_pc(); + } + } m_liveThreadCount[warpId]= liveThreadsCount; } void functionalCoreSim::execute(int inst_count, unsigned ctaid_cp) { - + cp_count= m_gpu->checkpoint_insn_Y; + cp_cta_resume= m_gpu->checkpoint_CTA_t; initializeCTA(ctaid_cp); - + int count=0; while(true){ bool someOneLive= false; bool allAtBarrier = true; for(unsigned i=0;i0 && count>inst_count && (m_kernel->get_uid()==m_gpu->checkpoint_kernel) && (ctaid_cp>=m_gpu->checkpoint_CTA) && (ctaid_cpcheckpoint_CTA_t) && m_gpu->checkpoint_option==1) + { + someOneLive=false; + break; + } if(!someOneLive) break; if(allAtBarrier){ for(unsigned i=0;ientry()->get_symtab(); + + + unsigned ctaid =m_kernel->get_next_cta_id_single(); + if(m_gpu->checkpoint_option==1 && (m_kernel->get_uid()==m_gpu->checkpoint_kernel) && (ctaid_cp>=m_gpu->checkpoint_CTA) && (ctaid_cpcheckpoint_CTA_t)) + { + char fname[2048]; + snprintf(fname,2048,"checkpoint_files/shared_mem_%d.txt",ctaid-1 ); + g_checkpoint->store_global_mem(m_thread[0]->m_shared_mem, fname , "%08x"); + for(int i=0; i<32*m_warp_count;i++) + { + char fname[2048]; + snprintf(fname,2048,"checkpoint_files/thread_%d_%d_reg.txt",i,ctaid-1 ); + m_thread[i]->print_reg_thread(fname); + char f1name[2048]; + snprintf(f1name,2048,"checkpoint_files/local_mem_thread_%d_%d_reg.txt",i,ctaid-1 ); + g_checkpoint->store_global_mem(m_thread[i]->m_local_mem, f1name , "%08x"); + m_thread[i]->set_done(); + m_thread[i]->exitCore(); + m_thread[i]->registerExit(); + } + + for(int i=0;iprint_checkpoint(fp); + fclose(fp); + } + } } diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index f57a3f7..31a33c6 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -183,7 +183,63 @@ void ptx_thread_info::set_reg( const symbol *reg, const ptx_reg_t &value ) m_last_set_operand_value = value; } +void ptx_thread_info::print_reg_thread(char * fname) +{ + + FILE *fp= fopen(fname,"w"); + assert(fp!=NULL); + + int size = m_regs.size(); + + if(size>0) + { + reg_map_t reg = m_regs.back(); + + typename reg_map_t::const_iterator it; + for (it = reg.begin(); it != reg.end(); ++it) + { + const std::string &name = it->first->name(); + const std::string &dec= it->first->decl_location(); + unsigned size = it->first->get_size_in_bytes(); + fprintf(fp,"%s %llu %s %d\n",name.c_str(),it->second, dec.c_str(),size ); + + } + //m_regs.pop_back(); + } + fclose(fp); + + } + +void ptx_thread_info::resume_reg_thread(char * fname, symbol_table * symtab) +{ + + FILE * fp2 = fopen(fname, "r"); + assert(fp2!=NULL); + //m_regs.push_back( reg_map_t() ); + char line [ 200 ]; + while ( fgets ( line, sizeof line, fp2 ) != NULL ) + { + symbol *reg; + char * pch; + unsigned size; + pch = strtok (line," "); + char * name =pch; + reg= symtab->lookup(name); + ptx_reg_t data; + pch = strtok (NULL," "); + data = atoi(pch); + pch = strtok (NULL," "); + char * decl= pch; + pch = strtok (NULL," "); + size = atoi(pch); + + + m_regs.back()[reg] = data; + } + fclose ( fp2 ); +} + ptx_reg_t ptx_thread_info::get_reg( const symbol *reg ) { diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index c706f23..79a6fcd 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -1414,19 +1414,38 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel ) // bind functional simulation state of threads to hardware resources (simulation) warp_set_t warps; unsigned nthreads_in_block= 0; - + function_info *kernel_func_info = kernel.entry(); + symbol_table * symtab= kernel_func_info->get_symtab(); + unsigned ctaid= kernel.get_next_cta_id_single(); + checkpoint *g_checkpoint= new checkpoint(); for (unsigned i = start_thread; iwarp_size; nthreads_in_block += ptx_sim_init_thread(kernel,&m_thread[i],m_sid,i,cta_size-(i-start_thread),m_config->n_thread_per_shader,this,free_cta_hw_id,warp_id,m_cluster->get_gpu()); m_threadState[i].m_active = true; - + // load thread local memory and register file + if(m_gpu->resume_option==1 && kernel.get_uid()==m_gpu->resume_kernel && ctaid>=m_gpu->resume_CTA && ctaidcheckpoint_CTA_t ) + { + char fname[2048]; + snprintf(fname,2048,"checkpoint_files/thread_%d_%d_reg.txt",i%cta_size,ctaid ); + m_thread[i]->resume_reg_thread(fname,symtab); + char f1name[2048]; + snprintf(f1name,2048,"checkpoint_files/local_mem_thread_%d_%d_reg.txt",i%cta_size,ctaid); + g_checkpoint->load_global_mem(m_thread[i]->m_local_mem, f1name); + } + // warps.set( warp_id ); } assert( nthreads_in_block > 0 && nthreads_in_block <= m_config->n_thread_per_shader); // should be at least one, but less than max m_cta_status[free_cta_hw_id]=nthreads_in_block; - + if(m_gpu->resume_option==1 && kernel.get_uid()==m_gpu->resume_kernel && ctaid>=m_gpu->resume_CTA && ctaidcheckpoint_CTA_t ) + { + char f1name[2048]; + snprintf(f1name,2048,"checkpoint_files/shared_mem_%d.txt", ctaid); + + g_checkpoint->load_global_mem(m_thread[start_thread]->m_shared_mem, f1name); + } // now that we know which warps are used in this CTA, we can allocate // resources for use in CTA-wide barrier operations m_barriers.allocate_barrier(free_cta_hw_id,warps); -- cgit v1.3