Making ATT work with Profilerv2

Change-Id: Ic9334aa80e40faaaf5c1a79ba37dbe52e8d31253 [ROCm/rocprofiler commit: 03c305dbd4]
2023-02-07 13:06:02 +05:30
parent f7772fb704
commit 4f742d346b
28 changed files with 3783 additions and 24 deletions
@@ -486,4 +486,146 @@ hsa_ven_amd_aqlprofile_profile_t* InitializeDeviceProfilingAqlPackets(
  return profile;
 }

+// ATT
+uint32_t g_output_buffer_size = 0x8000000;  // 128M x 16 = 2GB
+bool g_output_buffer_local = true;
+
+// Allocate system memory accessible by both CPU and GPU
+uint8_t* AllocateSysMemory(hsa_agent_t gpu_agent, size_t size, hsa_amd_memory_pool_t* cpu_pool) {
+  hsa_status_t status = HSA_STATUS_ERROR;
+  uint8_t* buffer = NULL;
+  size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
+  // if (!cpu_agents_.empty()) {
+  status = hsa_amd_memory_pool_allocate(*cpu_pool, size, 0, reinterpret_cast<void**>(&buffer));
+  // Both the CPU and GPU can access the memory
+  if (status == HSA_STATUS_SUCCESS) {
+    hsa_agent_t ag_list[1] = {gpu_agent};
+    status = hsa_amd_agents_allow_access(1, ag_list, NULL, buffer);
+  }
+  // }
+  uint8_t* ptr = (status == HSA_STATUS_SUCCESS) ? buffer : NULL;
+  return ptr;
+}
+
+// Allocate memory for use by a kernel of specified size
+uint8_t* AllocateLocalMemory(size_t size, hsa_amd_memory_pool_t* gpu_pool) {
+  hsa_status_t status = HSA_STATUS_ERROR;
+  uint8_t* buffer = NULL;
+  size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
+  status = hsa_amd_memory_pool_allocate(*gpu_pool, size, 0, reinterpret_cast<void**>(&buffer));
+  uint8_t* ptr = (status == HSA_STATUS_SUCCESS) ? buffer : NULL;
+  return ptr;
+}
+
+hsa_status_t Allocate(hsa_agent_t gpu_agent, hsa_ven_amd_aqlprofile_profile_t* profile,
+                      hsa_amd_memory_pool_t* cpu_pool, hsa_amd_memory_pool_t* gpu_pool) {
+  profile->command_buffer.ptr =
+      AllocateSysMemory(gpu_agent, profile->command_buffer.size, cpu_pool);
+  profile->output_buffer.size = g_output_buffer_size;
+  profile->output_buffer.ptr = (g_output_buffer_local)
+      ? AllocateLocalMemory(profile->output_buffer.size, gpu_pool)
+      : AllocateSysMemory(gpu_agent, profile->output_buffer.size, cpu_pool);
+  return (profile->command_buffer.ptr && profile->output_buffer.ptr) ? HSA_STATUS_SUCCESS
+                                                                     : HSA_STATUS_ERROR;
+}
+
+bool AllocateMemoryPools(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
+                         hsa_amd_memory_pool_t* cpu_pool, hsa_amd_memory_pool_t* gpu_pool) {
+  hsa_status_t status = hsa_amd_agent_iterate_memory_pools(cpu_agent, FindStandardPool, cpu_pool);
+  CHECK_HSA_STATUS("hsa_amd_agent_iterate_memory_pools(cpu_pool)", status);
+
+  status = hsa_amd_agent_iterate_memory_pools(gpu_agent, FindStandardPool, gpu_pool);
+  CHECK_HSA_STATUS("hsa_amd_agent_iterate_memory_pools(gpu_pool)", status);
+
+  return true;
+}
+
+// map between gpu agent handle and att_memory_pools_t
+typedef std::map<uint64_t, att_memory_pools_t*> att_mem_pools_map_t;
+
+att_mem_pools_map_t* agent_att_mem_pools_map = nullptr;
+std::atomic<bool> att_map_init{false};
+
+att_mem_pools_map_t* GetAttMemPoolsMap() {
+  if (!att_map_init.load(std::memory_order_relaxed)) {
+    agent_att_mem_pools_map = new att_mem_pools_map_t();
+    att_map_init.exchange(true, std::memory_order_release);
+  }
+
+  return agent_att_mem_pools_map;
+}
+
+
+att_memory_pools_t* GetAttMemPools(hsa_agent_t gpu_agent) {
+  auto it = GetAttMemPoolsMap()->find(gpu_agent.handle);
+  if (it != GetAttMemPoolsMap()->end()) {
+    return it->second;
+  }
+  printf("Error: att_memory_pools_t instance not found for given gpu agent handle: %lu\n",
+         gpu_agent.handle);
+
+  return nullptr;
+}
+
+// Generate start and stop packets for collecting ATT traces
+// Also generate and return the profile object which has the PM4
+// command buffer and the output buffer for retrieving the traces
+hsa_ven_amd_aqlprofile_profile_t* GenerateATTPackets(
+    hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
+    std::vector<hsa_ven_amd_aqlprofile_parameter_t>& att_params, packet_t* start_packet,
+    packet_t* stop_packet) {
+  att_memory_pools_t* att_mem_pools = NULL;
+  auto it = GetAttMemPoolsMap()->find(gpu_agent.handle);
+  if (it == GetAttMemPoolsMap()->end()) {
+    att_mem_pools = new att_memory_pools_t;
+
+    // Allocate memory pools for cpu and gpu
+    AllocateMemoryPools(cpu_agent, gpu_agent, &att_mem_pools->cpu_mem_pool,
+                        &att_mem_pools->gpu_mem_pool);
+
+    GetAttMemPoolsMap()->emplace(gpu_agent.handle, att_mem_pools);
+  } else
+    att_mem_pools = it->second;
+
+#pragma GCC diagnostic push
+#pragma GCC diagnostic ignored "-Wconversion-null"
+  // Preparing the profile structure to get the packets
+  hsa_ven_amd_aqlprofile_profile_t* profile =
+      new hsa_ven_amd_aqlprofile_profile_t{gpu_agent,
+                                           HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_TRACE,
+                                           nullptr,
+                                           0,
+                                           &att_params[0],
+                                           (uint32_t)att_params.size(),
+                                           NULL,
+                                           NULL};
+#pragma GCC diagnostic pop
+
+  // Check the profile buffer sizes
+  hsa_status_t status = hsa_ven_amd_aqlprofile_start(profile, NULL);
+  if (status != HSA_STATUS_SUCCESS) printf("Error: aqlprofile_start(NULL)");
+  // // Double output buffer size if concurrent
+  // if (is_concurrent) profile.output_buffer.size *= 2;
+
+  // TODO: create a separate class for memory allocations
+  // Maintain pools per device
+  // handle allocation and resource cleanup
+
+
+  // Allocate command and output buffers
+  // command buffer -> from CPU memory pool
+  // output buffer -> from GPU memory pool
+  status =
+      Allocate(gpu_agent, profile, &att_mem_pools->cpu_mem_pool, &att_mem_pools->gpu_mem_pool);
+  if (status != HSA_STATUS_SUCCESS) printf("Error: Allocate()");
+
+  // Generate start/stop/read profiling packets
+  status = hsa_ven_amd_aqlprofile_start(profile, start_packet);
+  if (status != HSA_STATUS_SUCCESS) printf("Error: aqlprofile_start");
+  status = hsa_ven_amd_aqlprofile_stop(profile, stop_packet);
+  if (status != HSA_STATUS_SUCCESS) printf("Error: aqlprofile_stop");
+  if (status == HSA_STATUS_ERROR) return nullptr;
+  return profile;
+}
+
 }  // namespace Packet
