platform/drm-tests: Add mali_stats test
Add mali_stats to drm-tests. This test includes a library for accessing
Mali performance counters, along with a simple test program
demonstrating how to use it.
Notes and technical specifications on these performance counters can be
found at go/mali-perf-metrics-notes.
BUG=b:221476799
TEST=Tested on asurada
Change-Id: I2127b04763f6200c6197cf91dac69fb4f4903136
Reviewed-on: https://chromium-review.googlesource.com/c/chromiumos/platform/drm-tests/+/3502554
Reviewed-by: Fritz Koenig <frkoenig@chromium.org>
Reviewed-by: Miguel Casas-Sanchez <mcasas@chromium.org>
Tested-by: Justin Green <greenjustin@google.com>
Commit-Queue: Justin Green <greenjustin@google.com>
diff --git a/Makefile b/Makefile
index 6a84647..e9227f0 100644
--- a/Makefile
+++ b/Makefile
@@ -22,6 +22,7 @@
CC_BINARY(drm_cursor_test) \
CC_BINARY(gamma_test) \
CC_BINARY(linear_bo_test) \
+ CC_BINARY(mali_stats) \
CC_BINARY(mapped_access_perf_test) \
CC_BINARY(mapped_texture_test) \
CC_BINARY(mmap_test) \
@@ -82,13 +83,17 @@
v4l2_stateful_decoder.o: $(SRC)/v4l2_stateful_decoder.c
v4l2_stateful_decoder.o: $(SRC)/bitstreams/bitstream_helper.h
-v4l2_stateful_decoder.o: $(SRC)/v4l2_macros.h
+v4l2_stateful_decoder.o: $(SRC)/v4l2_macros.h $(SRC)/logging.h
CC_BINARY(v4l2_stateful_decoder): v4l2_stateful_decoder.o
CC_BINARY(v4l2_stateful_decoder): CC_STATIC_LIBRARY(libbitstreams.pic.a)
CC_BINARY(v4l2_stateful_decoder): CC_STATIC_LIBRARY(libbsdrm.pic.a)
CC_BINARY(v4l2_stateful_decoder): LDLIBS += $(DRM_LIBS)
CC_BINARY(v4l2_stateful_decoder): LDLIBS += -lssl -lcrypto
+CC_BINARY(mali_stats): mali_stats.o
+CC_BINARY(mali_stats): $(SRC)/logging.h $(SRC)/mali/mali_ioctl.h
+CC_BINARY(mali_stats): CC_STATIC_LIBRARY(libmali.pic.a)
+
CC_BINARY(v4l2_stateful_encoder): v4l2_stateful_encoder.o
ifeq ($(USE_V4LPLUGIN),1)
CC_BINARY(v4l2_stateful_encoder): LDLIBS += -lv4l2
diff --git a/logging.h b/logging.h
new file mode 100644
index 0000000..d795a09
--- /dev/null
+++ b/logging.h
@@ -0,0 +1,41 @@
+/*
+ * Copyright 2022 The Chromium OS Authors. All rights reserved.
+ * Use of this source code is governed by a BSD-style license that can be
+ * found in the LICENSE file.
+ */
+
+#include <stdio.h>
+#include <stdlib.h>
+
+#ifndef __LOGGING_H__
+#define __LOGGING_H__
+
+enum logging_levels {
+ kLoggingDebug = -1,
+ kLoggingInfo = 0,
+ kLoggingError,
+ kLoggingFatal,
+ kLoggingLevelMax
+};
+
+#define DEFAULT_LOG_LEVEL kLoggingInfo
+
+#define LOG(level, stream, fmt, ...) \
+ do { \
+ if (level >= DEFAULT_LOG_LEVEL) { \
+ fprintf(stream, fmt, ##__VA_ARGS__); \
+ fprintf(stream, "\n"); \
+ fflush(stream); \
+ } \
+ } while (0)
+
+#define LOG_DEBUG(fmt, ...) LOG(kLoggingDebug, stderr, fmt, ##__VA_ARGS__)
+#define LOG_INFO(fmt, ...) LOG(kLoggingInfo, stderr, fmt, ##__VA_ARGS__)
+#define LOG_ERROR(fmt, ...) LOG(kLoggingError, stderr, fmt, ##__VA_ARGS__)
+#define LOG_FATAL(fmt, ...) \
+ do { \
+ LOG(kLoggingFatal, stderr, fmt, ##__VA_ARGS__); \
+ exit(EXIT_FAILURE); \
+ } while (0)
+
+#endif // __LOGGING_H__
diff --git a/mali/mali_gpu_perf_metrics.c b/mali/mali_gpu_perf_metrics.c
new file mode 100644
index 0000000..03d970d
--- /dev/null
+++ b/mali/mali_gpu_perf_metrics.c
@@ -0,0 +1,284 @@
+/*
+ * Copyright 2022 The Chromium OS Authors. All rights reserved.
+ * Use of this source code is governed by a BSD-style license that can be
+ * found in the LICENSE file.
+ */
+
+#include "mali/mali_gpu_perf_metrics.h"
+
+#include <assert.h>
+#include <errno.h>
+#include <fcntl.h>
+#include <stddef.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <sys/ioctl.h>
+#include <sys/mman.h>
+#include <unistd.h>
+
+#include "logging.h"
+#include "mali/mali_gpu_props.h"
+#include "mali/mali_ioctl.h"
+
+int num_shader_cores;
+int num_l2_caches;
+MaliGpuModel model;
+
+int gpufd = -1;
+int reader_fd = -1;
+
+struct kbase_hwcnt_reader_metadata reader_metadata;
+
+// The model logic here is also taken from gfx-pps.
+void populate_model(int product_id) {
+ int masked_product_id = product_id & 0xF00F;
+ switch (product_id) {
+ case 0x6956:
+ model = gpu_model_t60x;
+ break;
+ case 0x0620:
+ model = gpu_model_t62x;
+ break;
+ case 0x0720:
+ model = gpu_model_t72x;
+ break;
+ case 0x0750:
+ model = gpu_model_t76x;
+ break;
+ case 0x0820:
+ model = gpu_model_t82x;
+ break;
+ case 0x0830:
+ model = gpu_model_t83x;
+ break;
+ case 0x0860:
+ model = gpu_model_t86x;
+ break;
+ case 0x0880:
+ model = gpu_model_tfrx;
+ break;
+ default:
+ switch (masked_product_id) {
+ case 0x6000:
+ model = gpu_model_tmix;
+ break;
+ case 0x6001:
+ model = gpu_model_thex;
+ break;
+ case 0x7000:
+ model = gpu_model_tsix;
+ break;
+ case 0x7001:
+ model = gpu_model_tnox;
+ break;
+ case 0x7002:
+ model = gpu_model_tgox;
+ break;
+ case 0x7003:
+ model = gpu_model_tdvx;
+ break;
+ case 0x9000:
+ model = gpu_model_ttrx;
+ break;
+ case 0x9001:
+ case 0x9003:
+ model = gpu_model_tnax;
+ break;
+ default:
+ LOG_FATAL("Error: unsupported GPU with product ID 0x%x", product_id);
+ }
+ break;
+ }
+}
+
+void initialize_mali_perf_reader() {
+ int product_id = get_gpu_prop(gpu_prop_product_id);
+ populate_model(product_id);
+ num_shader_cores =
+ __builtin_popcount(get_gpu_prop(gpu_prop_shader_present_mask));
+ num_l2_caches = get_gpu_prop(gpu_prop_num_l2);
+
+ gpufd = open(kGpuDevice, O_RDWR | O_CLOEXEC);
+ if (gpufd < 0)
+ LOG_FATAL("Error opening GPU device! %s\n", strerror(errno));
+
+ struct kbase_ioctl_version_check version_check;
+ if (ioctl(gpufd, KBASE_IOCTL_VERSION_CHECK, &version_check) < 0)
+ LOG_FATAL("Error checking GPU version! %s\n", strerror(errno));
+ assert(version_check.major >= SUPPORTED_MAJOR_VERSION);
+ assert(version_check.minor >= SUPPORTED_MINOR_VERSION);
+
+ struct kbase_ioctl_set_flags init_flags;
+ init_flags.create_flags = BASE_CONTEXT_SYSTEM_MONITOR_SUBMIT_DISABLED;
+ if (ioctl(gpufd, KBASE_IOCTL_SET_FLAGS, &init_flags) < 0)
+ LOG_FATAL("Error initializing GPU context! %s\n", strerror(errno));
+
+ struct kbase_ioctl_hwcnt_reader_setup reader_setup;
+ reader_setup.num_buffers = 2;
+ reader_setup.job_manager_mask = 0xFFFFFFFF;
+ reader_setup.shader_mask = 0xFFFFFFFF;
+ reader_setup.tiler_mask = 0xFFFFFFFF;
+ reader_setup.mmu_l2_mask = 0xFFFFFFFF;
+ reader_fd = ioctl(gpufd, KBASE_IOCTL_HWCNT_READER_SETUP, &reader_setup);
+ if (reader_fd < 0)
+ LOG_FATAL("Error setting up hwcnt reader! %s\n", strerror(errno));
+
+ uint32_t api_version;
+ if (ioctl(reader_fd, KBASE_HWCNT_READER_GET_API_VERSION, &api_version) < 0)
+ LOG_FATAL("Error getting API version! %s\n", strerror(errno));
+ assert(api_version >= SUPPORTED_API_VERSION);
+
+ uint32_t hw_version;
+ if (ioctl(reader_fd, KBASE_HWCNT_READER_GET_HWVER, &hw_version) < 0)
+ LOG_FATAL("Error getting hardware version! %s\n", strerror(errno));
+ assert(hw_version >= SUPPORTED_HW_VERSION);
+}
+
+void cleanup_mali_perf_reader() {
+ close(gpufd);
+ close(reader_fd);
+}
+
+struct mali_counter_values get_counter_values_from_dump(
+ MaliGpuCounter counter,
+ uint32_t* dump,
+ size_t dump_size_bytes) {
+ struct mali_counter_values ret;
+ memset(&ret, 0, sizeof(struct mali_counter_values));
+
+ if (counter >> 8 != model) {
+ LOG_ERROR("Error: counter is of incorrect model type!\n");
+ return ret;
+ }
+
+ int lower_range = -1;
+ int upper_range = -1;
+ int counter_type = (counter >> 6) & 3;
+ switch (counter_type) {
+ case gpu_counter_job_manager:
+ lower_range = 0;
+ upper_range = 1;
+ break;
+ case gpu_counter_tiler:
+ lower_range = 1;
+ upper_range = 2;
+ break;
+ case gpu_counter_l2:
+ lower_range = 2;
+ upper_range = lower_range + num_l2_caches;
+ break;
+ case gpu_counter_shader:
+ lower_range = 2 + num_l2_caches;
+ upper_range = lower_range + num_shader_cores;
+ break;
+ default:
+ // This should never happen
+ break;
+ }
+
+ ret.counter = counter;
+ ret.num_values = upper_range - lower_range;
+ ret.values = (uint32_t*)malloc(ret.num_values * sizeof(uint32_t));
+
+ int index = 0;
+ int present_index = -1;
+ const int max_index = dump_size_bytes / (64 * sizeof(uint32_t));
+ while (index < max_index && present_index < upper_range) {
+ if (dump[index * 64 + 2])
+ present_index++;
+
+ if (present_index >= lower_range && present_index < upper_range) {
+ uint32_t counter_val = dump[index * 64 + (counter & 0x3F)];
+ ret.values[present_index - lower_range] = counter_val;
+ }
+
+ index++;
+ }
+
+ return ret;
+}
+
+size_t get_dump_size_bytes() {
+ uint32_t ret;
+ if (ioctl(reader_fd, KBASE_HWCNT_READER_GET_BUFFER_SIZE, &ret) < 0)
+ LOG_ERROR("Error reading dump buffer size! %s\n", strerror(errno));
+ return ret;
+}
+
+void initiate_dump() {
+ if (ioctl(reader_fd, KBASE_HWCNT_READER_CLEAR, NULL) < 0)
+ LOG_ERROR("Error clearing dump buffer! %s\n", strerror(errno));
+
+ if (ioctl(reader_fd, KBASE_HWCNT_READER_DUMP, NULL) < 0)
+ LOG_ERROR("Error dumping performance metrics! %s\n", strerror(errno));
+}
+
+void get_dump_buffer() {
+ if (ioctl(reader_fd, KBASE_HWCNT_READER_GET_BUFFER, &reader_metadata) < 0)
+ LOG_ERROR("Error getting dump buffer! %s\n", strerror(errno));
+}
+
+void put_dump_buffer() {
+ if (ioctl(reader_fd, KBASE_HWCNT_READER_PUT_BUFFER, &reader_metadata) < 0)
+ LOG_ERROR("Error putting dump buffer! %s\n", strerror(errno));
+}
+
+uint8_t* copy_dump_to_userspace(size_t dump_size_bytes) {
+ uint8_t* ret = (uint8_t*)malloc(dump_size_bytes);
+
+ int offset = 0;
+ const int kMaxWindowSize = 4096;
+ while (offset < dump_size_bytes) {
+ int curr_window_size = dump_size_bytes - offset < kMaxWindowSize
+ ? (dump_size_bytes - offset)
+ : kMaxWindowSize;
+
+ uint8_t* curr_window = (uint8_t*)mmap(NULL, curr_window_size, PROT_READ,
+ MAP_SHARED, reader_fd, offset);
+ if (curr_window == MAP_FAILED) {
+ LOG_ERROR("Error mapping dump buffer! %s\n", strerror(errno));
+ break;
+ }
+
+ memcpy(ret + offset, curr_window, curr_window_size);
+
+ munmap(curr_window, curr_window_size);
+
+ offset += kMaxWindowSize;
+ }
+
+ return ret;
+}
+
+struct mali_counter_response read_perf_metrics(MaliGpuCounter* counters,
+ size_t num_counters) {
+ initiate_dump();
+ get_dump_buffer();
+
+ size_t dump_size_bytes = get_dump_size_bytes();
+ uint8_t* dump_data = copy_dump_to_userspace(dump_size_bytes);
+
+ put_dump_buffer();
+
+ struct mali_counter_response ret;
+ ret.num_counters = num_counters;
+ ret.counter_values = (struct mali_counter_values*)malloc(
+ num_counters * sizeof(struct mali_counter_values));
+ for (int i = 0; i < num_counters; i++) {
+ ret.counter_values[i] = get_counter_values_from_dump(
+ counters[i], (uint32_t*)dump_data, dump_size_bytes);
+ }
+
+ free(dump_data);
+
+ return ret;
+}
+
+void free_counters(struct mali_counter_response counters) {
+ for (int i = 0; i < counters.num_counters; i++) {
+ free(counters.counter_values[i].values);
+ }
+ free(counters.counter_values);
+}
diff --git a/mali/mali_gpu_perf_metrics.h b/mali/mali_gpu_perf_metrics.h
new file mode 100644
index 0000000..3a6c813
--- /dev/null
+++ b/mali/mali_gpu_perf_metrics.h
@@ -0,0 +1,2864 @@
+/*
+ * Copyright 2022 The Chromium OS Authors. All rights reserved.
+ * Use of this source code is governed by a BSD-style license that can be
+ * found in the LICENSE file.
+ */
+
+#include <stddef.h>
+#include <stdint.h>
+
+#ifndef __MALI_GPU_PERF_METRICS_H__
+#define __MALI_GPU_PERF_METRICS_H__
+
+typedef enum {
+ gpu_counter_job_manager = 0,
+ gpu_counter_tiler = 1,
+ gpu_counter_shader = 2,
+ gpu_counter_l2 = 3
+} MaliGpuCounterType;
+
+typedef enum {
+ gpu_model_t60x = 0,
+ gpu_model_t62x = 1,
+ gpu_model_t72x = 2,
+ gpu_model_t76x = 3,
+ gpu_model_t82x = 4,
+ gpu_model_t83x = 5,
+ gpu_model_t86x = 6,
+ gpu_model_tfrx = 7,
+ gpu_model_tmix = 8,
+ gpu_model_thex = 9,
+ gpu_model_tsix = 10,
+ gpu_model_tnox = 11,
+ gpu_model_tgox = 12,
+ gpu_model_tdvx = 13,
+ gpu_model_ttrx = 14,
+ gpu_model_tnax = 15
+} MaliGpuModel;
+
+// These values are computed as:
+// index_in_counter_block | counter_type << 6 | model << 8
+// The indices were derived from the gfx-pps source code. More information at:
+// https://gitlab.freedesktop.org/Fahien/gfx-pps/-/blob/master/include/pps/gpu/panfrost/hwc_names.h # nocheck
+typedef enum {
+ t60x_messages_sent = 4,
+ t60x_messages_received = 5,
+ t60x_gpu_active = 6,
+ t60x_irq_active = 7,
+ t60x_js0_jobs = 8,
+ t60x_js0_tasks = 9,
+ t60x_js0_active = 10,
+ t60x_js0_wait_read = 12,
+ t60x_js0_wait_issue = 13,
+ t60x_js0_wait_depend = 14,
+ t60x_js0_wait_finish = 15,
+ t60x_js1_jobs = 16,
+ t60x_js1_tasks = 17,
+ t60x_js1_active = 18,
+ t60x_js1_wait_read = 20,
+ t60x_js1_wait_issue = 21,
+ t60x_js1_wait_depend = 22,
+ t60x_js1_wait_finish = 23,
+ t60x_js2_jobs = 24,
+ t60x_js2_tasks = 25,
+ t60x_js2_active = 26,
+ t60x_js2_wait_read = 28,
+ t60x_js2_wait_issue = 29,
+ t60x_js2_wait_depend = 30,
+ t60x_js2_wait_finish = 31,
+ t60x_ti_jobs_processed = 67,
+ t60x_ti_triangles = 68,
+ t60x_ti_quads = 69,
+ t60x_ti_polygons = 70,
+ t60x_ti_points = 71,
+ t60x_ti_lines = 72,
+ t60x_ti_vcache_hit = 73,
+ t60x_ti_vcache_miss = 74,
+ t60x_ti_front_facing = 75,
+ t60x_ti_back_facing = 76,
+ t60x_ti_prim_visible = 77,
+ t60x_ti_prim_culled = 78,
+ t60x_ti_prim_clipped = 79,
+ t60x_ti_level0 = 80,
+ t60x_ti_level1 = 81,
+ t60x_ti_level2 = 82,
+ t60x_ti_level3 = 83,
+ t60x_ti_level4 = 84,
+ t60x_ti_level5 = 85,
+ t60x_ti_level6 = 86,
+ t60x_ti_level7 = 87,
+ t60x_ti_command_1 = 88,
+ t60x_ti_command_2 = 89,
+ t60x_ti_command_3 = 90,
+ t60x_ti_command_4 = 91,
+ t60x_ti_command_4_7 = 92,
+ t60x_ti_command_8_15 = 93,
+ t60x_ti_command_16_63 = 94,
+ t60x_ti_command_64 = 95,
+ t60x_ti_compress_in = 96,
+ t60x_ti_compress_out = 97,
+ t60x_ti_compress_flush = 98,
+ t60x_ti_timestamps = 99,
+ t60x_ti_pcache_hit = 100,
+ t60x_ti_pcache_miss = 101,
+ t60x_ti_pcache_line = 102,
+ t60x_ti_pcache_stall = 103,
+ t60x_ti_wrbuf_hit = 104,
+ t60x_ti_wrbuf_miss = 105,
+ t60x_ti_wrbuf_line = 106,
+ t60x_ti_wrbuf_partial = 107,
+ t60x_ti_wrbuf_stall = 108,
+ t60x_ti_active = 109,
+ t60x_ti_loading_desc = 110,
+ t60x_ti_index_wait = 111,
+ t60x_ti_index_range_wait = 112,
+ t60x_ti_vertex_wait = 113,
+ t60x_ti_pcache_wait = 114,
+ t60x_ti_wrbuf_wait = 115,
+ t60x_ti_bus_read = 116,
+ t60x_ti_bus_write = 117,
+ t60x_ti_utlb_stall = 123,
+ t60x_ti_utlb_replay_miss = 124,
+ t60x_ti_utlb_replay_full = 125,
+ t60x_ti_utlb_new_miss = 126,
+ t60x_ti_utlb_hit = 127,
+ t60x_frag_active = 132,
+ t60x_frag_primitives = 133,
+ t60x_frag_primitives_dropped = 134,
+ t60x_frag_cycles_desc = 135,
+ t60x_frag_cycles_plr = 136,
+ t60x_frag_cycles_vert = 137,
+ t60x_frag_cycles_trisetup = 138,
+ t60x_frag_cycles_rast = 139,
+ t60x_frag_threads = 140,
+ t60x_frag_dummy_threads = 141, // # nocheck
+ t60x_frag_quads_rast = 142,
+ t60x_frag_quads_ezs_test = 143,
+ t60x_frag_quads_ezs_killed = 144,
+ t60x_frag_threads_lzs_test = 145,
+ t60x_frag_threads_lzs_killed = 146,
+ t60x_frag_cycles_no_tile = 147,
+ t60x_frag_num_tiles = 148,
+ t60x_frag_trans_elim = 149,
+ t60x_compute_active = 150,
+ t60x_compute_tasks = 151,
+ t60x_compute_threads = 152,
+ t60x_compute_cycles_desc = 153,
+ t60x_tripipe_active = 154,
+ t60x_arith_words = 155,
+ t60x_arith_cycles_reg = 156,
+ t60x_arith_cycles_l0 = 157,
+ t60x_arith_frag_depend = 158,
+ t60x_ls_words = 159,
+ t60x_ls_issues = 160,
+ t60x_ls_restarts = 161,
+ t60x_ls_reissues_miss = 162,
+ t60x_ls_reissues_vd = 163,
+ t60x_ls_reissue_attrib_miss = 164,
+ t60x_ls_no_wb = 165,
+ t60x_tex_words = 166,
+ t60x_tex_bubbles = 167,
+ t60x_tex_words_l0 = 168,
+ t60x_tex_words_desc = 169,
+ t60x_tex_issues = 170,
+ t60x_tex_recirc_fmiss = 171,
+ t60x_tex_recirc_desc = 172,
+ t60x_tex_recirc_multi = 173,
+ t60x_tex_recirc_pmiss = 174,
+ t60x_tex_recirc_conf = 175,
+ t60x_lsc_read_hits = 176,
+ t60x_lsc_read_misses = 177,
+ t60x_lsc_write_hits = 178,
+ t60x_lsc_write_misses = 179,
+ t60x_lsc_atomic_hits = 180,
+ t60x_lsc_atomic_misses = 181,
+ t60x_lsc_line_fetches = 182,
+ t60x_lsc_dirty_line = 183,
+ t60x_lsc_snoops = 184,
+ t60x_axi_tlb_stall = 185,
+ t60x_axi_tlb_miss = 186,
+ t60x_axi_tlb_transaction = 187,
+ t60x_ls_tlb_miss = 188,
+ t60x_ls_tlb_hit = 189,
+ t60x_axi_beats_read = 190,
+ t60x_axi_beats_written = 191,
+ t60x_mmu_hit = 196,
+ t60x_mmu_new_miss = 197,
+ t60x_mmu_replay_full = 198,
+ t60x_mmu_replay_miss = 199,
+ t60x_mmu_table_walk = 200,
+ t60x_utlb_hit = 208,
+ t60x_utlb_new_miss = 209,
+ t60x_utlb_replay_full = 210,
+ t60x_utlb_replay_miss = 211,
+ t60x_utlb_stall = 212,
+ t60x_l2_ext_write_beats = 222,
+ t60x_l2_ext_read_beats = 223,
+ t60x_l2_any_lookup = 224,
+ t60x_l2_read_lookup = 225,
+ t60x_l2_sread_lookup = 226,
+ t60x_l2_read_replay = 227,
+ t60x_l2_read_snoop = 228,
+ t60x_l2_read_hit = 229,
+ t60x_l2_clean_miss = 230,
+ t60x_l2_write_lookup = 231,
+ t60x_l2_swrite_lookup = 232,
+ t60x_l2_write_replay = 233,
+ t60x_l2_write_snoop = 234,
+ t60x_l2_write_hit = 235,
+ t60x_l2_ext_read_full = 236,
+ t60x_l2_ext_read_half = 237,
+ t60x_l2_ext_write_full = 238,
+ t60x_l2_ext_write_half = 239,
+ t60x_l2_ext_read = 240,
+ t60x_l2_ext_read_line = 241,
+ t60x_l2_ext_write = 242,
+ t60x_l2_ext_write_line = 243,
+ t60x_l2_ext_write_small = 244,
+ t60x_l2_ext_barrier = 245,
+ t60x_l2_ext_ar_stall = 246,
+ t60x_l2_ext_r_buf_full = 247,
+ t60x_l2_ext_rd_buf_full = 248,
+ t60x_l2_ext_r_raw = 249,
+ t60x_l2_ext_w_stall = 250,
+ t60x_l2_ext_w_buf_full = 251,
+ t60x_l2_ext_r_w_hazard = 252,
+ t60x_l2_tag_hazard = 253,
+ t60x_l2_snoop_full = 254,
+ t60x_l2_replay_full = 255,
+ t62x_messages_sent = 260,
+ t62x_messages_received = 261,
+ t62x_gpu_active = 262,
+ t62x_irq_active = 263,
+ t62x_js0_jobs = 264,
+ t62x_js0_tasks = 265,
+ t62x_js0_active = 266,
+ t62x_js0_wait_read = 268,
+ t62x_js0_wait_issue = 269,
+ t62x_js0_wait_depend = 270,
+ t62x_js0_wait_finish = 271,
+ t62x_js1_jobs = 272,
+ t62x_js1_tasks = 273,
+ t62x_js1_active = 274,
+ t62x_js1_wait_read = 276,
+ t62x_js1_wait_issue = 277,
+ t62x_js1_wait_depend = 278,
+ t62x_js1_wait_finish = 279,
+ t62x_js2_jobs = 280,
+ t62x_js2_tasks = 281,
+ t62x_js2_active = 282,
+ t62x_js2_wait_read = 284,
+ t62x_js2_wait_issue = 285,
+ t62x_js2_wait_depend = 286,
+ t62x_js2_wait_finish = 287,
+ t62x_ti_jobs_processed = 323,
+ t62x_ti_triangles = 324,
+ t62x_ti_quads = 325,
+ t62x_ti_polygons = 326,
+ t62x_ti_points = 327,
+ t62x_ti_lines = 328,
+ t62x_ti_vcache_hit = 329,
+ t62x_ti_vcache_miss = 330,
+ t62x_ti_front_facing = 331,
+ t62x_ti_back_facing = 332,
+ t62x_ti_prim_visible = 333,
+ t62x_ti_prim_culled = 334,
+ t62x_ti_prim_clipped = 335,
+ t62x_ti_level0 = 336,
+ t62x_ti_level1 = 337,
+ t62x_ti_level2 = 338,
+ t62x_ti_level3 = 339,
+ t62x_ti_level4 = 340,
+ t62x_ti_level5 = 341,
+ t62x_ti_level6 = 342,
+ t62x_ti_level7 = 343,
+ t62x_ti_command_1 = 344,
+ t62x_ti_command_2 = 345,
+ t62x_ti_command_3 = 346,
+ t62x_ti_command_4 = 347,
+ t62x_ti_command_5_7 = 348,
+ t62x_ti_command_8_15 = 349,
+ t62x_ti_command_16_63 = 350,
+ t62x_ti_command_64 = 351,
+ t62x_ti_compress_in = 352,
+ t62x_ti_compress_out = 353,
+ t62x_ti_compress_flush = 354,
+ t62x_ti_timestamps = 355,
+ t62x_ti_pcache_hit = 356,
+ t62x_ti_pcache_miss = 357,
+ t62x_ti_pcache_line = 358,
+ t62x_ti_pcache_stall = 359,
+ t62x_ti_wrbuf_hit = 360,
+ t62x_ti_wrbuf_miss = 361,
+ t62x_ti_wrbuf_line = 362,
+ t62x_ti_wrbuf_partial = 363,
+ t62x_ti_wrbuf_stall = 364,
+ t62x_ti_active = 365,
+ t62x_ti_loading_desc = 366,
+ t62x_ti_index_wait = 367,
+ t62x_ti_index_range_wait = 368,
+ t62x_ti_vertex_wait = 369,
+ t62x_ti_pcache_wait = 370,
+ t62x_ti_wrbuf_wait = 371,
+ t62x_ti_bus_read = 372,
+ t62x_ti_bus_write = 373,
+ t62x_ti_utlb_stall = 379,
+ t62x_ti_utlb_replay_miss = 380,
+ t62x_ti_utlb_replay_full = 381,
+ t62x_ti_utlb_new_miss = 382,
+ t62x_ti_utlb_hit = 383,
+ t62x_shader_core_active = 387,
+ t62x_frag_active = 388,
+ t62x_frag_primitives = 389,
+ t62x_frag_primitives_dropped = 390,
+ t62x_frag_cycles_desc = 391,
+ t62x_frag_cycles_fpkq_active = 392,
+ t62x_frag_cycles_vert = 393,
+ t62x_frag_cycles_trisetup = 394,
+ t62x_frag_cycles_ezs_active = 395,
+ t62x_frag_threads = 396,
+ t62x_frag_dummy_threads = 397, // # nocheck
+ t62x_frag_quads_rast = 398,
+ t62x_frag_quads_ezs_test = 399,
+ t62x_frag_quads_ezs_killed = 400,
+ t62x_frag_threads_lzs_test = 401,
+ t62x_frag_threads_lzs_killed = 402,
+ t62x_frag_cycles_no_tile = 403,
+ t62x_frag_num_tiles = 404,
+ t62x_frag_trans_elim = 405,
+ t62x_compute_active = 406,
+ t62x_compute_tasks = 407,
+ t62x_compute_threads = 408,
+ t62x_compute_cycles_desc = 409,
+ t62x_tripipe_active = 410,
+ t62x_arith_words = 411,
+ t62x_arith_cycles_reg = 412,
+ t62x_arith_cycles_l0 = 413,
+ t62x_arith_frag_depend = 414,
+ t62x_ls_words = 415,
+ t62x_ls_issues = 416,
+ t62x_ls_restarts = 417,
+ t62x_ls_reissues_miss = 418,
+ t62x_ls_reissues_vd = 419,
+ t62x_ls_reissue_attrib_miss = 420,
+ t62x_ls_no_wb = 421,
+ t62x_tex_words = 422,
+ t62x_tex_bubbles = 423,
+ t62x_tex_words_l0 = 424,
+ t62x_tex_words_desc = 425,
+ t62x_tex_issues = 426,
+ t62x_tex_recirc_fmiss = 427,
+ t62x_tex_recirc_desc = 428,
+ t62x_tex_recirc_multi = 429,
+ t62x_tex_recirc_pmiss = 430,
+ t62x_tex_recirc_conf = 431,
+ t62x_lsc_read_hits = 432,
+ t62x_lsc_read_misses = 433,
+ t62x_lsc_write_hits = 434,
+ t62x_lsc_write_misses = 435,
+ t62x_lsc_atomic_hits = 436,
+ t62x_lsc_atomic_misses = 437,
+ t62x_lsc_line_fetches = 438,
+ t62x_lsc_dirty_line = 439,
+ t62x_lsc_snoops = 440,
+ t62x_axi_tlb_stall = 441,
+ t62x_axi_tlb_miss = 442,
+ t62x_axi_tlb_transaction = 443,
+ t62x_ls_tlb_miss = 444,
+ t62x_ls_tlb_hit = 445,
+ t62x_axi_beats_read = 446,
+ t62x_axi_beats_written = 447,
+ t62x_mmu_hit = 452,
+ t62x_mmu_new_miss = 453,
+ t62x_mmu_replay_full = 454,
+ t62x_mmu_replay_miss = 455,
+ t62x_mmu_table_walk = 456,
+ t62x_utlb_hit = 464,
+ t62x_utlb_new_miss = 465,
+ t62x_utlb_replay_full = 466,
+ t62x_utlb_replay_miss = 467,
+ t62x_utlb_stall = 468,
+ t62x_l2_ext_write_beats = 478,
+ t62x_l2_ext_read_beats = 479,
+ t62x_l2_any_lookup = 480,
+ t62x_l2_read_lookup = 481,
+ t62x_l2_sread_lookup = 482,
+ t62x_l2_read_replay = 483,
+ t62x_l2_read_snoop = 484,
+ t62x_l2_read_hit = 485,
+ t62x_l2_clean_miss = 486,
+ t62x_l2_write_lookup = 487,
+ t62x_l2_swrite_lookup = 488,
+ t62x_l2_write_replay = 489,
+ t62x_l2_write_snoop = 490,
+ t62x_l2_write_hit = 491,
+ t62x_l2_ext_read_full = 492,
+ t62x_l2_ext_read_half = 493,
+ t62x_l2_ext_write_full = 494,
+ t62x_l2_ext_write_half = 495,
+ t62x_l2_ext_read = 496,
+ t62x_l2_ext_read_line = 497,
+ t62x_l2_ext_write = 498,
+ t62x_l2_ext_write_line = 499,
+ t62x_l2_ext_write_small = 500,
+ t62x_l2_ext_barrier = 501,
+ t62x_l2_ext_ar_stall = 502,
+ t62x_l2_ext_r_buf_full = 503,
+ t62x_l2_ext_rd_buf_full = 504,
+ t62x_l2_ext_r_raw = 505,
+ t62x_l2_ext_w_stall = 506,
+ t62x_l2_ext_w_buf_full = 507,
+ t62x_l2_ext_r_w_hazard = 508,
+ t62x_l2_tag_hazard = 509,
+ t62x_l2_snoop_full = 510,
+ t62x_l2_replay_full = 511,
+ t72x_gpu_active = 516,
+ t72x_irq_active = 517,
+ t72x_js0_jobs = 518,
+ t72x_js0_tasks = 519,
+ t72x_js0_active = 520,
+ t72x_js1_jobs = 521,
+ t72x_js1_tasks = 522,
+ t72x_js1_active = 523,
+ t72x_js2_jobs = 524,
+ t72x_js2_tasks = 525,
+ t72x_js2_active = 526,
+ t72x_ti_jobs_processed = 579,
+ t72x_ti_triangles = 580,
+ t72x_ti_quads = 581,
+ t72x_ti_polygons = 582,
+ t72x_ti_points = 583,
+ t72x_ti_lines = 584,
+ t72x_ti_front_facing = 585,
+ t72x_ti_back_facing = 586,
+ t72x_ti_prim_visible = 587,
+ t72x_ti_prim_culled = 588,
+ t72x_ti_prim_clipped = 589,
+ t72x_ti_active = 598,
+ t72x_frag_active = 644,
+ t72x_frag_primitives = 645,
+ t72x_frag_primitives_dropped = 646,
+ t72x_frag_threads = 647,
+ t72x_frag_dummy_threads = 648, // # nocheck
+ t72x_frag_quads_rast = 649,
+ t72x_frag_quads_ezs_test = 650,
+ t72x_frag_quads_ezs_killed = 651,
+ t72x_frag_threads_lzs_test = 652,
+ t72x_frag_threads_lzs_killed = 653,
+ t72x_frag_cycles_no_tile = 654,
+ t72x_frag_num_tiles = 655,
+ t72x_frag_trans_elim = 656,
+ t72x_compute_active = 657,
+ t72x_compute_tasks = 658,
+ t72x_compute_threads = 659,
+ t72x_tripipe_active = 660,
+ t72x_arith_words = 661,
+ t72x_arith_cycles_reg = 662,
+ t72x_ls_words = 663,
+ t72x_ls_issues = 664,
+ t72x_ls_restarts = 665,
+ t72x_ls_reissues_miss = 666,
+ t72x_tex_words = 667,
+ t72x_tex_bubbles = 668,
+ t72x_tex_issues = 669,
+ t72x_lsc_read_hits = 670,
+ t72x_lsc_read_misses = 671,
+ t72x_lsc_write_hits = 672,
+ t72x_lsc_write_misses = 673,
+ t72x_lsc_atomic_hits = 674,
+ t72x_lsc_atomic_misses = 675,
+ t72x_lsc_line_fetches = 676,
+ t72x_lsc_dirty_line = 677,
+ t72x_lsc_snoops = 678,
+ t72x_l2_ext_write_beat = 708,
+ t72x_l2_ext_read_beat = 709,
+ t72x_l2_read_snoop = 710,
+ t72x_l2_read_hit = 711,
+ t72x_l2_write_snoop = 712,
+ t72x_l2_write_hit = 713,
+ t72x_l2_ext_write_small = 714,
+ t72x_l2_ext_barrier = 715,
+ t72x_l2_ext_ar_stall = 716,
+ t72x_l2_ext_w_stall = 717,
+ t72x_l2_snoop_full = 718,
+ t76x_messages_sent = 772,
+ t76x_messages_received = 773,
+ t76x_gpu_active = 774,
+ t76x_irq_active = 775,
+ t76x_js0_jobs = 776,
+ t76x_js0_tasks = 777,
+ t76x_js0_active = 778,
+ t76x_js0_wait_read = 780,
+ t76x_js0_wait_issue = 781,
+ t76x_js0_wait_depend = 782,
+ t76x_js0_wait_finish = 783,
+ t76x_js1_jobs = 784,
+ t76x_js1_tasks = 785,
+ t76x_js1_active = 786,
+ t76x_js1_wait_read = 788,
+ t76x_js1_wait_issue = 789,
+ t76x_js1_wait_depend = 790,
+ t76x_js1_wait_finish = 791,
+ t76x_js2_jobs = 792,
+ t76x_js2_tasks = 793,
+ t76x_js2_active = 794,
+ t76x_js2_wait_read = 796,
+ t76x_js2_wait_issue = 797,
+ t76x_js2_wait_depend = 798,
+ t76x_js2_wait_finish = 799,
+ t76x_ti_jobs_processed = 835,
+ t76x_ti_triangles = 836,
+ t76x_ti_quads = 837,
+ t76x_ti_polygons = 838,
+ t76x_ti_points = 839,
+ t76x_ti_lines = 840,
+ t76x_ti_vcache_hit = 841,
+ t76x_ti_vcache_miss = 842,
+ t76x_ti_front_facing = 843,
+ t76x_ti_back_facing = 844,
+ t76x_ti_prim_visible = 845,
+ t76x_ti_prim_culled = 846,
+ t76x_ti_prim_clipped = 847,
+ t76x_ti_level0 = 848,
+ t76x_ti_level1 = 849,
+ t76x_ti_level2 = 850,
+ t76x_ti_level3 = 851,
+ t76x_ti_level4 = 852,
+ t76x_ti_level5 = 853,
+ t76x_ti_level6 = 854,
+ t76x_ti_level7 = 855,
+ t76x_ti_command_1 = 856,
+ t76x_ti_command_2 = 857,
+ t76x_ti_command_3 = 858,
+ t76x_ti_command_4 = 859,
+ t76x_ti_command_5_7 = 860,
+ t76x_ti_command_8_15 = 861,
+ t76x_ti_command_16_63 = 862,
+ t76x_ti_command_64 = 863,
+ t76x_ti_compress_in = 864,
+ t76x_ti_compress_out = 865,
+ t76x_ti_compress_flush = 866,
+ t76x_ti_timestamps = 867,
+ t76x_ti_pcache_hit = 868,
+ t76x_ti_pcache_miss = 869,
+ t76x_ti_pcache_line = 870,
+ t76x_ti_pcache_stall = 871,
+ t76x_ti_wrbuf_hit = 872,
+ t76x_ti_wrbuf_miss = 873,
+ t76x_ti_wrbuf_line = 874,
+ t76x_ti_wrbuf_partial = 875,
+ t76x_ti_wrbuf_stall = 876,
+ t76x_ti_active = 877,
+ t76x_ti_loading_desc = 878,
+ t76x_ti_index_wait = 879,
+ t76x_ti_index_range_wait = 880,
+ t76x_ti_vertex_wait = 881,
+ t76x_ti_pcache_wait = 882,
+ t76x_ti_wrbuf_wait = 883,
+ t76x_ti_bus_read = 884,
+ t76x_ti_bus_write = 885,
+ t76x_ti_utlb_hit = 891,
+ t76x_ti_utlb_new_miss = 892,
+ t76x_ti_utlb_replay_full = 893,
+ t76x_ti_utlb_replay_miss = 894,
+ t76x_ti_utlb_stall = 895,
+ t76x_frag_active = 900,
+ t76x_frag_primitives = 901,
+ t76x_frag_primitives_dropped = 902,
+ t76x_frag_cycles_desc = 903,
+ t76x_frag_cycles_fpkq_active = 904,
+ t76x_frag_cycles_vert = 905,
+ t76x_frag_cycles_trisetup = 906,
+ t76x_frag_cycles_ezs_active = 907,
+ t76x_frag_threads = 908,
+ t76x_frag_dummy_threads = 909, // # nocheck
+ t76x_frag_quads_rast = 910,
+ t76x_frag_quads_ezs_test = 911,
+ t76x_frag_quads_ezs_killed = 912,
+ t76x_frag_threads_lzs_test = 913,
+ t76x_frag_threads_lzs_killed = 914,
+ t76x_frag_cycles_no_tile = 915,
+ t76x_frag_num_tiles = 916,
+ t76x_frag_trans_elim = 917,
+ t76x_compute_active = 918,
+ t76x_compute_tasks = 919,
+ t76x_compute_threads = 920,
+ t76x_compute_cycles_desc = 921,
+ t76x_tripipe_active = 922,
+ t76x_arith_words = 923,
+ t76x_arith_cycles_reg = 924,
+ t76x_arith_cycles_l0 = 925,
+ t76x_arith_frag_depend = 926,
+ t76x_ls_words = 927,
+ t76x_ls_issues = 928,
+ t76x_ls_reissue_attr = 929,
+ t76x_ls_reissues_vary = 930,
+ t76x_ls_vary_rv_miss = 931,
+ t76x_ls_vary_rv_hit = 932,
+ t76x_ls_no_unpark = 933,
+ t76x_tex_words = 934,
+ t76x_tex_bubbles = 935,
+ t76x_tex_words_l0 = 936,
+ t76x_tex_words_desc = 937,
+ t76x_tex_issues = 938,
+ t76x_tex_recirc_fmiss = 939,
+ t76x_tex_recirc_desc = 940,
+ t76x_tex_recirc_multi = 941,
+ t76x_tex_recirc_pmiss = 942,
+ t76x_tex_recirc_conf = 943,
+ t76x_lsc_read_hits = 944,
+ t76x_lsc_read_op = 945,
+ t76x_lsc_write_hits = 946,
+ t76x_lsc_write_op = 947,
+ t76x_lsc_atomic_hits = 948,
+ t76x_lsc_atomic_op = 949,
+ t76x_lsc_line_fetches = 950,
+ t76x_lsc_dirty_line = 951,
+ t76x_lsc_snoops = 952,
+ t76x_axi_tlb_stall = 953,
+ t76x_axi_tlb_miss = 954,
+ t76x_axi_tlb_transaction = 955,
+ t76x_ls_tlb_miss = 956,
+ t76x_ls_tlb_hit = 957,
+ t76x_axi_beats_read = 958,
+ t76x_axi_beats_written = 959,
+ t76x_mmu_hit = 964,
+ t76x_mmu_new_miss = 965,
+ t76x_mmu_replay_full = 966,
+ t76x_mmu_replay_miss = 967,
+ t76x_mmu_table_walk = 968,
+ t76x_mmu_requests = 969,
+ t76x_utlb_hit = 972,
+ t76x_utlb_new_miss = 973,
+ t76x_utlb_replay_full = 974,
+ t76x_utlb_replay_miss = 975,
+ t76x_utlb_stall = 976,
+ t76x_l2_ext_write_beats = 990,
+ t76x_l2_ext_read_beats = 991,
+ t76x_l2_any_lookup = 992,
+ t76x_l2_read_lookup = 993,
+ t76x_l2_sread_lookup = 994,
+ t76x_l2_read_replay = 995,
+ t76x_l2_read_snoop = 996,
+ t76x_l2_read_hit = 997,
+ t76x_l2_clean_miss = 998,
+ t76x_l2_write_lookup = 999,
+ t76x_l2_swrite_lookup = 1000,
+ t76x_l2_write_replay = 1001,
+ t76x_l2_write_snoop = 1002,
+ t76x_l2_write_hit = 1003,
+ t76x_l2_ext_read_full = 1004,
+ t76x_l2_ext_write_full = 1006,
+ t76x_l2_ext_r_w_hazard = 1007,
+ t76x_l2_ext_read = 1008,
+ t76x_l2_ext_read_line = 1009,
+ t76x_l2_ext_write = 1010,
+ t76x_l2_ext_write_line = 1011,
+ t76x_l2_ext_write_small = 1012,
+ t76x_l2_ext_barrier = 1013,
+ t76x_l2_ext_ar_stall = 1014,
+ t76x_l2_ext_r_buf_full = 1015,
+ t76x_l2_ext_rd_buf_full = 1016,
+ t76x_l2_ext_r_raw = 1017,
+ t76x_l2_ext_w_stall = 1018,
+ t76x_l2_ext_w_buf_full = 1019,
+ t76x_l2_tag_hazard = 1021,
+ t76x_l2_snoop_full = 1022,
+ t76x_l2_replay_full = 1023,
+ t82x_messages_sent = 1028,
+ t82x_messages_received = 1029,
+ t82x_gpu_active = 1030,
+ t82x_irq_active = 1031,
+ t82x_js0_jobs = 1032,
+ t82x_js0_tasks = 1033,
+ t82x_js0_active = 1034,
+ t82x_js0_wait_read = 1036,
+ t82x_js0_wait_issue = 1037,
+ t82x_js0_wait_depend = 1038,
+ t82x_js0_wait_finish = 1039,
+ t82x_js1_jobs = 1040,
+ t82x_js1_tasks = 1041,
+ t82x_js1_active = 1042,
+ t82x_js1_wait_read = 1044,
+ t82x_js1_wait_issue = 1045,
+ t82x_js1_wait_depend = 1046,
+ t82x_js1_wait_finish = 1047,
+ t82x_js2_jobs = 1048,
+ t82x_js2_tasks = 1049,
+ t82x_js2_active = 1050,
+ t82x_js2_wait_read = 1052,
+ t82x_js2_wait_issue = 1053,
+ t82x_js2_wait_depend = 1054,
+ t82x_js2_wait_finish = 1055,
+ t82x_ti_jobs_processed = 1091,
+ t82x_ti_triangles = 1092,
+ t82x_ti_quads = 1093,
+ t82x_ti_polygons = 1094,
+ t82x_ti_points = 1095,
+ t82x_ti_lines = 1096,
+ t82x_ti_front_facing = 1097,
+ t82x_ti_back_facing = 1098,
+ t82x_ti_prim_visible = 1099,
+ t82x_ti_prim_culled = 1100,
+ t82x_ti_prim_clipped = 1101,
+ t82x_ti_active = 1110,
+ t82x_frag_active = 1156,
+ t82x_frag_primitives = 1157,
+ t82x_frag_primitives_dropped = 1158,
+ t82x_frag_cycles_desc = 1159,
+ t82x_frag_cycles_fpkq_active = 1160,
+ t82x_frag_cycles_vert = 1161,
+ t82x_frag_cycles_trisetup = 1162,
+ t82x_frag_cycles_ezs_active = 1163,
+ t82x_frag_threads = 1164,
+ t82x_frag_dummy_threads = 1165, // # nocheck
+ t82x_frag_quads_rast = 1166,
+ t82x_frag_quads_ezs_test = 1167,
+ t82x_frag_quads_ezs_killed = 1168,
+ t82x_frag_threads_lzs_test = 1169,
+ t82x_frag_threads_lzs_killed = 1170,
+ t82x_frag_cycles_no_tile = 1171,
+ t82x_frag_num_tiles = 1172,
+ t82x_frag_trans_elim = 1173,
+ t82x_compute_active = 1174,
+ t82x_compute_tasks = 1175,
+ t82x_compute_threads = 1176,
+ t82x_compute_cycles_desc = 1177,
+ t82x_tripipe_active = 1178,
+ t82x_arith_words = 1179,
+ t82x_arith_cycles_reg = 1180,
+ t82x_arith_cycles_l0 = 1181,
+ t82x_arith_frag_depend = 1182,
+ t82x_ls_words = 1183,
+ t82x_ls_issues = 1184,
+ t82x_ls_reissue_attr = 1185,
+ t82x_ls_reissues_vary = 1186,
+ t82x_ls_vary_rv_miss = 1187,
+ t82x_ls_vary_rv_hit = 1188,
+ t82x_ls_no_unpark = 1189,
+ t82x_tex_words = 1190,
+ t82x_tex_bubbles = 1191,
+ t82x_tex_words_l0 = 1192,
+ t82x_tex_words_desc = 1193,
+ t82x_tex_issues = 1194,
+ t82x_tex_recirc_fmiss = 1195,
+ t82x_tex_recirc_desc = 1196,
+ t82x_tex_recirc_multi = 1197,
+ t82x_tex_recirc_pmiss = 1198,
+ t82x_tex_recirc_conf = 1199,
+ t82x_lsc_read_hits = 1200,
+ t82x_lsc_read_op = 1201,
+ t82x_lsc_write_hits = 1202,
+ t82x_lsc_write_op = 1203,
+ t82x_lsc_atomic_hits = 1204,
+ t82x_lsc_atomic_op = 1205,
+ t82x_lsc_line_fetches = 1206,
+ t82x_lsc_dirty_line = 1207,
+ t82x_lsc_snoops = 1208,
+ t82x_axi_tlb_stall = 1209,
+ t82x_axi_tlb_miss = 1210,
+ t82x_axi_tlb_transaction = 1211,
+ t82x_ls_tlb_miss = 1212,
+ t82x_ls_tlb_hit = 1213,
+ t82x_axi_beats_read = 1214,
+ t82x_axi_beats_written = 1215,
+ t82x_mmu_hit = 1220,
+ t82x_mmu_new_miss = 1221,
+ t82x_mmu_replay_full = 1222,
+ t82x_mmu_replay_miss = 1223,
+ t82x_mmu_table_walk = 1224,
+ t82x_mmu_requests = 1225,
+ t82x_utlb_hit = 1228,
+ t82x_utlb_new_miss = 1229,
+ t82x_utlb_replay_full = 1230,
+ t82x_utlb_replay_miss = 1231,
+ t82x_utlb_stall = 1232,
+ t82x_l2_ext_write_beats = 1246,
+ t82x_l2_ext_read_beats = 1247,
+ t82x_l2_any_lookup = 1248,
+ t82x_l2_read_lookup = 1249,
+ t82x_l2_sread_lookup = 1250,
+ t82x_l2_read_replay = 1251,
+ t82x_l2_read_snoop = 1252,
+ t82x_l2_read_hit = 1253,
+ t82x_l2_clean_miss = 1254,
+ t82x_l2_write_lookup = 1255,
+ t82x_l2_swrite_lookup = 1256,
+ t82x_l2_write_replay = 1257,
+ t82x_l2_write_snoop = 1258,
+ t82x_l2_write_hit = 1259,
+ t82x_l2_ext_read_full = 1260,
+ t82x_l2_ext_write_full = 1262,
+ t82x_l2_ext_r_w_hazard = 1263,
+ t82x_l2_ext_read = 1264,
+ t82x_l2_ext_read_line = 1265,
+ t82x_l2_ext_write = 1266,
+ t82x_l2_ext_write_line = 1267,
+ t82x_l2_ext_write_small = 1268,
+ t82x_l2_ext_barrier = 1269,
+ t82x_l2_ext_ar_stall = 1270,
+ t82x_l2_ext_r_buf_full = 1271,
+ t82x_l2_ext_rd_buf_full = 1272,
+ t82x_l2_ext_r_raw = 1273,
+ t82x_l2_ext_w_stall = 1274,
+ t82x_l2_ext_w_buf_full = 1275,
+ t82x_l2_tag_hazard = 1277,
+ t82x_l2_snoop_full = 1278,
+ t82x_l2_replay_full = 1279,
+ t83x_messages_sent = 1284,
+ t83x_messages_received = 1285,
+ t83x_gpu_active = 1286,
+ t83x_irq_active = 1287,
+ t83x_js0_jobs = 1288,
+ t83x_js0_tasks = 1289,
+ t83x_js0_active = 1290,
+ t83x_js0_wait_read = 1292,
+ t83x_js0_wait_issue = 1293,
+ t83x_js0_wait_depend = 1294,
+ t83x_js0_wait_finish = 1295,
+ t83x_js1_jobs = 1296,
+ t83x_js1_tasks = 1297,
+ t83x_js1_active = 1298,
+ t83x_js1_wait_read = 1300,
+ t83x_js1_wait_issue = 1301,
+ t83x_js1_wait_depend = 1302,
+ t83x_js1_wait_finish = 1303,
+ t83x_js2_jobs = 1304,
+ t83x_js2_tasks = 1305,
+ t83x_js2_active = 1306,
+ t83x_js2_wait_read = 1308,
+ t83x_js2_wait_issue = 1309,
+ t83x_js2_wait_depend = 1310,
+ t83x_js2_wait_finish = 1311,
+ t83x_ti_jobs_processed = 1347,
+ t83x_ti_triangles = 1348,
+ t83x_ti_quads = 1349,
+ t83x_ti_polygons = 1350,
+ t83x_ti_points = 1351,
+ t83x_ti_lines = 1352,
+ t83x_ti_front_facing = 1353,
+ t83x_ti_back_facing = 1354,
+ t83x_ti_prim_visible = 1355,
+ t83x_ti_prim_culled = 1356,
+ t83x_ti_prim_clipped = 1357,
+ t83x_ti_active = 1366,
+ t83x_frag_active = 1412,
+ t83x_frag_primitives = 1413,
+ t83x_frag_primitives_dropped = 1414,
+ t83x_frag_cycles_desc = 1415,
+ t83x_frag_cycles_fpkq_active = 1416,
+ t83x_frag_cycles_vert = 1417,
+ t83x_frag_cycles_trisetup = 1418,
+ t83x_frag_cycles_ezs_active = 1419,
+ t83x_frag_threads = 1420,
+ t83x_frag_dummy_threads = 1421, // # nocheck
+ t83x_frag_quads_rast = 1422,
+ t83x_frag_quads_ezs_test = 1423,
+ t83x_frag_quads_ezs_killed = 1424,
+ t83x_frag_threads_lzs_test = 1425,
+ t83x_frag_threads_lzs_killed = 1426,
+ t83x_frag_cycles_no_tile = 1427,
+ t83x_frag_num_tiles = 1428,
+ t83x_frag_trans_elim = 1429,
+ t83x_compute_active = 1430,
+ t83x_compute_tasks = 1431,
+ t83x_compute_threads = 1432,
+ t83x_compute_cycles_desc = 1433,
+ t83x_tripipe_active = 1434,
+ t83x_arith_words = 1435,
+ t83x_arith_cycles_reg = 1436,
+ t83x_arith_cycles_l0 = 1437,
+ t83x_arith_frag_depend = 1438,
+ t83x_ls_words = 1439,
+ t83x_ls_issues = 1440,
+ t83x_ls_reissue_attr = 1441,
+ t83x_ls_reissues_vary = 1442,
+ t83x_ls_vary_rv_miss = 1443,
+ t83x_ls_vary_rv_hit = 1444,
+ t83x_ls_no_unpark = 1445,
+ t83x_tex_words = 1446,
+ t83x_tex_bubbles = 1447,
+ t83x_tex_words_l0 = 1448,
+ t83x_tex_words_desc = 1449,
+ t83x_tex_issues = 1450,
+ t83x_tex_recirc_fmiss = 1451,
+ t83x_tex_recirc_desc = 1452,
+ t83x_tex_recirc_multi = 1453,
+ t83x_tex_recirc_pmiss = 1454,
+ t83x_tex_recirc_conf = 1455,
+ t83x_lsc_read_hits = 1456,
+ t83x_lsc_read_op = 1457,
+ t83x_lsc_write_hits = 1458,
+ t83x_lsc_write_op = 1459,
+ t83x_lsc_atomic_hits = 1460,
+ t83x_lsc_atomic_op = 1461,
+ t83x_lsc_line_fetches = 1462,
+ t83x_lsc_dirty_line = 1463,
+ t83x_lsc_snoops = 1464,
+ t83x_axi_tlb_stall = 1465,
+ t83x_axi_tlb_miss = 1466,
+ t83x_axi_tlb_transaction = 1467,
+ t83x_ls_tlb_miss = 1468,
+ t83x_ls_tlb_hit = 1469,
+ t83x_axi_beats_read = 1470,
+ t83x_axi_beats_written = 1471,
+ t83x_mmu_hit = 1476,
+ t83x_mmu_new_miss = 1477,
+ t83x_mmu_replay_full = 1478,
+ t83x_mmu_replay_miss = 1479,
+ t83x_mmu_table_walk = 1480,
+ t83x_mmu_requests = 1481,
+ t83x_utlb_hit = 1484,
+ t83x_utlb_new_miss = 1485,
+ t83x_utlb_replay_full = 1486,
+ t83x_utlb_replay_miss = 1487,
+ t83x_utlb_stall = 1488,
+ t83x_l2_ext_write_beats = 1502,
+ t83x_l2_ext_read_beats = 1503,
+ t83x_l2_any_lookup = 1504,
+ t83x_l2_read_lookup = 1505,
+ t83x_l2_sread_lookup = 1506,
+ t83x_l2_read_replay = 1507,
+ t83x_l2_read_snoop = 1508,
+ t83x_l2_read_hit = 1509,
+ t83x_l2_clean_miss = 1510,
+ t83x_l2_write_lookup = 1511,
+ t83x_l2_swrite_lookup = 1512,
+ t83x_l2_write_replay = 1513,
+ t83x_l2_write_snoop = 1514,
+ t83x_l2_write_hit = 1515,
+ t83x_l2_ext_read_full = 1516,
+ t83x_l2_ext_write_full = 1518,
+ t83x_l2_ext_r_w_hazard = 1519,
+ t83x_l2_ext_read = 1520,
+ t83x_l2_ext_read_line = 1521,
+ t83x_l2_ext_write = 1522,
+ t83x_l2_ext_write_line = 1523,
+ t83x_l2_ext_write_small = 1524,
+ t83x_l2_ext_barrier = 1525,
+ t83x_l2_ext_ar_stall = 1526,
+ t83x_l2_ext_r_buf_full = 1527,
+ t83x_l2_ext_rd_buf_full = 1528,
+ t83x_l2_ext_r_raw = 1529,
+ t83x_l2_ext_w_stall = 1530,
+ t83x_l2_ext_w_buf_full = 1531,
+ t83x_l2_tag_hazard = 1533,
+ t83x_l2_snoop_full = 1534,
+ t83x_l2_replay_full = 1535,
+ t86x_messages_sent = 1540,
+ t86x_messages_received = 1541,
+ t86x_gpu_active = 1542,
+ t86x_irq_active = 1543,
+ t86x_js0_jobs = 1544,
+ t86x_js0_tasks = 1545,
+ t86x_js0_active = 1546,
+ t86x_js0_wait_read = 1548,
+ t86x_js0_wait_issue = 1549,
+ t86x_js0_wait_depend = 1550,
+ t86x_js0_wait_finish = 1551,
+ t86x_js1_jobs = 1552,
+ t86x_js1_tasks = 1553,
+ t86x_js1_active = 1554,
+ t86x_js1_wait_read = 1556,
+ t86x_js1_wait_issue = 1557,
+ t86x_js1_wait_depend = 1558,
+ t86x_js1_wait_finish = 1559,
+ t86x_js2_jobs = 1560,
+ t86x_js2_tasks = 1561,
+ t86x_js2_active = 1562,
+ t86x_js2_wait_read = 1564,
+ t86x_js2_wait_issue = 1565,
+ t86x_js2_wait_depend = 1566,
+ t86x_js2_wait_finish = 1567,
+ t86x_ti_jobs_processed = 1603,
+ t86x_ti_triangles = 1604,
+ t86x_ti_quads = 1605,
+ t86x_ti_polygons = 1606,
+ t86x_ti_points = 1607,
+ t86x_ti_lines = 1608,
+ t86x_ti_vcache_hit = 1609,
+ t86x_ti_vcache_miss = 1610,
+ t86x_ti_front_facing = 1611,
+ t86x_ti_back_facing = 1612,
+ t86x_ti_prim_visible = 1613,
+ t86x_ti_prim_culled = 1614,
+ t86x_ti_prim_clipped = 1615,
+ t86x_ti_level0 = 1616,
+ t86x_ti_level1 = 1617,
+ t86x_ti_level2 = 1618,
+ t86x_ti_level3 = 1619,
+ t86x_ti_level4 = 1620,
+ t86x_ti_level5 = 1621,
+ t86x_ti_level6 = 1622,
+ t86x_ti_level7 = 1623,
+ t86x_ti_command_1 = 1624,
+ t86x_ti_command_2 = 1625,
+ t86x_ti_command_3 = 1626,
+ t86x_ti_command_4 = 1627,
+ t86x_ti_command_5_7 = 1628,
+ t86x_ti_command_8_15 = 1629,
+ t86x_ti_command_16_63 = 1630,
+ t86x_ti_command_64 = 1631,
+ t86x_ti_compress_in = 1632,
+ t86x_ti_compress_out = 1633,
+ t86x_ti_compress_flush = 1634,
+ t86x_ti_timestamps = 1635,
+ t86x_ti_pcache_hit = 1636,
+ t86x_ti_pcache_miss = 1637,
+ t86x_ti_pcache_line = 1638,
+ t86x_ti_pcache_stall = 1639,
+ t86x_ti_wrbuf_hit = 1640,
+ t86x_ti_wrbuf_miss = 1641,
+ t86x_ti_wrbuf_line = 1642,
+ t86x_ti_wrbuf_partial = 1643,
+ t86x_ti_wrbuf_stall = 1644,
+ t86x_ti_active = 1645,
+ t86x_ti_loading_desc = 1646,
+ t86x_ti_index_wait = 1647,
+ t86x_ti_index_range_wait = 1648,
+ t86x_ti_vertex_wait = 1649,
+ t86x_ti_pcache_wait = 1650,
+ t86x_ti_wrbuf_wait = 1651,
+ t86x_ti_bus_read = 1652,
+ t86x_ti_bus_write = 1653,
+ t86x_ti_utlb_hit = 1659,
+ t86x_ti_utlb_new_miss = 1660,
+ t86x_ti_utlb_replay_full = 1661,
+ t86x_ti_utlb_replay_miss = 1662,
+ t86x_ti_utlb_stall = 1663,
+ t86x_frag_active = 1668,
+ t86x_frag_primitives = 1669,
+ t86x_frag_primitives_dropped = 1670,
+ t86x_frag_cycles_desc = 1671,
+ t86x_frag_cycles_fpkq_active = 1672,
+ t86x_frag_cycles_vert = 1673,
+ t86x_frag_cycles_trisetup = 1674,
+ t86x_frag_cycles_ezs_active = 1675,
+ t86x_frag_threads = 1676,
+ t86x_frag_dummy_threads = 1677, // # nocheck
+ t86x_frag_quads_rast = 1678,
+ t86x_frag_quads_ezs_test = 1679,
+ t86x_frag_quads_ezs_killed = 1680,
+ t86x_frag_threads_lzs_test = 1681,
+ t86x_frag_threads_lzs_killed = 1682,
+ t86x_frag_cycles_no_tile = 1683,
+ t86x_frag_num_tiles = 1684,
+ t86x_frag_trans_elim = 1685,
+ t86x_compute_active = 1686,
+ t86x_compute_tasks = 1687,
+ t86x_compute_threads = 1688,
+ t86x_compute_cycles_desc = 1689,
+ t86x_tripipe_active = 1690,
+ t86x_arith_words = 1691,
+ t86x_arith_cycles_reg = 1692,
+ t86x_arith_cycles_l0 = 1693,
+ t86x_arith_frag_depend = 1694,
+ t86x_ls_words = 1695,
+ t86x_ls_issues = 1696,
+ t86x_ls_reissue_attr = 1697,
+ t86x_ls_reissues_vary = 1698,
+ t86x_ls_vary_rv_miss = 1699,
+ t86x_ls_vary_rv_hit = 1700,
+ t86x_ls_no_unpark = 1701,
+ t86x_tex_words = 1702,
+ t86x_tex_bubbles = 1703,
+ t86x_tex_words_l0 = 1704,
+ t86x_tex_words_desc = 1705,
+ t86x_tex_issues = 1706,
+ t86x_tex_recirc_fmiss = 1707,
+ t86x_tex_recirc_desc = 1708,
+ t86x_tex_recirc_multi = 1709,
+ t86x_tex_recirc_pmiss = 1710,
+ t86x_tex_recirc_conf = 1711,
+ t86x_lsc_read_hits = 1712,
+ t86x_lsc_read_op = 1713,
+ t86x_lsc_write_hits = 1714,
+ t86x_lsc_write_op = 1715,
+ t86x_lsc_atomic_hits = 1716,
+ t86x_lsc_atomic_op = 1717,
+ t86x_lsc_line_fetches = 1718,
+ t86x_lsc_dirty_line = 1719,
+ t86x_lsc_snoops = 1720,
+ t86x_axi_tlb_stall = 1721,
+ t86x_axi_tlb_miss = 1722,
+ t86x_axi_tlb_transaction = 1723,
+ t86x_ls_tlb_miss = 1724,
+ t86x_ls_tlb_hit = 1725,
+ t86x_axi_beats_read = 1726,
+ t86x_axi_beats_written = 1727,
+ t86x_mmu_hit = 1732,
+ t86x_mmu_new_miss = 1733,
+ t86x_mmu_replay_full = 1734,
+ t86x_mmu_replay_miss = 1735,
+ t86x_mmu_table_walk = 1736,
+ t86x_mmu_requests = 1737,
+ t86x_utlb_hit = 1740,
+ t86x_utlb_new_miss = 1741,
+ t86x_utlb_replay_full = 1742,
+ t86x_utlb_replay_miss = 1743,
+ t86x_utlb_stall = 1744,
+ t86x_l2_ext_write_beats = 1758,
+ t86x_l2_ext_read_beats = 1759,
+ t86x_l2_any_lookup = 1760,
+ t86x_l2_read_lookup = 1761,
+ t86x_l2_sread_lookup = 1762,
+ t86x_l2_read_replay = 1763,
+ t86x_l2_read_snoop = 1764,
+ t86x_l2_read_hit = 1765,
+ t86x_l2_clean_miss = 1766,
+ t86x_l2_write_lookup = 1767,
+ t86x_l2_swrite_lookup = 1768,
+ t86x_l2_write_replay = 1769,
+ t86x_l2_write_snoop = 1770,
+ t86x_l2_write_hit = 1771,
+ t86x_l2_ext_read_full = 1772,
+ t86x_l2_ext_write_full = 1774,
+ t86x_l2_ext_r_w_hazard = 1775,
+ t86x_l2_ext_read = 1776,
+ t86x_l2_ext_read_line = 1777,
+ t86x_l2_ext_write = 1778,
+ t86x_l2_ext_write_line = 1779,
+ t86x_l2_ext_write_small = 1780,
+ t86x_l2_ext_barrier = 1781,
+ t86x_l2_ext_ar_stall = 1782,
+ t86x_l2_ext_r_buf_full = 1783,
+ t86x_l2_ext_rd_buf_full = 1784,
+ t86x_l2_ext_r_raw = 1785,
+ t86x_l2_ext_w_stall = 1786,
+ t86x_l2_ext_w_buf_full = 1787,
+ t86x_l2_tag_hazard = 1789,
+ t86x_l2_snoop_full = 1790,
+ t86x_l2_replay_full = 1791,
+ t88x_messages_sent = 1796,
+ t88x_messages_received = 1797,
+ t88x_gpu_active = 1798,
+ t88x_irq_active = 1799,
+ t88x_js0_jobs = 1800,
+ t88x_js0_tasks = 1801,
+ t88x_js0_active = 1802,
+ t88x_js0_wait_read = 1804,
+ t88x_js0_wait_issue = 1805,
+ t88x_js0_wait_depend = 1806,
+ t88x_js0_wait_finish = 1807,
+ t88x_js1_jobs = 1808,
+ t88x_js1_tasks = 1809,
+ t88x_js1_active = 1810,
+ t88x_js1_wait_read = 1812,
+ t88x_js1_wait_issue = 1813,
+ t88x_js1_wait_depend = 1814,
+ t88x_js1_wait_finish = 1815,
+ t88x_js2_jobs = 1816,
+ t88x_js2_tasks = 1817,
+ t88x_js2_active = 1818,
+ t88x_js2_wait_read = 1820,
+ t88x_js2_wait_issue = 1821,
+ t88x_js2_wait_depend = 1822,
+ t88x_js2_wait_finish = 1823,
+ t88x_ti_jobs_processed = 1859,
+ t88x_ti_triangles = 1860,
+ t88x_ti_quads = 1861,
+ t88x_ti_polygons = 1862,
+ t88x_ti_points = 1863,
+ t88x_ti_lines = 1864,
+ t88x_ti_vcache_hit = 1865,
+ t88x_ti_vcache_miss = 1866,
+ t88x_ti_front_facing = 1867,
+ t88x_ti_back_facing = 1868,
+ t88x_ti_prim_visible = 1869,
+ t88x_ti_prim_culled = 1870,
+ t88x_ti_prim_clipped = 1871,
+ t88x_ti_level0 = 1872,
+ t88x_ti_level1 = 1873,
+ t88x_ti_level2 = 1874,
+ t88x_ti_level3 = 1875,
+ t88x_ti_level4 = 1876,
+ t88x_ti_level5 = 1877,
+ t88x_ti_level6 = 1878,
+ t88x_ti_level7 = 1879,
+ t88x_ti_command_1 = 1880,
+ t88x_ti_command_2 = 1881,
+ t88x_ti_command_3 = 1882,
+ t88x_ti_command_4 = 1883,
+ t88x_ti_command_5_7 = 1884,
+ t88x_ti_command_8_15 = 1885,
+ t88x_ti_command_16_63 = 1886,
+ t88x_ti_command_64 = 1887,
+ t88x_ti_compress_in = 1888,
+ t88x_ti_compress_out = 1889,
+ t88x_ti_compress_flush = 1890,
+ t88x_ti_timestamps = 1891,
+ t88x_ti_pcache_hit = 1892,
+ t88x_ti_pcache_miss = 1893,
+ t88x_ti_pcache_line = 1894,
+ t88x_ti_pcache_stall = 1895,
+ t88x_ti_wrbuf_hit = 1896,
+ t88x_ti_wrbuf_miss = 1897,
+ t88x_ti_wrbuf_line = 1898,
+ t88x_ti_wrbuf_partial = 1899,
+ t88x_ti_wrbuf_stall = 1900,
+ t88x_ti_active = 1901,
+ t88x_ti_loading_desc = 1902,
+ t88x_ti_index_wait = 1903,
+ t88x_ti_index_range_wait = 1904,
+ t88x_ti_vertex_wait = 1905,
+ t88x_ti_pcache_wait = 1906,
+ t88x_ti_wrbuf_wait = 1907,
+ t88x_ti_bus_read = 1908,
+ t88x_ti_bus_write = 1909,
+ t88x_ti_utlb_hit = 1915,
+ t88x_ti_utlb_new_miss = 1916,
+ t88x_ti_utlb_replay_full = 1917,
+ t88x_ti_utlb_replay_miss = 1918,
+ t88x_ti_utlb_stall = 1919,
+ t88x_frag_active = 1924,
+ t88x_frag_primitives = 1925,
+ t88x_frag_primitives_dropped = 1926,
+ t88x_frag_cycles_desc = 1927,
+ t88x_frag_cycles_fpkq_active = 1928,
+ t88x_frag_cycles_vert = 1929,
+ t88x_frag_cycles_trisetup = 1930,
+ t88x_frag_cycles_ezs_active = 1931,
+ t88x_frag_threads = 1932,
+ t88x_frag_dummy_threads = 1933, // # nocheck
+ t88x_frag_quads_rast = 1934,
+ t88x_frag_quads_ezs_test = 1935,
+ t88x_frag_quads_ezs_killed = 1936,
+ t88x_frag_threads_lzs_test = 1937,
+ t88x_frag_threads_lzs_killed = 1938,
+ t88x_frag_cycles_no_tile = 1939,
+ t88x_frag_num_tiles = 1940,
+ t88x_frag_trans_elim = 1941,
+ t88x_compute_active = 1942,
+ t88x_compute_tasks = 1943,
+ t88x_compute_threads = 1944,
+ t88x_compute_cycles_desc = 1945,
+ t88x_tripipe_active = 1946,
+ t88x_arith_words = 1947,
+ t88x_arith_cycles_reg = 1948,
+ t88x_arith_cycles_l0 = 1949,
+ t88x_arith_frag_depend = 1950,
+ t88x_ls_words = 1951,
+ t88x_ls_issues = 1952,
+ t88x_ls_reissue_attr = 1953,
+ t88x_ls_reissues_vary = 1954,
+ t88x_ls_vary_rv_miss = 1955,
+ t88x_ls_vary_rv_hit = 1956,
+ t88x_ls_no_unpark = 1957,
+ t88x_tex_words = 1958,
+ t88x_tex_bubbles = 1959,
+ t88x_tex_words_l0 = 1960,
+ t88x_tex_words_desc = 1961,
+ t88x_tex_issues = 1962,
+ t88x_tex_recirc_fmiss = 1963,
+ t88x_tex_recirc_desc = 1964,
+ t88x_tex_recirc_multi = 1965,
+ t88x_tex_recirc_pmiss = 1966,
+ t88x_tex_recirc_conf = 1967,
+ t88x_lsc_read_hits = 1968,
+ t88x_lsc_read_op = 1969,
+ t88x_lsc_write_hits = 1970,
+ t88x_lsc_write_op = 1971,
+ t88x_lsc_atomic_hits = 1972,
+ t88x_lsc_atomic_op = 1973,
+ t88x_lsc_line_fetches = 1974,
+ t88x_lsc_dirty_line = 1975,
+ t88x_lsc_snoops = 1976,
+ t88x_axi_tlb_stall = 1977,
+ t88x_axi_tlb_miss = 1978,
+ t88x_axi_tlb_transaction = 1979,
+ t88x_ls_tlb_miss = 1980,
+ t88x_ls_tlb_hit = 1981,
+ t88x_axi_beats_read = 1982,
+ t88x_axi_beats_written = 1983,
+ t88x_mmu_hit = 1988,
+ t88x_mmu_new_miss = 1989,
+ t88x_mmu_replay_full = 1990,
+ t88x_mmu_replay_miss = 1991,
+ t88x_mmu_table_walk = 1992,
+ t88x_mmu_requests = 1993,
+ t88x_utlb_hit = 1996,
+ t88x_utlb_new_miss = 1997,
+ t88x_utlb_replay_full = 1998,
+ t88x_utlb_replay_miss = 1999,
+ t88x_utlb_stall = 2000,
+ t88x_l2_ext_write_beats = 2014,
+ t88x_l2_ext_read_beats = 2015,
+ t88x_l2_any_lookup = 2016,
+ t88x_l2_read_lookup = 2017,
+ t88x_l2_sread_lookup = 2018,
+ t88x_l2_read_replay = 2019,
+ t88x_l2_read_snoop = 2020,
+ t88x_l2_read_hit = 2021,
+ t88x_l2_clean_miss = 2022,
+ t88x_l2_write_lookup = 2023,
+ t88x_l2_swrite_lookup = 2024,
+ t88x_l2_write_replay = 2025,
+ t88x_l2_write_snoop = 2026,
+ t88x_l2_write_hit = 2027,
+ t88x_l2_ext_read_full = 2028,
+ t88x_l2_ext_write_full = 2030,
+ t88x_l2_ext_r_w_hazard = 2031,
+ t88x_l2_ext_read = 2032,
+ t88x_l2_ext_read_line = 2033,
+ t88x_l2_ext_write = 2034,
+ t88x_l2_ext_write_line = 2035,
+ t88x_l2_ext_write_small = 2036,
+ t88x_l2_ext_barrier = 2037,
+ t88x_l2_ext_ar_stall = 2038,
+ t88x_l2_ext_r_buf_full = 2039,
+ t88x_l2_ext_rd_buf_full = 2040,
+ t88x_l2_ext_r_raw = 2041,
+ t88x_l2_ext_w_stall = 2042,
+ t88x_l2_ext_w_buf_full = 2043,
+ t88x_l2_tag_hazard = 2045,
+ t88x_l2_snoop_full = 2046,
+ t88x_l2_replay_full = 2047,
+ thex_messages_sent = 2052,
+ thex_messages_received = 2053,
+ thex_gpu_active = 2054,
+ thex_irq_active = 2055,
+ thex_js0_jobs = 2056,
+ thex_js0_tasks = 2057,
+ thex_js0_active = 2058,
+ thex_js0_wait_read = 2060,
+ thex_js0_wait_issue = 2061,
+ thex_js0_wait_depend = 2062,
+ thex_js0_wait_finish = 2063,
+ thex_js1_jobs = 2064,
+ thex_js1_tasks = 2065,
+ thex_js1_active = 2066,
+ thex_js1_wait_read = 2068,
+ thex_js1_wait_issue = 2069,
+ thex_js1_wait_depend = 2070,
+ thex_js1_wait_finish = 2071,
+ thex_js2_jobs = 2072,
+ thex_js2_tasks = 2073,
+ thex_js2_active = 2074,
+ thex_js2_wait_read = 2076,
+ thex_js2_wait_issue = 2077,
+ thex_js2_wait_depend = 2078,
+ thex_js2_wait_finish = 2079,
+ thex_tiler_active = 2116,
+ thex_jobs_processed = 2117,
+ thex_triangles = 2118,
+ thex_lines = 2119,
+ thex_points = 2120,
+ thex_front_facing = 2121,
+ thex_back_facing = 2122,
+ thex_prim_visible = 2123,
+ thex_prim_culled = 2124,
+ thex_prim_clipped = 2125,
+ thex_prim_sat_culled = 2126,
+ thex_bus_read = 2129,
+ thex_bus_write = 2131,
+ thex_loading_desc = 2132,
+ thex_idvs_pos_shad_req = 2133,
+ thex_idvs_pos_shad_wait = 2134,
+ thex_idvs_pos_shad_stall = 2135,
+ thex_idvs_pos_fifo_full = 2136,
+ thex_prefetch_stall = 2137,
+ thex_vcache_hit = 2138,
+ thex_vcache_miss = 2139,
+ thex_vcache_line_wait = 2140,
+ thex_vfetch_pos_read_wait = 2141,
+ thex_vfetch_vertex_wait = 2142,
+ thex_vfetch_stall = 2143,
+ thex_primassy_stall = 2144,
+ thex_bbox_gen_stall = 2145,
+ thex_idvs_vbu_hit = 2146,
+ thex_idvs_vbu_miss = 2147,
+ thex_idvs_vbu_line_deallocate = 2148,
+ thex_idvs_var_shad_req = 2149,
+ thex_idvs_var_shad_stall = 2150,
+ thex_binner_stall = 2151,
+ thex_iter_stall = 2152,
+ thex_compress_miss = 2153,
+ thex_compress_stall = 2154,
+ thex_pcache_hit = 2155,
+ thex_pcache_miss = 2156,
+ thex_pcache_miss_stall = 2157,
+ thex_pcache_evict_stall = 2158,
+ thex_pmgr_ptr_wr_stall = 2159,
+ thex_pmgr_ptr_rd_stall = 2160,
+ thex_pmgr_cmd_wr_stall = 2161,
+ thex_wrbuf_active = 2162,
+ thex_wrbuf_hit = 2163,
+ thex_wrbuf_miss = 2164,
+ thex_wrbuf_no_free_line_stall = 2165,
+ thex_wrbuf_no_axi_id_stall = 2166,
+ thex_wrbuf_axi_stall = 2167,
+ thex_utlb_trans = 2171,
+ thex_utlb_trans_hit = 2172,
+ thex_utlb_trans_stall = 2173,
+ thex_utlb_trans_miss_delay = 2174,
+ thex_utlb_mmu_req = 2175,
+ thex_frag_active = 2180,
+ thex_frag_primitives = 2181,
+ thex_frag_prim_rast = 2182,
+ thex_frag_fpk_active = 2183,
+ thex_frag_starving = 2184,
+ thex_frag_warps = 2185,
+ thex_frag_partial_warps = 2186,
+ thex_frag_quads_rast = 2187,
+ thex_frag_quads_ezs_test = 2188,
+ thex_frag_quads_ezs_update = 2189,
+ thex_frag_quads_ezs_kill = 2190,
+ thex_frag_lzs_test = 2191,
+ thex_frag_lzs_kill = 2192,
+ thex_frag_ptiles = 2194,
+ thex_frag_trans_elim = 2195,
+ thex_quad_fpk_killer = 2196,
+ thex_compute_active = 2198,
+ thex_compute_tasks = 2199,
+ thex_compute_warps = 2200,
+ thex_compute_starving = 2201,
+ thex_exec_core_active = 2202,
+ thex_exec_active = 2203,
+ thex_exec_instr_count = 2204,
+ thex_exec_instr_diverged = 2205,
+ thex_exec_instr_starving = 2206,
+ thex_arith_instr_single_fma = 2207,
+ thex_arith_instr_double = 2208,
+ thex_arith_instr_msg = 2209,
+ thex_arith_instr_msg_only = 2210,
+ thex_tex_instr = 2211,
+ thex_tex_instr_mipmap = 2212,
+ thex_tex_instr_compressed = 2213,
+ thex_tex_instr_3d = 2214,
+ thex_tex_instr_trilinear = 2215,
+ thex_tex_coord_issue = 2216,
+ thex_tex_coord_stall = 2217,
+ thex_tex_starve_cache = 2218,
+ thex_tex_starve_filter = 2219,
+ thex_ls_mem_read_full = 2220,
+ thex_ls_mem_read_short = 2221,
+ thex_ls_mem_write_full = 2222,
+ thex_ls_mem_write_short = 2223,
+ thex_ls_mem_atomic = 2224,
+ thex_vary_instr = 2225,
+ thex_vary_slot_32 = 2226,
+ thex_vary_slot_16 = 2227,
+ thex_attr_instr = 2228,
+ thex_arith_instr_fp_mul = 2229,
+ thex_beats_rd_ftc = 2230,
+ thex_beats_rd_ftc_ext = 2231,
+ thex_beats_rd_lsc = 2232,
+ thex_beats_rd_lsc_ext = 2233,
+ thex_beats_rd_tex = 2234,
+ thex_beats_rd_tex_ext = 2235,
+ thex_beats_rd_other = 2236,
+ thex_beats_wr_lsc = 2237,
+ thex_beats_wr_tib = 2238,
+ thex_mmu_requests = 2244,
+ thex_l2_rd_msg_in = 2256,
+ thex_l2_rd_msg_in_stall = 2257,
+ thex_l2_wr_msg_in = 2258,
+ thex_l2_wr_msg_in_stall = 2259,
+ thex_l2_snp_msg_in = 2260,
+ thex_l2_snp_msg_in_stall = 2261,
+ thex_l2_rd_msg_out = 2262,
+ thex_l2_rd_msg_out_stall = 2263,
+ thex_l2_wr_msg_out = 2264,
+ thex_l2_any_lookup = 2265,
+ thex_l2_read_lookup = 2266,
+ thex_l2_write_lookup = 2267,
+ thex_l2_ext_snoop_lookup = 2268,
+ thex_l2_ext_read = 2269,
+ thex_l2_ext_read_nosnp = 2270,
+ thex_l2_ext_read_unique = 2271,
+ thex_l2_ext_read_beats = 2272,
+ thex_l2_ext_ar_stall = 2273,
+ thex_l2_ext_ar_cnt_q1 = 2274,
+ thex_l2_ext_ar_cnt_q2 = 2275,
+ thex_l2_ext_ar_cnt_q3 = 2276,
+ thex_l2_ext_rresp_0_127 = 2277,
+ thex_l2_ext_rresp_128_191 = 2278,
+ thex_l2_ext_rresp_192_255 = 2279,
+ thex_l2_ext_rresp_256_319 = 2280,
+ thex_l2_ext_rresp_320_383 = 2281,
+ thex_l2_ext_write = 2282,
+ thex_l2_ext_write_nosnp_full = 2283,
+ thex_l2_ext_write_nosnp_ptl = 2284,
+ thex_l2_ext_write_snp_full = 2285,
+ thex_l2_ext_write_snp_ptl = 2286,
+ thex_l2_ext_write_beats = 2287,
+ thex_l2_ext_w_stall = 2288,
+ thex_l2_ext_aw_cnt_q1 = 2289,
+ thex_l2_ext_aw_cnt_q2 = 2290,
+ thex_l2_ext_aw_cnt_q3 = 2291,
+ thex_l2_ext_snoop = 2292,
+ thex_l2_ext_snoop_stall = 2293,
+ thex_l2_ext_snoop_resp_clean = 2294,
+ thex_l2_ext_snoop_resp_data = 2295,
+ thex_l2_ext_snoop_internal = 2296,
+ tmix_messages_sent = 2308,
+ tmix_messages_received = 2309,
+ tmix_gpu_active = 2310,
+ tmix_irq_active = 2311,
+ tmix_js0_jobs = 2312,
+ tmix_js0_tasks = 2313,
+ tmix_js0_active = 2314,
+ tmix_js0_wait_read = 2316,
+ tmix_js0_wait_issue = 2317,
+ tmix_js0_wait_depend = 2318,
+ tmix_js0_wait_finish = 2319,
+ tmix_js1_jobs = 2320,
+ tmix_js1_tasks = 2321,
+ tmix_js1_active = 2322,
+ tmix_js1_wait_read = 2324,
+ tmix_js1_wait_issue = 2325,
+ tmix_js1_wait_depend = 2326,
+ tmix_js1_wait_finish = 2327,
+ tmix_js2_jobs = 2328,
+ tmix_js2_tasks = 2329,
+ tmix_js2_active = 2330,
+ tmix_js2_wait_read = 2332,
+ tmix_js2_wait_issue = 2333,
+ tmix_js2_wait_depend = 2334,
+ tmix_js2_wait_finish = 2335,
+ tmix_tiler_active = 2372,
+ tmix_jobs_processed = 2373,
+ tmix_triangles = 2374,
+ tmix_lines = 2375,
+ tmix_points = 2376,
+ tmix_front_facing = 2377,
+ tmix_back_facing = 2378,
+ tmix_prim_visible = 2379,
+ tmix_prim_culled = 2380,
+ tmix_prim_clipped = 2381,
+ tmix_prim_sat_culled = 2382,
+ tmix_bin_alloc_init = 2383,
+ tmix_bin_alloc_overflow = 2384,
+ tmix_bus_read = 2385,
+ tmix_bus_write = 2387,
+ tmix_loading_desc = 2388,
+ tmix_idvs_pos_shad_req = 2389,
+ tmix_idvs_pos_shad_wait = 2390,
+ tmix_idvs_pos_shad_stall = 2391,
+ tmix_idvs_pos_fifo_full = 2392,
+ tmix_prefetch_stall = 2393,
+ tmix_vcache_hit = 2394,
+ tmix_vcache_miss = 2395,
+ tmix_vcache_line_wait = 2396,
+ tmix_vfetch_pos_read_wait = 2397,
+ tmix_vfetch_vertex_wait = 2398,
+ tmix_vfetch_stall = 2399,
+ tmix_primassy_stall = 2400,
+ tmix_bbox_gen_stall = 2401,
+ tmix_idvs_vbu_hit = 2402,
+ tmix_idvs_vbu_miss = 2403,
+ tmix_idvs_vbu_line_deallocate = 2404,
+ tmix_idvs_var_shad_req = 2405,
+ tmix_idvs_var_shad_stall = 2406,
+ tmix_binner_stall = 2407,
+ tmix_iter_stall = 2408,
+ tmix_compress_miss = 2409,
+ tmix_compress_stall = 2410,
+ tmix_pcache_hit = 2411,
+ tmix_pcache_miss = 2412,
+ tmix_pcache_miss_stall = 2413,
+ tmix_pcache_evict_stall = 2414,
+ tmix_pmgr_ptr_wr_stall = 2415,
+ tmix_pmgr_ptr_rd_stall = 2416,
+ tmix_pmgr_cmd_wr_stall = 2417,
+ tmix_wrbuf_active = 2418,
+ tmix_wrbuf_hit = 2419,
+ tmix_wrbuf_miss = 2420,
+ tmix_wrbuf_no_free_line_stall = 2421,
+ tmix_wrbuf_no_axi_id_stall = 2422,
+ tmix_wrbuf_axi_stall = 2423,
+ tmix_utlb_trans = 2427,
+ tmix_utlb_trans_hit = 2428,
+ tmix_utlb_trans_stall = 2429,
+ tmix_utlb_trans_miss_delay = 2430,
+ tmix_utlb_mmu_req = 2431,
+ tmix_frag_active = 2436,
+ tmix_frag_primitives = 2437,
+ tmix_frag_prim_rast = 2438,
+ tmix_frag_fpk_active = 2439,
+ tmix_frag_starving = 2440,
+ tmix_frag_warps = 2441,
+ tmix_frag_partial_warps = 2442,
+ tmix_frag_quads_rast = 2443,
+ tmix_frag_quads_ezs_test = 2444,
+ tmix_frag_quads_ezs_update = 2445,
+ tmix_frag_quads_ezs_kill = 2446,
+ tmix_frag_lzs_test = 2447,
+ tmix_frag_lzs_kill = 2448,
+ tmix_frag_ptiles = 2450,
+ tmix_frag_trans_elim = 2451,
+ tmix_quad_fpk_killer = 2452,
+ tmix_compute_active = 2454,
+ tmix_compute_tasks = 2455,
+ tmix_compute_warps = 2456,
+ tmix_compute_starving = 2457,
+ tmix_exec_core_active = 2458,
+ tmix_exec_active = 2459,
+ tmix_exec_instr_count = 2460,
+ tmix_exec_instr_diverged = 2461,
+ tmix_exec_instr_starving = 2462,
+ tmix_arith_instr_single_fma = 2463,
+ tmix_arith_instr_double = 2464,
+ tmix_arith_instr_msg = 2465,
+ tmix_arith_instr_msg_only = 2466,
+ tmix_tex_instr = 2467,
+ tmix_tex_instr_mipmap = 2468,
+ tmix_tex_instr_compressed = 2469,
+ tmix_tex_instr_3d = 2470,
+ tmix_tex_instr_trilinear = 2471,
+ tmix_tex_coord_issue = 2472,
+ tmix_tex_coord_stall = 2473,
+ tmix_tex_starve_cache = 2474,
+ tmix_tex_starve_filter = 2475,
+ tmix_ls_mem_read_full = 2476,
+ tmix_ls_mem_read_short = 2477,
+ tmix_ls_mem_write_full = 2478,
+ tmix_ls_mem_write_short = 2479,
+ tmix_ls_mem_atomic = 2480,
+ tmix_vary_instr = 2481,
+ tmix_vary_slot_32 = 2482,
+ tmix_vary_slot_16 = 2483,
+ tmix_attr_instr = 2484,
+ tmix_arith_instr_fp_mul = 2485,
+ tmix_beats_rd_ftc = 2486,
+ tmix_beats_rd_ftc_ext = 2487,
+ tmix_beats_rd_lsc = 2488,
+ tmix_beats_rd_lsc_ext = 2489,
+ tmix_beats_rd_tex = 2490,
+ tmix_beats_rd_tex_ext = 2491,
+ tmix_beats_rd_other = 2492,
+ tmix_beats_wr_lsc = 2493,
+ tmix_beats_wr_tib = 2494,
+ tmix_mmu_requests = 2500,
+ tmix_l2_rd_msg_in = 2512,
+ tmix_l2_rd_msg_in_stall = 2513,
+ tmix_l2_wr_msg_in = 2514,
+ tmix_l2_wr_msg_in_stall = 2515,
+ tmix_l2_snp_msg_in = 2516,
+ tmix_l2_snp_msg_in_stall = 2517,
+ tmix_l2_rd_msg_out = 2518,
+ tmix_l2_rd_msg_out_stall = 2519,
+ tmix_l2_wr_msg_out = 2520,
+ tmix_l2_any_lookup = 2521,
+ tmix_l2_read_lookup = 2522,
+ tmix_l2_write_lookup = 2523,
+ tmix_l2_ext_snoop_lookup = 2524,
+ tmix_l2_ext_read = 2525,
+ tmix_l2_ext_read_nosnp = 2526,
+ tmix_l2_ext_read_unique = 2527,
+ tmix_l2_ext_read_beats = 2528,
+ tmix_l2_ext_ar_stall = 2529,
+ tmix_l2_ext_ar_cnt_q1 = 2530,
+ tmix_l2_ext_ar_cnt_q2 = 2531,
+ tmix_l2_ext_ar_cnt_q3 = 2532,
+ tmix_l2_ext_rresp_0_127 = 2533,
+ tmix_l2_ext_rresp_128_191 = 2534,
+ tmix_l2_ext_rresp_192_255 = 2535,
+ tmix_l2_ext_rresp_256_319 = 2536,
+ tmix_l2_ext_rresp_320_383 = 2537,
+ tmix_l2_ext_write = 2538,
+ tmix_l2_ext_write_nosnp_full = 2539,
+ tmix_l2_ext_write_nosnp_ptl = 2540,
+ tmix_l2_ext_write_snp_full = 2541,
+ tmix_l2_ext_write_snp_ptl = 2542,
+ tmix_l2_ext_write_beats = 2543,
+ tmix_l2_ext_w_stall = 2544,
+ tmix_l2_ext_aw_cnt_q1 = 2545,
+ tmix_l2_ext_aw_cnt_q2 = 2546,
+ tmix_l2_ext_aw_cnt_q3 = 2547,
+ tmix_l2_ext_snoop = 2548,
+ tmix_l2_ext_snoop_stall = 2549,
+ tmix_l2_ext_snoop_resp_clean = 2550,
+ tmix_l2_ext_snoop_resp_data = 2551,
+ tmix_l2_ext_snoop_internal = 2552,
+ tdvx_messages_sent = 2564,
+ tdvx_messages_received = 2565,
+ tdvx_gpu_active = 2566,
+ tdvx_irq_active = 2567,
+ tdvx_js0_jobs = 2568,
+ tdvx_js0_tasks = 2569,
+ tdvx_js0_active = 2570,
+ tdvx_js0_wait_flush = 2571,
+ tdvx_js0_wait_read = 2572,
+ tdvx_js0_wait_issue = 2573,
+ tdvx_js0_wait_depend = 2574,
+ tdvx_js0_wait_finish = 2575,
+ tdvx_js1_jobs = 2576,
+ tdvx_js1_tasks = 2577,
+ tdvx_js1_active = 2578,
+ tdvx_js1_wait_flush = 2579,
+ tdvx_js1_wait_read = 2580,
+ tdvx_js1_wait_issue = 2581,
+ tdvx_js1_wait_depend = 2582,
+ tdvx_js1_wait_finish = 2583,
+ tdvx_js2_jobs = 2584,
+ tdvx_js2_tasks = 2585,
+ tdvx_js2_active = 2586,
+ tdvx_js2_wait_flush = 2587,
+ tdvx_js2_wait_read = 2588,
+ tdvx_js2_wait_issue = 2589,
+ tdvx_js2_wait_depend = 2590,
+ tdvx_js2_wait_finish = 2591,
+ tdvx_cache_flush = 2623,
+ tdvx_tiler_active = 2628,
+ tdvx_jobs_processed = 2629,
+ tdvx_triangles = 2630,
+ tdvx_lines = 2631,
+ tdvx_points = 2632,
+ tdvx_front_facing = 2633,
+ tdvx_back_facing = 2634,
+ tdvx_prim_visible = 2635,
+ tdvx_prim_culled = 2636,
+ tdvx_prim_clipped = 2637,
+ tdvx_prim_sat_culled = 2638,
+ tdvx_bin_alloc_init = 2639,
+ tdvx_bin_alloc_overflow = 2640,
+ tdvx_bus_read = 2641,
+ tdvx_bus_write = 2643,
+ tdvx_loading_desc = 2644,
+ tdvx_idvs_pos_shad_req = 2645,
+ tdvx_idvs_pos_shad_wait = 2646,
+ tdvx_idvs_pos_shad_stall = 2647,
+ tdvx_idvs_pos_fifo_full = 2648,
+ tdvx_prefetch_stall = 2649,
+ tdvx_vcache_hit = 2650,
+ tdvx_vcache_miss = 2651,
+ tdvx_vcache_line_wait = 2652,
+ tdvx_vfetch_pos_read_wait = 2653,
+ tdvx_vfetch_vertex_wait = 2654,
+ tdvx_vfetch_stall = 2655,
+ tdvx_primassy_stall = 2656,
+ tdvx_bbox_gen_stall = 2657,
+ tdvx_idvs_vbu_hit = 2658,
+ tdvx_idvs_vbu_miss = 2659,
+ tdvx_idvs_vbu_line_deallocate = 2660,
+ tdvx_idvs_var_shad_req = 2661,
+ tdvx_idvs_var_shad_stall = 2662,
+ tdvx_binner_stall = 2663,
+ tdvx_iter_stall = 2664,
+ tdvx_compress_miss = 2665,
+ tdvx_compress_stall = 2666,
+ tdvx_pcache_hit = 2667,
+ tdvx_pcache_miss = 2668,
+ tdvx_pcache_miss_stall = 2669,
+ tdvx_pcache_evict_stall = 2670,
+ tdvx_pmgr_ptr_wr_stall = 2671,
+ tdvx_pmgr_ptr_rd_stall = 2672,
+ tdvx_pmgr_cmd_wr_stall = 2673,
+ tdvx_wrbuf_active = 2674,
+ tdvx_wrbuf_hit = 2675,
+ tdvx_wrbuf_miss = 2676,
+ tdvx_wrbuf_no_free_line_stall = 2677,
+ tdvx_wrbuf_no_axi_id_stall = 2678,
+ tdvx_wrbuf_axi_stall = 2679,
+ tdvx_utlb_trans = 2683,
+ tdvx_utlb_trans_hit = 2684,
+ tdvx_utlb_trans_stall = 2685,
+ tdvx_utlb_trans_miss_delay = 2686,
+ tdvx_utlb_mmu_req = 2687,
+ tdvx_frag_active = 2692,
+ tdvx_frag_primitives = 2693,
+ tdvx_frag_prim_rast = 2694,
+ tdvx_frag_fpk_active = 2695,
+ tdvx_frag_starving = 2696,
+ tdvx_frag_warps = 2697,
+ tdvx_frag_partial_warps = 2698,
+ tdvx_frag_quads_rast = 2699,
+ tdvx_frag_quads_ezs_test = 2700,
+ tdvx_frag_quads_ezs_update = 2701,
+ tdvx_frag_quads_ezs_kill = 2702,
+ tdvx_frag_lzs_test = 2703,
+ tdvx_frag_lzs_kill = 2704,
+ tdvx_frag_ptiles = 2706,
+ tdvx_frag_trans_elim = 2707,
+ tdvx_quad_fpk_killer = 2708,
+ tdvx_compute_active = 2710,
+ tdvx_compute_tasks = 2711,
+ tdvx_compute_warps = 2712,
+ tdvx_compute_starving = 2713,
+ tdvx_exec_core_active = 2714,
+ tdvx_exec_active = 2715,
+ tdvx_exec_instr_count = 2716,
+ tdvx_exec_instr_diverged = 2717,
+ tdvx_exec_instr_starving = 2718,
+ tdvx_arith_instr_single_fma = 2719,
+ tdvx_arith_instr_double = 2720,
+ tdvx_arith_instr_msg = 2721,
+ tdvx_arith_instr_msg_only = 2722,
+ tdvx_tex_msgi_num_quads = 2723,
+ tdvx_tex_dfch_num_passes = 2724,
+ tdvx_tex_dfch_num_passes_miss = 2725,
+ tdvx_tex_dfch_num_passes_mip_map = 2726,
+ tdvx_tex_tidx_num_split_mip_map = 2727,
+ tdvx_tex_tfch_num_lines_fetched = 2728,
+ tdvx_tex_tfch_num_lines_fetched_block_compressed = 2729,
+ tdvx_tex_tfch_num_operations = 2730,
+ tdvx_tex_filt_num_operations = 2731,
+ tdvx_ls_mem_read_full = 2732,
+ tdvx_ls_mem_read_short = 2733,
+ tdvx_ls_mem_write_full = 2734,
+ tdvx_ls_mem_write_short = 2735,
+ tdvx_ls_mem_atomic = 2736,
+ tdvx_vary_instr = 2737,
+ tdvx_vary_slot_32 = 2738,
+ tdvx_vary_slot_16 = 2739,
+ tdvx_attr_instr = 2740,
+ tdvx_arith_instr_fp_mul = 2741,
+ tdvx_beats_rd_ftc = 2742,
+ tdvx_beats_rd_ftc_ext = 2743,
+ tdvx_beats_rd_lsc = 2744,
+ tdvx_beats_rd_lsc_ext = 2745,
+ tdvx_beats_rd_tex = 2746,
+ tdvx_beats_rd_tex_ext = 2747,
+ tdvx_beats_rd_other = 2748,
+ tdvx_beats_wr_lsc_other = 2749,
+ tdvx_beats_wr_tib = 2750,
+ tdvx_beats_wr_lsc_wb = 2751,
+ tdvx_mmu_requests = 2756,
+ tdvx_mmu_table_reads_l3 = 2757,
+ tdvx_mmu_table_reads_l2 = 2758,
+ tdvx_mmu_hit_l3 = 2759,
+ tdvx_mmu_hit_l2 = 2760,
+ tdvx_mmu_s2_requests = 2761,
+ tdvx_mmu_s2_table_reads_l3 = 2762,
+ tdvx_mmu_s2_table_reads_l2 = 2763,
+ tdvx_mmu_s2_hit_l3 = 2764,
+ tdvx_mmu_s2_hit_l2 = 2765,
+ tdvx_l2_rd_msg_in = 2768,
+ tdvx_l2_rd_msg_in_stall = 2769,
+ tdvx_l2_wr_msg_in = 2770,
+ tdvx_l2_wr_msg_in_stall = 2771,
+ tdvx_l2_snp_msg_in = 2772,
+ tdvx_l2_snp_msg_in_stall = 2773,
+ tdvx_l2_rd_msg_out = 2774,
+ tdvx_l2_rd_msg_out_stall = 2775,
+ tdvx_l2_wr_msg_out = 2776,
+ tdvx_l2_any_lookup = 2777,
+ tdvx_l2_read_lookup = 2778,
+ tdvx_l2_write_lookup = 2779,
+ tdvx_l2_ext_snoop_lookup = 2780,
+ tdvx_l2_ext_read = 2781,
+ tdvx_l2_ext_read_nosnp = 2782,
+ tdvx_l2_ext_read_unique = 2783,
+ tdvx_l2_ext_read_beats = 2784,
+ tdvx_l2_ext_ar_stall = 2785,
+ tdvx_l2_ext_ar_cnt_q1 = 2786,
+ tdvx_l2_ext_ar_cnt_q2 = 2787,
+ tdvx_l2_ext_ar_cnt_q3 = 2788,
+ tdvx_l2_ext_rresp_0_127 = 2789,
+ tdvx_l2_ext_rresp_128_191 = 2790,
+ tdvx_l2_ext_rresp_192_255 = 2791,
+ tdvx_l2_ext_rresp_256_319 = 2792,
+ tdvx_l2_ext_rresp_320_383 = 2793,
+ tdvx_l2_ext_write = 2794,
+ tdvx_l2_ext_write_nosnp_full = 2795,
+ tdvx_l2_ext_write_nosnp_ptl = 2796,
+ tdvx_l2_ext_write_snp_full = 2797,
+ tdvx_l2_ext_write_snp_ptl = 2798,
+ tdvx_l2_ext_write_beats = 2799,
+ tdvx_l2_ext_w_stall = 2800,
+ tdvx_l2_ext_aw_cnt_q1 = 2801,
+ tdvx_l2_ext_aw_cnt_q2 = 2802,
+ tdvx_l2_ext_aw_cnt_q3 = 2803,
+ tdvx_l2_ext_snoop = 2804,
+ tdvx_l2_ext_snoop_stall = 2805,
+ tdvx_l2_ext_snoop_resp_clean = 2806,
+ tdvx_l2_ext_snoop_resp_data = 2807,
+ tdvx_l2_ext_snoop_internal = 2808,
+ tsix_messages_sent = 2820,
+ tsix_messages_received = 2821,
+ tsix_gpu_active = 2822,
+ tsix_irq_active = 2823,
+ tsix_js0_jobs = 2824,
+ tsix_js0_tasks = 2825,
+ tsix_js0_active = 2826,
+ tsix_js0_wait_flush = 2827,
+ tsix_js0_wait_read = 2828,
+ tsix_js0_wait_issue = 2829,
+ tsix_js0_wait_depend = 2830,
+ tsix_js0_wait_finish = 2831,
+ tsix_js1_jobs = 2832,
+ tsix_js1_tasks = 2833,
+ tsix_js1_active = 2834,
+ tsix_js1_wait_flush = 2835,
+ tsix_js1_wait_read = 2836,
+ tsix_js1_wait_issue = 2837,
+ tsix_js1_wait_depend = 2838,
+ tsix_js1_wait_finish = 2839,
+ tsix_js2_jobs = 2840,
+ tsix_js2_tasks = 2841,
+ tsix_js2_active = 2842,
+ tsix_js2_wait_flush = 2843,
+ tsix_js2_wait_read = 2844,
+ tsix_js2_wait_issue = 2845,
+ tsix_js2_wait_depend = 2846,
+ tsix_js2_wait_finish = 2847,
+ tsix_tiler_active = 2884,
+ tsix_jobs_processed = 2885,
+ tsix_triangles = 2886,
+ tsix_lines = 2887,
+ tsix_points = 2888,
+ tsix_front_facing = 2889,
+ tsix_back_facing = 2890,
+ tsix_prim_visible = 2891,
+ tsix_prim_culled = 2892,
+ tsix_prim_clipped = 2893,
+ tsix_prim_sat_culled = 2894,
+ tsix_bin_alloc_init = 2895,
+ tsix_bin_alloc_overflow = 2896,
+ tsix_bus_read = 2897,
+ tsix_bus_write = 2899,
+ tsix_loading_desc = 2900,
+ tsix_idvs_pos_shad_req = 2901,
+ tsix_idvs_pos_shad_wait = 2902,
+ tsix_idvs_pos_shad_stall = 2903,
+ tsix_idvs_pos_fifo_full = 2904,
+ tsix_prefetch_stall = 2905,
+ tsix_vcache_hit = 2906,
+ tsix_vcache_miss = 2907,
+ tsix_vcache_line_wait = 2908,
+ tsix_vfetch_pos_read_wait = 2909,
+ tsix_vfetch_vertex_wait = 2910,
+ tsix_vfetch_stall = 2911,
+ tsix_primassy_stall = 2912,
+ tsix_bbox_gen_stall = 2913,
+ tsix_idvs_vbu_hit = 2914,
+ tsix_idvs_vbu_miss = 2915,
+ tsix_idvs_vbu_line_deallocate = 2916,
+ tsix_idvs_var_shad_req = 2917,
+ tsix_idvs_var_shad_stall = 2918,
+ tsix_binner_stall = 2919,
+ tsix_iter_stall = 2920,
+ tsix_compress_miss = 2921,
+ tsix_compress_stall = 2922,
+ tsix_pcache_hit = 2923,
+ tsix_pcache_miss = 2924,
+ tsix_pcache_miss_stall = 2925,
+ tsix_pcache_evict_stall = 2926,
+ tsix_pmgr_ptr_wr_stall = 2927,
+ tsix_pmgr_ptr_rd_stall = 2928,
+ tsix_pmgr_cmd_wr_stall = 2929,
+ tsix_wrbuf_active = 2930,
+ tsix_wrbuf_hit = 2931,
+ tsix_wrbuf_miss = 2932,
+ tsix_wrbuf_no_free_line_stall = 2933,
+ tsix_wrbuf_no_axi_id_stall = 2934,
+ tsix_wrbuf_axi_stall = 2935,
+ tsix_utlb_trans = 2939,
+ tsix_utlb_trans_hit = 2940,
+ tsix_utlb_trans_stall = 2941,
+ tsix_utlb_trans_miss_delay = 2942,
+ tsix_utlb_mmu_req = 2943,
+ tsix_frag_active = 2948,
+ tsix_frag_primitives = 2949,
+ tsix_frag_prim_rast = 2950,
+ tsix_frag_fpk_active = 2951,
+ tsix_frag_starving = 2952,
+ tsix_frag_warps = 2953,
+ tsix_frag_partial_warps = 2954,
+ tsix_frag_quads_rast = 2955,
+ tsix_frag_quads_ezs_test = 2956,
+ tsix_frag_quads_ezs_update = 2957,
+ tsix_frag_quads_ezs_kill = 2958,
+ tsix_frag_lzs_test = 2959,
+ tsix_frag_lzs_kill = 2960,
+ tsix_frag_ptiles = 2962,
+ tsix_frag_trans_elim = 2963,
+ tsix_quad_fpk_killer = 2964,
+ tsix_compute_active = 2966,
+ tsix_compute_tasks = 2967,
+ tsix_compute_warps = 2968,
+ tsix_compute_starving = 2969,
+ tsix_exec_core_active = 2970,
+ tsix_exec_active = 2971,
+ tsix_exec_instr_count = 2972,
+ tsix_exec_instr_diverged = 2973,
+ tsix_exec_instr_starving = 2974,
+ tsix_arith_instr_single_fma = 2975,
+ tsix_arith_instr_double = 2976,
+ tsix_arith_instr_msg = 2977,
+ tsix_arith_instr_msg_only = 2978,
+ tsix_tex_msgi_num_quads = 2979,
+ tsix_tex_dfch_num_passes = 2980,
+ tsix_tex_dfch_num_passes_miss = 2981,
+ tsix_tex_dfch_num_passes_mip_map = 2982,
+ tsix_tex_tidx_num_split_mip_map = 2983,
+ tsix_tex_tfch_num_lines_fetched = 2984,
+ tsix_tex_tfch_num_lines_fetched_block_compressed = 2985,
+ tsix_tex_tfch_num_operations = 2986,
+ tsix_tex_filt_num_operations = 2987,
+ tsix_ls_mem_read_full = 2988,
+ tsix_ls_mem_read_short = 2989,
+ tsix_ls_mem_write_full = 2990,
+ tsix_ls_mem_write_short = 2991,
+ tsix_ls_mem_atomic = 2992,
+ tsix_vary_instr = 2993,
+ tsix_vary_slot_32 = 2994,
+ tsix_vary_slot_16 = 2995,
+ tsix_attr_instr = 2996,
+ tsix_arith_instr_fp_mul = 2997,
+ tsix_beats_rd_ftc = 2998,
+ tsix_beats_rd_ftc_ext = 2999,
+ tsix_beats_rd_lsc = 3000,
+ tsix_beats_rd_lsc_ext = 3001,
+ tsix_beats_rd_tex = 3002,
+ tsix_beats_rd_tex_ext = 3003,
+ tsix_beats_rd_other = 3004,
+ tsix_beats_wr_lsc_other = 3005,
+ tsix_beats_wr_tib = 3006,
+ tsix_beats_wr_lsc_wb = 3007,
+ tsix_mmu_requests = 3012,
+ tsix_mmu_table_reads_l3 = 3013,
+ tsix_mmu_table_reads_l2 = 3014,
+ tsix_mmu_hit_l3 = 3015,
+ tsix_mmu_hit_l2 = 3016,
+ tsix_mmu_s2_requests = 3017,
+ tsix_mmu_s2_table_reads_l3 = 3018,
+ tsix_mmu_s2_table_reads_l2 = 3019,
+ tsix_mmu_s2_hit_l3 = 3020,
+ tsix_mmu_s2_hit_l2 = 3021,
+ tsix_l2_rd_msg_in = 3024,
+ tsix_l2_rd_msg_in_stall = 3025,
+ tsix_l2_wr_msg_in = 3026,
+ tsix_l2_wr_msg_in_stall = 3027,
+ tsix_l2_snp_msg_in = 3028,
+ tsix_l2_snp_msg_in_stall = 3029,
+ tsix_l2_rd_msg_out = 3030,
+ tsix_l2_rd_msg_out_stall = 3031,
+ tsix_l2_wr_msg_out = 3032,
+ tsix_l2_any_lookup = 3033,
+ tsix_l2_read_lookup = 3034,
+ tsix_l2_write_lookup = 3035,
+ tsix_l2_ext_snoop_lookup = 3036,
+ tsix_l2_ext_read = 3037,
+ tsix_l2_ext_read_nosnp = 3038,
+ tsix_l2_ext_read_unique = 3039,
+ tsix_l2_ext_read_beats = 3040,
+ tsix_l2_ext_ar_stall = 3041,
+ tsix_l2_ext_ar_cnt_q1 = 3042,
+ tsix_l2_ext_ar_cnt_q2 = 3043,
+ tsix_l2_ext_ar_cnt_q3 = 3044,
+ tsix_l2_ext_rresp_0_127 = 3045,
+ tsix_l2_ext_rresp_128_191 = 3046,
+ tsix_l2_ext_rresp_192_255 = 3047,
+ tsix_l2_ext_rresp_256_319 = 3048,
+ tsix_l2_ext_rresp_320_383 = 3049,
+ tsix_l2_ext_write = 3050,
+ tsix_l2_ext_write_nosnp_full = 3051,
+ tsix_l2_ext_write_nosnp_ptl = 3052,
+ tsix_l2_ext_write_snp_full = 3053,
+ tsix_l2_ext_write_snp_ptl = 3054,
+ tsix_l2_ext_write_beats = 3055,
+ tsix_l2_ext_w_stall = 3056,
+ tsix_l2_ext_aw_cnt_q1 = 3057,
+ tsix_l2_ext_aw_cnt_q2 = 3058,
+ tsix_l2_ext_aw_cnt_q3 = 3059,
+ tsix_l2_ext_snoop = 3060,
+ tsix_l2_ext_snoop_stall = 3061,
+ tsix_l2_ext_snoop_resp_clean = 3062,
+ tsix_l2_ext_snoop_resp_data = 3063,
+ tsix_l2_ext_snoop_internal = 3064,
+ tnox_messages_sent = 3076,
+ tnox_messages_received = 3077,
+ tnox_gpu_active = 3078,
+ tnox_irq_active = 3079,
+ tnox_js0_jobs = 3080,
+ tnox_js0_tasks = 3081,
+ tnox_js0_active = 3082,
+ tnox_js0_wait_flush = 3083,
+ tnox_js0_wait_read = 3084,
+ tnox_js0_wait_issue = 3085,
+ tnox_js0_wait_depend = 3086,
+ tnox_js0_wait_finish = 3087,
+ tnox_js1_jobs = 3088,
+ tnox_js1_tasks = 3089,
+ tnox_js1_active = 3090,
+ tnox_js1_wait_flush = 3091,
+ tnox_js1_wait_read = 3092,
+ tnox_js1_wait_issue = 3093,
+ tnox_js1_wait_depend = 3094,
+ tnox_js1_wait_finish = 3095,
+ tnox_js2_jobs = 3096,
+ tnox_js2_tasks = 3097,
+ tnox_js2_active = 3098,
+ tnox_js2_wait_flush = 3099,
+ tnox_js2_wait_read = 3100,
+ tnox_js2_wait_issue = 3101,
+ tnox_js2_wait_depend = 3102,
+ tnox_js2_wait_finish = 3103,
+ tnox_cache_flush = 3135,
+ tnox_tiler_active = 3140,
+ tnox_jobs_processed = 3141,
+ tnox_triangles = 3142,
+ tnox_lines = 3143,
+ tnox_points = 3144,
+ tnox_front_facing = 3145,
+ tnox_back_facing = 3146,
+ tnox_prim_visible = 3147,
+ tnox_prim_culled = 3148,
+ tnox_prim_clipped = 3149,
+ tnox_prim_sat_culled = 3150,
+ tnox_bin_alloc_init = 3151,
+ tnox_bin_alloc_overflow = 3152,
+ tnox_bus_read = 3153,
+ tnox_bus_write = 3155,
+ tnox_loading_desc = 3156,
+ tnox_idvs_pos_shad_req = 3157,
+ tnox_idvs_pos_shad_wait = 3158,
+ tnox_idvs_pos_shad_stall = 3159,
+ tnox_idvs_pos_fifo_full = 3160,
+ tnox_prefetch_stall = 3161,
+ tnox_vcache_hit = 3162,
+ tnox_vcache_miss = 3163,
+ tnox_vcache_line_wait = 3164,
+ tnox_vfetch_pos_read_wait = 3165,
+ tnox_vfetch_vertex_wait = 3166,
+ tnox_vfetch_stall = 3167,
+ tnox_primassy_stall = 3168,
+ tnox_bbox_gen_stall = 3169,
+ tnox_idvs_vbu_hit = 3170,
+ tnox_idvs_vbu_miss = 3171,
+ tnox_idvs_vbu_line_deallocate = 3172,
+ tnox_idvs_var_shad_req = 3173,
+ tnox_idvs_var_shad_stall = 3174,
+ tnox_binner_stall = 3175,
+ tnox_iter_stall = 3176,
+ tnox_compress_miss = 3177,
+ tnox_compress_stall = 3178,
+ tnox_pcache_hit = 3179,
+ tnox_pcache_miss = 3180,
+ tnox_pcache_miss_stall = 3181,
+ tnox_pcache_evict_stall = 3182,
+ tnox_pmgr_ptr_wr_stall = 3183,
+ tnox_pmgr_ptr_rd_stall = 3184,
+ tnox_pmgr_cmd_wr_stall = 3185,
+ tnox_wrbuf_active = 3186,
+ tnox_wrbuf_hit = 3187,
+ tnox_wrbuf_miss = 3188,
+ tnox_wrbuf_no_free_line_stall = 3189,
+ tnox_wrbuf_no_axi_id_stall = 3190,
+ tnox_wrbuf_axi_stall = 3191,
+ tnox_utlb_trans = 3195,
+ tnox_utlb_trans_hit = 3196,
+ tnox_utlb_trans_stall = 3197,
+ tnox_utlb_trans_miss_delay = 3198,
+ tnox_utlb_mmu_req = 3199,
+ tnox_frag_active = 3204,
+ tnox_frag_primitives = 3205,
+ tnox_frag_prim_rast = 3206,
+ tnox_frag_fpk_active = 3207,
+ tnox_frag_starving = 3208,
+ tnox_frag_warps = 3209,
+ tnox_frag_partial_warps = 3210,
+ tnox_frag_quads_rast = 3211,
+ tnox_frag_quads_ezs_test = 3212,
+ tnox_frag_quads_ezs_update = 3213,
+ tnox_frag_quads_ezs_kill = 3214,
+ tnox_frag_lzs_test = 3215,
+ tnox_frag_lzs_kill = 3216,
+ tnox_warp_reg_size_64 = 3217,
+ tnox_frag_ptiles = 3218,
+ tnox_frag_trans_elim = 3219,
+ tnox_quad_fpk_killer = 3220,
+ tnox_full_quad_warps = 3221,
+ tnox_compute_active = 3222,
+ tnox_compute_tasks = 3223,
+ tnox_compute_warps = 3224,
+ tnox_compute_starving = 3225,
+ tnox_exec_core_active = 3226,
+ tnox_exec_active = 3227,
+ tnox_exec_instr_count = 3228,
+ tnox_exec_instr_diverged = 3229,
+ tnox_exec_instr_starving = 3230,
+ tnox_arith_instr_single_fma = 3231,
+ tnox_arith_instr_double = 3232,
+ tnox_arith_instr_msg = 3233,
+ tnox_arith_instr_msg_only = 3234,
+ tnox_tex_msgi_num_quads = 3235,
+ tnox_tex_dfch_num_passes = 3236,
+ tnox_tex_dfch_num_passes_miss = 3237,
+ tnox_tex_dfch_num_passes_mip_map = 3238,
+ tnox_tex_tidx_num_split_mip_map = 3239,
+ tnox_tex_tfch_num_lines_fetched = 3240,
+ tnox_tex_tfch_num_lines_fetched_block_compressed = 3241,
+ tnox_tex_tfch_num_operations = 3242,
+ tnox_tex_filt_num_operations = 3243,
+ tnox_ls_mem_read_full = 3244,
+ tnox_ls_mem_read_short = 3245,
+ tnox_ls_mem_write_full = 3246,
+ tnox_ls_mem_write_short = 3247,
+ tnox_ls_mem_atomic = 3248,
+ tnox_vary_instr = 3249,
+ tnox_vary_slot_32 = 3250,
+ tnox_vary_slot_16 = 3251,
+ tnox_attr_instr = 3252,
+ tnox_arith_instr_fp_mul = 3253,
+ tnox_beats_rd_ftc = 3254,
+ tnox_beats_rd_ftc_ext = 3255,
+ tnox_beats_rd_lsc = 3256,
+ tnox_beats_rd_lsc_ext = 3257,
+ tnox_beats_rd_tex = 3258,
+ tnox_beats_rd_tex_ext = 3259,
+ tnox_beats_rd_other = 3260,
+ tnox_beats_wr_lsc_other = 3261,
+ tnox_beats_wr_tib = 3262,
+ tnox_beats_wr_lsc_wb = 3263,
+ tnox_mmu_requests = 3268,
+ tnox_mmu_table_reads_l3 = 3269,
+ tnox_mmu_table_reads_l2 = 3270,
+ tnox_mmu_hit_l3 = 3271,
+ tnox_mmu_hit_l2 = 3272,
+ tnox_mmu_s2_requests = 3273,
+ tnox_mmu_s2_table_reads_l3 = 3274,
+ tnox_mmu_s2_table_reads_l2 = 3275,
+ tnox_mmu_s2_hit_l3 = 3276,
+ tnox_mmu_s2_hit_l2 = 3277,
+ tnox_l2_rd_msg_in = 3280,
+ tnox_l2_rd_msg_in_stall = 3281,
+ tnox_l2_wr_msg_in = 3282,
+ tnox_l2_wr_msg_in_stall = 3283,
+ tnox_l2_snp_msg_in = 3284,
+ tnox_l2_snp_msg_in_stall = 3285,
+ tnox_l2_rd_msg_out = 3286,
+ tnox_l2_rd_msg_out_stall = 3287,
+ tnox_l2_wr_msg_out = 3288,
+ tnox_l2_any_lookup = 3289,
+ tnox_l2_read_lookup = 3290,
+ tnox_l2_write_lookup = 3291,
+ tnox_l2_ext_snoop_lookup = 3292,
+ tnox_l2_ext_read = 3293,
+ tnox_l2_ext_read_nosnp = 3294,
+ tnox_l2_ext_read_unique = 3295,
+ tnox_l2_ext_read_beats = 3296,
+ tnox_l2_ext_ar_stall = 3297,
+ tnox_l2_ext_ar_cnt_q1 = 3298,
+ tnox_l2_ext_ar_cnt_q2 = 3299,
+ tnox_l2_ext_ar_cnt_q3 = 3300,
+ tnox_l2_ext_rresp_0_127 = 3301,
+ tnox_l2_ext_rresp_128_191 = 3302,
+ tnox_l2_ext_rresp_192_255 = 3303,
+ tnox_l2_ext_rresp_256_319 = 3304,
+ tnox_l2_ext_rresp_320_383 = 3305,
+ tnox_l2_ext_write = 3306,
+ tnox_l2_ext_write_nosnp_full = 3307,
+ tnox_l2_ext_write_nosnp_ptl = 3308,
+ tnox_l2_ext_write_snp_full = 3309,
+ tnox_l2_ext_write_snp_ptl = 3310,
+ tnox_l2_ext_write_beats = 3311,
+ tnox_l2_ext_w_stall = 3312,
+ tnox_l2_ext_aw_cnt_q1 = 3313,
+ tnox_l2_ext_aw_cnt_q2 = 3314,
+ tnox_l2_ext_aw_cnt_q3 = 3315,
+ tnox_l2_ext_snoop = 3316,
+ tnox_l2_ext_snoop_stall = 3317,
+ tnox_l2_ext_snoop_resp_clean = 3318,
+ tnox_l2_ext_snoop_resp_data = 3319,
+ tnox_l2_ext_snoop_internal = 3320,
+ tgox_messages_sent = 3332,
+ tgox_messages_received = 3333,
+ tgox_gpu_active = 3334,
+ tgox_irq_active = 3335,
+ tgox_js0_jobs = 3336,
+ tgox_js0_tasks = 3337,
+ tgox_js0_active = 3338,
+ tgox_js0_wait_flush = 3339,
+ tgox_js0_wait_read = 3340,
+ tgox_js0_wait_issue = 3341,
+ tgox_js0_wait_depend = 3342,
+ tgox_js0_wait_finish = 3343,
+ tgox_js1_jobs = 3344,
+ tgox_js1_tasks = 3345,
+ tgox_js1_active = 3346,
+ tgox_js1_wait_flush = 3347,
+ tgox_js1_wait_read = 3348,
+ tgox_js1_wait_issue = 3349,
+ tgox_js1_wait_depend = 3350,
+ tgox_js1_wait_finish = 3351,
+ tgox_js2_jobs = 3352,
+ tgox_js2_tasks = 3353,
+ tgox_js2_active = 3354,
+ tgox_js2_wait_flush = 3355,
+ tgox_js2_wait_read = 3356,
+ tgox_js2_wait_issue = 3357,
+ tgox_js2_wait_depend = 3358,
+ tgox_js2_wait_finish = 3359,
+ tgox_cache_flush = 3391,
+ tgox_tiler_active = 3396,
+ tgox_jobs_processed = 3397,
+ tgox_triangles = 3398,
+ tgox_lines = 3399,
+ tgox_points = 3400,
+ tgox_front_facing = 3401,
+ tgox_back_facing = 3402,
+ tgox_prim_visible = 3403,
+ tgox_prim_culled = 3404,
+ tgox_prim_clipped = 3405,
+ tgox_prim_sat_culled = 3406,
+ tgox_bin_alloc_init = 3407,
+ tgox_bin_alloc_overflow = 3408,
+ tgox_bus_read = 3409,
+ tgox_bus_write = 3411,
+ tgox_loading_desc = 3412,
+ tgox_idvs_pos_shad_req = 3413,
+ tgox_idvs_pos_shad_wait = 3414,
+ tgox_idvs_pos_shad_stall = 3415,
+ tgox_idvs_pos_fifo_full = 3416,
+ tgox_prefetch_stall = 3417,
+ tgox_vcache_hit = 3418,
+ tgox_vcache_miss = 3419,
+ tgox_vcache_line_wait = 3420,
+ tgox_vfetch_pos_read_wait = 3421,
+ tgox_vfetch_vertex_wait = 3422,
+ tgox_vfetch_stall = 3423,
+ tgox_primassy_stall = 3424,
+ tgox_bbox_gen_stall = 3425,
+ tgox_idvs_vbu_hit = 3426,
+ tgox_idvs_vbu_miss = 3427,
+ tgox_idvs_vbu_line_deallocate = 3428,
+ tgox_idvs_var_shad_req = 3429,
+ tgox_idvs_var_shad_stall = 3430,
+ tgox_binner_stall = 3431,
+ tgox_iter_stall = 3432,
+ tgox_compress_miss = 3433,
+ tgox_compress_stall = 3434,
+ tgox_pcache_hit = 3435,
+ tgox_pcache_miss = 3436,
+ tgox_pcache_miss_stall = 3437,
+ tgox_pcache_evict_stall = 3438,
+ tgox_pmgr_ptr_wr_stall = 3439,
+ tgox_pmgr_ptr_rd_stall = 3440,
+ tgox_pmgr_cmd_wr_stall = 3441,
+ tgox_wrbuf_active = 3442,
+ tgox_wrbuf_hit = 3443,
+ tgox_wrbuf_miss = 3444,
+ tgox_wrbuf_no_free_line_stall = 3445,
+ tgox_wrbuf_no_axi_id_stall = 3446,
+ tgox_wrbuf_axi_stall = 3447,
+ tgox_utlb_trans = 3451,
+ tgox_utlb_trans_hit = 3452,
+ tgox_utlb_trans_stall = 3453,
+ tgox_utlb_trans_miss_delay = 3454,
+ tgox_utlb_mmu_req = 3455,
+ tgox_frag_active = 3460,
+ tgox_frag_primitives = 3461,
+ tgox_frag_prim_rast = 3462,
+ tgox_frag_fpk_active = 3463,
+ tgox_frag_starving = 3464,
+ tgox_frag_warps = 3465,
+ tgox_frag_partial_warps = 3466,
+ tgox_frag_quads_rast = 3467,
+ tgox_frag_quads_ezs_test = 3468,
+ tgox_frag_quads_ezs_update = 3469,
+ tgox_frag_quads_ezs_kill = 3470,
+ tgox_frag_lzs_test = 3471,
+ tgox_frag_lzs_kill = 3472,
+ tgox_warp_reg_size_64 = 3473,
+ tgox_frag_ptiles = 3474,
+ tgox_frag_trans_elim = 3475,
+ tgox_quad_fpk_killer = 3476,
+ tgox_full_quad_warps = 3477,
+ tgox_compute_active = 3478,
+ tgox_compute_tasks = 3479,
+ tgox_compute_warps = 3480,
+ tgox_compute_starving = 3481,
+ tgox_exec_core_active = 3482,
+ tgox_exec_active = 3483,
+ tgox_exec_instr_count = 3484,
+ tgox_exec_instr_diverged = 3485,
+ tgox_exec_instr_starving = 3486,
+ tgox_arith_instr_single_fma = 3487,
+ tgox_arith_instr_double = 3488,
+ tgox_arith_instr_msg = 3489,
+ tgox_arith_instr_msg_only = 3490,
+ tgox_tex_msgi_num_quads = 3491,
+ tgox_tex_dfch_num_passes = 3492,
+ tgox_tex_dfch_num_passes_miss = 3493,
+ tgox_tex_dfch_num_passes_mip_map = 3494,
+ tgox_tex_tidx_num_split_mip_map = 3495,
+ tgox_tex_tfch_num_lines_fetched = 3496,
+ tgox_tex_tfch_num_lines_fetched_block_compressed = 3497,
+ tgox_tex_tfch_num_operations = 3498,
+ tgox_tex_filt_num_operations = 3499,
+ tgox_ls_mem_read_full = 3500,
+ tgox_ls_mem_read_short = 3501,
+ tgox_ls_mem_write_full = 3502,
+ tgox_ls_mem_write_short = 3503,
+ tgox_ls_mem_atomic = 3504,
+ tgox_vary_instr = 3505,
+ tgox_vary_slot_32 = 3506,
+ tgox_vary_slot_16 = 3507,
+ tgox_attr_instr = 3508,
+ tgox_arith_instr_fp_mul = 3509,
+ tgox_beats_rd_ftc = 3510,
+ tgox_beats_rd_ftc_ext = 3511,
+ tgox_beats_rd_lsc = 3512,
+ tgox_beats_rd_lsc_ext = 3513,
+ tgox_beats_rd_tex = 3514,
+ tgox_beats_rd_tex_ext = 3515,
+ tgox_beats_rd_other = 3516,
+ tgox_beats_wr_lsc_wb = 3517,
+ tgox_beats_wr_tib = 3518,
+ tgox_beats_wr_lsc_other = 3519,
+ tgox_mmu_requests = 3524,
+ tgox_mmu_table_reads_l3 = 3525,
+ tgox_mmu_table_reads_l2 = 3526,
+ tgox_mmu_hit_l3 = 3527,
+ tgox_mmu_hit_l2 = 3528,
+ tgox_mmu_s2_requests = 3529,
+ tgox_mmu_s2_table_reads_l3 = 3530,
+ tgox_mmu_s2_table_reads_l2 = 3531,
+ tgox_mmu_s2_hit_l3 = 3532,
+ tgox_mmu_s2_hit_l2 = 3533,
+ tgox_l2_rd_msg_in = 3536,
+ tgox_l2_rd_msg_in_stall = 3537,
+ tgox_l2_wr_msg_in = 3538,
+ tgox_l2_wr_msg_in_stall = 3539,
+ tgox_l2_snp_msg_in = 3540,
+ tgox_l2_snp_msg_in_stall = 3541,
+ tgox_l2_rd_msg_out = 3542,
+ tgox_l2_rd_msg_out_stall = 3543,
+ tgox_l2_wr_msg_out = 3544,
+ tgox_l2_any_lookup = 3545,
+ tgox_l2_read_lookup = 3546,
+ tgox_l2_write_lookup = 3547,
+ tgox_l2_ext_snoop_lookup = 3548,
+ tgox_l2_ext_read = 3549,
+ tgox_l2_ext_read_nosnp = 3550,
+ tgox_l2_ext_read_unique = 3551,
+ tgox_l2_ext_read_beats = 3552,
+ tgox_l2_ext_ar_stall = 3553,
+ tgox_l2_ext_ar_cnt_q1 = 3554,
+ tgox_l2_ext_ar_cnt_q2 = 3555,
+ tgox_l2_ext_ar_cnt_q3 = 3556,
+ tgox_l2_ext_rresp_0_127 = 3557,
+ tgox_l2_ext_rresp_128_191 = 3558,
+ tgox_l2_ext_rresp_192_255 = 3559,
+ tgox_l2_ext_rresp_256_319 = 3560,
+ tgox_l2_ext_rresp_320_383 = 3561,
+ tgox_l2_ext_write = 3562,
+ tgox_l2_ext_write_nosnp_full = 3563,
+ tgox_l2_ext_write_nosnp_ptl = 3564,
+ tgox_l2_ext_write_snp_full = 3565,
+ tgox_l2_ext_write_snp_ptl = 3566,
+ tgox_l2_ext_write_beats = 3567,
+ tgox_l2_ext_w_stall = 3568,
+ tgox_l2_ext_aw_cnt_q1 = 3569,
+ tgox_l2_ext_aw_cnt_q2 = 3570,
+ tgox_l2_ext_aw_cnt_q3 = 3571,
+ tgox_l2_ext_snoop = 3572,
+ tgox_l2_ext_snoop_stall = 3573,
+ tgox_l2_ext_snoop_resp_clean = 3574,
+ tgox_l2_ext_snoop_resp_data = 3575,
+ tgox_l2_ext_snoop_internal = 3576,
+ ttrx_messages_sent = 3588,
+ ttrx_messages_received = 3589,
+ ttrx_gpu_active = 3590,
+ ttrx_irq_active = 3591,
+ ttrx_js0_jobs = 3592,
+ ttrx_js0_tasks = 3593,
+ ttrx_js0_active = 3594,
+ ttrx_js0_wait_flush = 3595,
+ ttrx_js0_wait_read = 3596,
+ ttrx_js0_wait_issue = 3597,
+ ttrx_js0_wait_depend = 3598,
+ ttrx_js0_wait_finish = 3599,
+ ttrx_js1_jobs = 3600,
+ ttrx_js1_tasks = 3601,
+ ttrx_js1_active = 3602,
+ ttrx_js1_wait_flush = 3603,
+ ttrx_js1_wait_read = 3604,
+ ttrx_js1_wait_issue = 3605,
+ ttrx_js1_wait_depend = 3606,
+ ttrx_js1_wait_finish = 3607,
+ ttrx_js2_jobs = 3608,
+ ttrx_js2_tasks = 3609,
+ ttrx_js2_active = 3610,
+ ttrx_js2_wait_flush = 3611,
+ ttrx_js2_wait_read = 3612,
+ ttrx_js2_wait_issue = 3613,
+ ttrx_js2_wait_depend = 3614,
+ ttrx_js2_wait_finish = 3615,
+ ttrx_cache_flush = 3647,
+ ttrx_tiler_active = 3652,
+ ttrx_jobs_processed = 3653,
+ ttrx_triangles = 3654,
+ ttrx_lines = 3655,
+ ttrx_points = 3656,
+ ttrx_front_facing = 3657,
+ ttrx_back_facing = 3658,
+ ttrx_prim_visible = 3659,
+ ttrx_prim_culled = 3660,
+ ttrx_prim_clipped = 3661,
+ ttrx_prim_sat_culled = 3662,
+ ttrx_bin_alloc_init = 3663,
+ ttrx_bin_alloc_overflow = 3664,
+ ttrx_bus_read = 3665,
+ ttrx_bus_write = 3667,
+ ttrx_loading_desc = 3668,
+ ttrx_idvs_pos_shad_req = 3669,
+ ttrx_idvs_pos_shad_wait = 3670,
+ ttrx_idvs_pos_shad_stall = 3671,
+ ttrx_idvs_pos_fifo_full = 3672,
+ ttrx_prefetch_stall = 3673,
+ ttrx_vcache_hit = 3674,
+ ttrx_vcache_miss = 3675,
+ ttrx_vcache_line_wait = 3676,
+ ttrx_vfetch_pos_read_wait = 3677,
+ ttrx_vfetch_vertex_wait = 3678,
+ ttrx_vfetch_stall = 3679,
+ ttrx_primassy_stall = 3680,
+ ttrx_bbox_gen_stall = 3681,
+ ttrx_idvs_vbu_hit = 3682,
+ ttrx_idvs_vbu_miss = 3683,
+ ttrx_idvs_vbu_line_deallocate = 3684,
+ ttrx_idvs_var_shad_req = 3685,
+ ttrx_idvs_var_shad_stall = 3686,
+ ttrx_binner_stall = 3687,
+ ttrx_iter_stall = 3688,
+ ttrx_compress_miss = 3689,
+ ttrx_compress_stall = 3690,
+ ttrx_pcache_hit = 3691,
+ ttrx_pcache_miss = 3692,
+ ttrx_pcache_miss_stall = 3693,
+ ttrx_pcache_evict_stall = 3694,
+ ttrx_pmgr_ptr_wr_stall = 3695,
+ ttrx_pmgr_ptr_rd_stall = 3696,
+ ttrx_pmgr_cmd_wr_stall = 3697,
+ ttrx_wrbuf_active = 3698,
+ ttrx_wrbuf_hit = 3699,
+ ttrx_wrbuf_miss = 3700,
+ ttrx_wrbuf_no_free_line_stall = 3701,
+ ttrx_wrbuf_no_axi_id_stall = 3702,
+ ttrx_wrbuf_axi_stall = 3703,
+ ttrx_utlb_trans = 3707,
+ ttrx_utlb_trans_hit = 3708,
+ ttrx_utlb_trans_stall = 3709,
+ ttrx_utlb_trans_miss_delay = 3710,
+ ttrx_utlb_mmu_req = 3711,
+ ttrx_frag_active = 3716,
+ ttrx_frag_primitives_out = 3717,
+ ttrx_frag_prim_rast = 3718,
+ ttrx_frag_fpk_active = 3719,
+ ttrx_frag_starving = 3720,
+ ttrx_frag_warps = 3721,
+ ttrx_frag_partial_quads_rast = 3722,
+ ttrx_frag_quads_rast = 3723,
+ ttrx_frag_quads_ezs_test = 3724,
+ ttrx_frag_quads_ezs_update = 3725,
+ ttrx_frag_quads_ezs_kill = 3726,
+ ttrx_frag_lzs_test = 3727,
+ ttrx_frag_lzs_kill = 3728,
+ ttrx_warp_reg_size_64 = 3729,
+ ttrx_frag_ptiles = 3730,
+ ttrx_frag_trans_elim = 3731,
+ ttrx_quad_fpk_killer = 3732,
+ ttrx_full_quad_warps = 3733,
+ ttrx_compute_active = 3734,
+ ttrx_compute_tasks = 3735,
+ ttrx_compute_warps = 3736,
+ ttrx_compute_starving = 3737,
+ ttrx_exec_core_active = 3738,
+ ttrx_exec_instr_fma = 3739,
+ ttrx_exec_instr_cvt = 3740,
+ ttrx_exec_instr_sfu = 3741,
+ ttrx_exec_instr_msg = 3742,
+ ttrx_exec_instr_diverged = 3743,
+ ttrx_exec_icache_miss = 3744,
+ ttrx_exec_starve_arith = 3745,
+ ttrx_call_blend_shader = 3746,
+ ttrx_tex_msgi_num_flits = 3747,
+ ttrx_tex_dfch_clk_stalled = 3748,
+ ttrx_tex_tfch_clk_stalled = 3749,
+ ttrx_tex_tfch_starved_pending_data_fetch = 3750,
+ ttrx_tex_filt_num_operations = 3751,
+ ttrx_tex_filt_num_fxr_operations = 3752,
+ ttrx_tex_filt_num_fst_operations = 3753,
+ ttrx_tex_msgo_num_msg = 3754,
+ ttrx_tex_msgo_num_flits = 3755,
+ ttrx_ls_mem_read_full = 3756,
+ ttrx_ls_mem_read_short = 3757,
+ ttrx_ls_mem_write_full = 3758,
+ ttrx_ls_mem_write_short = 3759,
+ ttrx_ls_mem_atomic = 3760,
+ ttrx_vary_instr = 3761,
+ ttrx_vary_slot_32 = 3762,
+ ttrx_vary_slot_16 = 3763,
+ ttrx_attr_instr = 3764,
+ ttrx_arith_instr_fp_mul = 3765,
+ ttrx_beats_rd_ftc = 3766,
+ ttrx_beats_rd_ftc_ext = 3767,
+ ttrx_beats_rd_lsc = 3768,
+ ttrx_beats_rd_lsc_ext = 3769,
+ ttrx_beats_rd_tex = 3770,
+ ttrx_beats_rd_tex_ext = 3771,
+ ttrx_beats_rd_other = 3772,
+ ttrx_beats_wr_lsc_other = 3773,
+ ttrx_beats_wr_tib = 3774,
+ ttrx_beats_wr_lsc_wb = 3775,
+ ttrx_mmu_requests = 3780,
+ ttrx_mmu_table_reads_l3 = 3781,
+ ttrx_mmu_table_reads_l2 = 3782,
+ ttrx_mmu_hit_l3 = 3783,
+ ttrx_mmu_hit_l2 = 3784,
+ ttrx_mmu_s2_requests = 3785,
+ ttrx_mmu_s2_table_reads_l3 = 3786,
+ ttrx_mmu_s2_table_reads_l2 = 3787,
+ ttrx_mmu_s2_hit_l3 = 3788,
+ ttrx_mmu_s2_hit_l2 = 3789,
+ ttrx_l2_rd_msg_in = 3792,
+ ttrx_l2_rd_msg_in_stall = 3793,
+ ttrx_l2_wr_msg_in = 3794,
+ ttrx_l2_wr_msg_in_stall = 3795,
+ ttrx_l2_snp_msg_in = 3796,
+ ttrx_l2_snp_msg_in_stall = 3797,
+ ttrx_l2_rd_msg_out = 3798,
+ ttrx_l2_rd_msg_out_stall = 3799,
+ ttrx_l2_wr_msg_out = 3800,
+ ttrx_l2_any_lookup = 3801,
+ ttrx_l2_read_lookup = 3802,
+ ttrx_l2_write_lookup = 3803,
+ ttrx_l2_ext_snoop_lookup = 3804,
+ ttrx_l2_ext_read = 3805,
+ ttrx_l2_ext_read_nosnp = 3806,
+ ttrx_l2_ext_read_unique = 3807,
+ ttrx_l2_ext_read_beats = 3808,
+ ttrx_l2_ext_ar_stall = 3809,
+ ttrx_l2_ext_ar_cnt_q1 = 3810,
+ ttrx_l2_ext_ar_cnt_q2 = 3811,
+ ttrx_l2_ext_ar_cnt_q3 = 3812,
+ ttrx_l2_ext_rresp_0_127 = 3813,
+ ttrx_l2_ext_rresp_128_191 = 3814,
+ ttrx_l2_ext_rresp_192_255 = 3815,
+ ttrx_l2_ext_rresp_256_319 = 3816,
+ ttrx_l2_ext_rresp_320_383 = 3817,
+ ttrx_l2_ext_write = 3818,
+ ttrx_l2_ext_write_nosnp_full = 3819,
+ ttrx_l2_ext_write_nosnp_ptl = 3820,
+ ttrx_l2_ext_write_snp_full = 3821,
+ ttrx_l2_ext_write_snp_ptl = 3822,
+ ttrx_l2_ext_write_beats = 3823,
+ ttrx_l2_ext_w_stall = 3824,
+ ttrx_l2_ext_aw_cnt_q1 = 3825,
+ ttrx_l2_ext_aw_cnt_q2 = 3826,
+ ttrx_l2_ext_aw_cnt_q3 = 3827,
+ ttrx_l2_ext_snoop = 3828,
+ ttrx_l2_ext_snoop_stall = 3829,
+ ttrx_l2_ext_snoop_resp_clean = 3830,
+ ttrx_l2_ext_snoop_resp_data = 3831,
+ ttrx_l2_ext_snoop_internal = 3832,
+ tnax_messages_sent = 3844,
+ tnax_messages_received = 3845,
+ tnax_gpu_active = 3846,
+ tnax_irq_active = 3847,
+ tnax_js0_jobs = 3848,
+ tnax_js0_tasks = 3849,
+ tnax_js0_active = 3850,
+ tnax_js0_wait_flush = 3851,
+ tnax_js0_wait_read = 3852,
+ tnax_js0_wait_issue = 3853,
+ tnax_js0_wait_depend = 3854,
+ tnax_js0_wait_finish = 3855,
+ tnax_js1_jobs = 3856,
+ tnax_js1_tasks = 3857,
+ tnax_js1_active = 3858,
+ tnax_js1_wait_flush = 3859,
+ tnax_js1_wait_read = 3860,
+ tnax_js1_wait_issue = 3861,
+ tnax_js1_wait_depend = 3862,
+ tnax_js1_wait_finish = 3863,
+ tnax_js2_jobs = 3864,
+ tnax_js2_tasks = 3865,
+ tnax_js2_active = 3866,
+ tnax_js2_wait_flush = 3867,
+ tnax_js2_wait_read = 3868,
+ tnax_js2_wait_issue = 3869,
+ tnax_js2_wait_depend = 3870,
+ tnax_js2_wait_finish = 3871,
+ tnax_cache_flush = 3903,
+ tnax_tiler_active = 3908,
+ tnax_jobs_processed = 3909,
+ tnax_triangles = 3910,
+ tnax_lines = 3911,
+ tnax_points = 3912,
+ tnax_front_facing = 3913,
+ tnax_back_facing = 3914,
+ tnax_prim_visible = 3915,
+ tnax_prim_culled = 3916,
+ tnax_prim_clipped = 3917,
+ tnax_prim_sat_culled = 3918,
+ tnax_bin_alloc_init = 3919,
+ tnax_bin_alloc_overflow = 3920,
+ tnax_bus_read = 3921,
+ tnax_bus_write = 3923,
+ tnax_loading_desc = 3924,
+ tnax_idvs_pos_shad_req = 3925,
+ tnax_idvs_pos_shad_wait = 3926,
+ tnax_idvs_pos_shad_stall = 3927,
+ tnax_idvs_pos_fifo_full = 3928,
+ tnax_prefetch_stall = 3929,
+ tnax_vcache_hit = 3930,
+ tnax_vcache_miss = 3931,
+ tnax_vcache_line_wait = 3932,
+ tnax_vfetch_pos_read_wait = 3933,
+ tnax_vfetch_vertex_wait = 3934,
+ tnax_vfetch_stall = 3935,
+ tnax_primassy_stall = 3936,
+ tnax_bbox_gen_stall = 3937,
+ tnax_idvs_vbu_hit = 3938,
+ tnax_idvs_vbu_miss = 3939,
+ tnax_idvs_vbu_line_deallocate = 3940,
+ tnax_idvs_var_shad_req = 3941,
+ tnax_idvs_var_shad_stall = 3942,
+ tnax_binner_stall = 3943,
+ tnax_iter_stall = 3944,
+ tnax_compress_miss = 3945,
+ tnax_compress_stall = 3946,
+ tnax_pcache_hit = 3947,
+ tnax_pcache_miss = 3948,
+ tnax_pcache_miss_stall = 3949,
+ tnax_pcache_evict_stall = 3950,
+ tnax_pmgr_ptr_wr_stall = 3951,
+ tnax_pmgr_ptr_rd_stall = 3952,
+ tnax_pmgr_cmd_wr_stall = 3953,
+ tnax_wrbuf_active = 3954,
+ tnax_wrbuf_hit = 3955,
+ tnax_wrbuf_miss = 3956,
+ tnax_wrbuf_no_free_line_stall = 3957,
+ tnax_wrbuf_no_axi_id_stall = 3958,
+ tnax_wrbuf_axi_stall = 3959,
+ tnax_utlb_trans = 3963,
+ tnax_utlb_trans_hit = 3964,
+ tnax_utlb_trans_stall = 3965,
+ tnax_utlb_trans_miss_delay = 3966,
+ tnax_utlb_mmu_req = 3967,
+ tnax_frag_active = 3972,
+ tnax_frag_primitives_out = 3973,
+ tnax_frag_prim_rast = 3974,
+ tnax_frag_fpk_active = 3975,
+ tnax_frag_starving = 3976,
+ tnax_frag_warps = 3977,
+ tnax_frag_partial_quads_rast = 3978,
+ tnax_frag_quads_rast = 3979,
+ tnax_frag_quads_ezs_test = 3980,
+ tnax_frag_quads_ezs_update = 3981,
+ tnax_frag_quads_ezs_kill = 3982,
+ tnax_frag_lzs_test = 3983,
+ tnax_frag_lzs_kill = 3984,
+ tnax_warp_reg_size_64 = 3985,
+ tnax_frag_ptiles = 3986,
+ tnax_frag_trans_elim = 3987,
+ tnax_quad_fpk_killer = 3988,
+ tnax_full_quad_warps = 3989,
+ tnax_compute_active = 3990,
+ tnax_compute_tasks = 3991,
+ tnax_compute_warps = 3992,
+ tnax_compute_starving = 3993,
+ tnax_exec_core_active = 3994,
+ tnax_exec_instr_fma = 3995,
+ tnax_exec_instr_cvt = 3996,
+ tnax_exec_instr_sfu = 3997,
+ tnax_exec_instr_msg = 3998,
+ tnax_exec_instr_diverged = 3999,
+ tnax_exec_icache_miss = 4000,
+ tnax_exec_starve_arith = 4001,
+ tnax_call_blend_shader = 4002,
+ tnax_tex_msgi_num_flits = 4003,
+ tnax_tex_dfch_clk_stalled = 4004,
+ tnax_tex_tfch_clk_stalled = 4005,
+ tnax_tex_tfch_starved_pending_data_fetch = 4006,
+ tnax_tex_filt_num_operations = 4007,
+ tnax_tex_filt_num_fxr_operations = 4008,
+ tnax_tex_filt_num_fst_operations = 4009,
+ tnax_tex_msgo_num_msg = 4010,
+ tnax_tex_msgo_num_flits = 4011,
+ tnax_ls_mem_read_full = 4012,
+ tnax_ls_mem_read_short = 4013,
+ tnax_ls_mem_write_full = 4014,
+ tnax_ls_mem_write_short = 4015,
+ tnax_ls_mem_atomic = 4016,
+ tnax_vary_instr = 4017,
+ tnax_vary_slot_32 = 4018,
+ tnax_vary_slot_16 = 4019,
+ tnax_attr_instr = 4020,
+ tnax_arith_instr_fp_mul = 4021,
+ tnax_beats_rd_ftc = 4022,
+ tnax_beats_rd_ftc_ext = 4023,
+ tnax_beats_rd_lsc = 4024,
+ tnax_beats_rd_lsc_ext = 4025,
+ tnax_beats_rd_tex = 4026,
+ tnax_beats_rd_tex_ext = 4027,
+ tnax_beats_rd_other = 4028,
+ tnax_beats_wr_lsc_other = 4029,
+ tnax_beats_wr_tib = 4030,
+ tnax_beats_wr_lsc_wb = 4031,
+ tnax_mmu_requests = 4036,
+ tnax_mmu_table_reads_l3 = 4037,
+ tnax_mmu_table_reads_l2 = 4038,
+ tnax_mmu_hit_l3 = 4039,
+ tnax_mmu_hit_l2 = 4040,
+ tnax_mmu_s2_requests = 4041,
+ tnax_mmu_s2_table_reads_l3 = 4042,
+ tnax_mmu_s2_table_reads_l2 = 4043,
+ tnax_mmu_s2_hit_l3 = 4044,
+ tnax_mmu_s2_hit_l2 = 4045,
+ tnax_l2_rd_msg_in = 4048,
+ tnax_l2_rd_msg_in_stall = 4049,
+ tnax_l2_wr_msg_in = 4050,
+ tnax_l2_wr_msg_in_stall = 4051,
+ tnax_l2_snp_msg_in = 4052,
+ tnax_l2_snp_msg_in_stall = 4053,
+ tnax_l2_rd_msg_out = 4054,
+ tnax_l2_rd_msg_out_stall = 4055,
+ tnax_l2_wr_msg_out = 4056,
+ tnax_l2_any_lookup = 4057,
+ tnax_l2_read_lookup = 4058,
+ tnax_l2_write_lookup = 4059,
+ tnax_l2_ext_snoop_lookup = 4060,
+ tnax_l2_ext_read = 4061,
+ tnax_l2_ext_read_nosnp = 4062,
+ tnax_l2_ext_read_unique = 4063,
+ tnax_l2_ext_read_beats = 4064,
+ tnax_l2_ext_ar_stall = 4065,
+ tnax_l2_ext_ar_cnt_q1 = 4066,
+ tnax_l2_ext_ar_cnt_q2 = 4067,
+ tnax_l2_ext_ar_cnt_q3 = 4068,
+ tnax_l2_ext_rresp_0_127 = 4069,
+ tnax_l2_ext_rresp_128_191 = 4070,
+ tnax_l2_ext_rresp_192_255 = 4071,
+ tnax_l2_ext_rresp_256_319 = 4072,
+ tnax_l2_ext_rresp_320_383 = 4073,
+ tnax_l2_ext_write = 4074,
+ tnax_l2_ext_write_nosnp_full = 4075,
+ tnax_l2_ext_write_nosnp_ptl = 4076,
+ tnax_l2_ext_write_snp_full = 4077,
+ tnax_l2_ext_write_snp_ptl = 4078,
+ tnax_l2_ext_write_beats = 4079,
+ tnax_l2_ext_w_stall = 4080,
+ tnax_l2_ext_aw_cnt_q1 = 4081,
+ tnax_l2_ext_aw_cnt_q2 = 4082,
+ tnax_l2_ext_aw_cnt_q3 = 4083,
+ tnax_l2_ext_snoop = 4084,
+ tnax_l2_ext_snoop_stall = 4085,
+ tnax_l2_ext_snoop_resp_clean = 4086,
+ tnax_l2_ext_snoop_resp_data = 4087,
+ tnax_l2_ext_snoop_internal = 4088
+} MaliGpuCounter;
+
+struct mali_counter_values {
+ MaliGpuCounter counter;
+ uint32_t* values;
+ size_t num_values;
+};
+
+struct mali_counter_response {
+ struct mali_counter_values* counter_values;
+ size_t num_counters;
+};
+
+void initialize_mali_perf_reader();
+struct mali_counter_response read_perf_metrics(MaliGpuCounter* counters,
+ size_t num_counters);
+void free_counters(struct mali_counter_response counters);
+void cleanup_mali_perf_reader();
+
+#endif // __MALI_GPU_PERF_METRICS_H__
diff --git a/mali/mali_gpu_props.c b/mali/mali_gpu_props.c
new file mode 100644
index 0000000..a3190c9
--- /dev/null
+++ b/mali/mali_gpu_props.c
@@ -0,0 +1,106 @@
+/*
+ * Copyright 2022 The Chromium OS Authors. All rights reserved.
+ * Use of this source code is governed by a BSD-style license that can be
+ * found in the LICENSE file.
+ */
+
+#include "mali/mali_gpu_props.h"
+
+#include <assert.h>
+#include <errno.h>
+#include <fcntl.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <string.h>
+#include <sys/ioctl.h>
+#include <unistd.h>
+
+#include "logging.h"
+#include "mali/mali_ioctl.h"
+
+#define PROPS_BUFFER_SIZE 4096
+char props_buffer_cached = 0;
+uint8_t props_buffer[PROPS_BUFFER_SIZE];
+
+uint64_t get_cached_gpu_prop(MaliGpuProperty prop) {
+ int i = 0;
+ uint64_t ret;
+ while (i + 12 < PROPS_BUFFER_SIZE) {
+ uint32_t key = *(uint32_t*)(props_buffer + i);
+ if (!key)
+ break;
+
+ i += 4;
+ int size_type = key & 0x3;
+ int key_type = key >> 2;
+
+ switch (size_type) {
+ case 0:
+ ret = *(uint8_t*)(props_buffer + i);
+ i += 1;
+ break;
+ case 1:
+ ret = *(uint16_t*)(props_buffer + i);
+ i += 2;
+ break;
+ case 2:
+ ret = *(uint32_t*)(props_buffer + i);
+ i += 4;
+ break;
+ case 3:
+ ret = *(uint64_t*)(props_buffer + i);
+ i += 8;
+ break;
+ default:
+ LOG_ERROR(
+ "Error: GPU properties buffer contains malformed key at index "
+ "%d.\n",
+ i - 4);
+ return -1;
+ }
+
+ if (key_type == prop)
+ return ret;
+ }
+
+ LOG_ERROR("Error: GPU property %d not found.\n", prop);
+ return -1;
+}
+
+void get_gpu_property_buffer() {
+ int gpufd = open(kGpuDevice, O_RDWR | O_CLOEXEC);
+ if (gpufd < 0)
+ LOG_FATAL("Error opening GPU device! %s\n", strerror(errno));
+
+ // Note that even though we don't strictly speaking need this information for
+ // querying GPU properties, the GPU properties IOCTL will fail unless we check
+ // the version and set the flags.
+ struct kbase_ioctl_version_check version_check;
+ if (ioctl(gpufd, KBASE_IOCTL_VERSION_CHECK, &version_check) < 0)
+ LOG_FATAL("Error checking GPU version! %s\n", strerror(errno));
+ assert(version_check.major >= SUPPORTED_MAJOR_VERSION);
+ assert(version_check.minor >= SUPPORTED_MINOR_VERSION);
+
+ struct kbase_ioctl_set_flags init_flags;
+ init_flags.create_flags = BASE_CONTEXT_SYSTEM_MONITOR_SUBMIT_DISABLED;
+ if (ioctl(gpufd, KBASE_IOCTL_SET_FLAGS, &init_flags) < 0)
+ LOG_FATAL("Error initializing GPU context! %s\n", strerror(errno));
+
+ struct kbase_ioctl_get_gpuprops gpu_props_req;
+ gpu_props_req.buffer_ptr = (uint64_t)props_buffer;
+ gpu_props_req.size = PROPS_BUFFER_SIZE;
+ gpu_props_req.flags = 0;
+ if (ioctl(gpufd, KBASE_IOCTL_GET_GPUPROPS, &gpu_props_req) < 0)
+ LOG_FATAL("Error getting GPU properties! %s\n", strerror(errno));
+
+ props_buffer_cached = 1;
+
+ close(gpufd);
+}
+
+uint64_t get_gpu_prop(MaliGpuProperty prop) {
+ if (!props_buffer_cached)
+ get_gpu_property_buffer();
+
+ return get_cached_gpu_prop(prop);
+}
diff --git a/mali/mali_gpu_props.h b/mali/mali_gpu_props.h
new file mode 100644
index 0000000..eb3fd06
--- /dev/null
+++ b/mali/mali_gpu_props.h
@@ -0,0 +1,101 @@
+/*
+ * Copyright 2022 The Chromium OS Authors. All rights reserved.
+ * Use of this source code is governed by a BSD-style license that can be
+ * found in the LICENSE file.
+ */
+
+#include <stdint.h>
+
+#ifndef __MALI_GPU_PROPS_H__
+#define __MALI_GPU_PROPS_H__
+
+// Constant values found in
+// third_party/kernel/next/drivers/gpu/arm/valhall/mali_kbase_ioctl.h
+typedef enum {
+ gpu_prop_product_id = 1,
+ gpu_prop_version_status = 2,
+ gpu_prop_minor_revision = 3,
+ gpu_prop_major_revision = 4,
+ gpu_prop_freq_max_khz = 6,
+ gpu_prop_log2_program_counter_size = 8,
+ gpu_prop_texture_features_0 = 9,
+ gpu_prop_texture_features_1 = 10,
+ gpu_prop_texture_features_2 = 11,
+ gpu_prop_available_memory_size = 12,
+ gpu_prop_log2_l2_line_size = 13,
+ gpu_prop_log2_l2_cache_size = 14,
+ gpu_prop_num_l2 = 15,
+ gpu_prop_tiler_bin_size_bytes = 16,
+ gpu_prop_tiler_max_active_levels = 17,
+ gpu_prop_max_threads = 18,
+ gpu_prop_max_workgroup_size = 19,
+ gpu_prop_max_barrier_size = 20,
+ gpu_prop_max_registers = 21,
+ gpu_prop_max_task_queue = 22,
+ gpu_prop_max_thread_group_split = 23,
+ gpu_prop_impl_tech = 24,
+ gpu_prop_shader_present_mask = 25,
+ gpu_prop_tiler_present = 26,
+ gpu_prop_l2_present = 27,
+ gpu_prop_stack_present = 28,
+ gpu_prop_l2_features = 29,
+ gpu_prop_core_features = 30,
+ gpu_prop_mem_features = 31,
+ gpu_prop_mmu_features = 32,
+ gpu_prop_as_present = 33,
+ gpu_prop_js_present = 34,
+ gpu_prop_js_features_0 = 35,
+ gpu_prop_js_features_1 = 36,
+ gpu_prop_js_features_2 = 37,
+ gpu_prop_js_features_3 = 38,
+ gpu_prop_js_features_4 = 39,
+ gpu_prop_js_features_5 = 40,
+ gpu_prop_js_features_6 = 41,
+ gpu_prop_js_features_7 = 42,
+ gpu_prop_js_features_8 = 43,
+ gpu_prop_js_features_9 = 44,
+ gpu_prop_js_features_10 = 45,
+ gpu_prop_js_features_11 = 46,
+ gpu_prop_js_features_12 = 47,
+ gpu_prop_js_features_13 = 48,
+ gpu_prop_js_features_14 = 49,
+ gpu_prop_js_features_15 = 50,
+ gpu_prop_tiler_features = 51,
+ gpu_prop_raw_texture_features_0 = 52,
+ gpu_prop_raw_texture_features_1 = 53,
+ gpu_prop_raw_texture_features_2 = 54,
+ gpu_prop_gpu_id = 55,
+ gpu_prop_raw_max_threads = 56,
+ gpu_prop_raw_max_workgroupd_size = 57,
+ gpu_prop_raw_max_barrier_size = 58,
+ gpu_prop_thread_features = 59,
+ gpu_prop_coherency_mode = 60,
+ gpu_prop_num_coherency_groups = 61,
+ gpu_prop_num_core_coherency_groups = 62,
+ gpu_prop_coherency = 63,
+ gpu_prop_coherency_group_0 = 64,
+ gpu_prop_coherency_group_1 = 65,
+ gpu_prop_coherency_group_2 = 66,
+ gpu_prop_coherency_group_3 = 67,
+ gpu_prop_coherency_group_4 = 68,
+ gpu_prop_coherency_group_5 = 69,
+ gpu_prop_coherency_group_6 = 70,
+ gpu_prop_coherency_group_7 = 71,
+ gpu_prop_coherency_group_8 = 72,
+ gpu_prop_coherency_group_9 = 73,
+ gpu_prop_coherency_group_10 = 74,
+ gpu_prop_coherency_group_11 = 75,
+ gpu_prop_coherency_group_12 = 76,
+ gpu_prop_coherency_group_13 = 77,
+ gpu_prop_coherency_group_14 = 78,
+ gpu_prop_coherency_group_15 = 79,
+ gpu_prop_texture_features_3 = 80,
+ gpu_prop_raw_texture_features_3 = 81,
+ gpu_prop_num_exec_engines = 82,
+ gpu_prop_raw_thread_tls_alloc = 83,
+ gpu_prop_tls_alloc = 84
+} MaliGpuProperty;
+
+uint64_t get_gpu_prop(MaliGpuProperty prop);
+
+#endif // __MALI_GPU_PROPS_H__
diff --git a/mali/mali_ioctl.h b/mali/mali_ioctl.h
new file mode 100644
index 0000000..3546387
--- /dev/null
+++ b/mali/mali_ioctl.h
@@ -0,0 +1,84 @@
+/*
+ * Copyright 2022 The Chromium OS Authors. All rights reserved.
+ * Use of this source code is governed by a BSD-style license that can be
+ * found in the LICENSE file.
+ */
+
+#include <asm-generic/ioctl.h>
+#include <stdint.h>
+
+#ifndef __MALI_IOCTL_H__
+#define __MALI_IOCTL_H__
+
+static const char* kGpuDevice = "/dev/mali0";
+
+// Definitions courtesy of
+// third_party/kernel/v5.10/drivers/gpu/arm/valhall/mali_kbase_ioctl.h
+
+#define KBASE_IOCTL_TYPE 0x80
+#define KBASE_IOCTL_GET_GPUPROPS \
+ _IOW(KBASE_IOCTL_TYPE, 3, struct kbase_ioctl_get_gpuprops)
+#define KBASE_IOCTL_SET_FLAGS \
+ _IOW(KBASE_IOCTL_TYPE, 1, struct kbase_ioctl_set_flags)
+#define KBASE_IOCTL_VERSION_CHECK \
+ _IOWR(KBASE_IOCTL_TYPE, 0, struct kbase_ioctl_version_check)
+#define KBASE_IOCTL_HWCNT_READER_SETUP \
+ _IOW(KBASE_IOCTL_TYPE, 8, struct kbase_ioctl_hwcnt_reader_setup)
+
+#define BASE_CONTEXT_SYSTEM_MONITOR_SUBMIT_DISABLED ((uint32_t)1 << 1)
+
+struct __attribute__((packed)) kbase_ioctl_get_gpuprops {
+ uint64_t buffer_ptr;
+ uint32_t size;
+ uint32_t flags;
+};
+
+struct __attribute__((packed)) kbase_ioctl_set_flags {
+ uint32_t create_flags;
+};
+
+struct __attribute__((packed)) kbase_ioctl_version_check {
+ uint16_t major;
+ uint16_t minor;
+};
+
+struct __attribute__((packed)) kbase_ioctl_hwcnt_reader_setup {
+ uint32_t num_buffers;
+ uint32_t job_manager_mask;
+ uint32_t shader_mask;
+ uint32_t tiler_mask;
+ uint32_t mmu_l2_mask;
+};
+
+// Definitions courtesy of
+// third_party/kernel/v5.10/drivers/gpu/arm/valhall/mali_kbase_hwcnt_reader.h
+
+#define KBASE_HWCNT_READER 0xBE
+#define KBASE_HWCNT_READER_GET_HWVER _IOR(KBASE_HWCNT_READER, 0x00, uint32_t)
+#define KBASE_HWCNT_READER_GET_BUFFER_SIZE \
+ _IOR(KBASE_HWCNT_READER, 0x01, uint32_t)
+#define KBASE_HWCNT_READER_DUMP _IOW(KBASE_HWCNT_READER, 0x10, uint32_t)
+#define KBASE_HWCNT_READER_CLEAR _IOW(KBASE_HWCNT_READER, 0x11, uint32_t)
+#define KBASE_HWCNT_READER_GET_BUFFER \
+ _IOR(KBASE_HWCNT_READER, 0x20, struct kbase_hwcnt_reader_metadata)
+#define KBASE_HWCNT_READER_PUT_BUFFER \
+ _IOW(KBASE_HWCNT_READER, 0x21, struct kbase_hwcnt_reader_metadata)
+#define KBASE_HWCNT_READER_SET_INTERVAL _IOW(KBASE_HWCNT_READER, 0x30, uint32_t)
+#define KBASE_HWCNT_READER_ENABLE_EVENT _IOW(KBASE_HWCNT_READER, 0x40, uint32_t)
+#define KBASE_HWCNT_READER_DISABLE_EVENT \
+ _IOW(KBASE_HWCNT_READER, 0x41, uint32_t)
+#define KBASE_HWCNT_READER_GET_API_VERSION \
+ _IOW(KBASE_HWCNT_READER, 0xFF, uint32_t)
+
+struct kbase_hwcnt_reader_metadata {
+ uint64_t timestamp;
+ uint32_t event_id;
+ uint32_t buffer_idx;
+};
+
+#define SUPPORTED_MAJOR_VERSION 11
+#define SUPPORTED_MINOR_VERSION 21
+#define SUPPORTED_API_VERSION 1
+#define SUPPORTED_HW_VERSION 5
+
+#endif // __MALI_IOCTL_H__
diff --git a/mali/module.mk b/mali/module.mk
new file mode 100644
index 0000000..31ea7bb
--- /dev/null
+++ b/mali/module.mk
@@ -0,0 +1,11 @@
+# Copyright 2022 The Chromium OS Authors. All rights reserved.
+# Use of this source code is governed by a BSD-style license that can be
+# found in the LICENSE file.
+
+include common.mk
+
+CFLAGS += -std=gnu99 -I$(SRC)
+
+CC_STATIC_LIBRARY(libmali.pic.a): \
+ mali/mali_gpu_props.o \
+ mali/mali_gpu_perf_metrics.o
diff --git a/mali_stats.c b/mali_stats.c
new file mode 100644
index 0000000..3f2863b
--- /dev/null
+++ b/mali_stats.c
@@ -0,0 +1,51 @@
+/*
+ * Copyright 2022 The Chromium OS Authors. All rights reserved.
+ * Use of this source code is governed by a BSD-style license that can be
+ * found in the LICENSE file.
+ */
+
+#include <assert.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#include "mali/mali_gpu_perf_metrics.h"
+#include "mali/mali_gpu_props.h"
+
+int main(int argc, char** argv) {
+ initialize_mali_perf_reader();
+
+ MaliGpuCounter counters[] = {tnax_gpu_active, tnax_tiler_active,
+ tnax_compute_active, tnax_mmu_requests};
+ struct mali_counter_response response;
+ response = read_perf_metrics(counters, 4);
+
+ assert(response.num_counters == 4);
+
+ printf("Counter TNAx GPU_ACTIVE value(s):\n");
+ for (int j = 0; j < response.counter_values[0].num_values; j++) {
+ printf("%d\n", response.counter_values[0].values[j]);
+ }
+
+ printf("\n");
+ printf("Counter TNAx TILER_ACTIVE value(s):\n");
+ for (int j = 0; j < response.counter_values[1].num_values; j++) {
+ printf("%d\n", response.counter_values[1].values[j]);
+ }
+
+ printf("\n");
+ printf("Counter TNAx COMPUTE_ACTIVE value(s):\n");
+ for (int j = 0; j < response.counter_values[2].num_values; j++) {
+ printf("%d\n", response.counter_values[2].values[j]);
+ }
+
+ printf("\n");
+ printf("Counter TNAx MMU_REQUESTS value(s):\n");
+ for (int j = 0; j < response.counter_values[3].num_values; j++) {
+ printf("%d\n", response.counter_values[3].values[j]);
+ }
+
+ free_counters(response);
+
+ cleanup_mali_perf_reader();
+}
diff --git a/v4l2_macros.h b/v4l2_macros.h
index dd18698..6acb491 100644
--- a/v4l2_macros.h
+++ b/v4l2_macros.h
@@ -8,39 +8,11 @@
#define V4L2_MACROS_H
#include <linux/videodev2.h>
-#include <stdio.h>
-#include <stdlib.h>
+
+#include "logging.h"
#define FOURCC_SIZE 4
-enum logging_levels {
- kLoggingDebug = -1,
- kLoggingInfo = 0,
- kLoggingError,
- kLoggingFatal,
- kLoggingLevelMax
-};
-
-#define DEFAULT_LOG_LEVEL kLoggingInfo
-
-#define LOG(level, stream, fmt, ...) \
- do { \
- if (level >= DEFAULT_LOG_LEVEL) { \
- fprintf(stream, fmt, ##__VA_ARGS__); \
- fprintf(stream, "\n"); \
- fflush(stream); \
- } \
- } while (0)
-
-#define LOG_DEBUG(fmt, ...) LOG(kLoggingDebug, stderr, fmt, ##__VA_ARGS__)
-#define LOG_INFO(fmt, ...) LOG(kLoggingInfo, stderr, fmt, ##__VA_ARGS__)
-#define LOG_ERROR(fmt, ...) LOG(kLoggingError, stderr, fmt, ##__VA_ARGS__)
-#define LOG_FATAL(fmt, ...) \
- do { \
- LOG(kLoggingFatal, stderr, fmt, ##__VA_ARGS__); \
- exit(EXIT_FAILURE); \
- } while (0)
-
// TODO(frkoenig):
// P010 had not landed as an official V4L2 format yet. Once it has remove this.
#ifndef V4L2_PIX_FMT_P010
