Merge branch 'dev' of https://github.com/gpgpu-sim/gpgpu-sim_distribution into dev-private

7 files changed, 222 insertions, 6 deletions

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 97396ac..c70a570 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -1501,6 +1501,17 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun )
 		sscanf(mode,"%u", &g_ptx_sim_mode);
 	gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
 	kernel_config config = g_cuda_launch_stack.back();
+	{
+		dim3 gridDim = config.grid_dim();
+		dim3 blockDim = config.block_dim();
+		if (gridDim.x * gridDim.y * gridDim.z == 0 || blockDim.x * blockDim.y * blockDim.z == 0)
+		{
+			//can't launch
+			printf("can't launch a empty kernel\n");
+			g_cuda_launch_stack.pop_back();
+			return g_last_cudaError = cudaErrorInvalidConfiguration;
+		}
+	}
 	struct CUstream_st *stream = config.get_stream();
 	if(g_stream_manager->is_blocking())
 		stream = 0;
diff --git a/src/gpgpu-sim/gpu-cache.cc b/src/gpgpu-sim/gpu-cache.cc
index 565fae1..370f6e6 100644
--- a/src/gpgpu-sim/gpu-cache.cc
+++ b/src/gpgpu-sim/gpu-cache.cc
@@ -600,9 +600,11 @@ 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; 
@@ -623,6 +625,15 @@ void cache_stats::clear(){
     m_cache_fill_port_busy_cycles = 0; 
 }
 
+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);
+	}
+}
+
 void cache_stats::inc_stats(int access_type, int access_outcome){
     ///
     /// Increment the stat corresponding to (access_type, access_outcome) by 1.
@@ -632,6 +643,16 @@ void cache_stats::inc_stats(int access_type, int access_outcome){
 
     m_stats[access_type][access_outcome]++;
 }
+
+void cache_stats::inc_stats_pw(int access_type, int access_outcome){
+    ///
+    /// 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]++;
+}
+
 void cache_stats::inc_fail_stats(int access_type, int fail_outcome){
 
 	if(!check_fail_valid(access_type, fail_outcome))
@@ -718,6 +739,9 @@ cache_stats &cache_stats::operator+=(const cache_stats &cs){
         for(unsigned status=0; status<NUM_CACHE_REQUEST_STATUS; ++status){
             m_stats[type][status] += cs(type, status, false);
         }
+        for(unsigned status=0; status<NUM_CACHE_REQUEST_STATUS; ++status){
+            m_stats_pw[type][status] += cs(type, status, false);
+        }
         for(unsigned status=0; status<NUM_CACHE_RESERVATION_FAIL_STATUS; ++status){
             m_fail_stats[type][status] += cs(type, status, true);
 	   }
@@ -745,6 +769,7 @@ void cache_stats::print_stats(FILE *fout, const char *cache_name) const{
                 mem_access_type_str((enum mem_access_type)type),
                 cache_request_status_str((enum cache_request_status)status),
                 m_stats[type][status]);
+
             if(status != RESERVATION_FAIL)
             	 total_access[type]+= m_stats[type][status];
         }
@@ -804,6 +829,7 @@ unsigned long long cache_stats::get_stats(enum mem_access_type *access_type, uns
     }
     return total;
 }
+
 void cache_stats::get_sub_stats(struct cache_sub_stats &css) const{
     ///
     /// Overwrites "css" with the appropriate statistics from this cache.
@@ -834,6 +860,55 @@ void cache_stats::get_sub_stats(struct cache_sub_stats &css) const{
     css = t_css;
 }
 
+void cache_stats::get_sub_stats_pw(struct cache_sub_stats_pw &css) const{
+    ///
+    /// Overwrites "css" with the appropriate statistics from this cache.
+    ///
+    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];
+                }
+            }
+
+            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 == 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 == 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];
+                }
+            }
+        }
+    }
+
+    css = t_css;
+}
+
 bool cache_stats::check_valid(int type, int status) const{
     ///
     /// Verify a valid access_type/access_status
@@ -1501,6 +1576,7 @@ read_only_cache::access( new_addr_type addr,
     }
 
     m_stats.inc_stats(mf->get_access_type(), m_stats.select_stats_status(status, cache_status));
+    m_stats.inc_stats_pw(mf->get_access_type(), m_stats.select_stats_status(status, cache_status));
     return cache_status;
 }
 
@@ -1575,6 +1651,8 @@ data_cache::access( new_addr_type addr,
         = 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));
     return access_status;
 }
 
@@ -1639,6 +1717,7 @@ 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.inc_stats_pw(mf->get_access_type(), m_stats.select_stats_status(status, cache_status));
     return cache_status;
 }
 
diff --git a/src/gpgpu-sim/gpu-cache.h b/src/gpgpu-sim/gpu-cache.h
index 9725270..337f710 100644
--- a/src/gpgpu-sim/gpu-cache.h
+++ b/src/gpgpu-sim/gpu-cache.h
@@ -968,6 +968,66 @@ struct cache_sub_stats{
     void print_port_stats(FILE *fout, const char *cache_name) const; 
 };
 