@@ -55,5 +55,25 @@ hsa_amd_memory_pool_t& GetCommandPool();
 hsa_amd_memory_pool_t& GetOutputPool();


+hsa_ven_amd_aqlprofile_profile_t* GenerateATTPackets(
+    hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
+    std::vector<hsa_ven_amd_aqlprofile_parameter_t>& att_params, packet_t* start_packet,
+    packet_t* stop_packet);
+
+
+uint8_t* AllocateSysMemory(hsa_agent_t gpu_agent, size_t size, hsa_amd_memory_pool_t* cpu_pool);
+
+void get_command_buffer_map(std::map<size_t, uint8_t*> );
+void get_outbuffer_map(std::map<size_t, uint8_t*> );
+void initialize_pools(hsa_agent_t cpu_agent);
+
+typedef struct {
+  hsa_amd_memory_pool_t cpu_mem_pool;
+  hsa_amd_memory_pool_t gpu_mem_pool;
+} att_memory_pools_t;
+
+att_memory_pools_t* GetAttMemPools(hsa_agent_t gpu_agent);
+
+
 }  // namespace Packet
 #endif  // SRC_CORE_HSA_PACKETS_PACKETS_GENERATOR_H_
@@ -34,6 +34,8 @@
 #include "src/utils/helper.h"

 #define __NR_gettid 186
+#define MAX_ATT_PROFILES 16
+
 std::mutex sessions_pending_signal_lock;

 namespace rocmtools {
@@ -52,6 +54,8 @@ static inline bool IsEventMatch(const hsa_ven_amd_aqlprofile_event_t& event1,
      (event1.counter_id == event2.counter_id);
 }

+typedef std::vector<hsa_ven_amd_aqlprofile_info_data_t> att_trace_callback_data_t;
+
 static std::mutex ksymbol_map_lock;
 static std::map<uint64_t, std::string>* ksymbols;
 static std::atomic<bool> ksymbols_flag{true};
@@ -234,6 +238,17 @@ hsa_status_t pmcCallback(hsa_ven_amd_aqlprofile_info_type_t info_type,
  return status;
 }

+hsa_status_t attTraceDataCallback(hsa_ven_amd_aqlprofile_info_type_t info_type,
+                                   hsa_ven_amd_aqlprofile_info_data_t* info_data, void* data) {
+  hsa_status_t status = HSA_STATUS_SUCCESS;
+  att_trace_callback_data_t* passed_data = reinterpret_cast<att_trace_callback_data_t*>(data);
+  passed_data->push_back(*info_data);
+  // TODO: clear output buffers after copying
+  // either copy here or in AddattRecord
+
+  return status;
+}
+
 void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_signal_t& pending) {
  rocmtools::metrics::GetCounterData(pending.profile, pending.context->results_list);
  rocmtools::metrics::GetMetricsData(pending.context->results_map, pending.context->metrics_list);
@@ -260,6 +275,47 @@ void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_sign
  record->counters_count = rocprofiler_record_counters_instances_count_t{counters_vec.size()};
 }

