summaryrefslogtreecommitdiff
path: root/src
diff options
context:
space:
mode:
Diffstat (limited to 'src')
-rw-r--r--src/abstract_hardware_model.cc8
-rw-r--r--src/abstract_hardware_model.h28
-rw-r--r--src/accelwattch/core.cc2
-rw-r--r--src/accelwattch/logic.cc2
-rw-r--r--src/accelwattch/processor.cc18
-rw-r--r--src/accelwattch/xmlParser.cc2
-rw-r--r--src/cuda-sim/cuda-sim.cc4
-rw-r--r--src/cuda-sim/cuda-sim.h4
-rw-r--r--src/cuda-sim/cuda_device_printf.cc2
-rw-r--r--src/cuda-sim/half.h27
-rw-r--r--src/cuda-sim/ptx_sim.cc4
-rw-r--r--src/cuda-sim/ptx_sim.h4
-rw-r--r--src/gpgpu-sim/dram.cc8
-rw-r--r--src/gpgpu-sim/gpu-sim.h2
-rw-r--r--src/gpgpu-sim/l2cache.h4
-rw-r--r--src/gpgpu-sim/local_interconnect.cc50
-rw-r--r--src/gpgpu-sim/shader.cc2
-rw-r--r--src/gpgpu-sim/shader.h6
-rw-r--r--src/gpgpusim_entrypoint.cc20
-rw-r--r--src/stream_manager.cc39
-rw-r--r--src/stream_manager.h8
21 files changed, 150 insertions, 94 deletions
diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc
index 8743cc7..e8ddf95 100644
--- a/src/abstract_hardware_model.cc
+++ b/src/abstract_hardware_model.cc
@@ -216,20 +216,20 @@ new_addr_type line_size_based_tag_func(new_addr_type address,
return address & ~(line_size - 1);
}
-const char *mem_access_type_str(enum mem_access_type access_type){
+const char *mem_access_type_str(enum mem_access_type access_type) {
#define MA_TUP_BEGIN(X) static const char *access_type_str[] = {
#define MA_TUP(X) #X
#define MA_TUP_END(X) \
} \
;
- MEM_ACCESS_TYPE_TUP_DEF
+ MEM_ACCESS_TYPE_TUP_DEF
#undef MA_TUP_BEGIN
#undef MA_TUP
#undef MA_TUP_END
- assert(access_type < NUM_MEM_ACCESS_TYPE);
+ assert(access_type < NUM_MEM_ACCESS_TYPE);
-return access_type_str[access_type];
+ return access_type_str[access_type];
}
void warp_inst_t::clear_active(const active_mask_t &inactive) {
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index cddf523..06c3082 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -248,7 +248,9 @@ class kernel_info_t {
}
bool running() const { return m_num_cores_running > 0; }
bool done() const { return no_more_ctas_to_run() && !running(); }
- class function_info *entry() { return m_kernel_entry; }
+ class function_info *entry() {
+ return m_kernel_entry;
+ }
const class function_info *entry() const { return m_kernel_entry; }
size_t num_blocks() const {
@@ -298,7 +300,9 @@ class kernel_info_t {
std::list<class ptx_thread_info *> &active_threads() {
return m_active_threads;
}
- class memory_space *get_param_memory() { return m_param_mem; }
+ class memory_space *get_param_memory() {
+ return m_param_mem;
+ }
// The following functions access texture bindings present at the kernel's
// launch
@@ -605,9 +609,15 @@ class gpgpu_t {
void memcpy_from_gpu(void *dst, size_t src_start_addr, size_t count);
void memcpy_gpu_to_gpu(size_t dst, size_t src, size_t count);
- class memory_space *get_global_memory() { return m_global_mem; }
- class memory_space *get_tex_memory() { return m_tex_mem; }
- class memory_space *get_surf_memory() { return m_surf_mem; }
+ class memory_space *get_global_memory() {
+ return m_global_mem;
+ }
+ class memory_space *get_tex_memory() {
+ return m_tex_mem;
+ }
+ class memory_space *get_surf_memory() {
+ return m_surf_mem;
+ }
void gpgpu_ptx_sim_bindTextureToArray(const struct textureReference *texref,
const struct cudaArray *array);
@@ -1327,7 +1337,9 @@ class core_t {
virtual bool warp_waiting_at_barrier(unsigned warp_id) const = 0;
virtual void checkExecutionStatusAndUpdate(warp_inst_t &inst, unsigned t,
unsigned tid) = 0;
- class gpgpu_sim *get_gpu() { return m_gpu; }
+ class gpgpu_sim *get_gpu() {
+ return m_gpu;
+ }
void execute_warp_inst_t(warp_inst_t &inst, unsigned warpId = (unsigned)-1);
bool ptx_thread_done(unsigned hw_thread_id) const;
virtual void updateSIMTStack(unsigned warpId, warp_inst_t *inst);
@@ -1337,7 +1349,9 @@ class core_t {
void get_pdom_stack_top_info(unsigned warpId, unsigned *pc,
unsigned *rpc) const;
kernel_info_t *get_kernel_info() { return m_kernel; }
- class ptx_thread_info **get_thread_info() { return m_thread; }
+ class ptx_thread_info **get_thread_info() {
+ return m_thread;
+ }
unsigned get_warp_size() const { return m_warp_size; }
void and_reduction(unsigned ctaid, unsigned barid, bool value) {
reduction_storage[ctaid][barid] &= value;
diff --git a/src/accelwattch/core.cc b/src/accelwattch/core.cc
index fe11e05..556ed6c 100644
--- a/src/accelwattch/core.cc
+++ b/src/accelwattch/core.cc
@@ -2195,7 +2195,7 @@ EXECU::EXECU(ParseXML* XML_interface, int ithCore_,
bypass.area.set_area(bypass.area.get_area()
+fpTagBypass->area.get_area());
}*/
- } /* if (coredynp.core_ty==Inorder) */
+ } /* if (coredynp.core_ty==Inorder) */
else { // OOO
if (coredynp.scheu_ty == PhysicalRegFile) {
/* For physical register based OOO,
diff --git a/src/accelwattch/logic.cc b/src/accelwattch/logic.cc
index 3d3a1a4..7f40189 100644
--- a/src/accelwattch/logic.cc
+++ b/src/accelwattch/logic.cc
@@ -656,7 +656,7 @@ FunctionalUnit::FunctionalUnit(ParseXML *XML_interface, int ithCore_,
}
per_access_energy *= 0.5; // According to ARM data embedded processor has
// much lower per acc energy
- } /* if (XML->sys.Embedded) */
+ } /* if (XML->sys.Embedded) */
else {
if (fu_type == FPU) {
num_fu = coredynp.num_fpus;
diff --git a/src/accelwattch/processor.cc b/src/accelwattch/processor.cc
index a86b96d..d5c7cdd 100644
--- a/src/accelwattch/processor.cc
+++ b/src/accelwattch/processor.cc
@@ -664,22 +664,24 @@ void Processor::displayDeviceType(int device_type_, uint32_t indent) {
switch (device_type_) {
case 0:
- cout << indent_str
- << "Device Type= " << "ITRS high performance device type" << endl;
+ cout << indent_str << "Device Type= "
+ << "ITRS high performance device type" << endl;
break;
case 1:
- cout << indent_str
- << "Device Type= " << "ITRS low standby power device type" << endl;
+ cout << indent_str << "Device Type= "
+ << "ITRS low standby power device type" << endl;
break;
case 2:
- cout << indent_str
- << "Device Type= " << "ITRS low operating power device type" << endl;
+ cout << indent_str << "Device Type= "
+ << "ITRS low operating power device type" << endl;
break;
case 3:
- cout << indent_str << "Device Type= " << "LP-DRAM device type" << endl;
+ cout << indent_str << "Device Type= "
+ << "LP-DRAM device type" << endl;
break;
case 4:
- cout << indent_str << "Device Type= " << "COMM-DRAM device type" << endl;
+ cout << indent_str << "Device Type= "
+ << "COMM-DRAM device type" << endl;
break;
default: {
cout << indent_str << "Unknown Device Type" << endl;
diff --git a/src/accelwattch/xmlParser.cc b/src/accelwattch/xmlParser.cc
index 9f01ebe..084a50d 100644
--- a/src/accelwattch/xmlParser.cc
+++ b/src/accelwattch/xmlParser.cc
@@ -3241,7 +3241,7 @@ XMLSTR XMLParserBase64Tool::encode(unsigned char *inbuf, unsigned int inlen,
*(curr++) = base64EncodeTable[j >> 18];
*(curr++) = base64EncodeTable[(j >> 12) & 0x3f];
*(curr++) = base64EncodeTable[(j >> 6) & 0x3f];
- *(curr++) = base64EncodeTable[(j) & 0x3f];
+ *(curr++) = base64EncodeTable[(j)&0x3f];
if (formatted) {
if (!k) {
*(curr++) = _CXML('\n');
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 69d1eb7..14ae504 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -2782,7 +2782,9 @@ void print_ptxinfo() {
}
}
-struct gpgpu_ptx_sim_info get_ptxinfo() { return g_ptxinfo; }
+struct gpgpu_ptx_sim_info get_ptxinfo() {
+ return g_ptxinfo;
+}
std::map<unsigned, const char *> get_duplicate() { return g_duplicate; }
diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h
index b1caf0c..21e1ca0 100644
--- a/src/cuda-sim/cuda-sim.h
+++ b/src/cuda-sim/cuda-sim.h
@@ -101,8 +101,8 @@ class functionalCoreSim : public core_t {
bool *m_warpAtBarrier;
};
-#define RECONVERGE_RETURN_PC ((address_type) - 2)
-#define NO_BRANCH_DIVERGENCE ((address_type) - 1)
+#define RECONVERGE_RETURN_PC ((address_type)-2)
+#define NO_BRANCH_DIVERGENCE ((address_type)-1)
address_type get_return_pc(void *thd);
const char *get_ptxinfo_kname();
void print_ptxinfo();
diff --git a/src/cuda-sim/cuda_device_printf.cc b/src/cuda-sim/cuda_device_printf.cc
index 69ebc63..ef780e9 100644
--- a/src/cuda-sim/cuda_device_printf.cc
+++ b/src/cuda-sim/cuda_device_printf.cc
@@ -65,7 +65,7 @@ void my_cuda_printf(const char *fmtstr, const char *arg_list) {
fprintf(fp, buf, *((unsigned *)ptr));
arg_offset += 4;
} else if (c == 'f') {
- if (acc%2) {
+ if (acc % 2) {
arg_offset += 4;
ptr = (void *)&arg_list[arg_offset];
}
diff --git a/src/cuda-sim/half.h b/src/cuda-sim/half.h
index 67e607b..fab1a22 100644
--- a/src/cuda-sim/half.h
+++ b/src/cuda-sim/half.h
@@ -846,8 +846,10 @@ uint16 int2half_impl(T value) {
bits |= 0x7BFF + (R != std::round_toward_zero);
} else if (value) {
unsigned int m = value, exp = 24;
- for (; m < 0x400; m <<= 1, --exp);
- for (; m > 0x7FF; m >>= 1, ++exp);
+ for (; m < 0x400; m <<= 1, --exp)
+ ;
+ for (; m > 0x7FF; m >>= 1, ++exp)
+ ;
bits |= (exp << 10) + m;
if (exp > 24) {
if (R == std::round_to_nearest)
@@ -1272,7 +1274,8 @@ inline double half2float_impl(uint16 value, double, true_type) {
int abs = value & 0x7FFF;
if (abs) {
hi |= 0x3F000000 << static_cast<unsigned>(abs >= 0x7C00);
- for (; abs < 0x400; abs <<= 1, hi -= 0x100000);
+ for (; abs < 0x400; abs <<= 1, hi -= 0x100000)
+ ;
hi += static_cast<uint32>(abs) << 10;
}
uint64 bits = static_cast<uint64>(hi) << 32;
@@ -2113,7 +2116,8 @@ struct functions {
static half frexp(half arg, int *exp) {
int m = arg.data_ & 0x7FFF, e = -14;
if (m >= 0x7C00 || !m) return *exp = 0, arg;
- for (; m < 0x400; m <<= 1, --e);
+ for (; m < 0x400; m <<= 1, --e)
+ ;
return *exp = e + (m >> 10),
half(binary, (arg.data_ & 0x8000) | 0x3800 | (m & 0x3FF));
}
@@ -2131,7 +2135,8 @@ struct functions {
unsigned int mask = (1 << (25 - e)) - 1, m = arg.data_ & mask;
iptr->data_ = arg.data_ & ~mask;
if (!m) return half(binary, arg.data_ & 0x8000);
- for (; m < 0x400; m <<= 1, --e);
+ for (; m < 0x400; m <<= 1, --e)
+ ;
return half(binary, static_cast<uint16>((arg.data_ & 0x8000) | (e << 10) |
(m & 0x3FF)));
}
@@ -2143,7 +2148,8 @@ struct functions {
static half scalbln(half arg, long exp) {
unsigned int m = arg.data_ & 0x7FFF;
if (m >= 0x7C00 || !m) return arg;
- for (; m < 0x400; m <<= 1, --exp);
+ for (; m < 0x400; m <<= 1, --exp)
+ ;
exp += m >> 10;
uint16 value = arg.data_ & 0x8000;
if (exp > 30) {
@@ -2185,7 +2191,8 @@ struct functions {
if (abs < 0x7C00) {
int exp = (abs >> 10) - 15;
if (abs < 0x400)
- for (; abs < 0x200; abs <<= 1, --exp);
+ for (; abs < 0x200; abs <<= 1, --exp)
+ ;
return exp;
}
if (abs > 0x7C00) return FP_ILOGBNAN;
@@ -2201,11 +2208,13 @@ struct functions {
if (abs < 0x7C00) {
int exp = (abs >> 10) - 15;
if (abs < 0x400)
- for (; abs < 0x200; abs <<= 1, --exp);
+ for (; abs < 0x200; abs <<= 1, --exp)
+ ;
uint16 bits = (exp < 0) << 15;
if (exp) {
unsigned int m = std::abs(exp) << 6, e = 18;
- for (; m < 0x400; m <<= 1, --e);
+ for (; m < 0x400; m <<= 1, --e)
+ ;
bits |= (e << 10) + m;
}
return half(binary, bits);
diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc
index 9a22d02..d0de1ca 100644
--- a/src/cuda-sim/ptx_sim.cc
+++ b/src/cuda-sim/ptx_sim.cc
@@ -203,8 +203,8 @@ unsigned ptx_thread_info::get_builtin(int builtin_id, unsigned dim_mod) {
return (unsigned)(m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle);
case CLOCK64_REG:
// Change return value to unsigned long long?
- // Currently returns 32-bit unsigned, which may cause truncation for large values.
- // GPGPUSim clock is 4 times slower - multiply by 4
+ // Currently returns 32-bit unsigned, which may cause truncation for large
+ // values. GPGPUSim clock is 4 times slower - multiply by 4
return (m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle) * 4;
case HALFCLOCK_ID:
// GPGPUSim clock is 4 times slower - multiply by 4
diff --git a/src/cuda-sim/ptx_sim.h b/src/cuda-sim/ptx_sim.h
index 7128f8e..8eec922 100644
--- a/src/cuda-sim/ptx_sim.h
+++ b/src/cuda-sim/ptx_sim.h
@@ -340,7 +340,9 @@ class ptx_thread_info {
dim3 get_ctaid() const { return m_ctaid; }
dim3 get_tid() const { return m_tid; }
dim3 get_ntid() const { return m_ntid; }
- class gpgpu_sim *get_gpu() { return (gpgpu_sim *)m_gpu; }
+ class gpgpu_sim *get_gpu() {
+ return (gpgpu_sim *)m_gpu;
+ }
unsigned get_hw_tid() const { return m_hw_tid; }
unsigned get_hw_ctaid() const { return m_hw_ctaid; }
unsigned get_hw_wid() const { return m_hw_wid; }
diff --git a/src/gpgpu-sim/dram.cc b/src/gpgpu-sim/dram.cc
index 1e1ad3d..80e20d7 100644
--- a/src/gpgpu-sim/dram.cc
+++ b/src/gpgpu-sim/dram.cc
@@ -684,9 +684,13 @@ bool dram_t::issue_row_command(int j) {
}
// if mrq is being serviced by dram, gets popped after CL latency fulfilled
-class mem_fetch *dram_t::return_queue_pop() { return returnq->pop(); }
+class mem_fetch *dram_t::return_queue_pop() {
+ return returnq->pop();
+}
-class mem_fetch *dram_t::return_queue_top() { return returnq->top(); }
+class mem_fetch *dram_t::return_queue_top() {
+ return returnq->top();
+}
void dram_t::print(FILE *simFile) const {
unsigned i;
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index 68bdca7..98e52f4 100644
--- a/src/gpgpu-sim/gpu-sim.h
+++ b/src/gpgpu-sim/gpu-sim.h
@@ -879,7 +879,7 @@ class sst_gpgpu_sim : public gpgpu_sim {
* @param dst_start_addr
* @param count
*/
- void perf_memcpy_to_gpu(size_t dst_start_addr, size_t count) {};
+ void perf_memcpy_to_gpu(size_t dst_start_addr, size_t count){};
/**
* @brief Check if the SST config matches up with the
diff --git a/src/gpgpu-sim/l2cache.h b/src/gpgpu-sim/l2cache.h
index 9d164d7..65c9c38 100644
--- a/src/gpgpu-sim/l2cache.h
+++ b/src/gpgpu-sim/l2cache.h
@@ -106,7 +106,9 @@ class memory_partition_unit {
unsigned get_mpid() const { return m_id; }
- class gpgpu_sim *get_mgpu() const { return m_gpu; }
+ class gpgpu_sim *get_mgpu() const {
+ return m_gpu;
+ }
private:
unsigned m_id;
diff --git a/src/gpgpu-sim/local_interconnect.cc b/src/gpgpu-sim/local_interconnect.cc
index e4cdaa5..ce4415c 100644
--- a/src/gpgpu-sim/local_interconnect.cc
+++ b/src/gpgpu-sim/local_interconnect.cc
@@ -190,7 +190,7 @@ void xbar_router::iSLIP_Advance() {
input_nodes.insert(i);
unsigned out_node = in_buffers[i].front().output_deviceID;
- if(destination_set.find(out_node) != destination_set.end()) {
+ if (destination_set.find(out_node) != destination_set.end()) {
conflict_sub++;
}
destination_set.insert(out_node);
@@ -207,40 +207,40 @@ void xbar_router::iSLIP_Advance() {
for (auto dest : destination_set) {
if (Has_Buffer_Out(dest, 1)) {
unsigned start_node = next_node[dest];
- auto it = std::upper_bound(input_nodes.begin(), input_nodes.end(),
- start_node);
+ auto it =
+ std::upper_bound(input_nodes.begin(), input_nodes.end(), start_node);
for (unsigned j = 0; j < input_nodes.size(); j++, it++) {
if (it == input_nodes.end()) {
it = input_nodes.begin();
}
unsigned node_id = *it;
assert(!in_buffers[node_id].empty());
- Packet _packet = in_buffers[node_id].front();
- if (_packet.output_deviceID == dest) {
- out_buffers[_packet.output_deviceID].push(_packet);
- in_buffers[node_id].pop();
- input_nodes.erase(node_id); //can only be used once
- if (verbose)
- printf("%d : cycle %llu : send req from %d to %d\n", m_id, cycles,
- node_id, dest - _n_shader);
- if (grant_cycles_count == 1)
- next_node[dest] = (++node_id % total_nodes);
- if (verbose) {
- for (unsigned k = j + 1; k < total_nodes; ++k) {
- unsigned node_id2 = (k + next_node[dest]) % total_nodes;
- if (!in_buffers[node_id2].empty()) {
- Packet _packet2 = in_buffers[node_id2].front();
+ Packet _packet = in_buffers[node_id].front();
+ if (_packet.output_deviceID == dest) {
+ out_buffers[_packet.output_deviceID].push(_packet);
+ in_buffers[node_id].pop();
+ input_nodes.erase(node_id); // can only be used once
+ if (verbose)
+ printf("%d : cycle %llu : send req from %d to %d\n", m_id, cycles,
+ node_id, dest - _n_shader);
+ if (grant_cycles_count == 1)
+ next_node[dest] = (++node_id % total_nodes);
+ if (verbose) {
+ for (unsigned k = j + 1; k < total_nodes; ++k) {
+ unsigned node_id2 = (k + next_node[dest]) % total_nodes;
+ if (!in_buffers[node_id2].empty()) {
+ Packet _packet2 = in_buffers[node_id2].front();
- if (_packet2.output_deviceID == dest)
- printf("%d : cycle %llu : cannot send req from %d to %d\n",
- m_id, cycles, node_id2, dest - _n_shader);
- }
+ if (_packet2.output_deviceID == dest)
+ printf("%d : cycle %llu : cannot send req from %d to %d\n",
+ m_id, cycles, node_id2, dest - _n_shader);
}
}
-
- reqs++;
- break;
}
+
+ reqs++;
+ break;
+ }
}
} else {
out_buffer_full++;
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index 1ea011e..730cb6d 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -891,7 +891,7 @@ void shader_core_ctx::decode() {
// decode 1 or 2 instructions and place them into ibuffer
address_type pc = m_inst_fetch_buffer.m_pc;
const warp_inst_t *pI1 = get_next_inst(m_inst_fetch_buffer.m_warp_id, pc);
- if (pI1) {
+ if (pI1) {
m_warp[m_inst_fetch_buffer.m_warp_id]->ibuffer_fill(0, pI1);
m_warp[m_inst_fetch_buffer.m_warp_id]->inc_inst_in_pipeline();
m_stats->m_num_decoded_insn[m_sid]++;
diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h
index b90ce2f..5922d24 100644
--- a/src/gpgpu-sim/shader.h
+++ b/src/gpgpu-sim/shader.h
@@ -189,7 +189,7 @@ class shd_warp_t {
m_active_threads.reset(lane);
n_completed++;
}
- bool test_active(unsigned lane) {return m_active_threads.test(lane);}
+ bool test_active(unsigned lane) { return m_active_threads.test(lane); }
void set_last_fetch(unsigned long long sim_cycle) {
m_last_fetch = sim_cycle;
@@ -274,7 +274,9 @@ class shd_warp_t {
unsigned get_dynamic_warp_id() const { return m_dynamic_warp_id; }
unsigned get_warp_id() const { return m_warp_id; }
- class shader_core_ctx *get_shader() { return m_shader; }
+ class shader_core_ctx *get_shader() {
+ return m_shader;
+ }
private:
static const unsigned IBUFFER_SIZE = 2;
diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc
index be49229..4462dc3 100644
--- a/src/gpgpusim_entrypoint.cc
+++ b/src/gpgpusim_entrypoint.cc
@@ -58,7 +58,8 @@ class stream_manager *g_stream_manager() {
extern void SST_callback_cudaThreadSynchronize_done();
extern void SST_callback_cudaStreamSynchronize_done(cudaStream_t stream);
__attribute__((weak)) void SST_callback_cudaThreadSynchronize_done() {}
-__attribute__((weak)) void SST_callback_cudaStreamSynchronize_done(cudaStream_t stream) {}
+__attribute__((weak)) void SST_callback_cudaStreamSynchronize_done(
+ cudaStream_t stream) {}
void *gpgpu_sim_thread_sequential(void *ctx_ptr) {
gpgpu_context *ctx = (gpgpu_context *)ctx_ptr;
@@ -102,7 +103,8 @@ void *gpgpu_sim_thread_concurrent(void *ctx_ptr) {
fflush(stdout);
}
while (ctx->the_gpgpusim->g_stream_manager->empty_protected() &&
- !ctx->the_gpgpusim->g_sim_done);
+ !ctx->the_gpgpusim->g_sim_done)
+ ;
if (g_debug_execution >= 3) {
printf("GPGPU-Sim: ** START simulation thread (detected work) **\n");
ctx->the_gpgpusim->g_stream_manager->print(stdout);
@@ -191,7 +193,8 @@ bool SST_Cycle() {
// Check if Synchronize is done when SST previously requested
// cudaThreadSynchronize
if (GPGPU_Context()->requested_synchronize &&
- ((g_stream_manager()->empty_protected() && !GPGPUsim_ctx_ptr()->g_sim_active) ||
+ ((g_stream_manager()->empty_protected() &&
+ !GPGPUsim_ctx_ptr()->g_sim_active) ||
GPGPUsim_ctx_ptr()->g_sim_done)) {
SST_callback_cudaThreadSynchronize_done();
GPGPU_Context()->requested_synchronize = false;
@@ -199,19 +202,20 @@ bool SST_Cycle() {
// Polling to check for each stream if it is marked for requested with sync
if (g_stream_manager()->get_stream_zero()->requested_synchronize() &&
- ((g_stream_manager()->empty_protected() && !GPGPUsim_ctx_ptr()->g_sim_active) ||
- GPGPUsim_ctx_ptr()->g_sim_done)) {
+ ((g_stream_manager()->empty_protected() &&
+ !GPGPUsim_ctx_ptr()->g_sim_active) ||
+ GPGPUsim_ctx_ptr()->g_sim_done)) {
SST_callback_cudaStreamSynchronize_done(0);
g_stream_manager()->get_stream_zero()->reset_request_synchronize();
}
// Iterate through each stream to check if SST is waiting on
// it and it does not have any operation
- std::list<CUstream_st *>& streams = g_stream_manager()->get_concurrent_streams();
+ std::list<CUstream_st *> &streams =
+ g_stream_manager()->get_concurrent_streams();
for (auto it = streams.begin(); it != streams.end(); it++) {
CUstream_st *stream = *it;
- if (stream->requested_synchronize() &&
- stream->empty()) {
+ if (stream->requested_synchronize() && stream->empty()) {
// This stream is ready
SST_callback_cudaStreamSynchronize_done(stream);
stream->reset_request_synchronize();
diff --git a/src/stream_manager.cc b/src/stream_manager.cc
index d43964a..ab5b74f 100644
--- a/src/stream_manager.cc
+++ b/src/stream_manager.cc
@@ -34,21 +34,30 @@
unsigned CUstream_st::sm_next_stream_uid = 0;
-// SST memcpy callbacks, called after a stream operation is done via record_next_done()
-extern void SST_callback_memcpy_H2D_done(uint64_t dst, uint64_t src, size_t count, cudaStream_t stream);
-extern void SST_callback_memcpy_D2H_done(uint64_t dst, uint64_t src, size_t count, cudaStream_t stream);
+// SST memcpy callbacks, called after a stream operation is done via
+// record_next_done()
+extern void SST_callback_memcpy_H2D_done(uint64_t dst, uint64_t src,
+ size_t count, cudaStream_t stream);
+extern void SST_callback_memcpy_D2H_done(uint64_t dst, uint64_t src,
+ size_t count, cudaStream_t stream);
extern void SST_callback_memcpy_to_symbol_done();
extern void SST_callback_memcpy_from_symbol_done();
extern void SST_callback_cudaEventSynchronize_done(cudaEvent_t event);
extern void SST_callback_kernel_done(cudaStream_t stream);
-__attribute__((weak)) void SST_callback_memcpy_H2D_done(uint64_t dst, uint64_t src, size_t count, cudaStream_t stream) {}
-__attribute__((weak)) void SST_callback_memcpy_D2H_done(uint64_t dst, uint64_t src, size_t count, cudaStream_t stream) {}
+__attribute__((weak)) void SST_callback_memcpy_H2D_done(uint64_t dst,
+ uint64_t src,
+ size_t count,
+ cudaStream_t stream) {}
+__attribute__((weak)) void SST_callback_memcpy_D2H_done(uint64_t dst,
+ uint64_t src,
+ size_t count,
+ cudaStream_t stream) {}
__attribute__((weak)) void SST_callback_memcpy_to_symbol_done() {}
__attribute__((weak)) void SST_callback_memcpy_from_symbol_done() {}
-__attribute__((weak)) void SST_callback_cudaEventSynchronize_done(cudaEvent_t event);
+__attribute__((weak)) void SST_callback_cudaEventSynchronize_done(
+ cudaEvent_t event);
__attribute__((weak)) void SST_callback_kernel_done(cudaStream_t stream);
-
CUstream_st::CUstream_st() {
m_pending = false;
m_uid = sm_next_stream_uid++;
@@ -79,9 +88,7 @@ void CUstream_st::synchronize() {
} while (!done);
}
-bool CUstream_st::synchronize_check() {
- return m_operations.empty();
-}
+bool CUstream_st::synchronize_check() { return m_operations.empty(); }
void CUstream_st::push(const stream_operation &op) {
// called by host thread
@@ -142,14 +149,19 @@ bool stream_operation::do_operation(gpgpu_sim *gpu) {
gpu->memcpy_to_gpu(m_device_address_dst, m_host_address_src, m_cnt);
m_stream->record_next_done();
if (gpu->is_SST_mode()) {
- SST_callback_memcpy_H2D_done((uint64_t) m_device_address_dst, (uint64_t) m_host_address_src, m_cnt, m_stream->is_stream_zero_stream() ? 0 : m_stream);
+ SST_callback_memcpy_H2D_done(
+ (uint64_t)m_device_address_dst, (uint64_t)m_host_address_src, m_cnt,
+ m_stream->is_stream_zero_stream() ? 0 : m_stream);
}
break;
case stream_memcpy_device_to_host:
if (g_debug_execution >= 3) printf("memcpy device-to-host\n");
gpu->memcpy_from_gpu(m_host_address_dst, m_device_address_src, m_cnt);
m_stream->record_next_done();
- if (gpu->is_SST_mode()) SST_callback_memcpy_D2H_done((uint64_t) m_host_address_dst, (uint64_t) m_device_address_src, m_cnt, m_stream->is_stream_zero_stream() ? 0 : m_stream);
+ if (gpu->is_SST_mode())
+ SST_callback_memcpy_D2H_done(
+ (uint64_t)m_host_address_dst, (uint64_t)m_device_address_src, m_cnt,
+ m_stream->is_stream_zero_stream() ? 0 : m_stream);
break;
case stream_memcpy_device_to_device:
if (g_debug_execution >= 3) printf("memcpy device-to-device\n");
@@ -419,7 +431,8 @@ void stream_manager::add_stream(struct CUstream_st *stream) {
void stream_manager::destroy_stream(CUstream_st *stream) {
// called by host thread
pthread_mutex_lock(&m_lock);
- while (!stream->empty());
+ while (!stream->empty())
+ ;
std::list<CUstream_st *>::iterator s;
for (s = m_streams.begin(); s != m_streams.end(); s++) {
if (*s == stream) {
diff --git a/src/stream_manager.h b/src/stream_manager.h
index 55cfb8d..c6d2fc3 100644
--- a/src/stream_manager.h
+++ b/src/stream_manager.h
@@ -82,7 +82,7 @@ struct CUevent_st {
unsigned int m_issued;
time_t m_wallclock;
double m_gpu_tot_sim_cycle;
-
+
// SST related
bool m_requested_synchronize = false;
@@ -206,7 +206,9 @@ class stream_operation {
kernel_info_t *get_kernel() { return m_kernel; }
bool do_operation(gpgpu_sim *gpu);
void print(FILE *fp) const;
- struct CUstream_st *get_stream() { return m_stream; }
+ struct CUstream_st *get_stream() {
+ return m_stream;
+ }
void set_stream(CUstream_st *stream) { m_stream = stream; }
private:
@@ -284,7 +286,7 @@ class stream_manager {
unsigned size() { return m_streams.size(); };
bool is_blocking() { return m_cuda_launch_blocking; };
CUstream_st *get_stream_zero() { return &m_stream_zero; };
- std::list<CUstream_st *>& get_concurrent_streams() { return m_streams; };
+ std::list<CUstream_st *> &get_concurrent_streams() { return m_streams; };
private:
void print_impl(FILE *fp);
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 22:37:53 +0000