diff options
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 7 | ||||
| -rw-r--r-- | src/abstract_hardware_model.h | 3 | ||||
| -rw-r--r-- | src/cuda-sim/instructions.cc | 3 | ||||
| -rw-r--r-- | src/cuda-sim/opcodes.def | 2 |
4 files changed, 10 insertions, 5 deletions
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index e8a0e91..3eff4af 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -180,6 +180,9 @@ cudaError_t g_last_cudaError = cudaSuccess; extern stream_manager *g_stream_manager; +dim3 gridDim = (dim3){1,1,1}; +dim3 blockDim = (dim3){1,1,1}; + void register_ptx_function( const char *name, function_info *impl ) { // no longer need this @@ -959,8 +962,8 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 ); kernel_info_t *grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context); std::string kname = grid->name(); - dim3 gridDim = config.grid_dim(); - dim3 blockDim = config.block_dim(); + gridDim = config.grid_dim(); + blockDim = config.block_dim(); 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 6ed9b8e..46c3279 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -162,6 +162,9 @@ struct dim3 { }; #endif +extern dim3 gridDim; +extern dim3 blockDim; + void increment_x_then_y_then_z( dim3 &i, const dim3 &bound); class kernel_info_t { diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 8bdb94f..fd3b1fa 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -3772,7 +3772,7 @@ void sst_impl( const ptx_instruction *pI, ptx_thread_info *thread ) thread->m_last_dram_callback.function = bar_callback; thread->m_last_dram_callback.instruction = cpI; - int NUM_THREADS = 8; // (how do you get this dynamically?) + int NUM_THREADS = blockDim.x * blockDim.y * blockDim.z; if (src2_data.s64 == NUM_THREADS-1) { // pick only one thread to load all of the data back from sstarr memory unsigned offset = 0; @@ -3809,7 +3809,6 @@ void sst_impl( const ptx_instruction *pI, ptx_thread_info *thread ) // fill the rest of the array with zeros (dst should always have a 0 in it) while (offset < NUM_THREADS) { mem->write(addr+(offset*4),size/8,&dst_data.s64,thread,pI); - mem->write(addr+((NUM_THREADS+offset)*4),size/8,&dst_data.s64,thread,pI); offset++; } diff --git a/src/cuda-sim/opcodes.def b/src/cuda-sim/opcodes.def index 0c0eda9..1af04ea 100644 --- a/src/cuda-sim/opcodes.def +++ b/src/cuda-sim/opcodes.def @@ -103,7 +103,7 @@ OP_DEF(SHR_OP,shr_impl,"shr",1,1) OP_DEF(SIN_OP,sin_impl,"sin",1,4) OP_DEF(SLCT_OP,slct_impl,"slct",1,1) OP_DEF(SQRT_OP,sqrt_impl,"sqrt",1,4) -OP_DEF(SST_OP,sst_impl,"sst",1,1) +OP_DEF(SST_OP,sst_impl,"sst",1,5) OP_DEF(SSY_OP,ssy_impl,"ssy",0,3) OP_DEF(ST_OP,st_impl,"st",0,5) OP_DEF(SUB_OP,sub_impl,"sub",1,1) |