+void AddAttRecord(rocprofiler_record_att_tracer_t* record, hsa_agent_t gpu_agent,                  
+                   att_pending_signal_t& pending) {
+  att_trace_callback_data_t data;
+  hsa_ven_amd_aqlprofile_iterate_data(pending.profile, attTraceDataCallback, &data);
+
+  // Get CPU and GPU memory pools
+  Packet::att_memory_pools_t* att_mem_pools = Packet::GetAttMemPools(gpu_agent);
+
+  // Allocate memory for shader_engine_data
+  record->shader_engine_data = static_cast<rocprofiler_record_se_att_data_t*>(
+      malloc(data.size() * sizeof(rocprofiler_record_se_att_data_t)));
+
+  att_trace_callback_data_t::iterator trace_data_it;
+
+  uint32_t se_index = 0;
+  // iterate over the trace data collected from each shader engine
+  for (trace_data_it = data.begin(); trace_data_it != data.end(); trace_data_it++) {
+    const void* data_ptr = trace_data_it->trace_data.ptr;
+    const uint32_t data_size = trace_data_it->trace_data.size;
+    // fprintf(arg->file, "    SE(%u) size(%u)\n", data.sample_id, data_size);
+
+    void* buffer = NULL;
+    if (data_size != 0) {
+      // Allocate buffer on CPU to copy out trace data
+      buffer = Packet::AllocateSysMemory(gpu_agent, data_size, &att_mem_pools->cpu_mem_pool);
+      if (buffer == NULL) fatal("Trace data buffer allocation failed");
+
+      auto status =
+          rocmtools::hsa_support::GetCoreApiTable().hsa_memory_copy_fn(buffer, data_ptr, data_size);
+      if (status != HSA_STATUS_SUCCESS) fatal("Trace data memcopy to host failed");
+
+      record->shader_engine_data[se_index].buffer_ptr = buffer;
+      record->shader_engine_data[se_index].buffer_size = data_size;
+      ++se_index;
+
+      // TODO: clear output buffers after copying
+    }
+  }
+  record->shader_engine_data_count = data.size();
+}
+
 // static const size_t MEM_PAGE_BYTES = 0x1000;
 // static const size_t MEM_PAGE_MASK = MEM_PAGE_BYTES - 1;
 // static std::mutex begin_signal_lock;
@@ -418,6 +474,67 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
  return false;
 }

+bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) {
+  // TODO: finish implementation to iterate trace data and add it to rocmtools record
+  // and generic buffer
+
+  auto queue_info_session = static_cast<queue_info_session_t*>(data);
+  if (!queue_info_session || !GetROCMToolObj() ||
+      !GetROCMToolObj()->GetSession(queue_info_session->session_id) ||
+      !GetROCMToolObj()->GetSession(queue_info_session->session_id)->GetAttTracer())
+    return true;
+  rocmtools::Session* session = GetROCMToolObj()->GetSession(queue_info_session->session_id);
+  rocmtools::att::AttTracer* att_tracer = session->GetAttTracer();
+  std::vector<att_pending_signal_t>& pending_signals =
+      const_cast<std::vector<att_pending_signal_t>&>(
+          att_tracer->GetPendingSignals(queue_info_session->writer_id));
+
+  if (!pending_signals.empty()) {
+    for (auto it = pending_signals.begin(); it != pending_signals.end();
+         it = pending_signals.erase(it)) {
+      auto& pending = *it;
+      std::lock_guard<std::mutex> lock(session->GetSessionLock());
+      if (hsa_support::GetCoreApiTable().hsa_signal_load_relaxed_fn(pending.signal)) return true;
+      rocprofiler_record_att_tracer_t record{};
+      record.kernel_id = rocprofiler_kernel_id_t{pending.kernel_descriptor};
+      record.gpu_id = rocprofiler_agent_id_t{
+          (uint64_t)hsa_support::GetAgentInfo(queue_info_session->agent.handle).getIndex()};
+      record.kernel_properties = pending.kernel_properties;
+      record.thread_id = rocprofiler_thread_id_t{pending.thread_id};
+      record.queue_idx = rocprofiler_queue_index_t{pending.queue_index};
+      record.queue_id = rocprofiler_queue_id_t{queue_info_session->queue_id};
+      if (/*pending.counters_count > 0 && */ pending.profile) {
+        AddAttRecord(&record, queue_info_session->agent, pending);
+      }
+      record.header = {ROCPROFILER_ATT_TRACER_RECORD,
+                       rocprofiler_record_id_t{GetROCMToolObj()->GetUniqueRecordId()}};
+
+      if (pending.session_id.handle == 0) {
+        pending.session_id = GetROCMToolObj()->GetCurrentSessionId();
+      }
+      if (session->FindBuffer(pending.buffer_id)) {
+        Memory::GenericBuffer* buffer = session->GetBuffer(pending.buffer_id);
+        record.header.id = rocprofiler_record_id_t{GetROCMToolObj()->GetUniqueRecordId()};
+        buffer->AddRecord(record);
+      }
+      hsa_status_t status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(
+          (pending.profile->output_buffer.ptr));
+      if (status != HSA_STATUS_SUCCESS) {
+        printf("Error: Couldn't free output buffer memory\n");
+      }
+      status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(
+          (pending.profile->command_buffer.ptr));
+      if (status != HSA_STATUS_SUCCESS) {
+        printf("Error: Couldn't free command buffer memory\n");
+      }
+      delete pending.profile;
+    }
+  }
+  delete queue_info_session;
+
+  return false;
+}
+
 void CreateBarrierPacket(const hsa_signal_t& packet_completion_signal,
                         std::vector<Packet::packet_t>* transformed_packets) {
  hsa_barrier_and_packet_t barrier{0};
@@ -439,6 +556,12 @@ void SignalAsyncHandler(const hsa_signal_t& signal, void* data) {
  if (status != HSA_STATUS_SUCCESS) fatal("hsa_amd_signal_async_handler failed");
 }

+void signalAsyncHandlerATT(const hsa_signal_t& signal, void* data) {
+  hsa_status_t status = hsa_support::GetAmdExtTable().hsa_amd_signal_async_handler_fn(
+      signal, HSA_SIGNAL_CONDITION_EQ, 0, AsyncSignalHandlerATT, data);
+  if (status != HSA_STATUS_SUCCESS) fatal("hsa_amd_signal_async_handler failed");
+}
+
 void CreateSignal(uint32_t attribute, hsa_signal_t* signal) {
  hsa_status_t status =
      hsa_support::GetAmdExtTable().hsa_amd_signal_create_fn(1, 0, nullptr, attribute, signal);
@@ -459,6 +582,7 @@ template <typename Integral> constexpr Integral bit_extract(Integral x, int firs
  return (x >> first) & bit_mask<Integral>(0, last - first);
 }

+static int KernelInterceptCount = 0;
 std::atomic<uint32_t> WRITER_ID{0};
 /**
 * @brief This function is a queue write interceptor. It intercepts the
@@ -487,9 +611,12 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt

  bool is_counter_collection_mode = false;
  bool is_timestamp_collection_mode = false;
+  bool is_att_collection_mode = false;
  bool is_pc_sampling_collection_mode = false;
-
+  std::vector<rocprofiler_att_parameter_t> att_parameters_data;
  uint32_t replay_mode_count = 0;
+  std::vector<std::string> kernel_profile_names;
+  std::vector<std::string> att_counters_names;

  rocmtools::Session* session = nullptr;

@@ -509,6 +636,17 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
          session->GetFilterIdWithKind(ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION);
      rocmtools::Filter* filter = session->GetFilter(filter_id);
      buffer_id = filter->GetBufferId();
+    } else if (session && session->FindFilterWithKind(ROCPROFILER_ATT_TRACE_COLLECTION)) {
+      rocprofiler_filter_id_t filter_id =
+          session->GetFilterIdWithKind(ROCPROFILER_ATT_TRACE_COLLECTION);
+      rocmtools::Filter* filter = session->GetFilter(filter_id);
+      att_parameters_data = filter->GetAttParametersData();
+      is_att_collection_mode = true;
+      buffer_id = session->GetFilter(session->GetFilterIdWithKind(ROCPROFILER_ATT_TRACE_COLLECTION))
+                      ->GetBufferId();
+                      
+      att_counters_names = filter->GetCounterData();
+      kernel_profile_names = std::get<std::vector<std::string>>(filter->GetProperty(ROCPROFILER_FILTER_KERNEL_NAMES));
    } else if (session && session->FindFilterWithKind(ROCPROFILER_PC_SAMPLING_COLLECTION)) {
      is_pc_sampling_collection_mode = true;
    }
@@ -527,10 +665,12 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
    std::vector<std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>*
        profiles = nullptr;

+    
    // Searching accross all the packets given during this write
    for (size_t i = 0; i < pkt_count; ++i) {
      auto& original_packet = static_cast<const hsa_barrier_and_packet_t*>(packets)[i];

+      // +Skip kernel dispatch IDs not wanted
      // Skip packets other than kernel dispatch packets.
      if (bit_extract(original_packet.header, HSA_PACKET_HEADER_TYPE,
                      HSA_PACKET_HEADER_TYPE + HSA_PACKET_HEADER_WIDTH_TYPE - 1) !=
@@ -659,6 +799,185 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
    }
    /* Write the transformed packets to the hardware queue.  */
    writer(&transformed_packets[0], transformed_packets.size());
+  } else if (session_id.handle > 0 && pkt_count > 0 && is_att_collection_mode && session) {
+    // att start
+    // Getting Queue Data and Information
+    auto& queue_info = *static_cast<Queue*>(data);
+    std::lock_guard<std::mutex> lk(queue_info.qw_mutex);
+    Agent::AgentInfo* agentInfo = &(hsa_support::GetAgentInfo(queue_info.GetGPUAgent().handle));
+
+    if (agentInfo->getName().substr(0, 4) != "gfx9") {
+      printf("ATT collection is only supported for gfx9 at the moment!\n");
+      exit(1);
+    }
+
+    // Preparing att Packets
+    Packet::packet_t start_packet{};
+    Packet::packet_t stop_packet{};
+    hsa_ven_amd_aqlprofile_profile_t* profile = nullptr;
+
+    if (att_parameters_data.size() > 0 && is_att_collection_mode) {
+      // TODO sauverma: convert att_parameters_data to pass to generateattPackets
+      std::vector<hsa_ven_amd_aqlprofile_parameter_t> att_params;
+      int num_att_counters = 0;
+
+      for (rocprofiler_att_parameter_t& param : att_parameters_data) {
+        att_params.push_back({
+          static_cast<hsa_ven_amd_aqlprofile_parameter_name_t>(int(param.parameter_name)),
+          param.value
+        });
+        num_att_counters += param.parameter_name == ROCPROFILER_ATT_PERFCOUNTER;
+      }
+
+      if (att_counters_names.size() > 0) {
+        MetricsDict* metrics_dict_ = MetricsDict::Create(agentInfo);
+
+        for (const std::string& counter_name : att_counters_names) {
+          const Metric* metric = metrics_dict_->Get(counter_name);
+          const BaseMetric* base = dynamic_cast<const BaseMetric*>(metric);
+          if (!base) {
+            printf("Invalid base metric value: %s\n", counter_name.c_str());
+            exit(1);
+          }
+          std::vector<const counter_t*> counters;
+          base->GetCounters(counters);
+          hsa_ven_amd_aqlprofile_event_t event = counters[0]->event;
+          if (event.block_name != HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ) {
+            printf("Only events from the SQ block can be selected for ATT.");
+            exit(1);
+          }
+          att_params.push_back({
+            static_cast<hsa_ven_amd_aqlprofile_parameter_name_t>(int(ROCPROFILER_ATT_PERFCOUNTER)),
+            event.counter_id | (event.counter_id ? (0xF<<24) : 0)
+          });
+          num_att_counters += 1;
+        }
+
+        hsa_ven_amd_aqlprofile_parameter_t zero_perf = {
+          static_cast<hsa_ven_amd_aqlprofile_parameter_name_t>(int(ROCPROFILER_ATT_PERFCOUNTER)), 0};
+
+        // Fill other perfcounters with 0's
+        for(; num_att_counters<16; num_att_counters++) att_params.push_back(zero_perf);
+      }
+
+      // Get the PM4 Packets using packets_generator
+      profile = Packet::GenerateATTPackets(queue_info.GetCPUAgent(), queue_info.GetGPUAgent(),
+                                            att_params, &start_packet, &stop_packet);
+    }
+
+
+    // Searching across all the packets given during this write
+    for (size_t i = 0; i < pkt_count; ++i) {
+      auto& original_packet = static_cast<const hsa_barrier_and_packet_t*>(packets)[i];
+
+      // Skip packets other than kernel dispatch packets.
+      if (bit_extract(original_packet.header, HSA_PACKET_HEADER_TYPE,
+                      HSA_PACKET_HEADER_TYPE + HSA_PACKET_HEADER_WIDTH_TYPE - 1) !=
+          HSA_PACKET_TYPE_KERNEL_DISPATCH) {
+        transformed_packets.emplace_back(packets_arr[i]);
+        continue;
+      }
+
+      auto& kdispatch = static_cast<const hsa_kernel_dispatch_packet_s*>(packets)[i];
+      uint64_t kernel_object = kdispatch.kernel_object;
+      bool b_profile_this_object = false;
+
+      // Try to match the mangled kernel name with given matches in input.txt
+      try {
+        std::lock_guard<std::mutex> lock(ksymbol_map_lock);
+        assert(ksymbols);
+        const std::string& kernel_name = ksymbols->at(kernel_object);
+
+        // We want to initiate att profiling only if a match exists
+        for(const std::string& kernel_matches : kernel_profile_names) {
+          if (kernel_name.find(kernel_matches) != std::string::npos) {
+            b_profile_this_object = true;
+            break;
+          }
+        }
+        if (!b_profile_this_object) printf("Skipping: %s\n", kernel_name.c_str());
+      } catch (...) {
+        printf("Warning: Unknown name for object %lu\n", kernel_object);
+      }
+
+      // If no match was found or intercept count > maximum desired profiles, skip this kernel.
+      if (!b_profile_this_object || KernelInterceptCount >= MAX_ATT_PROFILES) {
+        printf("Skipping: %lu\n", kernel_object);
+        transformed_packets.emplace_back(packets_arr[i]);
+        continue;
+      }
+      KernelInterceptCount += 1;
+
+      uint32_t writer_id = WRITER_ID.fetch_add(1, std::memory_order_release);
+
+
+      if (att_parameters_data.size() > 0 && is_att_collection_mode && profile) {
+        // Adding start packet and its barrier with a dummy signal
+        hsa_signal_t dummy_signal{};
+        dummy_signal.handle = 0;
+        start_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
+        AddVendorSpecificPacket(&start_packet, &transformed_packets, dummy_signal);
+        CreateBarrierPacket(start_packet.completion_signal, &transformed_packets);
+      }
+
+      auto& packet = transformed_packets.emplace_back(packets_arr[i]);
+      auto& dispatch_packet = reinterpret_cast<hsa_kernel_dispatch_packet_t&>(packet);
+
+      CreateSignal(HSA_AMD_SIGNAL_AMD_GPU_ONLY, &packet.completion_signal);
+      // Adding the dispatch packet newly created signal to the pending signals
+      // list to be processed by the signal interrupt
+      rocprofiler_kernel_properties_t kernel_properties =
+          set_kernel_properties(dispatch_packet, queue_info.GetGPUAgent());
+      if (session && profile) {
+        session->GetAttTracer()->AddPendingSignals(
+            writer_id, dispatch_packet.kernel_object, dispatch_packet.completion_signal, session_id,
+            buffer_id, profile, kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index);
+      } else {
+        session->GetAttTracer()->AddPendingSignals(
+            writer_id, dispatch_packet.kernel_object, dispatch_packet.completion_signal, session_id,
+            buffer_id, nullptr, kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index);
+      }
+
+      // Make a copy of the original packet, adding its signal to a barrier packet
+      if (original_packet.completion_signal.handle) {
+        hsa_barrier_and_packet_t barrier{0};
+        barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
+        Packet::packet_t* __attribute__((__may_alias__)) pkt =
+            (reinterpret_cast<Packet::packet_t*>(&barrier));
+        transformed_packets.emplace_back(*pkt).completion_signal =
+            original_packet.completion_signal;
+      }
+
+      // Adding a barrier packet with the original packet's completion signal.
+      hsa_signal_t interrupt_signal;
+      CreateSignal(0, &interrupt_signal);
+
+      // Adding Stop PM4 Packets
+      if (att_parameters_data.size() > 0 && is_att_collection_mode && profile) {
+        stop_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
+        AddVendorSpecificPacket(&stop_packet, &transformed_packets, interrupt_signal);
+
+        // Added Interrupt Signal with barrier and provided handler for it
+        CreateBarrierPacket(interrupt_signal, &transformed_packets);
+      } else {
+        hsa_barrier_and_packet_t barrier{0};
+        barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
+        barrier.completion_signal = interrupt_signal;
+        Packet::packet_t* __attribute__((__may_alias__)) pkt =
+            (reinterpret_cast<Packet::packet_t*>(&barrier));
+        transformed_packets.emplace_back(*pkt);
+      }
+
+      // Creating Async Handler to be called every time the interrupt signal is
+      // marked complete
+      signalAsyncHandlerATT(
+          interrupt_signal,
+          new queue_info_session_t{queue_info.GetGPUAgent(), session_id, queue_info.GetQueueID(),
+                                   writer_id, interrupt_signal});
+    }
+    /* Write the transformed packets to the hardware queue.  */
+    writer(&transformed_packets[0], transformed_packets.size());
+    // ATT end
  } else {
    /* Write the original packets to the hardware queue if no profiling session
     * is active  */
@@ -43,7 +43,8 @@ GenericBuffer::GenericBuffer(rocprofiler_session_id_t session_id, rocprofiler_bu
    // pointer moves to the other buffer. Each buffer should be large enough to
    // hold at least 2 activity records, as record pairs may be written when
    // external correlation ids are used.
-    const size_t allocation_size = 2 * std::max(2 * sizeof(rocprofiler_record_header_t), buffer_size);
+    const size_t allocation_size =
+        2 * std::max(2 * sizeof(rocprofiler_record_header_t), buffer_size);
    pool_begin_ = nullptr;
    AllocateMemory(&pool_begin_, allocation_size);
    assert(pool_begin_ != nullptr && "pool allocator failed");
@@ -213,6 +214,12 @@ bool GetNextRecord(const rocprofiler_record_header_t* record,
      // break;
      *next = reinterpret_cast<const rocprofiler_record_header_t*>(tracer_record + 1);
    }
+    case ROCPROFILER_ATT_TRACER_RECORD: {
+      const rocprofiler_record_att_tracer_t* att_tracer_record =
+          reinterpret_cast<const rocprofiler_record_att_tracer_t*>(record);
+      *next = reinterpret_cast<const rocprofiler_record_header_t*>(att_tracer_record + 1);
+      break;
+    }
    default:
      const rocprofiler_record_tracer_t* tracer_record =
          reinterpret_cast<const rocprofiler_record_tracer_t*>(record);
@@ -0,0 +1,57 @@
+/* Copyright (c) 2022 Advanced Micro Devices, Inc.
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights
+ to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ copies of the Software, and to permit persons to whom the Software is
+ furnished to do so, subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE. */
+
+#include "att.h"
+#include <cassert>
+
+namespace rocmtools {
+
+namespace att {
+
+AttTracer::AttTracer(rocprofiler_buffer_id_t buffer_id, rocprofiler_filter_id_t filter_id,
+                       rocprofiler_session_id_t session_id)
+    : buffer_id_(buffer_id), filter_id_(filter_id), session_id_(session_id) {}
+AttTracer::~AttTracer() {}
+
+void AttTracer::AddPendingSignals(uint32_t writer_id, uint64_t kernel_object,
+                                   const hsa_signal_t& completion_signal,
+                                   rocprofiler_session_id_t session_id,
+                                   rocprofiler_buffer_id_t buffer_id,
+                                   hsa_ven_amd_aqlprofile_profile_t* profile,
+                                   rocprofiler_kernel_properties_t kernel_properties,
+                                   uint32_t thread_id, uint64_t queue_index) {
+  std::lock_guard<std::mutex> lock(sessions_pending_signals_lock_);
+  if (sessions_pending_signals_.find(writer_id) == sessions_pending_signals_.end())
+    sessions_pending_signals_.emplace(writer_id, std::vector<att_pending_signal_t>());
+  sessions_pending_signals_.at(writer_id).emplace_back(
+      att_pending_signal_t{kernel_object, completion_signal, session_id_, buffer_id, profile,
+                            kernel_properties, thread_id, queue_index});
+}
+
+const std::vector<att_pending_signal_t>& AttTracer::GetPendingSignals(uint32_t writer_id) {
+  std::lock_guard<std::mutex> lock(sessions_pending_signals_lock_);
+  assert(sessions_pending_signals_.find(writer_id) != sessions_pending_signals_.end() &&
+         "writer_id is not found in the pending_signals");
+  return sessions_pending_signals_.at(writer_id);
+}
+
+}  // namespace att
+
+}  // namespace rocmtools
@@ -0,0 +1,76 @@
+/* Copyright (c) 2022 Advanced Micro Devices, Inc.
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights
+ to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ copies of the Software, and to permit persons to whom the Software is
+ furnished to do so, subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE. */
+
+#ifndef SRC_CORE_SESSION_ATT_ATT_H_
+#define SRC_CORE_SESSION_ATT_ATT_H_
+
+#include <hsa/hsa_ven_amd_aqlprofile.h>
+
+#include <map>
+#include <mutex>
+#include <string>
+#include <vector>
+
+#include "inc/rocprofiler.h"
+
+namespace rocmtools {
+
+typedef struct {
+  uint64_t kernel_descriptor;
+  hsa_signal_t signal;
+  rocprofiler_session_id_t session_id;
+  rocprofiler_buffer_id_t buffer_id;
+  hsa_ven_amd_aqlprofile_profile_t* profile;
+  rocprofiler_kernel_properties_t kernel_properties;
+  uint32_t thread_id;
+  uint64_t queue_index;
+} att_pending_signal_t;
+
+namespace att {
+
+class AttTracer {
+ public:
+  AttTracer(rocprofiler_buffer_id_t buffer_id, rocprofiler_filter_id_t filter_id,
+             rocprofiler_session_id_t session_id);
+  ~AttTracer();
+
+  void AddPendingSignals(uint32_t writer_id, uint64_t kernel_object,
+                         const hsa_signal_t& completion_signal, rocprofiler_session_id_t session_id,
+                         rocprofiler_buffer_id_t buffer_id, hsa_ven_amd_aqlprofile_profile_t* profile,
+                         rocprofiler_kernel_properties_t kernel_properties, uint32_t thread_id,
+                         uint64_t queue_index);
+
+  const std::vector<att_pending_signal_t>& GetPendingSignals(uint32_t writer_id);
+
+ private:
+  rocprofiler_buffer_id_t buffer_id_;
+  rocprofiler_filter_id_t filter_id_;
+  rocprofiler_session_id_t session_id_;
+
+  std::mutex sessions_pending_signals_lock_;
+  std::map<uint32_t, std::vector<att_pending_signal_t>> sessions_pending_signals_;
+};
+
+}  // namespace att
+
+}  // namespace rocmtools
+
+
+#endif  // SRC_CORE_SESSION_ATT_ATT_H_
@@ -35,12 +35,25 @@ Filter::Filter(rocprofiler_filter_id_t id, rocprofiler_filter_kind_t filter_kind
    }
    case ROCPROFILER_COUNTERS_COLLECTION: {
      profiler_counter_names_.clear();
-      for (uint32_t j = 0; j < data_count; j++)
+      for (uint32_t j = 0; j < data_count; j++) {
        profiler_counter_names_.emplace_back(filter_data.counters_names[j]);
+      }
      break;
    }
-    case ROCPROFILER_PC_SAMPLING_COLLECTION:
-    case ROCPROFILER_ATT_TRACE: {
+    case ROCPROFILER_PC_SAMPLING_COLLECTION:{
+      break;
+    }
+    case ROCPROFILER_ATT_TRACE_COLLECTION: {
+      att_parameters_.clear();
+      profiler_counter_names_.clear();
+
+      for (uint32_t j = 0; j < data_count; j++) {
+        if (filter_data.att_parameters[j].parameter_name != ROCPROFILER_ATT_PERFCOUNTER_NAME) {
+          att_parameters_.emplace_back(filter_data.att_parameters[j]);
+        } else {
+          profiler_counter_names_.emplace_back(filter_data.att_parameters[j].counter_name);
+        }
+      }
      break;
    }
    case ROCPROFILER_SPM_COLLECTION: {
@@ -49,8 +62,9 @@ Filter::Filter(rocprofiler_filter_id_t id, rocprofiler_filter_kind_t filter_kind
    }
    case ROCPROFILER_API_TRACE: {
      tracer_apis_.clear();
-      for (uint32_t j = 0; j < data_count; j++)
-        tracer_apis_.emplace_back(filter_data.trace_apis[j]);
+      for (uint32_t j = 0; j < data_count; j++){
+          tracer_apis_.emplace_back(filter_data.trace_apis[j]);
+      }     
      break;
    }
    default: {
@@ -73,7 +87,7 @@ rocprofiler_filter_kind_t Filter::GetKind() { return kind_; }

 std::mutex counter_data_lock;
 std::vector<std::string> Filter::GetCounterData() {
-  if (kind_ == ROCPROFILER_COUNTERS_COLLECTION) {
+  if (kind_ == ROCPROFILER_COUNTERS_COLLECTION || kind_ == ROCPROFILER_ATT_TRACE_COLLECTION) {
    std::lock_guard<std::mutex> lock(counter_data_lock);
    return profiler_counter_names_;
  }
@@ -90,6 +104,16 @@ std::vector<rocprofiler_tracer_activity_domain_t> Filter::GetTraceData() {
      "Error: ROCMtools filter specified is not supported for "
      "profiler mode!\n");
 }
+
+std::vector<rocprofiler_att_parameter_t> Filter::GetAttParametersData() {
+  if (kind_ == ROCPROFILER_ATT_TRACE_COLLECTION) {
+    return att_parameters_;
+  }
+  fatal(
+      "Error: ROCMtools filter specified is not supported for "
+      "ATT tracing mode!\n");
+}
+
 rocprofiler_spm_parameter_t* Filter::GetSpmParameterData() {
  if (kind_ == ROCPROFILER_SPM_COLLECTION) {
    return spm_parameter_;
@@ -143,7 +167,8 @@ void Filter::SetProperty(rocprofiler_filter_property_t property) {
    }
    case ROCPROFILER_FILTER_KERNEL_NAMES: {
      if (kind_ == ROCPROFILER_COUNTERS_COLLECTION ||
-          kind_ == ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION) {
+          kind_ == ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION ||
+          kind_ == ROCPROFILER_ATT_TRACE_COLLECTION) {
        kernel_names_.clear();
        for (uint32_t j = 0; j < property.data_count; j++)
          kernel_names_.emplace_back(property.name_regex[j]);
@@ -166,23 +191,29 @@ std::variant<std::vector<std::string>, uint32_t*> Filter::GetProperty(
  switch (kind) {
    case ROCPROFILER_FILTER_GPU_NAME: {
      property = agent_names_;
+      break;
    }
    case ROCPROFILER_FILTER_RANGE: {
      property = static_cast<uint32_t*>(dispatch_range_);
+      break;
    }
    case ROCPROFILER_FILTER_KERNEL_NAMES: {
      property = kernel_names_;
+      break;
    }
    case ROCPROFILER_FILTER_HSA_TRACER_API_FUNCTIONS: {
      property = hsa_tracer_api_calls_;
+      break;
    }
    case ROCPROFILER_FILTER_HIP_TRACER_API_FUNCTIONS: {
      property = hip_tracer_api_calls_;
+      break;
    }
    default:
      fatal(
          "Error: ROCMtools filter specified is not supported for the given "
          "kind!");
+      break;
  }
  return property;
 }
@@ -47,7 +47,7 @@ class Filter {

  std::vector<std::string> GetCounterData();
  std::vector<rocprofiler_tracer_activity_domain_t> GetTraceData();
-
+  std::vector<rocprofiler_att_parameter_t> GetAttParametersData();
  void SetCallback(rocprofiler_sync_callback_t& callback);
  rocprofiler_sync_callback_t& GetCallback();

@@ -71,6 +71,7 @@ class Filter {
  std::vector<std::string> profiler_counter_names_;              // Counter Names to collect
  std::vector<rocprofiler_tracer_activity_domain_t> tracer_apis_;  // ROCTX/HIP/HSA API
  rocprofiler_spm_parameter_t* spm_parameter_;                     // spm parameter
+  std::vector<rocprofiler_att_parameter_t> att_parameters_;  // ATT Parameters

  rocprofiler_sync_callback_t callback_;
 };
@@ -54,6 +54,10 @@ Session::~Session() {
  //   delete tracer_;
  //   tracer_started_.exchange(false, std::memory_order_release);
  // }
+  if (att_tracer_started_.load(std::memory_order_release)) {
+    delete att_tracer_;
+    att_tracer_started_.exchange(false, std::memory_order_release);
+  }
  // {
  //   std::lock_guard<std::mutex> lock(filters_lock_);
  //   buffers_.clear();
@@ -99,6 +103,13 @@ void Session::Start() {
          GetFilter(GetFilterIdWithKind(ROCPROFILER_COUNTERS_COLLECTION))->GetId(), session_id_);
      profiler_started_.exchange(true, std::memory_order_release);
    }
+    if (FindFilterWithKind(ROCPROFILER_ATT_TRACE_COLLECTION)) {
+      if (att_tracer_started_.load(std::memory_order_release)) delete att_tracer_;
+      att_tracer_ = new att::AttTracer(
+          GetFilter(GetFilterIdWithKind(ROCPROFILER_ATT_TRACE_COLLECTION))->GetBufferId(),
+          GetFilter(GetFilterIdWithKind(ROCPROFILER_ATT_TRACE_COLLECTION))->GetId(), session_id_);
+      att_tracer_started_.exchange(true, std::memory_order_release);
+    }

    if (FindFilterWithKind(ROCPROFILER_SPM_COLLECTION)) {
      if (spm_started_.load(std::memory_order_release)) delete spmcounter_;
@@ -176,6 +187,7 @@ rocprofiler_session_id_t Session::GetId() { return session_id_; }
 bool Session::IsActive() { return is_active_; }

 profiler::Profiler* Session::GetProfiler() { return profiler_; }
+att::AttTracer* Session::GetAttTracer() { return att_tracer_; }
 tracer::Tracer* Session::GetTracer() { return tracer_; }
 spm::SpmCounters* Session::GetSpmCounter() { return spmcounter_; }
 pc_sampler::PCSampler* Session::GetPCSampler() { return pc_sampler_; }
@@ -37,6 +37,7 @@
 #include "src/core/session/filter.h"
 #include "profiler/profiler.h"
 #include "tracer/tracer.h"
+#include "att/att.h"
 #include "spm/spm.h"
 #include "src/pcsampler/session/pc_sampler.h"

@@ -58,6 +59,7 @@ class Session {

  profiler::Profiler* GetProfiler();
  tracer::Tracer* GetTracer();
+  att::AttTracer* GetAttTracer();
  spm::SpmCounters* GetSpmCounter();
  pc_sampler::PCSampler* GetPCSampler();

@@ -104,6 +106,8 @@ class Session {

  std::atomic<bool> profiler_started_{false};
  std::atomic<bool> tracer_started_{false};
+  std::atomic<bool> att_tracer_started_{false};
+  att::AttTracer* att_tracer_;
  std::atomic<bool> spm_started_{false};

  profiler::Profiler* profiler_;