+
+// Used for collecting AerialVision per-window statistics
+struct cache_sub_stats_pw{
+    unsigned accesses;
+    unsigned write_misses;
+    unsigned write_hits;
+    unsigned write_pending_hits;
+    unsigned write_res_fails;
+
+    unsigned read_misses;
+    unsigned read_hits;
+    unsigned read_pending_hits;
+    unsigned read_res_fails;
+
+    cache_sub_stats_pw(){
+        clear();
+    }
+    void clear(){
+        accesses = 0;
+        write_misses = 0;
+        write_hits = 0;
+        write_pending_hits = 0;
+        write_res_fails = 0;
+        read_misses = 0;
+        read_hits = 0;
+        read_pending_hits = 0;
+        read_res_fails = 0;
+    }
+    cache_sub_stats_pw &operator+=(const cache_sub_stats_pw &css){
+        ///
+        /// Overloading += operator to easily accumulate stats
+        ///
+        accesses += css.accesses;
+        write_misses += css.write_misses;
+        read_misses += css.read_misses;
+        write_pending_hits += css.write_pending_hits;
+        read_pending_hits += css.read_pending_hits;
+        write_res_fails += css.write_res_fails;
+        read_res_fails += css.read_res_fails;
+        return *this;
+    }
+
+    cache_sub_stats_pw operator+(const cache_sub_stats_pw &cs){
+        ///
+        /// Overloading + operator to easily accumulate stats
+        ///
+        cache_sub_stats_pw ret;
+        ret.accesses = accesses + cs.accesses;
+        ret.write_misses = write_misses + cs.write_misses;
+        ret.read_misses = read_misses + cs.read_misses;
+        ret.write_pending_hits = write_pending_hits + cs.write_pending_hits;
+        ret.read_pending_hits = read_pending_hits + cs.read_pending_hits;
+        ret.write_res_fails = write_res_fails + cs.write_res_fails;
+        ret.read_res_fails = read_res_fails + cs.read_res_fails;
+        return ret;
+    }
+
+};
+
+
 ///
 /// Cache_stats
 /// Used to record statistics for each cache.
@@ -978,7 +1038,11 @@ class cache_stats {
 public:
     cache_stats();
     void clear();
+    // Clear AerialVision cache stats after each window
+    void clear_pw();
     void inc_stats(int access_type, int access_outcome);
+    // Increment AerialVision cache stats
+    void inc_stats_pw(int access_type, int access_outcome);
     void inc_fail_stats(int access_type, int fail_outcome);
     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);
@@ -991,12 +1055,18 @@ public:
     unsigned long long get_stats(enum mem_access_type *access_type, unsigned num_access_type, enum cache_request_status *access_status, unsigned num_access_status)  const;
     void get_sub_stats(struct cache_sub_stats &css) const;
 
+    // Get per-window cache stats for AerialVision
+    void get_sub_stats_pw(struct cache_sub_stats_pw &css) const;
+
     void sample_cache_port_utility(bool data_port_busy, bool fill_port_busy); 
 private:
     bool check_valid(int type, int status) const;
     bool check_fail_valid(int type, int fail) const;
 
+
     std::vector< std::vector<unsigned long long> > m_stats;
+    // AerialVision cache stats (per-window)
+    std::vector< std::vector<unsigned long long> > m_stats_pw;
     std::vector< std::vector<unsigned long long> > m_fail_stats;
 
     unsigned long long m_cache_port_available_cycles; 
@@ -1082,6 +1152,14 @@ public:
     void get_sub_stats(struct cache_sub_stats &css) const {
         m_stats.get_sub_stats(css);
     }
+    // Clear per-window stats for AerialVision support
+    void clear_pw(){
+        m_stats.clear_pw();
+    }
+    // Per-window sub stats for AerialVision support
+    void get_sub_stats_pw(struct cache_sub_stats_pw &css) const {
+        m_stats.get_sub_stats_pw(css);
+    }
 
     // accessors for cache bandwidth availability 
     bool data_port_free() const { return m_bandwidth_management.data_port_free(); } 
diff --git a/src/gpgpu-sim/l2cache.cc b/src/gpgpu-sim/l2cache.cc
index 25da107..0edc3b7 100644
--- a/src/gpgpu-sim/l2cache.cc
+++ b/src/gpgpu-sim/l2cache.cc
@@ -519,14 +519,23 @@ void memory_sub_partition::print( FILE *fp ) const
 
 void memory_stats_t::visualizer_print( gzFile visualizer_file )
 {
-   // gzprintf(visualizer_file, "Ltwowritemiss: %d\n", L2_write_miss);
-   // gzprintf(visualizer_file, "Ltwowritehit: %d\n",  L2_write_access-L2_write_miss);
-   // gzprintf(visualizer_file, "Ltworeadmiss: %d\n", L2_read_miss);
-   // gzprintf(visualizer_file, "Ltworeadhit: %d\n", L2_read_access-L2_read_miss);
+   gzprintf(visualizer_file, "Ltwowritemiss: %d\n", L2_write_miss);
+   gzprintf(visualizer_file, "Ltwowritehit: %d\n",  L2_write_hit);
+   gzprintf(visualizer_file, "Ltworeadmiss: %d\n", L2_read_miss);
+   gzprintf(visualizer_file, "Ltworeadhit: %d\n", L2_read_hit);
+   clear_L2_stats_pw();
+
    if (num_mfs)
       gzprintf(visualizer_file, "averagemflatency: %lld\n", mf_total_lat/num_mfs);
 }
 
