GFX8 API

Change-Id: I9d0c430e4199f043226c8897f3320a7973cbdeda

runtime/hsa-ext-aql-profile/CMakeLists.txt

@@ -25,4 +25,4 @@ add_subdirectory ( ${TEST_DIR} ${PROJECT_BINARY_DIR}/test )
 #
 # Style format
 #
-execute_process ( COMMAND sh -xc "/usr/bin/find ${ROOT_DIR} -name '*.cpp' -o -name '*.hpp' -o -name '*.h' -exec /usr/bin/clang-format -i -style=file \{\} \;" )
+execute_process ( COMMAND sh -xc "/usr/bin/find ${PROJ_DIR} ${TEST_DIR} ${API_DIR} -name '*.cpp' -o -name '*.hpp' -o -name '*.h' -exec /usr/bin/clang-format -i -style=file \{\} \;" )

runtime/hsa-ext-aql-profile/inc/hsa_ext_amd_aql_profile.h

@@ -203,11 +203,11 @@ hsa_status_t hsa_ext_amd_aql_profile_stop(
     hsa_ext_amd_aql_pm4_packet_t* aql_stop_packet);    // [out] profile stop AQL packet
 
 // Legacy PM4 profiling packet size
-const unsigned HSA_EXT_AQL_PROFILE_LEGACY_PM4_PACKET_SIZE = 64;
-// Converting of the profiling AQL packet to PM4 packet, GFX8 support
+const unsigned HSA_EXT_AQL_PROFILE_LEGACY_PM4_PACKET_SIZE = 192;
+// GFX8 support, converting the profiling AQL packet to PM4 packet blob
 hsa_status_t hsa_ext_amd_aql_profile_legacy_get_pm4(
-    const hsa_ext_amd_aql_pm4_packet_t* aql_packet,  // AQL packet
-    void* pm4);                                      // PM4 packet blob
+    const hsa_ext_amd_aql_pm4_packet_t* aql_packet,  // [in] AQL packet
+    void* data);                                     // [out] PM4 packet blob
 
 //
 // Get profile info:

runtime/hsa-ext-aql-profile/src/aqlprofile/CMakeLists.txt

@@ -10,6 +10,7 @@ include_directories ( $ENV{ROCR_INC_DIR} )
 include_directories ( ${PROJ_DIR}/perfcounter )
 include_directories ( ${PROJ_DIR}/threadtrace )
 include_directories ( ${PROJ_DIR}/commandwriter )
+include_directories ( ${HSA_RUNTIME_OSC_DIR} )
 include_directories ( ${API_DIR} )
 
 #

runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile.cpp

@@ -258,10 +258,17 @@ PUBLIC_API hsa_status_t hsa_ext_amd_aql_profile_stop(
   return HSA_STATUS_SUCCESS;
 }
 
-// Converting of the profiling AQL packet to PM4 packet, GFX8 support
+// GFX8 support, converting of the profiling AQL packet to PM4 packet blob
 PUBLIC_API hsa_status_t hsa_ext_amd_aql_profile_legacy_get_pm4(
-    const aql_profile::packet_t* aql_packet, void* pm4) {
-  return HSA_STATUS_ERROR;
+    const aql_profile::packet_t* aql_packet, void* data) {
+  // Populate GFX8 pm4 packet blob
+  // Adding HSA barrier acquire packet
+  data = aql_profile::legacyAqlAcquire(aql_packet, data);
+  // Adding PM4 command packet
+  data = aql_profile::legacyPm4(aql_packet, data);
+  // Adding HSA barrier release packet
+  data = aql_profile::legacyAqlRelease(aql_packet, data);
+  return HSA_STATUS_SUCCESS;
 }
 
 // Method for getting the profile info

runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile.h

@@ -17,7 +17,10 @@ typedef hsa_ext_amd_aql_pm4_packet_t packet_t;
 typedef hsa_ext_amd_aql_profile_event_t event_t;
 
 void populateAql(void* cmdBuffer, uint32_t cmdSz, pm4_profile::CommandWriter* cmdWriter,
-                 packet_t* aqlPkt);
+                 packet_t* aql_packet);
+void* legacyAqlAcquire(const packet_t* aql_packet, void* data);
+void* legacyAqlRelease(const packet_t* aql_packet, void* data);
+void* legacyPm4(const packet_t* aql_packet, void* data);
 }
 
 #endif  // _AQL_PROFILE_H_

