diff options
| -rw-r--r-- | CHANGES | 5 | ||||
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 4 | ||||
| -rw-r--r-- | libcuda/cuobjdump.y | 2 | ||||
| -rw-r--r-- | src/cuda-sim/cuda-sim.cc | 7 | ||||
| -rw-r--r-- | src/cuda-sim/instructions.cc | 24 | ||||
| -rw-r--r-- | src/cuda-sim/ptx_loader.cc | 2 | ||||
| -rw-r--r-- | src/gpgpu-sim/dram_sched.cc | 2 | ||||
| -rw-r--r-- | src/gpgpu-sim/l2cache.h | 1 | ||||
| -rw-r--r-- | src/gpgpu-sim/shader.cc | 9 |
9 files changed, 47 insertions, 9 deletions
@@ -1,10 +1,5 @@ CHANGE LOG: -Version 3.1.2 versus 3.1.1 -- Bug Fixes: - - Fixed a minor compilation bug with newer gcc/g++ versions - - Removed compiler warnings with gcc 4.7.1 - Version 3.1.1 versus 3.1.0 - Add checks to top level makefile to ensure setup_environment is run and checks to setup_environment to ensure a compatible version of CUDA is installed. diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index cdfd99c..e189403 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -551,7 +551,9 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2D(void *dst, size_t dpitch, const void { CUctx_st *context = GPGPUSim_Context(); gpgpu_t *gpu = context->get_device()->get_gpgpu(); + struct cudaArray *cuArray_ptr; size_t size = spitch*height; + cuArray_ptr = (cudaArray*)dst; gpgpusim_ptx_assert( (dpitch==spitch), "different src and dst pitch not supported yet" ); if( kind == cudaMemcpyHostToDevice ) gpu->memcpy_to_gpu( (size_t)dst, src, size ); @@ -1799,7 +1801,7 @@ int CUDARTAPI __cudaSynchronizeThreads(void**, void*) //////// -//extern "C" int ptx_parse(); +extern "C" int ptx_parse(); extern "C" int ptx__scan_string(const char*); extern "C" FILE *ptx_in; diff --git a/libcuda/cuobjdump.y b/libcuda/cuobjdump.y index b07c754..914a3d6 100644 --- a/libcuda/cuobjdump.y +++ b/libcuda/cuobjdump.y @@ -103,7 +103,7 @@ ptxcode : ptxcode PTXLINE {fprintf(ptxfile, "%s", $2);} elfcode : elfcode ELFLINE {fprintf(elffile, "%s", $2);} | ; -sasscode : sasscode SASSLINE {fprintf(sassfile, "%s", $2);} +sasscode : sasscode SASSLINE {fprintf(sassfile, "%s", $2)} | ; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 89c240f..02c7dfd 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -438,6 +438,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 ) { + static unsigned size=1; std::map<unsigned,function_info*>::iterator f = g_pc_to_finfo.find(pc); if( f == g_pc_to_finfo.end() ) { fprintf(fp,"<no instruction at address 0x%x>", pc ); @@ -445,6 +446,7 @@ void ptx_print_insn( address_type pc, FILE *fp ) } function_info *finfo = f->second; assert( finfo ); + size = finfo->print_insn(pc,fp); } void ptx_instruction::set_opcode_and_latency() @@ -1088,6 +1090,8 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id) // Output instruction information to file and stdout if( config.get_ptx_inst_debug_to_file() != 0 && (config.get_ptx_inst_debug_thread_uid() == 0 || config.get_ptx_inst_debug_thread_uid() == get_uid()) ) { + dim3 ctaid = get_ctaid(); + dim3 tid = get_tid(); fprintf(m_gpu->get_ptx_inst_debug_file(), "[thd=%u] : (%s:%u - %s)\n", get_uid(), @@ -1453,6 +1457,7 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co printf("Execution error: No information for PTX symbol w/ hostVar=0x%Lx\n", (unsigned long long)hostVar ); abort(); } else printf("GPGPU-Sim PTX: gpgpu_ptx_sim_memcpy_symbol: Found PTX symbol w/ hostVar=0x%Lx\n", (unsigned long long)hostVar ); + const char *mem_name = NULL; memory_space *mem = NULL; std::map<std::string,symbol_table*>::iterator st = g_sym_name_to_symbol_table.find(sym_name.c_str()); @@ -1465,9 +1470,11 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co switch (mem_region.get_type()) { case const_space: mem = gpu->get_global_memory(); + mem_name = "global"; break; case global_space: mem = gpu->get_global_memory(); + mem_name = "global"; break; default: abort(); diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 8a737bc..6976ba0 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -882,6 +882,8 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread ) src2_data, // b op_result; // temp variable to hold operation result + bool data_ready = false; + // Get operand info of sources and destination const operand_info &dst = pI->dst(); // d const operand_info &src1 = pI->src1(); // a @@ -916,9 +918,11 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread ) case B32_TYPE: case U32_TYPE: op_result.u32 = data.u32 & src2_data.u32; + data_ready = true; break; case S32_TYPE: op_result.s32 = data.s32 & src2_data.s32; + data_ready = true; break; default: printf("Execution error: type mismatch (%x) with instruction\natom.AND only accepts b32\n", to_type); @@ -936,9 +940,11 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread ) case B32_TYPE: case U32_TYPE: op_result.u32 = data.u32 | src2_data.u32; + data_ready = true; break; case S32_TYPE: op_result.s32 = data.s32 | src2_data.s32; + data_ready = true; break; default: printf("Execution error: type mismatch (%x) with instruction\natom.OR only accepts b32\n", to_type); @@ -956,9 +962,11 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread ) case B32_TYPE: case U32_TYPE: op_result.u32 = data.u32 ^ src2_data.u32; + data_ready = true; break; case S32_TYPE: op_result.s32 = data.s32 ^ src2_data.s32; + data_ready = true; break; default: printf("Execution error: type mismatch (%x) with instruction\natom.XOR only accepts b32\n", to_type); @@ -980,13 +988,16 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread ) case B32_TYPE: case U32_TYPE: op_result.u32 = MY_CAS_I(data.u32, src2_data.u32, src3_data.u32); + data_ready = true; break; case B64_TYPE: case U64_TYPE: op_result.u64 = MY_CAS_I(data.u64, src2_data.u64, src3_data.u64); + data_ready = true; break; case S32_TYPE: op_result.s32 = MY_CAS_I(data.s32, src2_data.s32, src3_data.s32); + data_ready = true; break; default: printf("Execution error: type mismatch (%x) with instruction\natom.CAS only accepts b32 and b64\n", to_type); @@ -1003,13 +1014,16 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread ) case B32_TYPE: case U32_TYPE: op_result.u32 = MY_EXCH(data.u32, src2_data.u32); + data_ready = true; break; case B64_TYPE: case U64_TYPE: op_result.u64 = MY_EXCH(data.u64, src2_data.u64); + data_ready = true; break; case S32_TYPE: op_result.s32 = MY_EXCH(data.s32, src2_data.s32); + data_ready = true; break; default: printf("Execution error: type mismatch (%x) with instruction\natom.EXCH only accepts b32\n", to_type); @@ -1026,15 +1040,19 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread ) switch ( to_type ) { case U32_TYPE: op_result.u32 = data.u32 + src2_data.u32; + data_ready = true; break; case S32_TYPE: op_result.s32 = data.s32 + src2_data.s32; + data_ready = true; break; case U64_TYPE: op_result.u64 = data.u64 + src2_data.u64; + data_ready = true; break; case F32_TYPE: op_result.f32 = data.f32 + src2_data.f32; + data_ready = true; break; default: printf("Execution error: type mismatch with instruction\natom.ADD only accepts u32, s32, u64, and f32\n"); @@ -1050,6 +1068,7 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread ) switch ( to_type ) { case U32_TYPE: op_result.u32 = MY_INC_I(data.u32, src2_data.u32); + data_ready = true; break; default: printf("Execution error: type mismatch with instruction\natom.INC only accepts u32 and s32\n"); @@ -1065,6 +1084,7 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread ) switch ( to_type ) { case U32_TYPE: op_result.u32 = MY_DEC_I(data.u32, src2_data.u32); + data_ready = true; break; default: printf("Execution error: type mismatch with instruction\natom.DEC only accepts u32 and s32\n"); @@ -1080,9 +1100,11 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread ) switch ( to_type ) { case U32_TYPE: op_result.u32 = MY_MIN_I(data.u32, src2_data.u32); + data_ready = true; break; case S32_TYPE: op_result.s32 = MY_MIN_I(data.s32, src2_data.s32); + data_ready = true; break; default: printf("Execution error: type mismatch with instruction\natom.MIN only accepts u32 and s32\n"); @@ -1098,9 +1120,11 @@ void atom_callback( const inst_t* inst, ptx_thread_info* thread ) switch ( to_type ) { case U32_TYPE: op_result.u32 = MY_MAX_I(data.u32, src2_data.u32); + data_ready = true; break; case S32_TYPE: op_result.s32 = MY_MAX_I(data.s32, src2_data.s32); + data_ready = true; break; default: printf("Execution error: type mismatch with instruction\natom.MAX only accepts u32 and s32\n"); diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc index b5103df..56b9449 100644 --- a/src/cuda-sim/ptx_loader.cc +++ b/src/cuda-sim/ptx_loader.cc @@ -43,7 +43,7 @@ bool g_override_embedded_ptx = false; /// extern prototypes -//extern "C" int ptx_parse(); +extern "C" int ptx_parse(); extern "C" int ptx__scan_string(const char*); const char *g_ptxinfo_filename; diff --git a/src/gpgpu-sim/dram_sched.cc b/src/gpgpu-sim/dram_sched.cc index 48e96ee..e556edc 100644 --- a/src/gpgpu-sim/dram_sched.cc +++ b/src/gpgpu-sim/dram_sched.cc @@ -78,6 +78,7 @@ void frfcfs_scheduler::data_collection(unsigned int bank) dram_req_t *frfcfs_scheduler::schedule( unsigned bank, unsigned curr_row ) { + int row_hit = 0; if ( m_last_row[bank] == NULL ) { if ( m_queue[bank].empty() ) return NULL; @@ -94,6 +95,7 @@ dram_req_t *frfcfs_scheduler::schedule( unsigned bank, unsigned curr_row ) } } + row_hit=1; std::list<dram_req_t*>::iterator next = m_last_row[bank]->back(); dram_req_t *req = (*next); diff --git a/src/gpgpu-sim/l2cache.h b/src/gpgpu-sim/l2cache.h index 15c3806..22640b8 100644 --- a/src/gpgpu-sim/l2cache.h +++ b/src/gpgpu-sim/l2cache.h @@ -121,7 +121,6 @@ private: class L2interface : public mem_fetch_interface { public: L2interface( memory_partition_unit *unit ) { m_unit=unit; } - virtual ~L2interface(){} virtual bool full( unsigned size, bool write) const { // assume read and write packets all same size diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index a7daeaa..ddb544a 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -672,6 +672,10 @@ void TwoLevelScheduler::cycle() { //Do the scheduling only from activeWarps //If you schedule an instruction, move it to the end of the list + bool valid_inst = false; // there was one warp with a valid instruction to issue (didn't require flush due to control hazard) + bool ready_inst = false; // of the valid instructions, there was one not waiting for pending register writes + bool issued_inst = false; // of these we issued one + for ( std::list<int>::iterator warp_id = activeWarps.begin(); warp_id != activeWarps.end(); warp_id++) { @@ -691,13 +695,16 @@ void TwoLevelScheduler::cycle() { warp(*warp_id).set_next_pc(pc); warp(*warp_id).ibuffer_flush(); } else { + valid_inst = true; if ( !m_scoreboard->checkCollision(*warp_id, pI) ) { + ready_inst = true; const active_mask_t &active_mask = m_simt_stack[*warp_id]->get_active_mask(); assert( warp(*warp_id).inst_in_pipeline() ); if ( (pI->op == LOAD_OP) || (pI->op == STORE_OP) || (pI->op == MEMORY_BARRIER_OP) ) { if( m_mem_out->has_free() ) { m_shader->issue_warp(*m_mem_out,pI,active_mask,*warp_id); issued++; + issued_inst=true; warp_inst_issued = true; // Move it to pendingWarps unsigned currwarp = *warp_id; @@ -711,6 +718,7 @@ void TwoLevelScheduler::cycle() { // always prefer SP pipe for operations that can use both SP and SFU pipelines m_shader->issue_warp(*m_sp_out,pI,active_mask,*warp_id); issued++; + issued_inst=true; warp_inst_issued = true; //Move it to end of the activeWarps unsigned currwarp = *warp_id; @@ -720,6 +728,7 @@ void TwoLevelScheduler::cycle() { if( sfu_pipe_avail ) { m_shader->issue_warp(*m_sfu_out,pI,active_mask,*warp_id); issued++; + issued_inst=true; warp_inst_issued = true; //Move it to end of the activeWarps unsigned currwarp = *warp_id; |