+void memory_stats_t::clear_L2_stats_pw(){
+    L2_write_miss = 0;
+    L2_write_hit = 0;
+    L2_read_miss = 0;
+    L2_read_hit = 0;
+}
+
 void gpgpu_sim::print_dram_stats(FILE *fout) const
 {
 	unsigned cmd=0;
@@ -718,8 +727,29 @@ void memory_sub_partition::get_L2cache_sub_stats(struct cache_sub_stats &css) co
     }
 }
 
+void memory_sub_partition::get_L2cache_sub_stats_pw(struct cache_sub_stats_pw &css) const{
+    if (!m_config->m_L2_config.disabled()) {
+        m_L2cache->get_sub_stats_pw(css);
+    }
+}
+
+void memory_sub_partition::clear_L2cache_stats_pw() {
+    if (!m_config->m_L2_config.disabled()) {
+        m_L2cache->clear_pw();
+    }
+}
+
 void memory_sub_partition::visualizer_print( gzFile visualizer_file )
 {
-    // TODO: Add visualizer stats for L2 cache 
-}
+    // Support for L2 AerialVision stats
+    // Per-sub-partition stats would be trivial to extend from this
+    cache_sub_stats_pw temp_sub_stats;
+    get_L2cache_sub_stats_pw(temp_sub_stats);
+
+    m_stats->L2_read_miss += temp_sub_stats.read_misses;
+    m_stats->L2_write_miss += temp_sub_stats.write_misses;
+    m_stats->L2_read_hit += temp_sub_stats.read_hits;
+    m_stats->L2_write_hit += temp_sub_stats.write_hits;
 
+    clear_L2cache_stats_pw();
+}
diff --git a/src/gpgpu-sim/l2cache.h b/src/gpgpu-sim/l2cache.h
index 18c0a8b..beafdd3 100644
--- a/src/gpgpu-sim/l2cache.h
+++ b/src/gpgpu-sim/l2cache.h
@@ -180,6 +180,10 @@ public:
    void accumulate_L2cache_stats(class cache_stats &l2_stats) const;
    void get_L2cache_sub_stats(struct cache_sub_stats &css) const;
 
+   // Support for getting per-window L2 stats for AerialVision
+   void get_L2cache_sub_stats_pw(struct cache_sub_stats_pw &css) const;
+   void clear_L2cache_stats_pw();
+
    void force_l2_tag_update(new_addr_type addr, unsigned time, mem_access_sector_mask_t mask)
    {
         m_L2cache->force_tag_access( addr, m_memcpy_cycle_offset + time, mask );
diff --git a/src/gpgpu-sim/mem_latency_stat.cc b/src/gpgpu-sim/mem_latency_stat.cc
index 7f6cde9..11624f4 100644
--- a/src/gpgpu-sim/mem_latency_stat.cc
+++ b/src/gpgpu-sim/mem_latency_stat.cc
@@ -129,6 +129,12 @@ memory_stats_t::memory_stats_t( unsigned n_shader, const struct shader_core_conf
       }
    }
 
+   // AerialVision L2 stats
+   L2_read_miss = 0;
+   L2_write_miss = 0;
+   L2_read_hit = 0;
+   L2_write_hit = 0;
+
    L2_cbtoL2length = (unsigned int*) calloc(mem_config->m_n_mem, sizeof(unsigned int));
    L2_cbtoL2writelength = (unsigned int*) calloc(mem_config->m_n_mem, sizeof(unsigned int));
    L2_L2tocblength = (unsigned int*) calloc(mem_config->m_n_mem, sizeof(unsigned int));
diff --git a/src/gpgpu-sim/mem_latency_stat.h b/src/gpgpu-sim/mem_latency_stat.h
index 5b89202..982b9ae 100644
--- a/src/gpgpu-sim/mem_latency_stat.h
+++ b/src/gpgpu-sim/mem_latency_stat.h
@@ -47,6 +47,9 @@ public:
 
    void visualizer_print( gzFile visualizer_file );
 
+   // Reset local L2 stats that are aggregated each sampling window
+   void clear_L2_stats_pw();
+
    unsigned m_n_shader;
 
    const struct shader_core_config *m_shader_config;
@@ -84,6 +87,11 @@ public:
    
    unsigned ***mem_access_type_stats; // dram access type classification
 
+   // AerialVision L2 stats
+   unsigned L2_read_miss;
+   unsigned L2_write_miss;
+   unsigned L2_read_hit;
+   unsigned L2_write_hit;
 
    // L2 cache stats
    unsigned int *L2_cbtoL2length;