runtime/hsa-ext-aql-profile/src/aqlprofile/populate_aql.cpp

@@ -4,9 +4,92 @@
 #include "aql_profile.h"
 #include "cmdwriter.h"
 #include "amd_aql_pm4_ib_packet.h"
+#include "core/inc/amd_gpu_pm4.h"
 
 namespace aql_profile {
 
+typedef uint16_t aql_packet_header_t;
+
+void * legacyAqlAcquire(const packet_t* aql_packet, void * data) {
+  hsa_barrier_and_packet_t * aql_barrier =
+    reinterpret_cast<hsa_barrier_and_packet_t*>(data);
+  memset(aql_barrier, 0 , sizeof(hsa_barrier_and_packet_t));
+  const aql_packet_header_t aql_header_type =
+    HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
+  const aql_packet_header_t aql_header_barrier =
+    1ul << HSA_PACKET_HEADER_BARRIER;
+  aql_barrier->header |= aql_header_type;
+  aql_barrier->header |= aql_header_barrier;
+  return data + sizeof(hsa_barrier_and_packet_t);
+}
+
+void * legacyAqlRelease(const packet_t* aql_packet, void * data) {
+  hsa_barrier_and_packet_t * aql_barrier =
+    reinterpret_cast<hsa_barrier_and_packet_t*>(data);
+  memset(aql_barrier, 0 , sizeof(hsa_barrier_and_packet_t));
+  const aql_packet_header_t aql_header_type =
+    HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
+  const aql_packet_header_t aql_header_barrier =
+    1ul << HSA_PACKET_HEADER_BARRIER;
+  const aql_packet_header_t aql_header_release =
+    HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE;
+  aql_barrier->header |= aql_header_type;
+  aql_barrier->header |= aql_header_barrier;
+  aql_barrier->header |= aql_header_release;
+  aql_barrier->completion_signal = aql_packet->completion_signal;
+  return data + sizeof(hsa_barrier_and_packet_t);
+}
+
+void * legacyPm4(const packet_t* aql_packet, void *data) {
+  constexpr uint32_t major_version = 8;
+  constexpr uint32_t slot_size_b = 0x40;
+  constexpr uint32_t ib_jump_size_dw = 4;
+  constexpr uint32_t slot_size_dw = uint32_t(slot_size_b / sizeof(uint32_t));
+
+  const amd_aql_pm4_ib_packet_t* aql_pm4_ib =
+    reinterpret_cast<const amd_aql_pm4_ib_packet_t*>(aql_packet);
+  uint32_t * slot_data = (uint32_t*)data;
+  // Construct a set of PM4 to fit inside the AQL packet slot.
+  uint32_t slot_dw_idx = 0;
+
+  // Construct a no-op command to pad the queue slot.
+  constexpr uint32_t rel_mem_size_dw = 7;
+  constexpr uint32_t nop_pad_size_dw = slot_size_dw - (ib_jump_size_dw + rel_mem_size_dw);
+
+  uint32_t* nop_pad = &slot_data[slot_dw_idx];
+  slot_dw_idx += nop_pad_size_dw;
+
+  nop_pad[0] = PM4_HDR(PM4_HDR_IT_OPCODE_NOP, nop_pad_size_dw, major_version);
+
+  for (int i = 1; i < nop_pad_size_dw; ++i) {
+    nop_pad[i] = 0;
+  }
+
+  // Copy in command to execute the IB.
+  assert(slot_dw_idx + ib_jump_size_dw <= slot_size_dw && "PM4 exceeded queue slot size");
+  uint32_t* ib_jump = &slot_data[slot_dw_idx];
+  slot_dw_idx += ib_jump_size_dw;
+
+  memcpy(ib_jump, aql_pm4_ib->pm4_ib_command, sizeof(aql_pm4_ib->pm4_ib_command));
+
+  // Construct a command to advance the read index and invalidate the packet
+  // header. This must be the last command since this releases the queue slot
+  // for writing.
+  assert(slot_dw_idx + rel_mem_size_dw <= slot_size_dw && "PM4 exceeded queue slot size");
+  uint32_t* rel_mem = &slot_data[slot_dw_idx];
+
+  rel_mem[0] =
+      PM4_HDR(PM4_HDR_IT_OPCODE_RELEASE_MEM, rel_mem_size_dw, major_version);
+  rel_mem[1] = PM4_RELEASE_MEM_DW1_EVENT_INDEX(PM4_RELEASE_MEM_EVENT_INDEX_AQL);
+  rel_mem[2] = 0;
+  rel_mem[3] = 0;
+  rel_mem[4] = 0;
+  rel_mem[5] = 0;
+  rel_mem[6] = 0;
+
+  return data + slot_size_b;
+}
+
 void populateAql(uint32_t* ib_packet, packet_t* aql_packet) {
   // Populate relevant fields of Aql pkt
   // Size of IB pkt is four DWords
@@ -32,10 +115,10 @@ void populateAql(uint32_t* ib_packet, packet_t* aql_packet) {
 }
 
 void populateAql(void* cmd_buffer, uint32_t cmd_size,
-                 pm4_profile::CommandWriter* cmd_writer, packet_t* ppt_packet) {
+                 pm4_profile::CommandWriter* cmd_writer, packet_t* aql_packet) {
   pm4_profile::DefaultCmdBuf ib_buffer;
   cmd_writer->BuildIndirectBufferCmd(&ib_buffer, cmd_buffer, (size_t)cmd_size);
-  uint32_t* ib_cmds = (uint32_t*)ib_buffer.Base();
-  populateAql(ib_cmds, ppt_packet);
+  uint32_t* ib_packet = (uint32_t*)ib_buffer.Base();
+  populateAql(ib_packet, aql_packet);
 }
 }

runtime/hsa-ext-aql-profile/test/CMakeLists.txt

@@ -3,9 +3,6 @@
 #
 include_directories ( $ENV{ROCR_INC_DIR} )
 include_directories ( ${API_DIR} )
-include_directories ( ${PROJ_DIR}/cmdwriter )
-include_directories ( ${PROJ_DIR}/perfcounter )
-include_directories ( ${PROJ_DIR}/threadtrace )
 include_directories ( ${PROJ_DIR}/aqlprofile )
 include_directories ( ${TEST_DIR}/common )
 include_directories ( ${TEST_DIR}/ctrl )
@@ -45,4 +42,4 @@ set ( SRC_LIST ${SRC_LIST} ${TEST_DIR}/ctrl/test_hsa.cpp )
 set ( LIB_LIST ${TEST_LIBS} ${COMMON_LIB} ${CORE_UTILS_LIB} ${ROCR_LIB} ${TARGET_LIB} )
 set ( EXE_NAME "ctrl" )
 add_executable ( ${EXE_NAME} ${SRC_LIST} )
-target_link_libraries( ${EXE_NAME} ${LIB_LIST} c stdc++ dl pthread rt )
+target_link_libraries( ${EXE_NAME} ${LIB_LIST} c stdc++ dl pthread rt atomic )

runtime/hsa-ext-aql-profile/test/ctrl/test_hsa.cpp

@@ -144,10 +144,11 @@ bool TestHSA::run() {
   // Initialize the dispatch packet.
   hsa_kernel_dispatch_packet_t aql;
   memset(&aql, 0, sizeof(aql));
-  // Set the packet's type, acquire and release fences
+  // Set the packet's type, barrier bit, acquire and release fences
   aql.header = HSA_PACKET_TYPE_KERNEL_DISPATCH;
-  aql.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE;
-  aql.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE;
+  aql.header |= 1ul << HSA_PACKET_HEADER_BARRIER;
+  aql.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE;
+  aql.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE;
   // Populate Aql packet with default values
   aql.setup = 1;
   aql.grid_size_x = work_grid_size;

runtime/hsa-ext-aql-profile/test/ctrl/test_pmgr.cpp

@@ -30,7 +30,7 @@ OF THE POSSIBILITY OF SUCH DAMAGE.
 
 #include "test_pmgr.h"
 
-bool TestPMgr::addPacket(const packet_t* packet) {
+bool TestPMgr::addPacketGfx9(const packet_t* packet) {
   packet_t aql_packet = *packet;
 
   // Compute the write index of queue and copy Aql packet into it
@@ -57,6 +57,36 @@ bool TestPMgr::addPacket(const packet_t* packet) {
   return true;
 }
 
+bool TestPMgr::addPacketGfx8(const packet_t* packet) {
+  // Create Legacy PM4 data
+  const hsa_ext_amd_aql_pm4_packet_t * aql_packet =
+    (const hsa_ext_amd_aql_pm4_packet_t *) packet;
+  slot_pm4_s data;
+  hsa_ext_amd_aql_profile_legacy_get_pm4(aql_packet, (void*)data.words);
+
+  // Compute the write index of queue and copy Aql packet into it
+  uint64_t que_idx = hsa_queue_load_write_index_relaxed(getQueue());
+  const uint32_t mask = getQueue()->size - 1;
+
+  // Copy Aql packet into queue buffer
+  packet_t* ptr = ((packet_t*)(getQueue()->base_address)) + (que_idx & mask);
+  slot_pm4_t * slot_pm4 = (slot_pm4_t*)ptr;
+  slot_pm4->store(data, std::memory_order_relaxed);
+
+  // Increment the write index and ring the doorbell to dispatch the kernel.
+  hsa_queue_store_write_index_relaxed(getQueue(), (que_idx + SLOT_PM4_SIZE_AQLP));
+  hsa_signal_store_relaxed(getQueue()->doorbell_signal, que_idx + SLOT_PM4_SIZE_AQLP - 1);
+
+  return true;
+}
+
+bool TestPMgr::addPacket(const packet_t* packet) {
+  const char * agent_name = getAgentInfo()->name;
+  return (strncmp(agent_name, "gfx8", 4) == 0) ?
+    addPacketGfx8(packet) :
+    addPacketGfx9(packet);
+}
+
 bool TestPMgr::run() {
   // Build Aql Pkts
   const bool active = buildPackets();

runtime/hsa-ext-aql-profile/test/ctrl/test_pmgr.h

@@ -28,6 +28,8 @@ OF THE POSSIBILITY OF SUCH DAMAGE.
 #ifndef _TEST_SMGR_H_
 #define _TEST_SMGR_H_
 
+#include <atomic>
+
 #include "test_aql.h"
 #include "amd_aql_pm4_ib_packet.h"
 
@@ -35,9 +37,8 @@ OF THE POSSIBILITY OF SUCH DAMAGE.
 class TestPMgr : public TestAql {
  public:
   typedef amd_aql_pm4_ib_packet_t packet_t;
-
- private:
-  bool addPacket(const packet_t* packet);
+  TestPMgr(TestAql* t);
+  bool run();
 
  protected:
   packet_t prePacket;
@@ -49,9 +50,19 @@ class TestPMgr : public TestAql {
   virtual bool dumpData() { return false; }
   virtual bool initialize(int argc, char** argv);
 
- public:
-  TestPMgr(TestAql* t);
-  bool run();
+ private:
+  enum {
+    SLOT_PM4_SIZE_DW = HSA_EXT_AQL_PROFILE_LEGACY_PM4_PACKET_SIZE / sizeof(uint32_t),
+    SLOT_PM4_SIZE_AQLP = HSA_EXT_AQL_PROFILE_LEGACY_PM4_PACKET_SIZE / 64
+  };
+  struct slot_pm4_s {
+    uint32_t words[SLOT_PM4_SIZE_DW];
+  };
+  typedef std::atomic<slot_pm4_s> slot_pm4_t;
+
+  bool addPacket(const packet_t* packet);
+  bool addPacketGfx8(const packet_t* packet);
+  bool addPacketGfx9(const packet_t* packet);
 };
 
 #endif  // _TEST_SMGR_H_