From 21b18be3832b80b00d4890f041818bac431866e7 Mon Sep 17 00:00:00 2001 From: JRPan <25518778+JRPan@users.noreply.github.com> Date: Mon, 18 Aug 2025 23:09:47 -0700 Subject: running formatter (#126) * running formatter * consolidate CI runs * use cluster to run formatter * use cluster to run formatter * Add a CI-Success step --- src/abstract_hardware_model.cc | 8 +++--- src/abstract_hardware_model.h | 28 +++++++++++++++----- src/accelwattch/core.cc | 2 +- src/accelwattch/logic.cc | 2 +- src/accelwattch/processor.cc | 18 +++++++------ src/accelwattch/xmlParser.cc | 2 +- src/cuda-sim/cuda-sim.cc | 4 ++- src/cuda-sim/cuda-sim.h | 4 +-- src/cuda-sim/cuda_device_printf.cc | 2 +- src/cuda-sim/half.h | 27 ++++++++++++------- src/cuda-sim/ptx_sim.cc | 4 +-- src/cuda-sim/ptx_sim.h | 4 ++- src/gpgpu-sim/dram.cc | 8 ++++-- src/gpgpu-sim/gpu-sim.h | 2 +- src/gpgpu-sim/l2cache.h | 4 ++- src/gpgpu-sim/local_interconnect.cc | 52 ++++++++++++++++++------------------- src/gpgpu-sim/shader.cc | 2 +- src/gpgpu-sim/shader.h | 6 +++-- src/gpgpusim_entrypoint.cc | 20 ++++++++------ src/stream_manager.cc | 39 ++++++++++++++++++---------- src/stream_manager.h | 8 +++--- 21 files changed, 151 insertions(+), 95 deletions(-) (limited to 'src') 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 &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 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(abs >= 0x7C00); - for (; abs < 0x400; abs <<= 1, hi -= 0x100000); + for (; abs < 0x400; abs <<= 1, hi -= 0x100000) + ; hi += static_cast(abs) << 10; } uint64 bits = static_cast(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((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(); - - 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); - } + 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); } } - - 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& streams = g_stream_manager()->get_concurrent_streams(); + std::list &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::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& get_concurrent_streams() { return m_streams; }; + std::list &get_concurrent_streams() { return m_streams; }; private: void print_impl(FILE *fp); -- cgit v1.3