summaryrefslogtreecommitdiff
path: root/src/cuda-sim/instructions.cc
diff options
context:
space:
mode:
authorsspenst <[email protected]>2016-07-08 14:51:44 -0700
committersspenst <[email protected]>2016-07-08 14:51:44 -0700
commit6c1fb702e17b00fd7de72ac7dd4a31584d5978b9 (patch)
treed9ec9e297afabe0878c19a520f01e4ca123a75f1 /src/cuda-sim/instructions.cc
parent877cbd077ffaf112b68973fdb7db8f10505303ee (diff)
Made gridDim and blockDim global variables so that they can be accessed from sst_impl
Diffstat (limited to 'src/cuda-sim/instructions.cc')
-rw-r--r--src/cuda-sim/instructions.cc3
1 files changed, 1 insertions, 2 deletions
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++;
}