From 68a91076b2aab8f60bae551d6df6b3a8aa411463 Mon Sep 17 00:00:00 2001 From: Ahmad Alawneh Date: Tue, 20 Jun 2023 14:42:00 -0400 Subject: fix most c warnings --- src/gpgpu-sim/mem_fetch.cc | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) (limited to 'src/gpgpu-sim/mem_fetch.cc') diff --git a/src/gpgpu-sim/mem_fetch.cc b/src/gpgpu-sim/mem_fetch.cc index 456d891..0d86046 100644 --- a/src/gpgpu-sim/mem_fetch.cc +++ b/src/gpgpu-sim/mem_fetch.cc @@ -84,10 +84,10 @@ mem_fetch::~mem_fetch() { m_status = MEM_FETCH_DELETED; } #undef MF_TUP_END void mem_fetch::print(FILE *fp, bool print_inst) const { - if (this == NULL) { - fprintf(fp, " \n"); - return; - } + // if (this == NULL) { // doenst make sense! + // fprintf(fp, " \n"); + // return; + // } fprintf(fp, " mf: uid=%6u, sid%02u:w%02u, part=%u, ", m_request_uid, m_sid, m_wid, m_raw_addr.chip); m_access.print(fp); -- cgit v1.3 From 38b4df5653ecbd9907a3d39b125640cd4fb7d012 Mon Sep 17 00:00:00 2001 From: JRPan <25518778+JRPan@users.noreply.github.com> Date: Tue, 20 Aug 2024 20:43:32 -0400 Subject: Stream stats (#71) * Temp commit for Justin and Cassie to sync on code changes for adding per-stream status. * Resolved compile errors. * Removed redundant parameter * Passed cuda_stream_id from accelsim to gpgpusim * Cleaned up unused changes * Changed vector to map, having operator problems. * StreamID defaults to zero * Implemented streams to inc_stats and so on * Fixed TOTAL_ACCESS counts * Implemented GLOBAL_TIMER. * Fixed m_shader->get_kernel SEGFAULT issue in shader.cc. * Use warp_init to track streamID instead of issue_warp * Removed temp debug print * Modified cache_stats to only print data from latest finished stream Added optional arg to cache_stats::print_stats, cache_stats::print_fail_stats and their upstream functions. When streamID is specified, print stats from that stream. When not specified, print all stats. NOTE: current implementation depending on streamid never equals -1 * Removed default arg values of streamID * modified constructor of mem_fetch to pass in streamID * changed get_streamid to get_streamID * Added TODO to gpgpusim_entrypoint.cc and power_stat.cc * Only collect power stats when enabled * print last finished stream in PTX mode using last_streamID * take out additional printf * Add a field to baseline cache to indicate cache level * save gpu object in cache * Print stream ID only once per kernel * rm test print * use -1 for default stream id * cleanup debug prints * remove GLOABL_TIMER * Automated clang-format * Should be correct to print everything in power model * addressing concerns & errors * Automated clang-format * add m_stats_pw in operator+ * Automated Format --------- Co-authored-by: Justin Qiao Co-authored-by: Justin Qiao <71228724+ShichenQiao@users.noreply.github.com> Co-authored-by: Tim Rogers Co-authored-by: JRPan Co-authored-by: purdue-jenkins --- libcuda/cuda_runtime_api.cc | 4 +- src/abstract_hardware_model.cc | 8 +- src/abstract_hardware_model.h | 21 +- src/gpgpu-sim/gpu-cache.cc | 533 +++++++++++++++++++++++++++++------------ src/gpgpu-sim/gpu-cache.h | 77 ++++-- src/gpgpu-sim/gpu-sim.cc | 72 ++++-- src/gpgpu-sim/gpu-sim.h | 18 +- src/gpgpu-sim/l2cache.cc | 29 +-- src/gpgpu-sim/l2cache.h | 7 +- src/gpgpu-sim/mem_fetch.cc | 9 +- src/gpgpu-sim/mem_fetch.h | 10 +- src/gpgpu-sim/power_stat.cc | 6 +- src/gpgpu-sim/shader.cc | 42 ++-- src/gpgpu-sim/shader.h | 19 +- src/gpgpusim_entrypoint.cc | 6 +- src/stream_manager.cc | 12 +- 16 files changed, 613 insertions(+), 260 deletions(-) (limited to 'src/gpgpu-sim/mem_fetch.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index b64c3d9..b540ffd 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1421,7 +1421,9 @@ cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlagsInternal( dim3 gridDim(context->get_device()->get_gpgpu()->max_cta_per_core() * context->get_device()->get_gpgpu()->get_config().num_shader()); dim3 blockDim(blockSize); - kernel_info_t result(gridDim, blockDim, entry); + // because this fuction is only checking for resource requirements, we do + // not care which stream this kernel runs at, just picked -1 + kernel_info_t result(gridDim, blockDim, entry, -1); // if(entry == NULL){ // *numBlocks = 1; // return g_last_cudaError = cudaErrorUnknown; diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index fd056c6..e8ddf95 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -49,12 +49,14 @@ void mem_access_t::init(gpgpu_context *ctx) { m_addr = 0; m_req_size = 0; } + void warp_inst_t::issue(const active_mask_t &mask, unsigned warp_id, unsigned long long cycle, int dynamic_warp_id, - int sch_id) { + int sch_id, unsigned long long streamID) { m_warp_active_mask = mask; m_warp_issued_mask = mask; m_uid = ++(m_config->gpgpu_ctx->warp_inst_sm_next_uid); + m_streamID = streamID; m_warp_id = warp_id; m_dynamic_warp_id = dynamic_warp_id; issue_cycle = cycle; @@ -755,7 +757,8 @@ void warp_inst_t::completed(unsigned long long cycle) const { } kernel_info_t::kernel_info_t(dim3 gridDim, dim3 blockDim, - class function_info *entry) { + class function_info *entry, + unsigned long long streamID) { m_kernel_entry = entry; m_grid_dim = gridDim; m_block_dim = blockDim; @@ -765,6 +768,7 @@ kernel_info_t::kernel_info_t(dim3 gridDim, dim3 blockDim, m_next_tid = m_next_cta; m_num_cores_running = 0; m_uid = (entry->gpgpu_ctx->kernel_info_m_next_uid)++; + m_streamID = streamID; m_param_mem = new memory_space_impl<8192>("param", 64 * 1024); // Jin: parent and child kernel management for CDP diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index e5f3b78..98a4039 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -233,7 +233,8 @@ class kernel_info_t { // m_num_cores_running=0; // m_param_mem=NULL; // } - kernel_info_t(dim3 gridDim, dim3 blockDim, class function_info *entry); + kernel_info_t(dim3 gridDim, dim3 blockDim, class function_info *entry, + unsigned long long streamID); kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, @@ -292,6 +293,7 @@ class kernel_info_t { m_next_tid.x < m_block_dim.x; } unsigned get_uid() const { return m_uid; } + unsigned long long get_streamID() const { return m_streamID; } std::string get_name() const { return name(); } std::string name() const; @@ -325,7 +327,8 @@ class kernel_info_t { class function_info *m_kernel_entry; - unsigned m_uid; + unsigned m_uid; // Kernel ID + unsigned long long m_streamID; // These maps contain the snapshot of the texture mappings at kernel launch std::map m_NameToCudaArray; @@ -900,8 +903,8 @@ class mem_fetch_interface { class mem_fetch_allocator { public: virtual mem_fetch *alloc(new_addr_type addr, mem_access_type type, - unsigned size, bool wr, - unsigned long long cycle) const = 0; + unsigned size, bool wr, unsigned long long cycle, + unsigned long long streamID) const = 0; virtual mem_fetch *alloc(const class warp_inst_t &inst, const mem_access_t &access, unsigned long long cycle) const = 0; @@ -911,7 +914,8 @@ class mem_fetch_allocator { const mem_access_sector_mask_t §or_mask, unsigned size, bool wr, unsigned long long cycle, unsigned wid, unsigned sid, unsigned tpc, - mem_fetch *original_mf) const = 0; + mem_fetch *original_mf, + unsigned long long streamID) const = 0; }; // the maximum number of destination, source, or address uarch operands in a @@ -1059,6 +1063,7 @@ class warp_inst_t : public inst_t { // constructors warp_inst_t() { m_uid = 0; + m_streamID = (unsigned long long)-1; m_empty = true; m_config = NULL; @@ -1071,6 +1076,7 @@ class warp_inst_t : public inst_t { } warp_inst_t(const core_config *config) { m_uid = 0; + m_streamID = (unsigned long long)-1; assert(config->warp_size <= MAX_WARP_SIZE); m_config = config; m_empty = true; @@ -1098,7 +1104,8 @@ class warp_inst_t : public inst_t { void clear() { m_empty = true; } void issue(const active_mask_t &mask, unsigned warp_id, - unsigned long long cycle, int dynamic_warp_id, int sch_id); + unsigned long long cycle, int dynamic_warp_id, int sch_id, + unsigned long long streamID); const active_mask_t &get_active_mask() const { return m_warp_active_mask; } void completed(unsigned long long cycle) @@ -1226,11 +1233,13 @@ class warp_inst_t : public inst_t { void print(FILE *fout) const; unsigned get_uid() const { return m_uid; } + unsigned long long get_streamID() const { return m_streamID; } unsigned get_schd_id() const { return m_scheduler_id; } active_mask_t get_warp_active_mask() const { return m_warp_active_mask; } protected: unsigned m_uid; + unsigned long long m_streamID; bool m_empty; bool m_cache_hit; unsigned long long issue_cycle; diff --git a/src/gpgpu-sim/gpu-cache.cc b/src/gpgpu-sim/gpu-cache.cc index 32cc56b..cd3c880 100644 --- a/src/gpgpu-sim/gpu-cache.cc +++ b/src/gpgpu-sim/gpu-cache.cc @@ -634,14 +634,6 @@ void mshr_table::display(FILE *fp) const { /***************************************************************** Caches * *****************************************************************/ cache_stats::cache_stats() { - m_stats.resize(NUM_MEM_ACCESS_TYPE); - m_stats_pw.resize(NUM_MEM_ACCESS_TYPE); - m_fail_stats.resize(NUM_MEM_ACCESS_TYPE); - for (unsigned i = 0; i < NUM_MEM_ACCESS_TYPE; ++i) { - m_stats[i].resize(NUM_CACHE_REQUEST_STATUS, 0); - m_stats_pw[i].resize(NUM_CACHE_REQUEST_STATUS, 0); - m_fail_stats[i].resize(NUM_CACHE_RESERVATION_FAIL_STATUS, 0); - } m_cache_port_available_cycles = 0; m_cache_data_port_busy_cycles = 0; m_cache_fill_port_busy_cycles = 0; @@ -651,11 +643,10 @@ void cache_stats::clear() { /// /// Zero out all current cache statistics /// - for (unsigned i = 0; i < NUM_MEM_ACCESS_TYPE; ++i) { - std::fill(m_stats[i].begin(), m_stats[i].end(), 0); - std::fill(m_stats_pw[i].begin(), m_stats_pw[i].end(), 0); - std::fill(m_fail_stats[i].begin(), m_fail_stats[i].end(), 0); - } + m_stats.clear(); + m_stats_pw.clear(); + m_fail_stats.clear(); + m_cache_port_available_cycles = 0; m_cache_data_port_busy_cycles = 0; m_cache_fill_port_busy_cycles = 0; @@ -665,35 +656,67 @@ void cache_stats::clear_pw() { /// /// Zero out per-window cache statistics /// - for (unsigned i = 0; i < NUM_MEM_ACCESS_TYPE; ++i) { - std::fill(m_stats_pw[i].begin(), m_stats_pw[i].end(), 0); - } + m_stats_pw.clear(); } -void cache_stats::inc_stats(int access_type, int access_outcome) { +void cache_stats::inc_stats(int access_type, int access_outcome, + unsigned long long streamID) { /// /// Increment the stat corresponding to (access_type, access_outcome) by 1. /// if (!check_valid(access_type, access_outcome)) assert(0 && "Unknown cache access type or access outcome"); - m_stats[access_type][access_outcome]++; + if (m_stats.find(streamID) == m_stats.end()) { + std::vector> new_val; + new_val.resize(NUM_MEM_ACCESS_TYPE); + for (unsigned j = 0; j < NUM_MEM_ACCESS_TYPE; ++j) { + new_val[j].resize(NUM_CACHE_REQUEST_STATUS, 0); + } + m_stats.insert(std::pair>>( + streamID, new_val)); + } + m_stats.at(streamID)[access_type][access_outcome]++; } -void cache_stats::inc_stats_pw(int access_type, int access_outcome) { +void cache_stats::inc_stats_pw(int access_type, int access_outcome, + unsigned long long streamID) { /// /// Increment the corresponding per-window cache stat /// if (!check_valid(access_type, access_outcome)) assert(0 && "Unknown cache access type or access outcome"); - m_stats_pw[access_type][access_outcome]++; + + if (m_stats_pw.find(streamID) == m_stats_pw.end()) { + std::vector> new_val; + new_val.resize(NUM_MEM_ACCESS_TYPE); + for (unsigned j = 0; j < NUM_MEM_ACCESS_TYPE; ++j) { + new_val[j].resize(NUM_CACHE_REQUEST_STATUS, 0); + } + m_stats_pw.insert(std::pair>>( + streamID, new_val)); + } + m_stats_pw.at(streamID)[access_type][access_outcome]++; } -void cache_stats::inc_fail_stats(int access_type, int fail_outcome) { +void cache_stats::inc_fail_stats(int access_type, int fail_outcome, + unsigned long long streamID) { if (!check_fail_valid(access_type, fail_outcome)) assert(0 && "Unknown cache access type or access fail"); - m_fail_stats[access_type][fail_outcome]++; + if (m_fail_stats.find(streamID) == m_fail_stats.end()) { + std::vector> new_val; + new_val.resize(NUM_MEM_ACCESS_TYPE); + for (unsigned j = 0; j < NUM_MEM_ACCESS_TYPE; ++j) { + new_val[j].resize(NUM_CACHE_RESERVATION_FAIL_STATUS, 0); + } + m_fail_stats.insert(std::pair>>( + streamID, new_val)); + } + m_fail_stats.at(streamID)[access_type][fail_outcome]++; } enum cache_request_status cache_stats::select_stats_status( @@ -712,7 +735,8 @@ enum cache_request_status cache_stats::select_stats_status( } unsigned long long &cache_stats::operator()(int access_type, int access_outcome, - bool fail_outcome) { + bool fail_outcome, + unsigned long long streamID) { /// /// Simple method to read/modify the stat corresponding to (access_type, /// access_outcome) Used overloaded () to avoid the need for separate @@ -722,17 +746,18 @@ unsigned long long &cache_stats::operator()(int access_type, int access_outcome, if (!check_fail_valid(access_type, access_outcome)) assert(0 && "Unknown cache access type or fail outcome"); - return m_fail_stats[access_type][access_outcome]; + return m_fail_stats.at(streamID)[access_type][access_outcome]; } else { if (!check_valid(access_type, access_outcome)) assert(0 && "Unknown cache access type or access outcome"); - return m_stats[access_type][access_outcome]; + return m_stats.at(streamID)[access_type][access_outcome]; } } unsigned long long cache_stats::operator()(int access_type, int access_outcome, - bool fail_outcome) const { + bool fail_outcome, + unsigned long long streamID) const { /// /// Const accessor into m_stats. /// @@ -740,12 +765,12 @@ unsigned long long cache_stats::operator()(int access_type, int access_outcome, if (!check_fail_valid(access_type, access_outcome)) assert(0 && "Unknown cache access type or fail outcome"); - return m_fail_stats[access_type][access_outcome]; + return m_fail_stats.at(streamID)[access_type][access_outcome]; } else { if (!check_valid(access_type, access_outcome)) assert(0 && "Unknown cache access type or access outcome"); - return m_stats[access_type][access_outcome]; + return m_stats.at(streamID)[access_type][access_outcome]; } } @@ -754,15 +779,74 @@ cache_stats cache_stats::operator+(const cache_stats &cs) { /// Overloaded + operator to allow for simple stat accumulation /// cache_stats ret; - for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { - for (unsigned status = 0; status < NUM_CACHE_REQUEST_STATUS; ++status) { - ret(type, status, false) = - m_stats[type][status] + cs(type, status, false); + for (auto iter = m_stats.begin(); iter != m_stats.end(); ++iter) { + unsigned long long streamID = iter->first; + ret.m_stats.insert(std::pair>>( + streamID, m_stats.at(streamID))); + } + for (auto iter = m_stats_pw.begin(); iter != m_stats_pw.end(); ++iter) { + unsigned long long streamID = iter->first; + ret.m_stats_pw.insert( + std::pair>>( + streamID, m_stats_pw.at(streamID))); + } + for (auto iter = m_fail_stats.begin(); iter != m_fail_stats.end(); ++iter) { + unsigned long long streamID = iter->first; + ret.m_fail_stats.insert( + std::pair>>( + streamID, m_fail_stats.at(streamID))); + } + for (auto iter = cs.m_stats.begin(); iter != cs.m_stats.end(); ++iter) { + unsigned long long streamID = iter->first; + if (ret.m_stats.find(streamID) == ret.m_stats.end()) { + ret.m_stats.insert( + std::pair>>( + streamID, cs.m_stats.at(streamID))); + } else { + for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { + for (unsigned status = 0; status < NUM_CACHE_REQUEST_STATUS; ++status) { + ret.m_stats.at(streamID)[type][status] += + cs(type, status, false, streamID); + } + } } - for (unsigned status = 0; status < NUM_CACHE_RESERVATION_FAIL_STATUS; - ++status) { - ret(type, status, true) = - m_fail_stats[type][status] + cs(type, status, true); + } + for (auto iter = cs.m_stats_pw.begin(); iter != cs.m_stats_pw.end(); ++iter) { + unsigned long long streamID = iter->first; + if (ret.m_stats_pw.find(streamID) == ret.m_stats_pw.end()) { + ret.m_stats_pw.insert( + std::pair>>( + streamID, cs.m_stats_pw.at(streamID))); + } else { + for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { + for (unsigned status = 0; status < NUM_CACHE_REQUEST_STATUS; ++status) { + ret.m_stats_pw.at(streamID)[type][status] += + cs(type, status, false, streamID); + } + } + } + } + for (auto iter = cs.m_fail_stats.begin(); iter != cs.m_fail_stats.end(); + ++iter) { + unsigned long long streamID = iter->first; + if (ret.m_fail_stats.find(streamID) == ret.m_fail_stats.end()) { + ret.m_fail_stats.insert( + std::pair>>( + streamID, cs.m_fail_stats.at(streamID))); + } else { + for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { + for (unsigned status = 0; status < NUM_CACHE_RESERVATION_FAIL_STATUS; + ++status) { + ret.m_fail_stats.at(streamID)[type][status] += + cs(type, status, true, streamID); + } + } } } ret.m_cache_port_available_cycles = @@ -778,16 +862,52 @@ cache_stats &cache_stats::operator+=(const cache_stats &cs) { /// /// Overloaded += operator to allow for simple stat accumulation /// - for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { - for (unsigned status = 0; status < NUM_CACHE_REQUEST_STATUS; ++status) { - m_stats[type][status] += cs(type, status, false); + for (auto iter = cs.m_stats.begin(); iter != cs.m_stats.end(); ++iter) { + unsigned long long streamID = iter->first; + if (m_stats.find(streamID) == m_stats.end()) { + m_stats.insert(std::pair>>( + streamID, cs.m_stats.at(streamID))); + } else { + for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { + for (unsigned status = 0; status < NUM_CACHE_REQUEST_STATUS; ++status) { + m_stats.at(streamID)[type][status] += + cs(type, status, false, streamID); + } + } } - for (unsigned status = 0; status < NUM_CACHE_REQUEST_STATUS; ++status) { - m_stats_pw[type][status] += cs(type, status, false); + } + for (auto iter = cs.m_stats_pw.begin(); iter != cs.m_stats_pw.end(); ++iter) { + unsigned long long streamID = iter->first; + if (m_stats_pw.find(streamID) == m_stats_pw.end()) { + m_stats_pw.insert(std::pair>>( + streamID, cs.m_stats_pw.at(streamID))); + } else { + for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { + for (unsigned status = 0; status < NUM_CACHE_REQUEST_STATUS; ++status) { + m_stats_pw.at(streamID)[type][status] += + cs(type, status, false, streamID); + } + } } - for (unsigned status = 0; status < NUM_CACHE_RESERVATION_FAIL_STATUS; - ++status) { - m_fail_stats[type][status] += cs(type, status, true); + } + for (auto iter = cs.m_fail_stats.begin(); iter != cs.m_fail_stats.end(); + ++iter) { + unsigned long long streamID = iter->first; + if (m_fail_stats.find(streamID) == m_fail_stats.end()) { + m_fail_stats.insert( + std::pair>>( + streamID, cs.m_fail_stats.at(streamID))); + } else { + for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { + for (unsigned status = 0; status < NUM_CACHE_RESERVATION_FAIL_STATUS; + ++status) { + m_fail_stats.at(streamID)[type][status] += + cs(type, status, true, streamID); + } + } } } m_cache_port_available_cycles += cs.m_cache_port_available_cycles; @@ -796,46 +916,65 @@ cache_stats &cache_stats::operator+=(const cache_stats &cs) { return *this; } -void cache_stats::print_stats(FILE *fout, const char *cache_name) const { +void cache_stats::print_stats(FILE *fout, unsigned long long streamID, + const char *cache_name) const { /// - /// Print out each non-zero cache statistic for every memory access type and - /// status "cache_name" defaults to "Cache_stats" when no argument is - /// provided, otherwise the provided name is used. The printed format is + /// For a given CUDA stream, print out each non-zero cache statistic for every + /// memory access type and status "cache_name" defaults to "Cache_stats" when + /// no argument is provided, otherwise the provided name is used. The printed + /// format is /// "[][] = " - /// + /// Specify streamID to be -1 to print every stream. + std::vector total_access; - total_access.resize(NUM_MEM_ACCESS_TYPE, 0); std::string m_cache_name = cache_name; - for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { - for (unsigned status = 0; status < NUM_CACHE_REQUEST_STATUS; ++status) { - fprintf(fout, "\t%s[%s][%s] = %llu\n", m_cache_name.c_str(), - mem_access_type_str((enum mem_access_type)type), - cache_request_status_str((enum cache_request_status)status), - m_stats[type][status]); + for (auto iter = m_stats.begin(); iter != m_stats.end(); ++iter) { + unsigned long long streamid = iter->first; + // when streamID is specified, skip stats for all other streams, otherwise, + // print stats from all streams + if ((streamID != -1) && (streamid != streamID)) continue; + total_access.clear(); + total_access.resize(NUM_MEM_ACCESS_TYPE, 0); + for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { + for (unsigned status = 0; status < NUM_CACHE_REQUEST_STATUS; ++status) { + fprintf(fout, "\t%s[%s][%s] = %llu\n", m_cache_name.c_str(), + mem_access_type_str((enum mem_access_type)type), + cache_request_status_str((enum cache_request_status)status), + m_stats.at(streamid)[type][status]); - if (status != RESERVATION_FAIL && status != MSHR_HIT) - // MSHR_HIT is a special type of SECTOR_MISS - // so its already included in the SECTOR_MISS - total_access[type] += m_stats[type][status]; + if (status != RESERVATION_FAIL && status != MSHR_HIT) + // MSHR_HIT is a special type of SECTOR_MISS + // so its already included in the SECTOR_MISS + total_access[type] += m_stats.at(streamid)[type][status]; + } + } + for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { + if (total_access[type] > 0) + fprintf(fout, "\t%s[%s][%s] = %u\n", m_cache_name.c_str(), + mem_access_type_str((enum mem_access_type)type), "TOTAL_ACCESS", + total_access[type]); } - } - for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { - if (total_access[type] > 0) - fprintf(fout, "\t%s[%s][%s] = %u\n", m_cache_name.c_str(), - mem_access_type_str((enum mem_access_type)type), "TOTAL_ACCESS", - total_access[type]); } } -void cache_stats::print_fail_stats(FILE *fout, const char *cache_name) const { +void cache_stats::print_fail_stats(FILE *fout, unsigned long long streamID, + const char *cache_name) const { std::string m_cache_name = cache_name; - for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { - for (unsigned fail = 0; fail < NUM_CACHE_RESERVATION_FAIL_STATUS; ++fail) { - if (m_fail_stats[type][fail] > 0) { - fprintf(fout, "\t%s[%s][%s] = %llu\n", m_cache_name.c_str(), - mem_access_type_str((enum mem_access_type)type), - cache_fail_status_str((enum cache_reservation_fail_reason)fail), - m_fail_stats[type][fail]); + for (auto iter = m_fail_stats.begin(); iter != m_fail_stats.end(); ++iter) { + unsigned long long streamid = iter->first; + // when streamID is specified, skip stats for all other streams, otherwise, + // print stats from all streams + if ((streamID != -1) && (streamid != streamID)) continue; + for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { + for (unsigned fail = 0; fail < NUM_CACHE_RESERVATION_FAIL_STATUS; + ++fail) { + if (m_fail_stats.at(streamid)[type][fail] > 0) { + fprintf( + fout, "\t%s[%s][%s] = %llu\n", m_cache_name.c_str(), + mem_access_type_str((enum mem_access_type)type), + cache_fail_status_str((enum cache_reservation_fail_reason)fail), + m_fail_stats.at(streamid)[type][fail]); + } } } } @@ -866,11 +1005,14 @@ unsigned long long cache_stats::get_stats( /// cache_request_statuses. /// unsigned long long total = 0; - for (unsigned type = 0; type < num_access_type; ++type) { - for (unsigned status = 0; status < num_access_status; ++status) { - if (!check_valid((int)access_type[type], (int)access_status[status])) - assert(0 && "Unknown cache access type or access outcome"); - total += m_stats[access_type[type]][access_status[status]]; + for (auto iter = m_stats.begin(); iter != m_stats.end(); ++iter) { + unsigned long long streamID = iter->first; + for (unsigned type = 0; type < num_access_type; ++type) { + for (unsigned status = 0; status < num_access_status; ++status) { + if (!check_valid((int)access_type[type], (int)access_status[status])) + assert(0 && "Unknown cache access type or access outcome"); + total += m_stats.at(streamID)[access_type[type]][access_status[status]]; + } } } return total; @@ -883,18 +1025,23 @@ void cache_stats::get_sub_stats(struct cache_sub_stats &css) const { struct cache_sub_stats t_css; t_css.clear(); - for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { - for (unsigned status = 0; status < NUM_CACHE_REQUEST_STATUS; ++status) { - if (status == HIT || status == MISS || status == SECTOR_MISS || - status == HIT_RESERVED) - t_css.accesses += m_stats[type][status]; + for (auto iter = m_stats.begin(); iter != m_stats.end(); ++iter) { + unsigned long long streamID = iter->first; + for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { + for (unsigned status = 0; status < NUM_CACHE_REQUEST_STATUS; ++status) { + if (status == HIT || status == MISS || status == SECTOR_MISS || + status == HIT_RESERVED) + t_css.accesses += m_stats.at(streamID)[type][status]; - if (status == MISS || status == SECTOR_MISS) - t_css.misses += m_stats[type][status]; + if (status == MISS || status == SECTOR_MISS) + t_css.misses += m_stats.at(streamID)[type][status]; - if (status == HIT_RESERVED) t_css.pending_hits += m_stats[type][status]; + if (status == HIT_RESERVED) + t_css.pending_hits += m_stats.at(streamID)[type][status]; - if (status == RESERVATION_FAIL) t_css.res_fails += m_stats[type][status]; + if (status == RESERVATION_FAIL) + t_css.res_fails += m_stats.at(streamID)[type][status]; + } } } @@ -912,41 +1059,48 @@ void cache_stats::get_sub_stats_pw(struct cache_sub_stats_pw &css) const { struct cache_sub_stats_pw t_css; t_css.clear(); - for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { - for (unsigned status = 0; status < NUM_CACHE_REQUEST_STATUS; ++status) { - if (status == HIT || status == MISS || status == SECTOR_MISS || - status == HIT_RESERVED) - t_css.accesses += m_stats_pw[type][status]; - - if (status == HIT) { - if (type == GLOBAL_ACC_R || type == CONST_ACC_R || type == INST_ACC_R) { - t_css.read_hits += m_stats_pw[type][status]; - } else if (type == GLOBAL_ACC_W) { - t_css.write_hits += m_stats_pw[type][status]; + for (auto iter = m_stats_pw.begin(); iter != m_stats_pw.end(); ++iter) { + unsigned long long streamID = iter->first; + for (unsigned type = 0; type < NUM_MEM_ACCESS_TYPE; ++type) { + for (unsigned status = 0; status < NUM_CACHE_REQUEST_STATUS; ++status) { + if (status == HIT || status == MISS || status == SECTOR_MISS || + status == HIT_RESERVED) + t_css.accesses += m_stats_pw.at(streamID)[type][status]; + + if (status == HIT) { + if (type == GLOBAL_ACC_R || type == CONST_ACC_R || + type == INST_ACC_R) { + t_css.read_hits += m_stats_pw.at(streamID)[type][status]; + } else if (type == GLOBAL_ACC_W) { + t_css.write_hits += m_stats_pw.at(streamID)[type][status]; + } } - } - if (status == MISS || status == SECTOR_MISS) { - if (type == GLOBAL_ACC_R || type == CONST_ACC_R || type == INST_ACC_R) { - t_css.read_misses += m_stats_pw[type][status]; - } else if (type == GLOBAL_ACC_W) { - t_css.write_misses += m_stats_pw[type][status]; + if (status == MISS || status == SECTOR_MISS) { + if (type == GLOBAL_ACC_R || type == CONST_ACC_R || + type == INST_ACC_R) { + t_css.read_misses += m_stats_pw.at(streamID)[type][status]; + } else if (type == GLOBAL_ACC_W) { + t_css.write_misses += m_stats_pw.at(streamID)[type][status]; + } } - } - if (status == HIT_RESERVED) { - if (type == GLOBAL_ACC_R || type == CONST_ACC_R || type == INST_ACC_R) { - t_css.read_pending_hits += m_stats_pw[type][status]; - } else if (type == GLOBAL_ACC_W) { - t_css.write_pending_hits += m_stats_pw[type][status]; + if (status == HIT_RESERVED) { + if (type == GLOBAL_ACC_R || type == CONST_ACC_R || + type == INST_ACC_R) { + t_css.read_pending_hits += m_stats_pw.at(streamID)[type][status]; + } else if (type == GLOBAL_ACC_W) { + t_css.write_pending_hits += m_stats_pw.at(streamID)[type][status]; + } } - } - if (status == RESERVATION_FAIL) { - if (type == GLOBAL_ACC_R || type == CONST_ACC_R || type == INST_ACC_R) { - t_css.read_res_fails += m_stats_pw[type][status]; - } else if (type == GLOBAL_ACC_W) { - t_css.write_res_fails += m_stats_pw[type][status]; + if (status == RESERVATION_FAIL) { + if (type == GLOBAL_ACC_R || type == CONST_ACC_R || + type == INST_ACC_R) { + t_css.read_res_fails += m_stats_pw.at(streamID)[type][status]; + } else if (type == GLOBAL_ACC_W) { + t_css.write_res_fails += m_stats_pw.at(streamID)[type][status]; + } } } } @@ -1139,6 +1293,50 @@ void baseline_cache::display_state(FILE *fp) const { fprintf(fp, "\n"); } +void baseline_cache::inc_aggregated_stats(cache_request_status status, + cache_request_status cache_status, + mem_fetch *mf, + enum cache_gpu_level level) { + if (level == L1_GPU_CACHE) { + m_gpu->aggregated_l1_stats.inc_stats( + mf->get_streamID(), mf->get_access_type(), + m_gpu->aggregated_l1_stats.select_stats_status(status, cache_status)); + } else if (level == L2_GPU_CACHE) { + m_gpu->aggregated_l2_stats.inc_stats( + mf->get_streamID(), mf->get_access_type(), + m_gpu->aggregated_l2_stats.select_stats_status(status, cache_status)); + } +} + +void baseline_cache::inc_aggregated_fail_stats( + cache_request_status status, cache_request_status cache_status, + mem_fetch *mf, enum cache_gpu_level level) { + if (level == L1_GPU_CACHE) { + m_gpu->aggregated_l1_stats.inc_fail_stats( + mf->get_streamID(), mf->get_access_type(), + m_gpu->aggregated_l1_stats.select_stats_status(status, cache_status)); + } else if (level == L2_GPU_CACHE) { + m_gpu->aggregated_l2_stats.inc_fail_stats( + mf->get_streamID(), mf->get_access_type(), + m_gpu->aggregated_l2_stats.select_stats_status(status, cache_status)); + } +} + +void baseline_cache::inc_aggregated_stats_pw(cache_request_status status, + cache_request_status cache_status, + mem_fetch *mf, + enum cache_gpu_level level) { + if (level == L1_GPU_CACHE) { + m_gpu->aggregated_l1_stats.inc_stats_pw( + mf->get_streamID(), mf->get_access_type(), + m_gpu->aggregated_l1_stats.select_stats_status(status, cache_status)); + } else if (level == L2_GPU_CACHE) { + m_gpu->aggregated_l2_stats.inc_stats_pw( + mf->get_streamID(), mf->get_access_type(), + m_gpu->aggregated_l2_stats.select_stats_status(status, cache_status)); + } +} + /// Read miss handler without writeback void baseline_cache::send_read_request(new_addr_type addr, new_addr_type block_addr, @@ -1170,7 +1368,7 @@ void baseline_cache::send_read_request(new_addr_type addr, m_tag_array->access(block_addr, time, cache_index, wb, evicted, mf); m_mshrs.add(mshr_addr, mf); - m_stats.inc_stats(mf->get_access_type(), MSHR_HIT); + m_stats.inc_stats(mf->get_access_type(), MSHR_HIT, mf->get_streamID()); do_miss = true; } else if (!mshr_hit && mshr_avail && @@ -1191,9 +1389,11 @@ void baseline_cache::send_read_request(new_addr_type addr, do_miss = true; } else if (mshr_hit && !mshr_avail) - m_stats.inc_fail_stats(mf->get_access_type(), MSHR_MERGE_ENRTY_FAIL); + m_stats.inc_fail_stats(mf->get_access_type(), MSHR_MERGE_ENRTY_FAIL, + mf->get_streamID()); else if (!mshr_hit && !mshr_avail) - m_stats.inc_fail_stats(mf->get_access_type(), MSHR_ENRTY_FAIL); + m_stats.inc_fail_stats(mf->get_access_type(), MSHR_ENRTY_FAIL, + mf->get_streamID()); else assert(0); } @@ -1253,7 +1453,8 @@ cache_request_status data_cache::wr_hit_wt(new_addr_type addr, std::list &events, enum cache_request_status status) { if (miss_queue_full(0)) { - m_stats.inc_fail_stats(mf->get_access_type(), MISS_QUEUE_FULL); + m_stats.inc_fail_stats(mf->get_access_type(), MISS_QUEUE_FULL, + mf->get_streamID()); return RESERVATION_FAIL; // cannot handle request this cycle } @@ -1281,7 +1482,8 @@ cache_request_status data_cache::wr_hit_we(new_addr_type addr, std::list &events, enum cache_request_status status) { if (miss_queue_full(0)) { - m_stats.inc_fail_stats(mf->get_access_type(), MISS_QUEUE_FULL); + m_stats.inc_fail_stats(mf->get_access_type(), MISS_QUEUE_FULL, + mf->get_streamID()); return RESERVATION_FAIL; // cannot handle request this cycle } @@ -1330,11 +1532,14 @@ enum cache_request_status data_cache::wr_miss_wa_naive( (m_miss_queue.size() < m_config.m_miss_queue_size)))) { // check what is the exactly the failure reason if (miss_queue_full(2)) - m_stats.inc_fail_stats(mf->get_access_type(), MISS_QUEUE_FULL); + m_stats.inc_fail_stats(mf->get_access_type(), MISS_QUEUE_FULL, + mf->get_streamID()); else if (mshr_hit && !mshr_avail) - m_stats.inc_fail_stats(mf->get_access_type(), MSHR_MERGE_ENRTY_FAIL); + m_stats.inc_fail_stats(mf->get_access_type(), MSHR_MERGE_ENRTY_FAIL, + mf->get_streamID()); else if (!mshr_hit && !mshr_avail) - m_stats.inc_fail_stats(mf->get_access_type(), MSHR_ENRTY_FAIL); + m_stats.inc_fail_stats(mf->get_access_type(), MSHR_ENRTY_FAIL, + mf->get_streamID()); else assert(0); @@ -1353,10 +1558,10 @@ enum cache_request_status data_cache::wr_miss_wa_naive( mf->get_access_warp_mask(), mf->get_access_byte_mask(), mf->get_access_sector_mask(), m_gpu->gpgpu_ctx); - mem_fetch *n_mf = - new mem_fetch(*ma, NULL, mf->get_ctrl_size(), mf->get_wid(), - mf->get_sid(), mf->get_tpc(), mf->get_mem_config(), - m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle); + mem_fetch *n_mf = new mem_fetch( + *ma, NULL, mf->get_streamID(), mf->get_ctrl_size(), mf->get_wid(), + mf->get_sid(), mf->get_tpc(), mf->get_mem_config(), + m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle); bool do_miss = false; bool wb = false; @@ -1378,7 +1583,7 @@ enum cache_request_status data_cache::wr_miss_wa_naive( evicted.m_block_addr, m_wrbk_type, mf->get_access_warp_mask(), evicted.m_byte_mask, evicted.m_sector_mask, evicted.m_modified_size, true, m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle, -1, -1, -1, - NULL); + NULL, mf->get_streamID()); // the evicted block may have wrong chip id when advanced L2 hashing is // used, so set the right chip address from the original mf wb->set_chip(mf->get_tlx_addr().chip); @@ -1404,7 +1609,8 @@ enum cache_request_status data_cache::wr_miss_wa_fetch_on_write( // reserve mshr if (miss_queue_full(0)) { - m_stats.inc_fail_stats(mf->get_access_type(), MISS_QUEUE_FULL); + m_stats.inc_fail_stats(mf->get_access_type(), MISS_QUEUE_FULL, + mf->get_streamID()); return RESERVATION_FAIL; // cannot handle request this cycle } @@ -1431,7 +1637,7 @@ enum cache_request_status data_cache::wr_miss_wa_fetch_on_write( evicted.m_block_addr, m_wrbk_type, mf->get_access_warp_mask(), evicted.m_byte_mask, evicted.m_sector_mask, evicted.m_modified_size, true, m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle, -1, -1, -1, - NULL); + NULL, mf->get_streamID()); // the evicted block may have wrong chip id when advanced L2 hashing is // used, so set the right chip address from the original mf wb->set_chip(mf->get_tlx_addr().chip); @@ -1451,11 +1657,14 @@ enum cache_request_status data_cache::wr_miss_wa_fetch_on_write( (m_miss_queue.size() < m_config.m_miss_queue_size)))) { // check what is the exactly the failure reason if (miss_queue_full(1)) - m_stats.inc_fail_stats(mf->get_access_type(), MISS_QUEUE_FULL); + m_stats.inc_fail_stats(mf->get_access_type(), MISS_QUEUE_FULL, + mf->get_streamID()); else if (mshr_hit && !mshr_avail) - m_stats.inc_fail_stats(mf->get_access_type(), MSHR_MERGE_ENRTY_FAIL); + m_stats.inc_fail_stats(mf->get_access_type(), MSHR_MERGE_ENRTY_FAIL, + mf->get_streamID()); else if (!mshr_hit && !mshr_avail) - m_stats.inc_fail_stats(mf->get_access_type(), MSHR_ENRTY_FAIL); + m_stats.inc_fail_stats(mf->get_access_type(), MSHR_ENRTY_FAIL, + mf->get_streamID()); else assert(0); @@ -1468,7 +1677,8 @@ enum cache_request_status data_cache::wr_miss_wa_fetch_on_write( if (m_mshrs.probe(mshr_addr) && m_mshrs.is_read_after_write_pending(mshr_addr) && mf->is_write()) { // assert(0); - m_stats.inc_fail_stats(mf->get_access_type(), MSHR_RW_PENDING); + m_stats.inc_fail_stats(mf->get_access_type(), MSHR_RW_PENDING, + mf->get_streamID()); return RESERVATION_FAIL; } @@ -1479,8 +1689,8 @@ enum cache_request_status data_cache::wr_miss_wa_fetch_on_write( mf->get_access_sector_mask(), m_gpu->gpgpu_ctx); mem_fetch *n_mf = new mem_fetch( - *ma, NULL, mf->get_ctrl_size(), mf->get_wid(), mf->get_sid(), - mf->get_tpc(), mf->get_mem_config(), + *ma, NULL, mf->get_streamID(), mf->get_ctrl_size(), mf->get_wid(), + mf->get_sid(), mf->get_tpc(), mf->get_mem_config(), m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle, NULL, mf); new_addr_type block_addr = m_config.block_addr(addr); @@ -1504,7 +1714,7 @@ enum cache_request_status data_cache::wr_miss_wa_fetch_on_write( evicted.m_block_addr, m_wrbk_type, mf->get_access_warp_mask(), evicted.m_byte_mask, evicted.m_sector_mask, evicted.m_modified_size, true, m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle, -1, -1, -1, - NULL); + NULL, mf->get_streamID()); // the evicted block may have wrong chip id when advanced L2 hashing is // used, so set the right chip address from the original mf wb->set_chip(mf->get_tlx_addr().chip); @@ -1528,7 +1738,8 @@ enum cache_request_status data_cache::wr_miss_wa_lazy_fetch_on_read( // mshr if (miss_queue_full(0)) { - m_stats.inc_fail_stats(mf->get_access_type(), MISS_QUEUE_FULL); + m_stats.inc_fail_stats(mf->get_access_type(), MISS_QUEUE_FULL, + mf->get_streamID()); return RESERVATION_FAIL; // cannot handle request this cycle } @@ -1571,7 +1782,7 @@ enum cache_request_status data_cache::wr_miss_wa_lazy_fetch_on_read( evicted.m_block_addr, m_wrbk_type, mf->get_access_warp_mask(), evicted.m_byte_mask, evicted.m_sector_mask, evicted.m_modified_size, true, m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle, -1, -1, -1, - NULL); + NULL, mf->get_streamID()); // the evicted block may have wrong chip id when advanced L2 hashing is // used, so set the right chip address from the original mf wb->set_chip(mf->get_tlx_addr().chip); @@ -1589,7 +1800,8 @@ enum cache_request_status data_cache::wr_miss_no_wa( new_addr_type addr, unsigned cache_index, mem_fetch *mf, unsigned time, std::list &events, enum cache_request_status status) { if (miss_queue_full(0)) { - m_stats.inc_fail_stats(mf->get_access_type(), MISS_QUEUE_FULL); + m_stats.inc_fail_stats(mf->get_access_type(), MISS_QUEUE_FULL, + mf->get_streamID()); return RESERVATION_FAIL; // cannot handle request this cycle } @@ -1634,7 +1846,8 @@ enum cache_request_status data_cache::rd_miss_base( if (miss_queue_full(1)) { // cannot handle request this cycle // (might need to generate two requests) - m_stats.inc_fail_stats(mf->get_access_type(), MISS_QUEUE_FULL); + m_stats.inc_fail_stats(mf->get_access_type(), MISS_QUEUE_FULL, + mf->get_streamID()); return RESERVATION_FAIL; } @@ -1653,7 +1866,7 @@ enum cache_request_status data_cache::rd_miss_base( evicted.m_block_addr, m_wrbk_type, mf->get_access_warp_mask(), evicted.m_byte_mask, evicted.m_sector_mask, evicted.m_modified_size, true, m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle, -1, -1, -1, - NULL); + NULL, mf->get_streamID()); // the evicted block may have wrong chip id when advanced L2 hashing is // used, so set the right chip address from the original mf wb->set_chip(mf->get_tlx_addr().chip); @@ -1693,16 +1906,20 @@ enum cache_request_status read_only_cache::access( cache_status = RESERVATION_FAIL; } else { cache_status = RESERVATION_FAIL; - m_stats.inc_fail_stats(mf->get_access_type(), MISS_QUEUE_FULL); + m_stats.inc_fail_stats(mf->get_access_type(), MISS_QUEUE_FULL, + mf->get_streamID()); } } else { - m_stats.inc_fail_stats(mf->get_access_type(), LINE_ALLOC_FAIL); + m_stats.inc_fail_stats(mf->get_access_type(), LINE_ALLOC_FAIL, + mf->get_streamID()); } m_stats.inc_stats(mf->get_access_type(), - m_stats.select_stats_status(status, cache_status)); + m_stats.select_stats_status(status, cache_status), + mf->get_streamID()); m_stats.inc_stats_pw(mf->get_access_type(), - m_stats.select_stats_status(status, cache_status)); + m_stats.select_stats_status(status, cache_status), + mf->get_streamID()); return cache_status; } @@ -1730,7 +1947,8 @@ enum cache_request_status data_cache::process_tag_probe( } else { // the only reason for reservation fail here is LINE_ALLOC_FAIL (i.e all // lines are reserved) - m_stats.inc_fail_stats(mf->get_access_type(), LINE_ALLOC_FAIL); + m_stats.inc_fail_stats(mf->get_access_type(), LINE_ALLOC_FAIL, + mf->get_streamID()); } } else { // Read if (probe_status == HIT) { @@ -1742,7 +1960,8 @@ enum cache_request_status data_cache::process_tag_probe( } else { // the only reason for reservation fail here is LINE_ALLOC_FAIL (i.e all // lines are reserved) - m_stats.inc_fail_stats(mf->get_access_type(), LINE_ALLOC_FAIL); + m_stats.inc_fail_stats(mf->get_access_type(), LINE_ALLOC_FAIL, + mf->get_streamID()); } } @@ -1767,9 +1986,11 @@ enum cache_request_status data_cache::access(new_addr_type addr, mem_fetch *mf, enum cache_request_status access_status = process_tag_probe(wr, probe_status, addr, cache_index, mf, time, events); m_stats.inc_stats(mf->get_access_type(), - m_stats.select_stats_status(probe_status, access_status)); - m_stats.inc_stats_pw(mf->get_access_type(), m_stats.select_stats_status( - probe_status, access_status)); + m_stats.select_stats_status(probe_status, access_status), + mf->get_streamID()); + m_stats.inc_stats_pw(mf->get_access_type(), + m_stats.select_stats_status(probe_status, access_status), + mf->get_streamID()); return access_status; } @@ -1831,9 +2052,11 @@ enum cache_request_status tex_cache::access(new_addr_type addr, mem_fetch *mf, cache_status = HIT_RESERVED; } m_stats.inc_stats(mf->get_access_type(), - m_stats.select_stats_status(status, cache_status)); + m_stats.select_stats_status(status, cache_status), + mf->get_streamID()); m_stats.inc_stats_pw(mf->get_access_type(), - m_stats.select_stats_status(status, cache_status)); + m_stats.select_stats_status(status, cache_status), + mf->get_streamID()); return cache_status; } diff --git a/src/gpgpu-sim/gpu-cache.h b/src/gpgpu-sim/gpu-cache.h index 5fd40a9..c07695f 100644 --- a/src/gpgpu-sim/gpu-cache.h +++ b/src/gpgpu-sim/gpu-cache.h @@ -72,6 +72,13 @@ enum cache_event_type { WRITE_ALLOCATE_SENT }; +enum cache_gpu_level { + L1_GPU_CACHE = 0, + L2_GPU_CACHE, + OTHER_GPU_CACHE, + NUM_CACHE_GPU_LEVELS +}; + struct evicted_block_info { new_addr_type m_block_addr; unsigned m_modified_size; @@ -1200,20 +1207,26 @@ class cache_stats { void clear(); // Clear AerialVision cache stats after each window void clear_pw(); - void inc_stats(int access_type, int access_outcome); + void inc_stats(int access_type, int access_outcome, + unsigned long long streamID); // Increment AerialVision cache stats - void inc_stats_pw(int access_type, int access_outcome); - void inc_fail_stats(int access_type, int fail_outcome); + void inc_stats_pw(int access_type, int access_outcome, + unsigned long long streamID); + void inc_fail_stats(int access_type, int fail_outcome, + unsigned long long streamID); enum cache_request_status select_stats_status( enum cache_request_status probe, enum cache_request_status access) const; unsigned long long &operator()(int access_type, int access_outcome, - bool fail_outcome); + bool fail_outcome, + unsigned long long streamID); unsigned long long operator()(int access_type, int access_outcome, - bool fail_outcome) const; + bool fail_outcome, + unsigned long long streamID) const; cache_stats operator+(const cache_stats &cs); cache_stats &operator+=(const cache_stats &cs); - void print_stats(FILE *fout, const char *cache_name = "Cache_stats") const; - void print_fail_stats(FILE *fout, + void print_stats(FILE *fout, unsigned long long streamID, + const char *cache_name = "Cache_stats") const; + void print_fail_stats(FILE *fout, unsigned long long streamID, const char *cache_name = "Cache_fail_stats") const; unsigned long long get_stats(enum mem_access_type *access_type, @@ -1231,10 +1244,14 @@ class cache_stats { bool check_valid(int type, int status) const; bool check_fail_valid(int type, int fail) const; - std::vector > m_stats; + // CUDA streamID -> cache stats[NUM_MEM_ACCESS_TYPE] + std::map>> + m_stats; // AerialVision cache stats (per-window) - std::vector > m_stats_pw; - std::vector > m_fail_stats; + std::map>> + m_stats_pw; + std::map>> + m_fail_stats; unsigned long long m_cache_port_available_cycles; unsigned long long m_cache_data_port_busy_cycles; @@ -1264,11 +1281,14 @@ class baseline_cache : public cache_t { public: baseline_cache(const char *name, cache_config &config, int core_id, int type_id, mem_fetch_interface *memport, - enum mem_fetch_status status) + enum mem_fetch_status status, enum cache_gpu_level level, + gpgpu_sim *gpu) : m_config(config), m_tag_array(new tag_array(config, core_id, type_id)), m_mshrs(config.m_mshr_entries, config.m_mshr_max_merge), - m_bandwidth_management(config) { + m_bandwidth_management(config), + m_level(level), + m_gpu(gpu) { init(name, config, memport, status); } @@ -1336,6 +1356,15 @@ class baseline_cache : public cache_t { bool fill_port_free() const { return m_bandwidth_management.fill_port_free(); } + void inc_aggregated_stats(cache_request_status status, + cache_request_status cache_status, mem_fetch *mf, + enum cache_gpu_level level); + void inc_aggregated_fail_stats(cache_request_status status, + cache_request_status cache_status, + mem_fetch *mf, enum cache_gpu_level level); + void inc_aggregated_stats_pw(cache_request_status status, + cache_request_status cache_status, mem_fetch *mf, + enum cache_gpu_level level); // This is a gapping hole we are poking in the system to quickly handle // filling the cache on cudamemcopies. We don't care about anything other than @@ -1367,6 +1396,8 @@ class baseline_cache : public cache_t { std::list m_miss_queue; enum mem_fetch_status m_miss_queue_status; mem_fetch_interface *m_memport; + cache_gpu_level m_level; + gpgpu_sim *m_gpu; struct extra_mf_fields { extra_mf_fields() { m_valid = false; } @@ -1453,8 +1484,10 @@ class read_only_cache : public baseline_cache { public: read_only_cache(const char *name, cache_config &config, int core_id, int type_id, mem_fetch_interface *memport, - enum mem_fetch_status status) - : baseline_cache(name, config, core_id, type_id, memport, status) {} + enum mem_fetch_status status, enum cache_gpu_level level, + gpgpu_sim *gpu) + : baseline_cache(name, config, core_id, type_id, memport, status, level, + gpu) {} /// Access cache for read_only_cache: returns RESERVATION_FAIL if request /// could not be accepted (for any reason) @@ -1478,8 +1511,10 @@ class data_cache : public baseline_cache { data_cache(const char *name, cache_config &config, int core_id, int type_id, mem_fetch_interface *memport, mem_fetch_allocator *mfcreator, enum mem_fetch_status status, mem_access_type wr_alloc_type, - mem_access_type wrbk_type, class gpgpu_sim *gpu) - : baseline_cache(name, config, core_id, type_id, memport, status) { + mem_access_type wrbk_type, class gpgpu_sim *gpu, + enum cache_gpu_level level) + : baseline_cache(name, config, core_id, type_id, memport, status, level, + gpu) { init(mfcreator); m_wr_alloc_type = wr_alloc_type; m_wrbk_type = wrbk_type; @@ -1668,9 +1703,10 @@ class l1_cache : public data_cache { public: l1_cache(const char *name, cache_config &config, int core_id, int type_id, mem_fetch_interface *memport, mem_fetch_allocator *mfcreator, - enum mem_fetch_status status, class gpgpu_sim *gpu) + enum mem_fetch_status status, class gpgpu_sim *gpu, + enum cache_gpu_level level) : data_cache(name, config, core_id, type_id, memport, mfcreator, status, - L1_WR_ALLOC_R, L1_WRBK_ACC, gpu) {} + L1_WR_ALLOC_R, L1_WRBK_ACC, gpu, level) {} virtual ~l1_cache() {} @@ -1693,9 +1729,10 @@ class l2_cache : public data_cache { public: l2_cache(const char *name, cache_config &config, int core_id, int type_id, mem_fetch_interface *memport, mem_fetch_allocator *mfcreator, - enum mem_fetch_status status, class gpgpu_sim *gpu) + enum mem_fetch_status status, class gpgpu_sim *gpu, + enum cache_gpu_level level) : data_cache(name, config, core_id, type_id, memport, mfcreator, status, - L2_WR_ALLOC_R, L2_WRBK_ACC, gpu) {} + L2_WR_ALLOC_R, L2_WRBK_ACC, gpu, level) {} virtual ~l2_cache() {} diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 1cb8a25..0c922bd 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -788,6 +788,22 @@ void increment_x_then_y_then_z(dim3 &i, const dim3 &bound) { } void gpgpu_sim::launch(kernel_info_t *kinfo) { + unsigned kernelID = kinfo->get_uid(); + unsigned long long streamID = kinfo->get_streamID(); + + kernel_time_t kernel_time = {gpu_tot_sim_cycle + gpu_sim_cycle, 0}; + if (gpu_kernel_time.find(streamID) == gpu_kernel_time.end()) { + std::map new_val; + new_val.insert(std::pair(kernelID, kernel_time)); + gpu_kernel_time.insert( + std::pair>( + streamID, new_val)); + } else { + gpu_kernel_time.at(streamID).insert( + std::pair(kernelID, kernel_time)); + ////////// assume same kernel ID do not appear more than once + } + unsigned cta_size = kinfo->threads_per_cta(); if (cta_size > m_shader_config->n_thread_per_shader) { printf( @@ -893,7 +909,10 @@ kernel_info_t *gpgpu_sim::select_kernel() { } unsigned gpgpu_sim::finished_kernel() { - if (m_finished_kernel.empty()) return 0; + if (m_finished_kernel.empty()) { + last_streamID = -1; + return 0; + } unsigned result = m_finished_kernel.front(); m_finished_kernel.pop_front(); return result; @@ -901,6 +920,11 @@ unsigned gpgpu_sim::finished_kernel() { void gpgpu_sim::set_kernel_done(kernel_info_t *kernel) { unsigned uid = kernel->get_uid(); + last_uid = uid; + unsigned long long streamID = kernel->get_streamID(); + last_streamID = streamID; + gpu_kernel_time.at(streamID).at(uid).end_cycle = + gpu_tot_sim_cycle + gpu_sim_cycle; m_finished_kernel.push_back(uid); std::vector::iterator k; for (k = m_running_kernels.begin(); k != m_running_kernels.end(); k++) { @@ -971,6 +995,9 @@ gpgpu_sim::gpgpu_sim(const gpgpu_sim_config &config, gpgpu_context *ctx) gpu_tot_sim_cycle_parition_util = 0; partiton_replys_in_parallel = 0; partiton_replys_in_parallel_total = 0; + last_streamID = -1; + + gpu_kernel_time.clear(); m_memory_partition_unit = new memory_partition_unit *[m_memory_config->m_n_mem]; @@ -1178,9 +1205,9 @@ PowerscalingCoefficients *gpgpu_sim::get_scaling_coeffs() { return m_gpgpusim_wrapper->get_scaling_coeffs(); } -void gpgpu_sim::print_stats() { +void gpgpu_sim::print_stats(unsigned long long streamID) { gpgpu_ctx->stats->ptx_file_line_stats_write_file(); - gpu_print_stat(); + gpu_print_stat(streamID); if (g_network_mode) { printf( @@ -1363,12 +1390,15 @@ void gpgpu_sim::clear_executed_kernel_info() { m_executed_kernel_names.clear(); m_executed_kernel_uids.clear(); } -void gpgpu_sim::gpu_print_stat() { + +void gpgpu_sim::gpu_print_stat(unsigned long long streamID) { FILE *statfout = stdout; std::string kernel_info_str = executed_kernel_info_string(); fprintf(statfout, "%s", kernel_info_str.c_str()); + printf("kernel_stream_id = %llu\n", streamID); + printf("gpu_sim_cycle = %lld\n", gpu_sim_cycle); printf("gpu_sim_insn = %lld\n", gpu_sim_insn); printf("gpu_ipc = %12.4f\n", (float)gpu_sim_insn / gpu_sim_cycle); @@ -1440,9 +1470,10 @@ void gpgpu_sim::gpu_print_stat() { m_cluster[i]->get_cache_stats(core_cache_stats); } printf("\nTotal_core_cache_stats:\n"); - core_cache_stats.print_stats(stdout, "Total_core_cache_stats_breakdown"); + core_cache_stats.print_stats(stdout, streamID, + "Total_core_cache_stats_breakdown"); printf("\nTotal_core_cache_fail_stats:\n"); - core_cache_stats.print_fail_stats(stdout, + core_cache_stats.print_fail_stats(stdout, streamID, "Total_core_cache_fail_stats_breakdown"); shader_print_scheduler_stat(stdout, false); @@ -1510,9 +1541,10 @@ void gpgpu_sim::gpu_print_stat() { printf("L2_total_cache_reservation_fails = %llu\n", total_l2_css.res_fails); printf("L2_total_cache_breakdown:\n"); - l2_stats.print_stats(stdout, "L2_cache_stats_breakdown"); + l2_stats.print_stats(stdout, streamID, "L2_cache_stats_breakdown"); printf("L2_total_cache_reservation_fail_breakdown:\n"); - l2_stats.print_fail_stats(stdout, "L2_cache_stats_fail_breakdown"); + l2_stats.print_fail_stats(stdout, streamID, + "L2_cache_stats_fail_breakdown"); total_l2_css.print_port_stats(stdout, "L2_cache"); } } @@ -1955,8 +1987,10 @@ void gpgpu_sim::cycle() { if (mf) partiton_reqs_in_parallel_per_cycle++; } m_memory_sub_partition[i]->cache_cycle(gpu_sim_cycle + gpu_tot_sim_cycle); - m_memory_sub_partition[i]->accumulate_L2cache_stats( - m_power_stats->pwr_mem_stat->l2_cache_stats[CURRENT_STAT_IDX]); + if (m_config.g_power_simulation_enabled) { + m_memory_sub_partition[i]->accumulate_L2cache_stats( + m_power_stats->pwr_mem_stat->l2_cache_stats[CURRENT_STAT_IDX]); + } } } partiton_reqs_in_parallel += partiton_reqs_in_parallel_per_cycle; @@ -1978,14 +2012,16 @@ void gpgpu_sim::cycle() { *active_sms += m_cluster[i]->get_n_active_sms(); } // Update core icnt/cache stats for AccelWattch - m_cluster[i]->get_icnt_stats( - m_power_stats->pwr_mem_stat->n_simt_to_mem[CURRENT_STAT_IDX][i], - m_power_stats->pwr_mem_stat->n_mem_to_simt[CURRENT_STAT_IDX][i]); - m_cluster[i]->get_cache_stats( - m_power_stats->pwr_mem_stat->core_cache_stats[CURRENT_STAT_IDX]); - m_cluster[i]->get_current_occupancy( - gpu_occupancy.aggregate_warp_slot_filled, - gpu_occupancy.aggregate_theoretical_warp_slots); + if (m_config.g_power_simulation_enabled) { + m_cluster[i]->get_icnt_stats( + m_power_stats->pwr_mem_stat->n_simt_to_mem[CURRENT_STAT_IDX][i], + m_power_stats->pwr_mem_stat->n_mem_to_simt[CURRENT_STAT_IDX][i]); + m_cluster[i]->get_cache_stats( + m_power_stats->pwr_mem_stat->core_cache_stats[CURRENT_STAT_IDX]); + m_cluster[i]->get_current_occupancy( + gpu_occupancy.aggregate_warp_slot_filled, + gpu_occupancy.aggregate_theoretical_warp_slots); + } } float temp = 0; for (unsigned i = 0; i < m_shader_config->num_shader(); i++) { diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index d43b399..8e81451 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -539,7 +539,7 @@ class gpgpu_sim : public gpgpu_t { (m_config.gpu_max_completed_cta_opt && (gpu_completed_cta >= m_config.gpu_max_completed_cta_opt)); } - void print_stats(); + void print_stats(unsigned long long streamID); void update_stats(); void deadlock_check(); void inc_completed_cta() { gpu_completed_cta++; } @@ -568,7 +568,7 @@ class gpgpu_sim : public gpgpu_t { void decrement_kernel_latency(); const gpgpu_sim_config &get_config() const { return m_config; } - void gpu_print_stat(); + void gpu_print_stat(unsigned long long streamID); void dump_pipeline(int mask, int s, int m) const; void perf_memcpy_to_gpu(size_t dst_start_addr, size_t count); @@ -685,6 +685,17 @@ class gpgpu_sim : public gpgpu_t { occupancy_stats gpu_occupancy; occupancy_stats gpu_tot_occupancy; + typedef struct { + unsigned long long start_cycle; + unsigned long long end_cycle; + } kernel_time_t; + std::map> + gpu_kernel_time; + unsigned long long last_streamID; + unsigned long long last_uid; + cache_stats aggregated_l1_stats; + cache_stats aggregated_l2_stats; + // performance counter for stalls due to congestion. unsigned int gpu_stall_dramfull; unsigned int gpu_stall_icnt2sh; @@ -712,6 +723,9 @@ class gpgpu_sim : public gpgpu_t { public: bool is_functional_sim() { return m_functional_sim; } kernel_info_t *get_functional_kernel() { return m_functional_sim_kernel; } + std::vector get_running_kernels() { + return m_running_kernels; + } void functional_launch(kernel_info_t *k) { m_functional_sim = true; m_functional_sim_kernel = k; diff --git a/src/gpgpu-sim/l2cache.cc b/src/gpgpu-sim/l2cache.cc index 8469453..52eed0e 100644 --- a/src/gpgpu-sim/l2cache.cc +++ b/src/gpgpu-sim/l2cache.cc @@ -51,12 +51,12 @@ mem_fetch *partition_mf_allocator::alloc(new_addr_type addr, mem_access_type type, unsigned size, - bool wr, - unsigned long long cycle) const { + bool wr, unsigned long long cycle, + unsigned long long streamID) const { assert(wr); mem_access_t access(type, addr, size, wr, m_memory_config->gpgpu_ctx); - mem_fetch *mf = new mem_fetch(access, NULL, WRITE_PACKET_SIZE, -1, -1, -1, - m_memory_config, cycle); + mem_fetch *mf = new mem_fetch(access, NULL, streamID, WRITE_PACKET_SIZE, -1, + -1, -1, m_memory_config, cycle); return mf; } @@ -65,12 +65,12 @@ mem_fetch *partition_mf_allocator::alloc( const mem_access_byte_mask_t &byte_mask, const mem_access_sector_mask_t §or_mask, unsigned size, bool wr, unsigned long long cycle, unsigned wid, unsigned sid, unsigned tpc, - mem_fetch *original_mf) const { + mem_fetch *original_mf, unsigned long long streamID) const { mem_access_t access(type, addr, size, wr, active_mask, byte_mask, sector_mask, m_memory_config->gpgpu_ctx); - mem_fetch *mf = - new mem_fetch(access, NULL, wr ? WRITE_PACKET_SIZE : READ_PACKET_SIZE, - wid, sid, tpc, m_memory_config, cycle, original_mf); + mem_fetch *mf = new mem_fetch(access, NULL, streamID, + wr ? WRITE_PACKET_SIZE : READ_PACKET_SIZE, wid, + sid, tpc, m_memory_config, cycle, original_mf); return mf; } memory_partition_unit::memory_partition_unit(unsigned partition_id, @@ -436,9 +436,9 @@ memory_sub_partition::memory_sub_partition(unsigned sub_partition_id, m_mf_allocator = new partition_mf_allocator(config); if (!m_config->m_L2_config.disabled()) - m_L2cache = - new l2_cache(L2c_name, m_config->m_L2_config, -1, -1, m_L2interface, - m_mf_allocator, IN_PARTITION_L2_MISS_QUEUE, gpu); + m_L2cache = new l2_cache(L2c_name, m_config->m_L2_config, -1, -1, + m_L2interface, m_mf_allocator, + IN_PARTITION_L2_MISS_QUEUE, gpu, L2_GPU_CACHE); unsigned int icnt_L2; unsigned int L2_dram; @@ -733,7 +733,7 @@ memory_sub_partition::breakdown_request_to_sector_requests(mem_fetch *mf) { mf->get_access_warp_mask(), mf->get_access_byte_mask() & mask, std::bitset().set(i), SECTOR_SIZE, mf->is_write(), m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle, mf->get_wid(), - mf->get_sid(), mf->get_tpc(), mf); + mf->get_sid(), mf->get_tpc(), mf, mf->get_streamID()); result.push_back(n_mf); } @@ -756,7 +756,7 @@ memory_sub_partition::breakdown_request_to_sector_requests(mem_fetch *mf) { mf->get_access_byte_mask() & mask, std::bitset().set(i), SECTOR_SIZE, mf->is_write(), m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle, mf->get_wid(), - mf->get_sid(), mf->get_tpc(), mf); + mf->get_sid(), mf->get_tpc(), mf, mf->get_streamID()); result.push_back(n_mf); } @@ -772,7 +772,8 @@ memory_sub_partition::breakdown_request_to_sector_requests(mem_fetch *mf) { mf->get_access_warp_mask(), mf->get_access_byte_mask() & mask, std::bitset().set(i), SECTOR_SIZE, mf->is_write(), m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle, - mf->get_wid(), mf->get_sid(), mf->get_tpc(), mf); + mf->get_wid(), mf->get_sid(), mf->get_tpc(), mf, + mf->get_streamID()); result.push_back(n_mf); } diff --git a/src/gpgpu-sim/l2cache.h b/src/gpgpu-sim/l2cache.h index ccf9b70..65c9c38 100644 --- a/src/gpgpu-sim/l2cache.h +++ b/src/gpgpu-sim/l2cache.h @@ -52,15 +52,16 @@ class partition_mf_allocator : public mem_fetch_allocator { return NULL; } virtual mem_fetch *alloc(new_addr_type addr, mem_access_type type, - unsigned size, bool wr, - unsigned long long cycle) const; + unsigned size, bool wr, unsigned long long cycle, + unsigned long long streamID) const; virtual mem_fetch *alloc(new_addr_type addr, mem_access_type type, const active_mask_t &active_mask, const mem_access_byte_mask_t &byte_mask, const mem_access_sector_mask_t §or_mask, unsigned size, bool wr, unsigned long long cycle, unsigned wid, unsigned sid, unsigned tpc, - mem_fetch *original_mf) const; + mem_fetch *original_mf, + unsigned long long streamID) const; private: const memory_config *m_memory_config; diff --git a/src/gpgpu-sim/mem_fetch.cc b/src/gpgpu-sim/mem_fetch.cc index 0d86046..7211a7d 100644 --- a/src/gpgpu-sim/mem_fetch.cc +++ b/src/gpgpu-sim/mem_fetch.cc @@ -35,10 +35,10 @@ unsigned mem_fetch::sm_next_mf_request_uid = 1; mem_fetch::mem_fetch(const mem_access_t &access, const warp_inst_t *inst, - unsigned ctrl_size, unsigned wid, unsigned sid, - unsigned tpc, const memory_config *config, - unsigned long long cycle, mem_fetch *m_original_mf, - mem_fetch *m_original_wr_mf) + unsigned long long streamID, unsigned ctrl_size, + unsigned wid, unsigned sid, unsigned tpc, + const memory_config *config, unsigned long long cycle, + mem_fetch *m_original_mf, mem_fetch *m_original_wr_mf) : m_access(access) { @@ -48,6 +48,7 @@ mem_fetch::mem_fetch(const mem_access_t &access, const warp_inst_t *inst, m_inst = *inst; assert(wid == m_inst.warp_id()); } + m_streamID = streamID; m_data_size = access.get_size(); m_ctrl_size = ctrl_size; m_sid = sid; diff --git a/src/gpgpu-sim/mem_fetch.h b/src/gpgpu-sim/mem_fetch.h index 283fe80..7704218 100644 --- a/src/gpgpu-sim/mem_fetch.h +++ b/src/gpgpu-sim/mem_fetch.h @@ -54,9 +54,10 @@ class memory_config; class mem_fetch { public: mem_fetch(const mem_access_t &access, const warp_inst_t *inst, - unsigned ctrl_size, unsigned wid, unsigned sid, unsigned tpc, - const memory_config *config, unsigned long long cycle, - mem_fetch *original_mf = NULL, mem_fetch *original_wr_mf = NULL); + unsigned long long streamID, unsigned ctrl_size, unsigned wid, + unsigned sid, unsigned tpc, const memory_config *config, + unsigned long long cycle, mem_fetch *original_mf = NULL, + mem_fetch *original_wr_mf = NULL); ~mem_fetch(); void set_status(enum mem_fetch_status status, unsigned long long cycle); @@ -105,6 +106,7 @@ class mem_fetch { unsigned get_timestamp() const { return m_timestamp; } unsigned get_return_timestamp() const { return m_timestamp2; } unsigned get_icnt_receive_time() const { return m_icnt_receive_time; } + unsigned long long get_streamID() const { return m_streamID; } enum mem_access_type get_access_type() const { return m_access.get_type(); } const active_mask_t &get_access_warp_mask() const { @@ -163,6 +165,8 @@ class mem_fetch { // requesting instruction (put last so mem_fetch prints nicer in gdb) warp_inst_t m_inst; + unsigned long long m_streamID; + static unsigned sm_next_mf_request_uid; const memory_config *m_mem_config; diff --git a/src/gpgpu-sim/power_stat.cc b/src/gpgpu-sim/power_stat.cc index dead4a0..764652b 100644 --- a/src/gpgpu-sim/power_stat.cc +++ b/src/gpgpu-sim/power_stat.cc @@ -181,11 +181,11 @@ void power_mem_stat_t::print(FILE *fout) const { total_mem_reads + total_mem_writes); fprintf(fout, "Total memory controller reads: %u\n", total_mem_reads); fprintf(fout, "Total memory controller writes: %u\n", total_mem_writes); - + // TODO: print_stats(require stream ID input) fprintf(fout, "Core cache stats:\n"); - core_cache_stats->print_stats(fout); + core_cache_stats->print_stats(fout, -1); fprintf(fout, "L2 cache stats:\n"); - l2_cache_stats->print_stats(fout); + l2_cache_stats->print_stats(fout, -1); } power_core_stat_t::power_core_stat_t(const shader_core_config *shader_config, diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 9fe4c09..4d4f112 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -57,11 +57,11 @@ mem_fetch *shader_core_mem_fetch_allocator::alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr, - unsigned long long cycle) const { + unsigned long long cycle, unsigned long long streamID) const { mem_access_t access(type, addr, size, wr, m_memory_config->gpgpu_ctx); - mem_fetch *mf = - new mem_fetch(access, NULL, wr ? WRITE_PACKET_SIZE : READ_PACKET_SIZE, -1, - m_core_id, m_cluster_id, m_memory_config, cycle); + mem_fetch *mf = new mem_fetch( + access, NULL, streamID, wr ? WRITE_PACKET_SIZE : READ_PACKET_SIZE, -1, + m_core_id, m_cluster_id, m_memory_config, cycle); return mf; } @@ -70,12 +70,12 @@ mem_fetch *shader_core_mem_fetch_allocator::alloc( const mem_access_byte_mask_t &byte_mask, const mem_access_sector_mask_t §or_mask, unsigned size, bool wr, unsigned long long cycle, unsigned wid, unsigned sid, unsigned tpc, - mem_fetch *original_mf) const { + mem_fetch *original_mf, unsigned long long streamID) const { mem_access_t access(type, addr, size, wr, active_mask, byte_mask, sector_mask, m_memory_config->gpgpu_ctx); mem_fetch *mf = new mem_fetch( - access, NULL, wr ? WRITE_PACKET_SIZE : READ_PACKET_SIZE, wid, m_core_id, - m_cluster_id, m_memory_config, cycle, original_mf); + access, NULL, streamID, wr ? WRITE_PACKET_SIZE : READ_PACKET_SIZE, wid, + m_core_id, m_cluster_id, m_memory_config, cycle, original_mf); return mf; } ///////////////////////////////////////////////////////////////////////////// @@ -178,7 +178,7 @@ void shader_core_ctx::create_front_pipeline() { snprintf(name, STRSIZE, "L1I_%03d", m_sid); m_L1I = new read_only_cache(name, m_config->m_L1I_config, m_sid, get_shader_instruction_cache_id(), m_icnt, - IN_L1I_MISS_QUEUE); + IN_L1I_MISS_QUEUE, OTHER_GPU_CACHE, m_gpu); } void shader_core_ctx::create_schedulers() { @@ -447,7 +447,7 @@ void shader_core_ctx::create_exec_pipeline() { m_ldst_unit = new ldst_unit(m_icnt, m_mem_fetch_allocator, this, &m_operand_collector, m_scoreboard, m_config, - m_memory_config, m_stats, m_sid, m_tpc); + m_memory_config, m_stats, m_sid, m_tpc, m_gpu); m_fu.push_back(m_ldst_unit); m_dispatch_port.push_back(ID_OC_MEM); m_issue_port.push_back(OC_EX_MEM); @@ -567,7 +567,8 @@ void shader_core_ctx::init_warps(unsigned cta_id, unsigned start_thread, start_pc = pc; } - m_warp[i]->init(start_pc, cta_id, i, active_threads, m_dynamic_warp_id); + m_warp[i]->init(start_pc, cta_id, i, active_threads, m_dynamic_warp_id, + kernel.get_streamID()); ++m_dynamic_warp_id; m_not_completed += n_active; ++m_active_warps; @@ -985,8 +986,8 @@ void shader_core_ctx::fetch() { // mem_fetch *mf = m_mem_fetch_allocator->alloc() mem_access_t acc(INST_ACC_R, ppc, nbytes, false, m_gpu->gpgpu_ctx); mem_fetch *mf = new mem_fetch( - acc, NULL /*we don't have an instruction yet*/, READ_PACKET_SIZE, - warp_id, m_sid, m_tpc, m_memory_config, + acc, NULL, m_warp[warp_id]->get_kernel_info()->get_streamID(), + READ_PACKET_SIZE, warp_id, m_sid, m_tpc, m_memory_config, m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle); std::list events; enum cache_request_status status; @@ -1040,10 +1041,10 @@ void shader_core_ctx::issue_warp(register_set &pipe_reg_set, m_warp[warp_id]->ibuffer_free(); assert(next_inst->valid()); **pipe_reg = *next_inst; // static instruction information - (*pipe_reg)->issue(active_mask, warp_id, - m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle, - m_warp[warp_id]->get_dynamic_warp_id(), - sch_id); // dynamic instruction information + (*pipe_reg)->issue( + active_mask, warp_id, m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle, + m_warp[warp_id]->get_dynamic_warp_id(), sch_id, + m_warp[warp_id]->get_streamID()); // dynamic instruction information m_stats->shader_cycle_distro[2 + (*pipe_reg)->active_count()]++; func_exec_inst(**pipe_reg); @@ -2597,7 +2598,7 @@ void ldst_unit::init(mem_fetch_interface *icnt, IN_SHADER_L1T_ROB); m_L1C = new read_only_cache(L1C_name, m_config->m_L1C_config, m_sid, get_shader_constant_cache_id(), icnt, - IN_L1C_MISS_QUEUE); + IN_L1C_MISS_QUEUE, OTHER_GPU_CACHE, m_gpu); m_L1D = NULL; m_mem_rc = NO_RC_FAIL; m_num_writeback_clients = @@ -2613,9 +2614,10 @@ ldst_unit::ldst_unit(mem_fetch_interface *icnt, shader_core_ctx *core, opndcoll_rfu_t *operand_collector, Scoreboard *scoreboard, const shader_core_config *config, const memory_config *mem_config, shader_core_stats *stats, - unsigned sid, unsigned tpc) + unsigned sid, unsigned tpc, gpgpu_sim *gpu) : pipelined_simd_unit(NULL, config, config->smem_latency, core, 0), - m_next_wb(config) { + m_next_wb(config), + m_gpu(gpu) { assert(config->smem_latency > 1); init(icnt, mf_allocator, core, operand_collector, scoreboard, config, mem_config, stats, sid, tpc); @@ -2624,7 +2626,7 @@ ldst_unit::ldst_unit(mem_fetch_interface *icnt, snprintf(L1D_name, STRSIZE, "L1D_%03d", m_sid); m_L1D = new l1_cache(L1D_name, m_config->m_L1D_config, m_sid, get_shader_normal_cache_id(), m_icnt, m_mf_allocator, - IN_L1D_MISS_QUEUE, core->get_gpu()); + IN_L1D_MISS_QUEUE, core->get_gpu(), L1_GPU_CACHE); l1_latency_queue.resize(m_config->m_L1D_config.l1_banks); assert(m_config->m_L1D_config.l1_latency > 0); diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index 92691d3..e658a14 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -120,6 +120,7 @@ class shd_warp_t { m_done_exit = true; m_last_fetch = 0; m_next = 0; + m_streamID = (unsigned long long)-1; // Jin: cdp support m_cdp_latency = 0; @@ -140,8 +141,9 @@ class shd_warp_t { m_ldgdepbar_buf.clear(); } void init(address_type start_pc, unsigned cta_id, unsigned wid, - const std::bitset &active, - unsigned dynamic_warp_id) { + const std::bitset &active, unsigned dynamic_warp_id, + unsigned long long streamID) { + m_streamID = streamID; m_cta_id = cta_id; m_warp_id = wid; m_dynamic_warp_id = dynamic_warp_id; @@ -265,6 +267,7 @@ class shd_warp_t { m_inst_in_pipeline--; } + unsigned long long get_streamID() const { return m_streamID; } unsigned get_cta_id() const { return m_cta_id; } unsigned get_dynamic_warp_id() const { return m_dynamic_warp_id; } @@ -277,6 +280,7 @@ class shd_warp_t { private: static const unsigned IBUFFER_SIZE = 2; class shader_core_ctx *m_shader; + unsigned long long m_streamID; unsigned m_cta_id; unsigned m_warp_id; unsigned m_warp_size; @@ -1345,7 +1349,7 @@ class ldst_unit : public pipelined_simd_unit { shader_core_ctx *core, opndcoll_rfu_t *operand_collector, Scoreboard *scoreboard, const shader_core_config *config, const memory_config *mem_config, class shader_core_stats *stats, - unsigned sid, unsigned tpc); + unsigned sid, unsigned tpc, gpgpu_sim *gpu); // Add a structure to record the LDGSTS instructions, // similar to m_pending_writes, but since LDGSTS does not have a output @@ -1435,6 +1439,7 @@ class ldst_unit : public pipelined_simd_unit { warp_inst_t &inst); mem_stage_stall_type process_memory_access_queue_l1cache(l1_cache *cache, warp_inst_t &inst); + gpgpu_sim *m_gpu; const memory_config *m_memory_config; class mem_fetch_interface *m_icnt; @@ -2025,18 +2030,20 @@ class shader_core_mem_fetch_allocator : public mem_fetch_allocator { m_memory_config = config; } mem_fetch *alloc(new_addr_type addr, mem_access_type type, unsigned size, - bool wr, unsigned long long cycle) const; + bool wr, unsigned long long cycle, + unsigned long long streamID) const; mem_fetch *alloc(new_addr_type addr, mem_access_type type, const active_mask_t &active_mask, const mem_access_byte_mask_t &byte_mask, const mem_access_sector_mask_t §or_mask, unsigned size, bool wr, unsigned long long cycle, unsigned wid, - unsigned sid, unsigned tpc, mem_fetch *original_mf) const; + unsigned sid, unsigned tpc, mem_fetch *original_mf, + unsigned long long streamID) const; mem_fetch *alloc(const warp_inst_t &inst, const mem_access_t &access, unsigned long long cycle) const { warp_inst_t inst_copy = inst; mem_fetch *mf = new mem_fetch( - access, &inst_copy, + access, &inst_copy, inst.get_streamID(), access.is_write() ? WRITE_PACKET_SIZE : READ_PACKET_SIZE, inst.warp_id(), m_core_id, m_cluster_id, m_memory_config, cycle); return mf; diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index f4287d8..42c6981 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -57,7 +57,8 @@ void *gpgpu_sim_thread_sequential(void *ctx_ptr) { ctx->the_gpgpusim->g_the_gpu->cycle(); ctx->the_gpgpusim->g_the_gpu->deadlock_check(); } - ctx->the_gpgpusim->g_the_gpu->print_stats(); + ctx->the_gpgpusim->g_the_gpu->print_stats( + ctx->the_gpgpusim->g_the_gpu->last_streamID); ctx->the_gpgpusim->g_the_gpu->update_stats(); ctx->print_simulation_time(); } @@ -144,7 +145,8 @@ void *gpgpu_sim_thread_concurrent(void *ctx_ptr) { fflush(stdout); } if (sim_cycles) { - ctx->the_gpgpusim->g_the_gpu->print_stats(); + ctx->the_gpgpusim->g_the_gpu->print_stats( + ctx->the_gpgpusim->g_the_gpu->last_streamID); ctx->the_gpgpusim->g_the_gpu->update_stats(); ctx->print_simulation_time(); } diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 0ce3c6a..72f8bb0 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -302,6 +302,14 @@ bool stream_manager::register_finished_kernel(unsigned grid_uid) { void stream_manager::stop_all_running_kernels() { pthread_mutex_lock(&m_lock); + std::vector finished_streams; + std::vector running_kernels = m_gpu->get_running_kernels(); + for (kernel_info_t *k : running_kernels) { + if (k != NULL) { + finished_streams.push_back(k->get_streamID()); + } + } + // Signal m_gpu to stop all running kernels m_gpu->stop_all_running_kernels(); @@ -312,7 +320,9 @@ void stream_manager::stop_all_running_kernels() { } // If any kernels completed, print out the current stats - if (count > 0) m_gpu->print_stats(); + for (unsigned long long streamID : finished_streams) { + m_gpu->print_stats(streamID); + } pthread_mutex_unlock(&m_lock); } -- cgit v1.3 From 752d4e5bf622b0d7c730e3eb2f1b3e3cf91e81fa Mon Sep 17 00:00:00 2001 From: WilliamMTK Date: Wed, 11 Dec 2024 16:52:40 -0500 Subject: Add SST integration into gpgpusim (#44) * Add accommodations to run gpgpusim with SST simulation framework through balar * Output setup_environment options when sourcing * Add SST directive check when creating sim thread * Add sst side test for jenkins * sst-integration: update Jenkinsfile with offical sst-elements repo and fix bugs in pipeline script * sst-integration: direct jenkins to rebuild gpgpusim before testing for sst * sst-integration: fix bugs in sst repos config * sst-integration: let Jenkins rebuilds simulator Since the simulator needs to be configured with both normal mode and sst mode, need to rebuild make target to clean prior runs. * sst-integration: Update Jenkinsfile to source env vars when running balar test * sst-integration: refactor code to remove __SST__ flag * sst-integration: fix a bug that init cluster twice for sst * sst-integration: fix a bug of not sending mem packets to SST * sst-integration: remove sst flags from makefiles and setup_env * sst-integration: add comments to SST changes * sst-integration: remove rebuilding simulator in jenkins when testing for SST * sst-integration: revert simulator build script * Add a function to support querying function argument info for SST * sst-integration: add version detection for vanadis binary * Automated Format * add version detection support for gcc 10+ * sst-integration: add cudaMallocHost for SST * sst-integration: fix a compilation bug * sst-integration: add sst balar unittest CI * sst-integration: specify GPU_ARCH for CI test * sst-integration: use bash for github actions * sst-integration: use https links for sst repos * sst-integration: add SST dependencies to CI config * sst-integration: remove sudo * sst-integration: default to yes for apt install * sst-integration: add manual trigger for github action * sst-integration: remove wrong on event * sst-integration: limit CPU usage for compilation * sst-integration: fix wrong path * sst-integration: use personal repo for testing * sst-integration: remove sst-core source in CI to free space * sst-integration: SST_Cycle use print stats with stream id * Automated Format * sst-integration: check for diskspace and try to clean it * sst-integration: move out of docker image * sst-integration: testing for ci path * sst-integration: fix syntax * sst-integration: pass env vars * sst-integration: set env properly * sst-integration: merge LLVM build and test into same job * sst-integration: fix step order * sst-integration: checkout correct branch for env-setup * sst-integration: remove resourcing gpu apps * sst-integration: revert back to docker github action * sst-integration: enable debug trace for sst testing * sst-integration: resourcing gpu app for env vars * sst-integration: use GPUAPPS_ROOT for path for gpu app * sst-integration: use GPUAPPS_ROOT for path for gpu app * sst-integration: enable parallel ci tests and fix not returning with cudaMallocHostSST * sst-integration: using debug flag for CI run * sst-integration: revert debug ci run * sst-integration: CI skips cuda sdk download and launch multiple jobs * sst-integration: reenable parallel tests * sst-integration: reduce concurrent test thread count * sst-integration: skip long test for github runner * sst-integration: try running CI with single core * sst-integrtion: add callback to SST to check thread sync is done in SST_Cycle() * sst-integration: ignore lookup if already found and add callbacks to SST * Automated Format * sst-integration: add support for indirect texture access * Automated Format * sste-integration: fix up for PR * Automated Format --------- Co-authored-by: purdue-jenkins --- .github/workflows/sst_integration.yml | 80 ++++++++++ Jenkinsfile | 35 +++++ Makefile | 2 + libcuda/cuda_api_object.h | 17 +++ libcuda/cuda_runtime_api.cc | 274 +++++++++++++++++++++++++++++++--- libcuda/gpgpu_context.h | 4 + setup_environment | 18 ++- src/cuda-sim/cuda-sim.cc | 45 ++++-- src/cuda-sim/instructions.cc | 11 ++ src/cuda-sim/ptx_ir.cc | 16 ++ src/cuda-sim/ptx_ir.h | 2 + src/cuda-sim/ptx_sim.h | 3 + src/gpgpu-sim/gpu-cache.cc | 1 + src/gpgpu-sim/gpu-sim.cc | 200 ++++++++++++++++++++++--- src/gpgpu-sim/gpu-sim.h | 139 ++++++++++++++++- src/gpgpu-sim/mem_fetch.cc | 12 +- src/gpgpu-sim/mem_latency_stat.cc | 17 ++- src/gpgpu-sim/shader.cc | 135 +++++++++++++++-- src/gpgpu-sim/shader.h | 101 +++++++++++++ src/gpgpusim_entrypoint.cc | 138 +++++++++++++++-- src/stream_manager.cc | 12 +- 21 files changed, 1176 insertions(+), 86 deletions(-) create mode 100644 .github/workflows/sst_integration.yml (limited to 'src/gpgpu-sim/mem_fetch.cc') diff --git a/.github/workflows/sst_integration.yml b/.github/workflows/sst_integration.yml new file mode 100644 index 0000000..03635db --- /dev/null +++ b/.github/workflows/sst_integration.yml @@ -0,0 +1,80 @@ +# Workflow with cmake build system +name: SST Integration Test + +# Controls when the workflow will run +on: + # Triggers the workflow on push or pull request events but only for the mydev branch + push: + branches-ignore: + - "gh-readonly-queue**" + pull_request: + + # Allows you to run this workflow manually from the Actions tab + workflow_dispatch: + +# A workflow run is made up of one or more jobs that can run sequentially or in parallel +jobs: + build-QV100: + runs-on: ubuntu-latest + defaults: + run: + shell: bash + strategy: + matrix: + # test_type: [simple, medium, long] + test_type: [simple, medium] + container: + image: tgrogers/accel-sim_regress:SST-Integration-Ubuntu-22.04-cuda-11.7-llvm-18.1.8-riscv-gnu-2024.08.06-nightly + env: + CONFIG: QV100 + GPU_ARCH: sm_70 + + # Steps represent a sequence of tasks that will be executed as part of the job + steps: + # Checks-out your repository under $GITHUB_WORKSPACE, so your job can access it + - uses: actions/checkout@v4 + - name: Build GPGPU-Sim in SST mode + run: | + source ./setup_environment + make -j4 + - name: Prepare SST dependencies + run: | + apt install -y openmpi-bin openmpi-common libtool libtool-bin autoconf python3 python3-dev automake build-essential git + # Use personal repo for now + - name: Build SST-Core + run: | + git clone https://github.com/William-An/sst-core.git + cd sst-core + git pull + git checkout devel + ./autogen.sh + ./configure --prefix=`realpath ../sstcore-install` --disable-mpi --disable-mem-pools + make -j4 + make install + cd .. + rm -rf ./sst-core + # Use personal repo for now + - name: Build SST-Elements + run: | + git clone https://github.com/William-An/sst-elements.git + source ./setup_environment + cd sst-elements + git pull + git checkout balar-mmio-vanadis-llvm + ./autogen.sh + ./configure --prefix=`realpath ../sstelements-install` --with-sst-core=`realpath ../sstcore-install` --with-cuda=$CUDA_INSTALL_PATH --with-gpgpusim=$GPGPUSIM_ROOT + make -j4 + make install + # Have to resource the gpu app + # Also fake a SDK since rodinia 2.0 does not need this, speed things up on github + - name: Balar Test + run: | + pip install testtools blessings pygments + source ./setup_environment + mkdir 4.2 + mkdir fake_sdk + export NVIDIA_COMPUTE_SDK_LOCATION=$(readlink -f ./fake_sdk) + source $GPUAPPS_ROOT/src/setup_environment sst + rm -rf 4.2 + rm -f gpucomputingsdk_4.2.9_linux.run + ./sstcore-install/bin/sst-test-elements -w "*balar*${{ matrix.test_type }}*" \ No newline at end of file diff --git a/Jenkinsfile b/Jenkinsfile index f6676bf..4ef467b 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -87,6 +87,41 @@ pipeline { ssh tgrogers@dynamo.ecn.purdue.edu "cd $PLOTDIR && rm -rf latest && cp -r ${BUILD_NUMBER} latest"' } } + stage('sst-core-build') { + steps { + sh 'rm -rf sstcore-install' + sh 'rm -rf sst-core && git clone git@github.com:sstsimulator/sst-core.git' + sh '''#!/bin/bash + cd sst-core + ./autogen.sh + ./configure --prefix=`realpath ../sstcore-install` --disable-mpi --disable-mem-pools + make -j 10 + make install''' + } + } + stage('sst-elements-build') { + steps { + sh 'rm -rf sstelements-install' + sh 'rm -rf sst-elements && git clone git@github.com:sstsimulator/sst-elements.git' + // First sourcing the env_setup and setup_environment script for env vars + sh '''#!/bin/bash + source ./env-setup/11.0_env_setup.sh + source `pwd`/setup_environment + cd sst-elements + ./autogen.sh + ./configure --prefix=`realpath ../sstelements-install` --with-sst-core=`realpath ../sstcore-install` --with-cuda=$CUDA_INSTALL_PATH --with-gpgpusim=$GPGPUSIM_ROOT + make -j 10 + make install''' + } + } + stage('sst balar test') { + steps { + sh '''#!/bin/bash + source ./env-setup/11.0_env_setup.sh + source `pwd`/setup_environment sst + ./sstcore-install/bin/sst-test-elements -p ./sst-elements/src/sst/elements/balar/tests''' + } + } } post { success { diff --git a/Makefile b/Makefile index 82ea399..37dba01 100644 --- a/Makefile +++ b/Makefile @@ -34,6 +34,7 @@ INTERSIM ?= intersim2 include version_detection.mk +# Check for debug ifeq ($(GPGPUSIM_CONFIG), gcc-$(CC_VERSION)/cuda-$(CUDART_VERSION)/debug) export DEBUG=1 else @@ -168,6 +169,7 @@ $(SIM_LIB_DIR)/libcudart.so: makedirs $(LIBS) cudalib if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.0; fi if [ ! -f $(SIM_LIB_DIR)/libcudart.so.10.1 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.10.1; fi if [ ! -f $(SIM_LIB_DIR)/libcudart.so.11.0 ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart.so.11.0; fi + if [ ! -f $(SIM_LIB_DIR)/libcudart_mod.so ]; then ln -s libcudart.so $(SIM_LIB_DIR)/libcudart_mod.so; fi $(SIM_LIB_DIR)/libcudart.dylib: makedirs $(LIBS) cudalib g++ -dynamiclib -Wl,-headerpad_max_install_names,-undefined,dynamic_lookup,-compatibility_version,1.1,-current_version,1.1\ diff --git a/libcuda/cuda_api_object.h b/libcuda/cuda_api_object.h index d292e22..e620e57 100644 --- a/libcuda/cuda_api_object.h +++ b/libcuda/cuda_api_object.h @@ -1,6 +1,7 @@ #ifndef __cuda_api_object_h__ #define __cuda_api_object_h__ +#include #include #include #include @@ -193,9 +194,25 @@ class cuda_runtime_api { // backward pointer class gpgpu_context *gpgpu_ctx; // member function list + + // For SST and other potential simulator interface + void cuobjdumpInit(const char *fn); + void extract_code_using_cuobjdump(const char *fn); + void extract_ptx_files_using_cuobjdump(CUctx_st *context, const char *fn); + + // For running GPGPUSim alone void cuobjdumpInit(); void extract_code_using_cuobjdump(); void extract_ptx_files_using_cuobjdump(CUctx_st *context); + + // Internal functions for the above public methods + void cuobjdumpInit_internal(std::function ctx_extract_code_func); + void extract_code_using_cuobjdump_internal( + CUctx_st *context, std::string &app_binary, + std::function ctx_extract_ptx_func); + void extract_ptx_files_using_cuobjdump_internal(CUctx_st *context, + std::string &app_binary); + std::list pruneSectionList(CUctx_st *context); std::list mergeMatchingSections(std::string identifier); std::list mergeSections(); diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index b540ffd..5dfd3fc 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -109,6 +109,7 @@ #include #include #include +#include #include #include #include @@ -151,6 +152,9 @@ #include #endif +// SST cycle +extern bool SST_Cycle(); + /*DEVICE_BUILTIN*/ struct cudaArray { void *devPtr; @@ -412,6 +416,13 @@ void setCuobjdumpsassfilename( //! processes (e.g. cuobjdump) reading /proc//exe will see the emulator //! executable instead of the application binary. //! +// In SST need the string to pass the binary information +// as we cannot get it from /proc/self/exe +std::string get_app_binary(const char *fn) { + printf("self exe links to: %s\n", fn); + return fn; +} + std::string get_app_binary() { char self_exe_path[1025]; #ifdef __APPLE__ @@ -453,19 +464,25 @@ char *get_app_binary_name(std::string abs_path) { return self_exe_path; } -static int get_app_cuda_version() { +static int get_app_cuda_version_internal(std::string app_binary) { int app_cuda_version = 0; char fname[1024]; snprintf(fname, 1024, "_app_cuda_version_XXXXXX"); int fd = mkstemp(fname); close(fd); + // Weili: Add way to extract CUDA version information from Balar Vanadis + // binary (stored as a const string) std::string app_cuda_version_command = - "ldd " + get_app_binary() + + "ldd " + app_binary + " | grep libcudart.so | sed 's/.*libcudart.so.\\(.*\\) =>.*/\\1/' > " + + fname + " && strings " + app_binary + + " | grep libcudart_vanadis.a | sed " + "'s/.*libcudart_vanadis.a.\\(.*\\)/\\1/' >> " + fname; int res = system(app_cuda_version_command.c_str()); if (res == -1) { - printf("Error - Cannot detect the app's CUDA version.\n"); + printf("Error - Cannot detect the app's CUDA version. Command: %s\n", + app_cuda_version_command.c_str()); exit(1); } FILE *cmd = fopen(fname, "r"); @@ -476,12 +493,24 @@ static int get_app_cuda_version() { } fclose(cmd); if (app_cuda_version == 0) { - printf("Error - Cannot detect the app's CUDA version.\n"); + printf("Error - Cannot detect the app's CUDA version. Command: %s\n", + app_cuda_version_command.c_str()); exit(1); } return app_cuda_version; } +static int get_app_cuda_version(const char *fn) { + // Use for other simulator integration + std::string app_binary = get_app_binary(fn); + return get_app_cuda_version_internal(app_binary); +} + +static int get_app_cuda_version() { + std::string app_binary = get_app_binary(); + return get_app_cuda_version_internal(app_binary); +} + //! Keep track of the association between filename and cubin handle void cuda_runtime_api::cuobjdumpRegisterFatBinary(unsigned int handle, const char *filename, @@ -574,8 +603,11 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimitInternal( return g_last_cudaError = cudaSuccess; } -void **cudaRegisterFatBinaryInternal(void *fatCubin, - gpgpu_context *gpgpu_ctx = NULL) { +// Internal implementation for cudaRegisterFatBiaryInternal +void **cudaRegisterFatBiaryInternal_impl( + void *fatCubin, gpgpu_context *gpgpu_ctx, std::string &app_binary_path, + int app_cuda_version, + std::function ctx_cuobjdumpInit_func) { gpgpu_context *ctx; if (gpgpu_ctx) { ctx = gpgpu_ctx; @@ -606,11 +638,9 @@ void **cudaRegisterFatBinaryInternal(void *fatCubin, // compiled with a newer version of CUDA to run apps compiled with older // versions of CUDA. This is especially useful for PTXPLUS execution. // Skip cuda version check for pytorch application - std::string app_binary_path = get_app_binary(); int pos = app_binary_path.find("python"); if (pos == std::string::npos) { // Not pytorch app : checking cuda version - int app_cuda_version = get_app_cuda_version(); assert( app_cuda_version == CUDART_VERSION / 1000 && "The app must be compiled with same major version as the simulator."); @@ -661,7 +691,7 @@ void **cudaRegisterFatBinaryInternal(void *fatCubin, * then for next calls, only returns the appropriate number */ assert(fat_cubin_handle >= 1); - if (fat_cubin_handle == 1) ctx->api->cuobjdumpInit(); + if (fat_cubin_handle == 1) ctx_cuobjdumpInit_func(ctx); ctx->api->cuobjdumpRegisterFatBinary(fat_cubin_handle, filename, context); return (void **)fat_cubin_handle; @@ -753,6 +783,28 @@ void **cudaRegisterFatBinaryInternal(void *fatCubin, #endif } +void **cudaRegisterFatBinaryInternal(const char *fn, void *fatCubin, + gpgpu_context *gpgpu_ctx = NULL) { + std::string app_binary_path = get_app_binary(fn); + int app_cuda_version = get_app_cuda_version(fn); + auto ctx_cuobjdumpInit = [=](gpgpu_context *ctx) { + ctx->api->cuobjdumpInit(fn); + }; + return cudaRegisterFatBiaryInternal_impl(fatCubin, gpgpu_ctx, app_binary_path, + app_cuda_version, ctx_cuobjdumpInit); +} + +void **cudaRegisterFatBinaryInternal(void *fatCubin, + gpgpu_context *gpgpu_ctx = NULL) { + std::string app_binary_path = get_app_binary(); + int app_cuda_version = get_app_cuda_version(); + auto ctx_cuobjdumpInit = [](gpgpu_context *ctx) { + ctx->api->cuobjdumpInit(); + }; + return cudaRegisterFatBiaryInternal_impl(fatCubin, gpgpu_ctx, app_binary_path, + app_cuda_version, ctx_cuobjdumpInit); +} + void cudaRegisterFunctionInternal(void **fatCubinHandle, const char *hostFun, char *deviceFun, const char *deviceName, int thread_limit, uint3 *tid, uint3 *bid, @@ -1057,6 +1109,24 @@ cudaError_t cudaMallocHostInternal(void **ptr, size_t size, } } +// SST malloc done by vanadis, we just need to record the memory addr +cudaError_t CUDARTAPI cudaMallocHostSSTInternal( + void *addr, size_t size, gpgpu_context *gpgpu_ctx = NULL) { + gpgpu_context *ctx; + if (gpgpu_ctx) { + ctx = gpgpu_ctx; + } else { + ctx = GPGPU_Context(); + } + if (g_debug_execution >= 3) { + announce_call(__my_func__); + } + // track pinned memory size allocated in the host so that same amount of + // memory is also allocated in GPU. + ctx->api->pinned_memory_size[addr] = size; + return g_last_cudaError = cudaSuccess; +} + __host__ cudaError_t CUDARTAPI cudaMallocPitchInternal(void **devPtr, size_t *pitch, size_t width, size_t height, gpgpu_context *gpgpu_ctx = NULL) { @@ -2301,13 +2371,77 @@ cudaDeviceSynchronizeInternal(gpgpu_context *gpgpu_ctx = NULL) { * * *******************************************************************************/ -extern "C" { - /******************************************************************************* * * - * * + * SST Specific functions, used by Balar * * * *******************************************************************************/ + +/** + * @brief Custom function to get CUDA function parameter size and offset + * from PTX parsing result + * + * @param hostFun + * @param index + * @return std::tuple + */ +std::tuple SST_cudaGetParamConfig( + uint64_t hostFun, unsigned index) { + if (g_debug_execution >= 3) { + announce_call(__my_func__); + } + CUctx_st *context = GPGPUSim_Context(GPGPU_Context()); + function_info *entry = context->get_kernel((char *)hostFun); + cudaError_t result = cudaSuccess; + size_t size = 0; + unsigned alignment = 0; + if (index >= entry->num_args()) { + result = cudaErrorAssert; + } else { + std::pair p = entry->get_param_config(index); + size = p.first; + alignment = p.second; + } + return std::tuple(result, size, alignment); +} + +extern "C" { +void SST_receive_mem_reply(unsigned core_id, void *mem_req) { + CUctx_st *context = GPGPUSim_Context(GPGPU_Context()); + static_cast(context->get_device()->get_gpgpu()) + ->SST_receive_mem_reply(core_id, mem_req); + // printf("GPGPU-sim: Recived Request\n"); +} + +bool SST_gpu_core_cycle() { return SST_Cycle(); } + +void SST_gpgpusim_numcores_equal_check(unsigned sst_numcores) { + CUctx_st *context = GPGPUSim_Context(GPGPU_Context()); + static_cast(context->get_device()->get_gpgpu()) + ->SST_gpgpusim_numcores_equal_check(sst_numcores); +} + +uint64_t cudaMallocSST(void **devPtr, size_t size) { + if (g_debug_execution >= 3) { + announce_call(__my_func__); + } + void *test_malloc; + test_malloc = (void *)malloc(size); + void **test_malloc2 = &test_malloc; + CUctx_st *context = GPGPUSim_Context(GPGPU_Context()); + *test_malloc2 = context->get_device()->get_gpgpu()->gpu_malloc(size); + printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n", size, + (unsigned long long)*test_malloc2); + if (g_debug_execution >= 3) + printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n", + size, (unsigned long long)*test_malloc2); + return (uint64_t)*test_malloc2; +} + +__host__ cudaError_t CUDARTAPI cudaMallocHostSST(void *addr, size_t size) { + return cudaMallocHostSSTInternal(addr, size); +} + cudaError_t cudaPeekAtLastError(void) { return g_last_cudaError; } __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) { @@ -2534,6 +2668,7 @@ cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( * * * * *******************************************************************************/ + __host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count) { return cudaMemsetInternal(mem, c, count); } @@ -2754,11 +2889,32 @@ __host__ const char *CUDARTAPI cudaGetErrorString(cudaError_t error) { return strdup(buf); } +// SST specific cuda apis +__host__ cudaError_t CUDARTAPI cudaSetupArgumentSST(uint64_t arg, + uint8_t value[200], + size_t size, + size_t offset) { + void *local_value; + local_value = (void *)malloc(size); + + if (arg) { + memcpy(local_value, (void *)&arg, size); + } else { + memcpy(local_value, value, size); + } + return cudaSetupArgumentInternal(local_value, size, offset); +} + __host__ cudaError_t CUDARTAPI cudaSetupArgument(const void *arg, size_t size, size_t offset) { return cudaSetupArgumentInternal(arg, size, offset); } +// SST specific cuda apis +__host__ cudaError_t CUDARTAPI cudaLaunchSST(uint64_t hostFun) { + return cudaLaunchInternal((char *)hostFun); +} + __host__ cudaError_t CUDARTAPI cudaLaunch(const char *hostFun) { return cudaLaunchInternal(hostFun); } @@ -2933,6 +3089,27 @@ __host__ cudaError_t CUDARTAPI cudaThreadSynchronize(void) { return cudaThreadSynchronizeInternal(); } +__host__ cudaError_t CUDARTAPI cudaThreadSynchronizeSST(void) { + // For SST, perform a one-time check and let SST_Cycle() + // do the polling test and invoke callback to SST + // to signal ThreadSynchonize done + gpgpu_context *ctx = GPGPU_Context(); + if (g_debug_execution >= 3) { + announce_call(__my_func__); + } + + // Called on host side + bool thread_sync_done = ctx->synchronize_check(); + g_last_cudaError = cudaSuccess; + if (thread_sync_done) { + // We are already done, so no need to poll for sync done + ctx->requested_synchronize = false; + return cudaSuccess; + } else { + return cudaErrorNotReady; + } +} + int CUDARTAPI __cudaSynchronizeThreads(void **, void *) { if (g_debug_execution >= 3) { announce_call(__my_func__); @@ -2992,10 +3169,10 @@ __host__ cudaError_t CUDARTAPI cudaGetExportTable( // extracts all ptx files from binary and dumps into // prog_name.unique_no.sm_<>.ptx files -void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context) { +void cuda_runtime_api::extract_ptx_files_using_cuobjdump_internal( + CUctx_st *context, std::string &app_binary) { char command[1000]; char *pytorch_bin = getenv("PYTORCH_BIN"); - std::string app_binary = get_app_binary(); char ptx_list_file_name[1024]; snprintf(ptx_list_file_name, 1024, "_cuobjdump_list_ptx_XXXXXX"); @@ -3062,6 +3239,17 @@ void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context) { } } +void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context, + const char *fn) { + std::string app_binary = get_app_binary(fn); + this->extract_ptx_files_using_cuobjdump_internal(context, app_binary); +} + +void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context) { + std::string app_binary = get_app_binary(); + this->extract_ptx_files_using_cuobjdump_internal(context, app_binary); +} + //! Call cuobjdump to extract everything (-elf -sass -ptx) /*! * This Function extract the whole PTX (for all the files) using cuobjdump @@ -3069,13 +3257,12 @@ void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context) { *with each binary in its own file It is also responsible for extracting the *libraries linked to the binary if the option is enabled * */ -void cuda_runtime_api::extract_code_using_cuobjdump() { - CUctx_st *context = GPGPUSim_Context(gpgpu_ctx); - +void cuda_runtime_api::extract_code_using_cuobjdump_internal( + CUctx_st *context, std::string &app_binary, + std::function ctx_extract_ptx_func) { // prevent the dumping by cuobjdump everytime we execute the code! const char *override_cuobjdump = getenv("CUOBJDUMP_SIM_FILE"); char command[1000]; - std::string app_binary = get_app_binary(); // Running cuobjdump using dynamic link to current process snprintf(command, 1000, "md5sum %s ", app_binary.c_str()); printf("Running md5sum using \"%s\"\n", command); @@ -3090,7 +3277,7 @@ void cuda_runtime_api::extract_code_using_cuobjdump() { // used by ptxas. int result = 0; #if (CUDART_VERSION >= 6000) - extract_ptx_files_using_cuobjdump(context); + ctx_extract_ptx_func(context); return; #endif // TODO: redundant to dump twice. how can it be prevented? @@ -3222,6 +3409,26 @@ void cuda_runtime_api::extract_code_using_cuobjdump() { } } +void cuda_runtime_api::extract_code_using_cuobjdump(const char *fn) { + CUctx_st *context = GPGPUSim_Context(gpgpu_ctx); + std::string app_binary = get_app_binary(fn); + auto ctx_extract_ptx_func = [=](CUctx_st *context) { + extract_ptx_files_using_cuobjdump(context, fn); + }; + extract_code_using_cuobjdump_internal(context, app_binary, + ctx_extract_ptx_func); +} + +void cuda_runtime_api::extract_code_using_cuobjdump() { + CUctx_st *context = GPGPUSim_Context(gpgpu_ctx); + std::string app_binary = get_app_binary(); + auto ctx_extract_ptx_func = [=](CUctx_st *context) { + extract_ptx_files_using_cuobjdump(context); + }; + extract_code_using_cuobjdump_internal(context, app_binary, + ctx_extract_ptx_func); +} + //! Read file into char* // TODO: convert this to C++ streams, will be way cleaner char *readfile(const std::string filename) { @@ -3466,10 +3673,11 @@ cuobjdumpPTXSection *cuda_runtime_api::findPTXSection( } //! Extract the code using cuobjdump and remove unnecessary sections -void cuda_runtime_api::cuobjdumpInit() { +void cuda_runtime_api::cuobjdumpInit_internal( + std::function ctx_extract_code_func) { CUctx_st *context = GPGPUSim_Context(gpgpu_ctx); - extract_code_using_cuobjdump(); // extract all the output of cuobjdump to - // _cuobjdump_*.* + ctx_extract_code_func(); // extract all the output of cuobjdump to + // _cuobjdump_*.* const char *pre_load = getenv("CUOBJDUMP_SIM_FILE"); if (pre_load == NULL || strlen(pre_load) == 0) { cuobjdumpSectionList = pruneSectionList(context); @@ -3477,6 +3685,16 @@ void cuda_runtime_api::cuobjdumpInit() { } } +void cuda_runtime_api::cuobjdumpInit(const char *fn) { + auto ctx_extract_code_func = [=]() { extract_code_using_cuobjdump(fn); }; + cuobjdumpInit_internal(ctx_extract_code_func); +} + +void cuda_runtime_api::cuobjdumpInit() { + auto ctx_extract_code_func = [=]() { extract_code_using_cuobjdump(); }; + cuobjdumpInit_internal(ctx_extract_code_func); +} + //! Either submit PTX for simulation or convert SASS to PTXPlus and submit it void gpgpu_context::cuobjdumpParseBinary(unsigned int handle) { CUctx_st *context = GPGPUSim_Context(this); @@ -3587,6 +3805,10 @@ void gpgpu_context::cuobjdumpParseBinary(unsigned int handle) { extern "C" { +void **CUDARTAPI __cudaRegisterFatBinarySST(const char *fn) { + return cudaRegisterFatBinaryInternal(fn, NULL); +} + void **CUDARTAPI __cudaRegisterFatBinary(void *fatCubin) { if (g_debug_execution >= 3) { announce_call(__my_func__); @@ -3619,6 +3841,14 @@ cudaError_t CUDARTAPI __cudaPopCallConfiguration(dim3 *gridDim, dim3 *blockDim, return g_last_cudaError = cudaSuccess; } +void CUDARTAPI __cudaRegisterFunctionSST(unsigned fatCubinHandle, + uint64_t hostFun, + char deviceFun[512]) { + cudaRegisterFunctionInternal((void **)fatCubinHandle, (const char *)hostFun, + (char *)deviceFun, NULL, NULL, NULL, NULL, NULL, + NULL); +} + void CUDARTAPI __cudaRegisterFunction(void **fatCubinHandle, const char *hostFun, char *deviceFun, const char *deviceName, int thread_limit, diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h index d0cd7c4..5ef2117 100644 --- a/libcuda/gpgpu_context.h +++ b/libcuda/gpgpu_context.h @@ -44,6 +44,9 @@ class gpgpu_context { s_g_pc_to_insn; // a direct mapping from PC to instruction bool debug_tensorcore; + // SST related + bool requested_synchronize = false; + // objects pointers for each file cuda_runtime_api *api; ptxinfo_data *ptxinfo; @@ -54,6 +57,7 @@ class gpgpu_context { ptx_stats *stats; // member function list void synchronize(); + bool synchronize_check(); void exit_simulation(); void print_simulation_time(); int gpgpu_opencl_ptx_sim_main_perf(kernel_info_t *grid); diff --git a/setup_environment b/setup_environment index 3428101..2fac1b9 100644 --- a/setup_environment +++ b/setup_environment @@ -46,7 +46,6 @@ fi CC_VERSION=$(gcc --version | head -1 | awk '{for(i=1;i<=NF;i++){ if(match($i,/^[0-9]+\.[0-9]+\.[0-9]+$/)) {print $i; exit 0}}}') - CUDA_VERSION_STRING=`$CUDA_INSTALL_PATH/bin/nvcc --version | awk '/release/ {print $5;}' | sed 's/,//'`; export CUDA_VERSION_NUMBER=`echo $CUDA_VERSION_STRING | sed 's/\./ /' | awk '{printf("%02u%02u", 10*int($1), 10*$2);}'` if [ $CUDA_VERSION_NUMBER -gt 11100 -o $CUDA_VERSION_NUMBER -lt 2030 ]; then @@ -61,13 +60,18 @@ if [ $CUDA_VERSION_NUMBER -ge 6000 ]; then export CUOBJDUMP_SIM_FILE=jj fi +# Simple configure, loop through all positional arguments +# Default config +export GPGPUSIM_CONFIG=gcc-$CC_VERSION/cuda-$CUDA_VERSION_NUMBER/release -if [ $# = '1' ] ; -then - export GPGPUSIM_CONFIG=gcc-$CC_VERSION/cuda-$CUDA_VERSION_NUMBER/$1 -else - export GPGPUSIM_CONFIG=gcc-$CC_VERSION/cuda-$CUDA_VERSION_NUMBER/release -fi +for opt in $@ +do + if [[ $opt == 'debug' ]] ; then + # Debug mode + echo -n "enabled debug mode " + export GPGPUSIM_CONFIG=gcc-$CC_VERSION/cuda-$CUDA_VERSION_NUMBER/$1 + fi +done export QTINC=/usr/include diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 833d33f..2fd90c0 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1305,7 +1305,12 @@ void function_info::add_param_name_type_size(unsigned index, std::string name, void function_info::add_param_data(unsigned argn, struct gpgpu_ptx_sim_arg *args) { const void *data = args->m_start; - + if (g_debug_execution >= 3) { + if (args->m_nbytes == 4) + printf("ADD_PARAM_DATA %d\n", *((uint32_t *)data)); + else + printf("ADD_PARAM_DATA %p\n", *((void **)data)); + } bool scratchpad_memory_param = false; // Is this parameter in CUDA shared memory or OpenCL local memory @@ -1746,6 +1751,17 @@ static unsigned get_tex_datasize(const ptx_instruction *pI, ptx_thread_info *thread) { const operand_info &src1 = pI->src1(); // the name of the texture std::string texname = src1.name(); + // If indirect access, use register's value as address + // to find the symbol + if (src1.is_reg()) { + const operand_info &dst = pI->dst(); + ptx_reg_t src1_data = + thread->get_operand_value(src1, dst, pI->get_type(), thread, 1); + addr_t sym_addr = src1_data.u64; + symbol *texRef = thread->get_symbol_table()->lookup_by_addr(sym_addr); + assert(texRef != NULL); + texname = texRef->name(); + } /* For programs with many streams, textures can be bound and unbound @@ -2285,15 +2301,24 @@ void cuda_sim::gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, sym_name = g->second; mem_region = global_space; } - if (g_globals.find(hostVar) != g_globals.end()) { - found_sym = true; - sym_name = hostVar; - mem_region = global_space; - } - if (g_constants.find(hostVar) != g_constants.end()) { - found_sym = true; - sym_name = hostVar; - mem_region = const_space; + + // Weili: Only attempt to find symbol as it is a string + // if we could not find it in previously registered variable. + // This will avoid constructing std::string() from hostVar address + // where it is not a string as + // Use of a string naming a variable as the symbol parameter was deprecated in + // CUDA 4.1 and removed in CUDA 5.0. + if (!found_sym) { + if (g_globals.find(hostVar) != g_globals.end()) { + found_sym = true; + sym_name = hostVar; + mem_region = global_space; + } + if (g_constants.find(hostVar) != g_constants.end()) { + found_sym = true; + sym_name = hostVar; + mem_region = const_space; + } } if (!found_sym) { diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index 108de97..843bf0b 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -6055,6 +6055,17 @@ void tex_impl(const ptx_instruction *pI, ptx_thread_info *thread) { // to be fetched std::string texname = src1.name(); + // If indirect access, use register's value as address + // to find the symbol + if (src1.is_reg()) { + ptx_reg_t src1_data = + thread->get_operand_value(src1, dst, pI->get_type(), thread, 1); + addr_t sym_addr = src1_data.u64; + symbol *texRef = thread->get_symbol_table()->lookup_by_addr(sym_addr); + assert(texRef != NULL); + texname = texRef->name(); + } + unsigned to_type = pI->get_type(); unsigned c_type = pI->get_type2(); fflush(stdout); diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc index 1399209..4e500cc 100644 --- a/src/cuda-sim/ptx_ir.cc +++ b/src/cuda-sim/ptx_ir.cc @@ -139,6 +139,22 @@ symbol *symbol_table::lookup(const char *identifier) { return NULL; } +symbol *symbol_table::lookup_by_addr(addr_t addr) { + for (auto it = m_symbols.begin(); it != m_symbols.end(); ++it) { + symbol *sym = it->second; + + // check if symbol has the addr to be found + if ((!sym->is_reg()) && (sym->has_valid_address()) && + (sym->get_address() == addr)) { + return sym; + } + } + if (m_parent) { + return m_parent->lookup_by_addr(addr); + } + return NULL; +} + symbol *symbol_table::add_variable(const char *identifier, const type_info *type, unsigned size, const char *filename, unsigned line) { diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h index d253866..b08a692 100644 --- a/src/cuda-sim/ptx_ir.h +++ b/src/cuda-sim/ptx_ir.h @@ -205,6 +205,7 @@ class symbol { const std::string &name() const { return m_name; } const std::string &decl_location() const { return m_decl_location; } const type_info *type() const { return m_type; } + bool has_valid_address() const { return m_address_valid; } addr_t get_address() const { assert(m_is_label || !m_type->get_key().is_reg()); // todo : other assertions @@ -310,6 +311,7 @@ class symbol_table { void set_ptx_version(float ver, unsigned ext); void set_sm_target(const char *target, const char *ext, const char *ext2); symbol *lookup(const char *identifier); + symbol *lookup_by_addr(addr_t addr); std::string get_scope_name() const { return m_scope_name; } symbol *add_variable(const char *identifier, const type_info *type, unsigned size, const char *filename, unsigned line); diff --git a/src/cuda-sim/ptx_sim.h b/src/cuda-sim/ptx_sim.h index f0c26ef..8eec922 100644 --- a/src/cuda-sim/ptx_sim.h +++ b/src/cuda-sim/ptx_sim.h @@ -459,6 +459,9 @@ class ptx_thread_info { // Jin: get corresponding kernel grid for CDP purpose kernel_info_t &get_kernel() { return m_kernel; } + // Weili: access symbol_table + symbol_table *get_symbol_table() { return m_symbol_table; } + public: addr_t m_last_effective_address; bool m_branch_taken; diff --git a/src/gpgpu-sim/gpu-cache.cc b/src/gpgpu-sim/gpu-cache.cc index cd3c880..0ea9ff6 100644 --- a/src/gpgpu-sim/gpu-cache.cc +++ b/src/gpgpu-sim/gpu-cache.cc @@ -2062,6 +2062,7 @@ enum cache_request_status tex_cache::access(new_addr_type addr, mem_fetch *mf, void tex_cache::cycle() { // send next request to lower level of memory + // TODO: Use different full() for sst_mem_interface? if (!m_request_fifo.empty()) { mem_fetch *mf = m_request_fifo.peek(); if (!m_memport->full(mf->get_ctrl_size(), false)) { diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 5bd4180..b92494b 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -319,6 +319,9 @@ void memory_config::reg_options(class OptionParser *opp) { "elimnate_rw_turnaround i.e set tWTR and tRTW = 0", "0"); option_parser_register(opp, "-icnt_flit_size", OPT_UINT32, &icnt_flit_size, "icnt_flit_size", "32"); + // SST mode activate + option_parser_register(opp, "-SST_mode", OPT_BOOL, &SST_mode, "SST mode", + "0"); m_address_mapping.addrdec_setoption(opp); } @@ -955,6 +958,16 @@ void exec_gpgpu_sim::createSIMTCluster() { m_shader_stats, m_memory_stats); } +// SST get its own simt_cluster +void sst_gpgpu_sim::createSIMTCluster() { + m_cluster = new simt_core_cluster *[m_shader_config->n_simt_clusters]; + for (unsigned i = 0; i < m_shader_config->n_simt_clusters; i++) + m_cluster[i] = + new sst_simt_core_cluster(this, i, m_shader_config, m_memory_config, + m_shader_stats, m_memory_stats); + SST_gpgpu_reply_buffer.resize(m_shader_config->n_simt_clusters); +} + gpgpu_sim::gpgpu_sim(const gpgpu_sim_config &config, gpgpu_context *ctx) : gpgpu_t(config, ctx), m_config(config) { gpgpu_ctx = ctx; @@ -999,26 +1012,29 @@ gpgpu_sim::gpgpu_sim(const gpgpu_sim_config &config, gpgpu_context *ctx) gpu_kernel_time.clear(); - m_memory_partition_unit = - new memory_partition_unit *[m_memory_config->m_n_mem]; - m_memory_sub_partition = - new memory_sub_partition *[m_memory_config->m_n_mem_sub_partition]; - for (unsigned i = 0; i < m_memory_config->m_n_mem; i++) { - m_memory_partition_unit[i] = - new memory_partition_unit(i, m_memory_config, m_memory_stats, this); - for (unsigned p = 0; - p < m_memory_config->m_n_sub_partition_per_memory_channel; p++) { - unsigned submpid = - i * m_memory_config->m_n_sub_partition_per_memory_channel + p; - m_memory_sub_partition[submpid] = - m_memory_partition_unit[i]->get_sub_partition(p); + // TODO: somehow move this logic to the sst_gpgpu_sim constructor? + if (!m_config.is_SST_mode()) { + // Init memory if not in SST mode + m_memory_partition_unit = + new memory_partition_unit *[m_memory_config->m_n_mem]; + m_memory_sub_partition = + new memory_sub_partition *[m_memory_config->m_n_mem_sub_partition]; + for (unsigned i = 0; i < m_memory_config->m_n_mem; i++) { + m_memory_partition_unit[i] = + new memory_partition_unit(i, m_memory_config, m_memory_stats, this); + for (unsigned p = 0; + p < m_memory_config->m_n_sub_partition_per_memory_channel; p++) { + unsigned submpid = + i * m_memory_config->m_n_sub_partition_per_memory_channel + p; + m_memory_sub_partition[submpid] = + m_memory_partition_unit[i]->get_sub_partition(p); + } } - } - - icnt_wrapper_init(); - icnt_create(m_shader_config->n_simt_clusters, - m_memory_config->m_n_mem_sub_partition); + icnt_wrapper_init(); + icnt_create(m_shader_config->n_simt_clusters, + m_memory_config->m_n_mem_sub_partition); + } time_vector_create(NUM_MEM_REQ_STAT); fprintf(stdout, "GPGPU-Sim uArch: performance model initialization complete.\n"); @@ -1037,6 +1053,22 @@ gpgpu_sim::gpgpu_sim(const gpgpu_sim_config &config, gpgpu_context *ctx) m_functional_sim_kernel = NULL; } +void sst_gpgpu_sim::SST_receive_mem_reply(unsigned core_id, void *mem_req) { + assert(core_id < m_shader_config->n_simt_clusters); + mem_fetch *mf = (mem_fetch *)mem_req; + + (SST_gpgpu_reply_buffer[core_id]).push_back(mf); +} + +mem_fetch *sst_gpgpu_sim::SST_pop_mem_reply(unsigned core_id) { + if (SST_gpgpu_reply_buffer[core_id].size() > 0) { + mem_fetch *temp = SST_gpgpu_reply_buffer[core_id].front(); + SST_gpgpu_reply_buffer[core_id].pop_front(); + return temp; + } else + return NULL; +} + int gpgpu_sim::shared_mem_size() const { return m_shader_config->gpgpu_shmem_size; } @@ -1132,6 +1164,26 @@ bool gpgpu_sim::active() { return false; } +bool sst_gpgpu_sim::active() { + if (m_config.gpu_max_cycle_opt && + (gpu_tot_sim_cycle + gpu_sim_cycle) >= m_config.gpu_max_cycle_opt) + return false; + if (m_config.gpu_max_insn_opt && + (gpu_tot_sim_insn + gpu_sim_insn) >= m_config.gpu_max_insn_opt) + return false; + if (m_config.gpu_max_cta_opt && + (gpu_tot_issued_cta >= m_config.gpu_max_cta_opt)) + return false; + if (m_config.gpu_max_completed_cta_opt && + (gpu_completed_cta >= m_config.gpu_max_completed_cta_opt)) + return false; + if (m_config.gpu_deadlock_detect && gpu_deadlock) return false; + for (unsigned i = 0; i < m_shader_config->n_simt_clusters; i++) + if (m_cluster[i]->get_not_completed() > 0) return true; + if (get_more_cta_left()) return true; + return false; +} + void gpgpu_sim::init() { // run a CUDA grid on the GPU microarchitecture simulator gpu_sim_cycle = 0; @@ -2157,6 +2209,11 @@ void gpgpu_sim::cycle() { } } +void sst_gpgpu_sim::cycle() { + SST_cycle(); + return; +} + void shader_core_ctx::dump_warp_state(FILE *fout) const { fprintf(fout, "\n"); fprintf(fout, "per warp functional simulation status:\n"); @@ -2236,3 +2293,110 @@ const shader_core_config *gpgpu_sim::getShaderCoreConfig() { const memory_config *gpgpu_sim::getMemoryConfig() { return m_memory_config; } simt_core_cluster *gpgpu_sim::getSIMTCluster() { return *m_cluster; } + +void sst_gpgpu_sim::SST_gpgpusim_numcores_equal_check(unsigned sst_numcores) { + if (m_shader_config->n_simt_clusters != sst_numcores) { + assert( + "\nSST core is not equal the GPGPU-sim cores. Open gpgpu-sim.config " + "file and ensure n_simt_clusters" + "is the same as SST gpu cores.\n" && + 0); + } else { + printf("\nSST GPU core is equal the GPGPU-sim cores = %d\n", sst_numcores); + } +} + +void sst_gpgpu_sim::SST_cycle() { + // shader core loading (pop from ICNT into core) follows CORE clock + for (unsigned i = 0; i < m_shader_config->n_simt_clusters; i++) + static_cast(m_cluster[i])->icnt_cycle_SST(); + + // L1 cache + shader core pipeline stages + m_power_stats->pwr_mem_stat->core_cache_stats[CURRENT_STAT_IDX].clear(); + for (unsigned i = 0; i < m_shader_config->n_simt_clusters; i++) { + if (m_cluster[i]->get_not_completed() || get_more_cta_left()) { + m_cluster[i]->core_cycle(); + *active_sms += m_cluster[i]->get_n_active_sms(); + } + // Update core icnt/cache stats for GPUWattch + m_cluster[i]->get_icnt_stats( + m_power_stats->pwr_mem_stat->n_simt_to_mem[CURRENT_STAT_IDX][i], + m_power_stats->pwr_mem_stat->n_mem_to_simt[CURRENT_STAT_IDX][i]); + m_cluster[i]->get_cache_stats( + m_power_stats->pwr_mem_stat->core_cache_stats[CURRENT_STAT_IDX]); + } + float temp = 0; + for (unsigned i = 0; i < m_shader_config->num_shader(); i++) { + temp += m_shader_stats->m_pipeline_duty_cycle[i]; + } + temp = temp / m_shader_config->num_shader(); + *average_pipeline_duty_cycle = ((*average_pipeline_duty_cycle) + temp); + // cout<<"Average pipeline duty cycle: "<<*average_pipeline_duty_cycle<= g_single_step)) { + asm("int $03"); + } + gpu_sim_cycle++; + if (g_interactive_debugger_enabled) gpgpu_debug(); + + // McPAT main cycle (interface with McPAT) +#ifdef GPGPUSIM_POWER_MODEL + if (m_config.g_power_simulation_enabled) { + mcpat_cycle(m_config, getShaderCoreConfig(), m_gpgpusim_wrapper, + m_power_stats, m_config.gpu_stat_sample_freq, gpu_tot_sim_cycle, + gpu_sim_cycle, gpu_tot_sim_insn, gpu_sim_insn, + m_config.g_dvfs_enabled); + } +#endif + + issue_block2core(); + + if (!(gpu_sim_cycle % m_config.gpu_stat_sample_freq)) { + time_t days, hrs, minutes, sec; + time_t curr_time; + time(&curr_time); + unsigned long long elapsed_time = + MAX(curr_time - gpgpu_ctx->the_gpgpusim->g_simulation_starttime, 1); + if ((elapsed_time - last_liveness_message_time) >= + m_config.liveness_message_freq) { + days = elapsed_time / (3600 * 24); + hrs = elapsed_time / 3600 - 24 * days; + minutes = elapsed_time / 60 - 60 * (hrs + 24 * days); + sec = elapsed_time - 60 * (minutes + 60 * (hrs + 24 * days)); + + last_liveness_message_time = elapsed_time; + } + visualizer_printstat(); + m_memory_stats->memlatstat_lat_pw(); + if (m_config.gpgpu_runtime_stat && (m_config.gpu_runtime_stat_flag != 0)) { + if (m_config.gpu_runtime_stat_flag & GPU_RSTAT_BW_STAT) { + for (unsigned i = 0; i < m_memory_config->m_n_mem; i++) + m_memory_partition_unit[i]->print_stat(stdout); + printf("maxmrqlatency = %d \n", m_memory_stats->max_mrq_latency); + printf("maxmflatency = %d \n", m_memory_stats->max_mf_latency); + } + if (m_config.gpu_runtime_stat_flag & GPU_RSTAT_SHD_INFO) + shader_print_runtime_stat(stdout); + if (m_config.gpu_runtime_stat_flag & GPU_RSTAT_L1MISS) + shader_print_l1_miss_stat(stdout); + if (m_config.gpu_runtime_stat_flag & GPU_RSTAT_SCHED) + shader_print_scheduler_stat(stdout, false); + } + } + + if (!(gpu_sim_cycle % 20000)) { + // deadlock detection + if (m_config.gpu_deadlock_detect && gpu_sim_insn == last_gpu_sim_insn) { + gpu_deadlock = true; + } else { + last_gpu_sim_insn = gpu_sim_insn; + } + } + try_snap_shot(gpu_sim_cycle); + spill_log_to_file(stdout, 0, gpu_sim_cycle); + +#if (CUDART_VERSION >= 5000) + // launch device kernel + gpgpu_ctx->device_runtime->launch_one_device_kernel(); +#endif +} diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 8e81451..d0c2a17 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -69,6 +69,38 @@ class gpgpu_context; extern tr1_hash_map address_random_interleaving; +// SST communication functions +/** + * @brief Check if SST requests buffer is full + * + * @param core_id + * @return true + * @return false + */ +extern bool is_SST_buffer_full(unsigned core_id); + +/** + * @brief Send loads to SST memory backend + * + * @param core_id + * @param address + * @param size + * @param mem_req + */ +extern void send_read_request_SST(unsigned core_id, uint64_t address, + size_t size, void *mem_req); + +/** + * @brief Send stores to SST memory backend + * + * @param core_id + * @param address + * @param size + * @param mem_req + */ +extern void send_write_request_SST(unsigned core_id, uint64_t address, + size_t size, void *mem_req); + enum dram_ctrl_t { DRAM_FIFO = 0, DRAM_FRFCFS = 1 }; enum hw_perf_t { @@ -274,6 +306,14 @@ class memory_config { } void reg_options(class OptionParser *opp); + /** + * @brief Check if the config script is in SST mode + * + * @return true + * @return false + */ + bool is_SST_mode() const { return SST_mode; } + bool m_valid; mutable l2_cache_config m_L2_config; bool m_L2_texure_only; @@ -351,7 +391,7 @@ class memory_config { unsigned write_low_watermark; bool m_perf_sim_memcpy; bool simple_dram_model; - + bool SST_mode; gpgpu_context *gpgpu_ctx; }; @@ -398,6 +438,15 @@ class gpgpu_sim_config : public power_config, unsigned num_shader() const { return m_shader_config.num_shader(); } unsigned num_cluster() const { return m_shader_config.n_simt_clusters; } unsigned get_max_concurrent_kernel() const { return max_concurrent_kernel; } + + /** + * @brief Check if we are in SST mode + * + * @return true + * @return false + */ + bool is_SST_mode() const { return m_memory_config.SST_mode; } + unsigned checkpoint_option; size_t stack_limit() const { return stack_size_limit; } @@ -462,6 +511,7 @@ class gpgpu_sim_config : public power_config, unsigned long long liveness_message_freq; friend class gpgpu_sim; + friend class sst_gpgpu_sim; }; struct occupancy_stats { @@ -600,10 +650,18 @@ class gpgpu_sim : public gpgpu_t { void hit_watchpoint(unsigned watchpoint_num, ptx_thread_info *thd, const ptx_instruction *pI); + /** + * @brief Check if we are in SST mode + * + * @return true + * @return false + */ + bool is_SST_mode() { return m_config.is_SST_mode(); } + // backward pointer class gpgpu_context *gpgpu_ctx; - private: + protected: // clocks void reinit_clock_domains(void); int next_clock_domain(void); @@ -715,7 +773,7 @@ class gpgpu_sim : public gpgpu_t { void set_cache_config(std::string kernel_name); // Jin: functional simulation for CDP - private: + protected: // set by stream operation every time a functoinal simulation is done bool m_functional_sim; kernel_info_t *m_functional_sim_kernel; @@ -748,4 +806,79 @@ class exec_gpgpu_sim : public gpgpu_sim { virtual void createSIMTCluster(); }; +/** + * @brief A GPGPUSim class customized to SST Balar interfacing + * + */ +class sst_gpgpu_sim : public gpgpu_sim { + public: + sst_gpgpu_sim(const gpgpu_sim_config &config, gpgpu_context *ctx) + : gpgpu_sim(config, ctx) { + createSIMTCluster(); + } + + // SST memory handling + std::vector> + SST_gpgpu_reply_buffer; /** SST mem response queue */ + + /** + * @brief Receive mem request's response from SST and put + * it in a buffer (SST_gpgpu_reply_buffer) + * + * @param core_id + * @param mem_req + */ + void SST_receive_mem_reply(unsigned core_id, void *mem_req); + + /** + * @brief Pop the head of the buffer queue to get the + * memory response + * + * @param core_id + * @return mem_fetch* + */ + mem_fetch *SST_pop_mem_reply(unsigned core_id); + + virtual void createSIMTCluster(); + + // SST Balar interfacing + /** + * @brief Advance core and collect stats + * + */ + void SST_cycle(); + + /** + * @brief Wrapper of SST_cycle() + * + */ + void cycle(); + + /** + * @brief Whether the GPU is active, removed test for + * memory system since that is handled in SST + * + * @return true + * @return false + */ + bool active(); + + /** + * @brief SST mode use SST memory system instead, so the memcpy + * is empty here + * + * @param dst_start_addr + * @param count + */ + void perf_memcpy_to_gpu(size_t dst_start_addr, size_t count){}; + + /** + * @brief Check if the SST config matches up with the + * gpgpusim.config in core number + * + * @param sst_numcores SST core count + */ + void SST_gpgpusim_numcores_equal_check(unsigned sst_numcores); +}; + #endif diff --git a/src/gpgpu-sim/mem_fetch.cc b/src/gpgpu-sim/mem_fetch.cc index 7211a7d..809c920 100644 --- a/src/gpgpu-sim/mem_fetch.cc +++ b/src/gpgpu-sim/mem_fetch.cc @@ -54,9 +54,15 @@ mem_fetch::mem_fetch(const mem_access_t &access, const warp_inst_t *inst, m_sid = sid; m_tpc = tpc; m_wid = wid; - config->m_address_mapping.addrdec_tlx(access.get_addr(), &m_raw_addr); - m_partition_addr = - config->m_address_mapping.partition_address(access.get_addr()); + + if (!config->is_SST_mode()) { + // In SST memory model, the SST memory hierarchy is + // responsible to generate the correct address mapping + config->m_address_mapping.addrdec_tlx(access.get_addr(), &m_raw_addr); + m_partition_addr = + config->m_address_mapping.partition_address(access.get_addr()); + } + m_type = m_access.is_write() ? WRITE_REQUEST : READ_REQUEST; m_timestamp = cycle; m_timestamp2 = 0; diff --git a/src/gpgpu-sim/mem_latency_stat.cc b/src/gpgpu-sim/mem_latency_stat.cc index 63d7ee8..c77a686 100644 --- a/src/gpgpu-sim/mem_latency_stat.cc +++ b/src/gpgpu-sim/mem_latency_stat.cc @@ -203,7 +203,15 @@ unsigned memory_stats_t::memlatstat_done(mem_fetch *mf) { } void memory_stats_t::memlatstat_read_done(mem_fetch *mf) { - if (m_memory_config->gpgpu_memlatency_stat) { + if (m_memory_config->SST_mode) { + // in SST mode, we just calculate mem latency + unsigned mf_latency; + mf_latency = + (m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle) - mf->get_timestamp(); + num_mfs++; + mf_total_lat += mf_latency; + if (mf_latency > max_mf_latency) max_mf_latency = mf_latency; + } else if (m_memory_config->gpgpu_memlatency_stat) { unsigned mf_latency = memlatstat_done(mf); if (mf_latency > mf_max_lat_table[mf->get_tlx_addr().chip][mf->get_tlx_addr().bk]) @@ -273,7 +281,12 @@ void memory_stats_t::memlatstat_print(unsigned n_mem, unsigned gpu_mem_n_bk) { unsigned max_bank_accesses, min_bank_accesses, max_chip_accesses, min_chip_accesses; - if (m_memory_config->gpgpu_memlatency_stat) { + if (m_memory_config->SST_mode) { + // in SST mode, we just calculate mem latency + printf("max_mem_SST_latency = %d \n", max_mf_latency); + if (num_mfs) + printf("average_mf_SST_latency = %lld \n", mf_total_lat / num_mfs); + } else if (m_memory_config->gpgpu_memlatency_stat) { printf("maxmflatency = %d \n", max_mf_latency); printf("max_icnt2mem_latency = %d \n", max_icnt2mem_latency); printf("maxmrqlatency = %d \n", max_mrq_latency); diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 4d4f112..7482e0e 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -162,7 +162,10 @@ void shader_core_ctx::create_front_pipeline() { } // m_icnt = new shader_memory_interface(this,cluster); - if (m_config->gpgpu_perfect_mem) { + if (m_memory_config->SST_mode) { + m_icnt = new sst_memory_interface( + this, static_cast(m_cluster)); + } else if (m_config->gpgpu_perfect_mem) { m_icnt = new perfect_memory_interface(this, m_cluster); } else { m_icnt = new shader_memory_interface(this, m_cluster); @@ -2281,7 +2284,15 @@ bool ldst_unit::memory_cycle(warp_inst_t &inst, inst.is_store() ? WRITE_PACKET_SIZE : READ_PACKET_SIZE; unsigned size = access.get_size() + control_size; // printf("Interconnect:Addr: %x, size=%d\n",access.get_addr(),size); - if (m_icnt->full(size, inst.is_store() || inst.isatomic())) { + if (m_memory_config->SST_mode && + (static_cast(m_icnt)->full( + size, inst.is_store() || inst.isatomic(), access.get_type()))) { + // SST need mf type here + // Cast it to sst_memory_interface pointer first as this full() method + // is not a virtual method in parent class + stall_cond = ICNT_RC_FAIL; + } else if (!m_memory_config->SST_mode && + (m_icnt->full(size, inst.is_store() || inst.isatomic()))) { stall_cond = ICNT_RC_FAIL; } else { mem_fetch *mf = @@ -2846,7 +2857,10 @@ void ldst_unit::cycle() { } } else { if (mf->get_type() == WRITE_ACK || - (m_config->gpgpu_perfect_mem && mf->get_is_write())) { + ((m_config->gpgpu_perfect_mem || m_memory_config->SST_mode) && + mf->get_is_write())) { + // SST memory is handled by SST mem hierarchy + // Perfect mem m_core->store_ack(mf); m_response_fifo.pop_front(); delete mf; @@ -4020,7 +4034,8 @@ void shader_core_ctx::accept_ldst_unit_response(mem_fetch *mf) { void shader_core_ctx::store_ack(class mem_fetch *mf) { assert(mf->get_type() == WRITE_ACK || - (m_config->gpgpu_perfect_mem && mf->get_is_write())); + ((m_config->gpgpu_perfect_mem || m_memory_config->SST_mode) && + mf->get_is_write())); unsigned warp_id = mf->get_wid(); m_warp[warp_id]->dec_store_req(); } @@ -4573,7 +4588,46 @@ bool simt_core_cluster::icnt_injection_buffer_full(unsigned size, bool write) { return !::icnt_has_buffer(m_cluster_id, request_size); } +bool sst_simt_core_cluster::SST_injection_buffer_full(unsigned size, bool write, + mem_access_type type) { + switch (type) { + case CONST_ACC_R: + case INST_ACC_R: { + return response_queue_full(); + break; + } + default: { + return ::is_SST_buffer_full(m_cluster_id); + break; + } + } +} + void simt_core_cluster::icnt_inject_request_packet(class mem_fetch *mf) { + // Update stats based on mf type + update_icnt_stats(mf); + + // The packet size varies depending on the type of request: + // - For write request and atomic request, the packet contains the data + // - For read request (i.e. not write nor atomic), the packet only has control + // metadata + unsigned int packet_size = mf->size(); + if (!mf->get_is_write() && !mf->isatomic()) { + packet_size = mf->get_ctrl_size(); + } + m_stats->m_outgoing_traffic_stats->record_traffic(mf, packet_size); + unsigned destination = mf->get_sub_partition_id(); + mf->set_status(IN_ICNT_TO_MEM, + m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle); + if (!mf->get_is_write() && !mf->isatomic()) + ::icnt_push(m_cluster_id, m_config->mem2device(destination), (void *)mf, + mf->get_ctrl_size()); + else + ::icnt_push(m_cluster_id, m_config->mem2device(destination), (void *)mf, + mf->size()); +} + +void simt_core_cluster::update_icnt_stats(class mem_fetch *mf) { // stats if (mf->get_is_write()) m_stats->made_write_mfs++; @@ -4618,6 +4672,12 @@ void simt_core_cluster::icnt_inject_request_packet(class mem_fetch *mf) { default: assert(0); } +} + +void sst_simt_core_cluster::icnt_inject_request_packet_to_SST( + class mem_fetch *mf) { + // Update stats + update_icnt_stats(mf); // The packet size varies depending on the type of request: // - For write request and atomic request, the packet contains the data @@ -4628,15 +4688,25 @@ void simt_core_cluster::icnt_inject_request_packet(class mem_fetch *mf) { packet_size = mf->get_ctrl_size(); } m_stats->m_outgoing_traffic_stats->record_traffic(mf, packet_size); - unsigned destination = mf->get_sub_partition_id(); mf->set_status(IN_ICNT_TO_MEM, m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle); - if (!mf->get_is_write() && !mf->isatomic()) - ::icnt_push(m_cluster_id, m_config->mem2device(destination), (void *)mf, - mf->get_ctrl_size()); - else - ::icnt_push(m_cluster_id, m_config->mem2device(destination), (void *)mf, - mf->size()); + switch (mf->get_access_type()) { + case CONST_ACC_R: + case INST_ACC_R: { + push_response_fifo(mf); + break; + } + default: { + if (!mf->get_is_write() && !mf->isatomic()) + ::send_read_request_SST(m_cluster_id, mf->get_addr(), + mf->get_data_size(), (void *)mf); + else + ::send_write_request_SST(m_cluster_id, mf->get_addr(), + mf->get_data_size(), (void *)mf); + + break; + } + } } void simt_core_cluster::icnt_cycle() { @@ -4678,6 +4748,49 @@ void simt_core_cluster::icnt_cycle() { } } +void sst_simt_core_cluster::icnt_cycle_SST() { + if (!m_response_fifo.empty()) { + mem_fetch *mf = m_response_fifo.front(); + unsigned cid = m_config->sid_to_cid(mf->get_sid()); + if (mf->get_access_type() == INST_ACC_R) { + // instruction fetch response + if (!m_core[cid]->fetch_unit_response_buffer_full()) { + m_response_fifo.pop_front(); + m_core[cid]->accept_fetch_response(mf); + } + } else { + // data response + if (!m_core[cid]->ldst_unit_response_buffer_full()) { + m_response_fifo.pop_front(); + m_memory_stats->memlatstat_read_done(mf); + m_core[cid]->accept_ldst_unit_response(mf); + } + } + } + + // pop from SST buffers + if (m_response_fifo.size() < m_config->n_simt_ejection_buffer_size) { + mem_fetch *mf = (mem_fetch *)(static_cast(get_gpu()) + ->SST_pop_mem_reply(m_cluster_id)); + if (!mf) return; + assert(mf->get_tpc() == m_cluster_id); + + // do atomic here + // For now, we execute atomic when the mem reply comes back + // This needs to be validated + if (mf && mf->isatomic()) mf->do_atomic(); + + unsigned int packet_size = + (mf->get_is_write()) ? mf->get_ctrl_size() : mf->size(); + m_stats->m_incoming_traffic_stats->record_traffic(mf, packet_size); + mf->set_status(IN_CLUSTER_TO_SHADER_QUEUE, + m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle); + // m_memory_stats->memlatstat_read_done(mf,m_shader_config->max_warps_per_shader); + m_response_fifo.push_back(mf); + m_stats->n_mem_to_simt[m_cluster_id] += mf->get_num_flits(false); + } +} + void simt_core_cluster::get_pdom_stack_top_info(unsigned sid, unsigned tid, unsigned *pc, unsigned *rpc) const { diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index e658a14..ee10af6 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -2015,6 +2015,7 @@ class shader_core_stats : public shader_core_stats_pod { friend class shader_core_ctx; friend class ldst_unit; friend class simt_core_cluster; + friend class sst_simt_core_cluster; friend class scheduler_unit; friend class TwoLevelScheduler; friend class LooseRoundRobbinScheduler; @@ -2624,6 +2625,7 @@ class simt_core_cluster { void cache_invalidate(); bool icnt_injection_buffer_full(unsigned size, bool write); void icnt_inject_request_packet(class mem_fetch *mf); + void update_icnt_stats(class mem_fetch *mf); // for perfect memory interface bool response_queue_full() { @@ -2685,6 +2687,50 @@ class exec_simt_core_cluster : public simt_core_cluster { virtual void create_shader_core_ctx(); }; +/** + * @brief SST cluster class + * + */ +class sst_simt_core_cluster : public exec_simt_core_cluster { + public: + sst_simt_core_cluster(class gpgpu_sim *gpu, unsigned cluster_id, + const shader_core_config *config, + const memory_config *mem_config, + class shader_core_stats *stats, + class memory_stats_t *mstats) + : exec_simt_core_cluster(gpu, cluster_id, config, mem_config, stats, + mstats) {} + + /** + * @brief Check if SST memory request injection + * buffer is full by using extern + * function is_SST_buffer_full() + * defined in Balar + * + * @param size + * @param write + * @param type + * @return true + * @return false + */ + bool SST_injection_buffer_full(unsigned size, bool write, + mem_access_type type); + + /** + * @brief Send memory request packets to SST + * memory + * + * @param mf + */ + void icnt_inject_request_packet_to_SST(class mem_fetch *mf); + + /** + * @brief Advance ICNT between core and SST + * + */ + void icnt_cycle_SST(); +}; + class shader_memory_interface : public mem_fetch_interface { public: shader_memory_interface(shader_core_ctx *core, simt_core_cluster *cluster) { @@ -2725,6 +2771,61 @@ class perfect_memory_interface : public mem_fetch_interface { simt_core_cluster *m_cluster; }; +/** + * @brief SST memory interface + * + */ +class sst_memory_interface : public mem_fetch_interface { + public: + sst_memory_interface(shader_core_ctx *core, sst_simt_core_cluster *cluster) { + m_core = core; + m_cluster = cluster; + } + /** + * @brief For constant, inst, tex cache access + * + * @param size + * @param write + * @return true + * @return false + */ + virtual bool full(unsigned size, bool write) const { + assert(false && "Use the full() method with access type instead!"); + return true; + } + + /** + * @brief With SST, the core will direct all mem access except for + * constant, tex, and inst reads to SST mem system + * (i.e. not modeling constant mem right now), thus + * requiring the mem_access_type information to be passed in + * + * @param size + * @param write + * @param type + * @return true + * @return false + */ + bool full(unsigned size, bool write, mem_access_type type) const { + return m_cluster->SST_injection_buffer_full(size, write, type); + } + + /** + * @brief Push memory request to SST memory system and + * update stats + * + * @param mf + */ + virtual void push(mem_fetch *mf) { + m_core->inc_simt_to_mem(mf->get_num_flits(true)); + m_cluster->icnt_inject_request_packet_to_SST(mf); + } + + private: + shader_core_ctx *m_core; + sst_simt_core_cluster *m_cluster; +}; + inline int scheduler_unit::get_sid() const { return m_shader->get_sid(); } #endif /* SHADER_H */ diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 42c6981..839fef6 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -43,6 +43,20 @@ static int sg_argc = 3; static const char *sg_argv[] = {"", "-config", "gpgpusim.config"}; +// Help funcs to avoid multiple '->' for SST +GPGPUsim_ctx *GPGPUsim_ctx_ptr() { return GPGPU_Context()->the_gpgpusim; } + +class sst_gpgpu_sim *g_the_gpu() { + return static_cast(GPGPUsim_ctx_ptr()->g_the_gpu); +} + +class stream_manager *g_stream_manager() { + return GPGPUsim_ctx_ptr()->g_stream_manager; +} + +// SST callback +extern void SST_callback_cudaThreadSynchronize_done(); + void *gpgpu_sim_thread_sequential(void *ctx_ptr) { gpgpu_context *ctx = (gpgpu_context *)ctx_ptr; // at most one kernel running at a time @@ -169,6 +183,75 @@ void *gpgpu_sim_thread_concurrent(void *ctx_ptr) { return NULL; } +bool sst_sim_cycles = false; + +bool SST_Cycle() { + // Check if Synchronize is done when SST previously requested + // cudaThreadSynchronize + if (GPGPU_Context()->requested_synchronize && + ((g_stream_manager()->empty() && !GPGPUsim_ctx_ptr()->g_sim_active) || + GPGPUsim_ctx_ptr()->g_sim_done)) { + SST_callback_cudaThreadSynchronize_done(); + GPGPU_Context()->requested_synchronize = false; + } + + if (g_stream_manager()->empty_protected() && + !GPGPUsim_ctx_ptr()->g_sim_done && !g_the_gpu()->active()) { + GPGPUsim_ctx_ptr()->g_sim_active = false; + // printf("stream is empty %d \n", g_stream_manager->empty()); + return false; + } + + if (g_stream_manager()->operation(&sst_sim_cycles) && + !g_the_gpu()->active()) { + if (sst_sim_cycles) { + sst_sim_cycles = false; + } + return false; + } + + // printf("GPGPU-Sim: Give GPU Cycle\n"); + GPGPUsim_ctx_ptr()->g_sim_active = true; + + // functional simulation + if (g_the_gpu()->is_functional_sim()) { + kernel_info_t *kernel = g_the_gpu()->get_functional_kernel(); + assert(kernel); + GPGPUsim_ctx_ptr()->gpgpu_ctx->func_sim->gpgpu_cuda_ptx_sim_main_func( + *kernel); + g_the_gpu()->finish_functional_sim(kernel); + } + + // performance simulation + if (g_the_gpu()->active()) { + g_the_gpu()->SST_cycle(); + sst_sim_cycles = true; + g_the_gpu()->deadlock_check(); + } else { + if (g_the_gpu()->cycle_insn_cta_max_hit()) { + g_stream_manager()->stop_all_running_kernels(); + GPGPUsim_ctx_ptr()->g_sim_done = true; + GPGPUsim_ctx_ptr()->g_sim_active = false; + GPGPUsim_ctx_ptr()->break_limit = true; + } + } + + if (!g_the_gpu()->active()) { + g_the_gpu()->print_stats(GPGPUsim_ctx_ptr()->g_the_gpu->last_streamID); + g_the_gpu()->update_stats(); + GPGPU_Context()->print_simulation_time(); + } + + if (GPGPUsim_ctx_ptr()->break_limit) { + printf( + "GPGPU-Sim: ** break due to reaching the maximum cycles (or " + "instructions) **\n"); + return true; + } + + return false; +} + void gpgpu_context::synchronize() { printf("GPGPU-Sim: synchronize waiting for inactive GPU simulation\n"); the_gpgpusim->g_stream_manager->print(stdout); @@ -187,6 +270,27 @@ void gpgpu_context::synchronize() { // sem_post(&g_sim_signal_start); } +bool gpgpu_context::synchronize_check() { + // printf("GPGPU-Sim: synchronize checking for inactive GPU simulation\n"); + requested_synchronize = true; + the_gpgpusim->g_stream_manager->print(stdout); + fflush(stdout); + // sem_wait(&g_sim_signal_finish); + bool done = false; + pthread_mutex_lock(&(the_gpgpusim->g_sim_lock)); + done = (the_gpgpusim->g_stream_manager->empty() && + !the_gpgpusim->g_sim_active) || + the_gpgpusim->g_sim_done; + pthread_mutex_unlock(&(the_gpgpusim->g_sim_lock)); + if (done) { + printf( + "GPGPU-Sim: synchronize checking: detected inactive GPU simulation " + "thread\n"); + } + fflush(stdout); + return done; +} + void gpgpu_context::exit_simulation() { the_gpgpusim->g_sim_done = true; printf("GPGPU-Sim: exit_simulation called\n"); @@ -220,8 +324,14 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf() { assert(setlocale(LC_NUMERIC, "C")); the_gpgpusim->g_the_gpu_config->init(); - the_gpgpusim->g_the_gpu = - new exec_gpgpu_sim(*(the_gpgpusim->g_the_gpu_config), this); + if (the_gpgpusim->g_the_gpu_config->is_SST_mode()) { + // Create SST specific GPGPUSim + the_gpgpusim->g_the_gpu = + new sst_gpgpu_sim(*(the_gpgpusim->g_the_gpu_config), this); + } else { + the_gpgpusim->g_the_gpu = + new exec_gpgpu_sim(*(the_gpgpusim->g_the_gpu_config), this); + } the_gpgpusim->g_stream_manager = new stream_manager( (the_gpgpusim->g_the_gpu), func_sim->g_cuda_launch_blocking); @@ -237,12 +347,17 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf() { void gpgpu_context::start_sim_thread(int api) { if (the_gpgpusim->g_sim_done) { the_gpgpusim->g_sim_done = false; - if (api == 1) { - pthread_create(&(the_gpgpusim->g_simulation_thread), NULL, - gpgpu_sim_thread_concurrent, (void *)this); + if (the_gpgpusim->g_the_gpu_config->is_SST_mode()) { + // Do not create concurrent thread in SST mode + g_the_gpu()->init(); } else { - pthread_create(&(the_gpgpusim->g_simulation_thread), NULL, - gpgpu_sim_thread_sequential, (void *)this); + if (api == 1) { + pthread_create(&(the_gpgpusim->g_simulation_thread), NULL, + gpgpu_sim_thread_concurrent, (void *)this); + } else { + pthread_create(&(the_gpgpusim->g_simulation_thread), NULL, + gpgpu_sim_thread_sequential, (void *)this); + } } } } @@ -266,8 +381,13 @@ void gpgpu_context::print_simulation_time() { const unsigned cycles_per_sec = (unsigned)(the_gpgpusim->g_the_gpu->gpu_tot_sim_cycle / difference); printf("gpgpu_simulation_rate = %u (cycle/sec)\n", cycles_per_sec); - printf("gpgpu_silicon_slowdown = %ux\n", - the_gpgpusim->g_the_gpu->shader_clock() * 1000 / cycles_per_sec); + + if (cycles_per_sec == 0) { + printf("gpgpu_silicon_slowdown = Nan\n"); + } else { + printf("gpgpu_silicon_slowdown = %ux\n", + the_gpgpusim->g_the_gpu->shader_clock() * 1000 / cycles_per_sec); + } fflush(stdout); } diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 72f8bb0..b974791 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -34,6 +34,12 @@ unsigned CUstream_st::sm_next_stream_uid = 0; +// SST memcpy callbacks +extern void SST_callback_memcpy_H2D_done(); +extern void SST_callback_memcpy_D2H_done(); +extern void SST_callback_memcpy_to_symbol_done(); +extern void SST_callback_memcpy_from_symbol_done(); + CUstream_st::CUstream_st() { m_pending = false; m_uid = sm_next_stream_uid++; @@ -122,11 +128,13 @@ bool stream_operation::do_operation(gpgpu_sim *gpu) { if (g_debug_execution >= 3) printf("memcpy host-to-device\n"); 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(); 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(); break; case stream_memcpy_device_to_device: if (g_debug_execution >= 3) printf("memcpy device-to-device\n"); @@ -138,12 +146,14 @@ bool stream_operation::do_operation(gpgpu_sim *gpu) { gpu->gpgpu_ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( m_symbol, m_host_address_src, m_cnt, m_offset, 1, gpu); m_stream->record_next_done(); + if (gpu->is_SST_mode()) SST_callback_memcpy_to_symbol_done(); break; case stream_memcpy_from_symbol: if (g_debug_execution >= 3) printf("memcpy from symbol\n"); gpu->gpgpu_ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( m_symbol, m_host_address_dst, m_cnt, m_offset, 0, gpu); m_stream->record_next_done(); + if (gpu->is_SST_mode()) SST_callback_memcpy_from_symbol_done(); break; case stream_kernel_launch: if (m_sim_mode) { // Functional Sim @@ -472,7 +482,7 @@ void stream_manager::push(stream_operation op) { } if (g_debug_execution >= 3) print_impl(stdout); pthread_mutex_unlock(&m_lock); - if (m_cuda_launch_blocking || stream == NULL) { + if (!m_gpu->is_SST_mode() && (m_cuda_launch_blocking || stream == NULL)) { unsigned int wait_amount = 100; unsigned int wait_cap = 100000; // 100ms while (!empty()) { -- cgit v1.3