From 35b376e2ee943331d5ec7b4a54260456779f2ad5 Mon Sep 17 00:00:00 2001 From: Evgeny Date: Tue, 13 Jun 2017 16:27:18 -0500 Subject: [PATCH] GFX8 API Change-Id: I9d0c430e4199f043226c8897f3320a7973cbdeda --- runtime/hsa-ext-aql-profile/CMakeLists.txt | 2 +- .../inc/hsa_ext_amd_aql_profile.h | 8 +- .../src/aqlprofile/CMakeLists.txt | 1 + .../aqlprofile}/amd_aql_pm4_ib_packet.h | 0 .../src/aqlprofile/aql_profile.cpp | 13 ++- .../src/aqlprofile/aql_profile.h | 5 +- .../src/aqlprofile/populate_aql.cpp | 89 ++++++++++++++++++- .../hsa-ext-aql-profile/test/CMakeLists.txt | 5 +- .../test/ctrl/test_hsa.cpp | 7 +- .../test/ctrl/test_pmgr.cpp | 32 ++++++- .../hsa-ext-aql-profile/test/ctrl/test_pmgr.h | 23 +++-- 11 files changed, 159 insertions(+), 26 deletions(-) rename runtime/hsa-ext-aql-profile/{inc => src/aqlprofile}/amd_aql_pm4_ib_packet.h (100%) diff --git a/runtime/hsa-ext-aql-profile/CMakeLists.txt b/runtime/hsa-ext-aql-profile/CMakeLists.txt index ceb476dca3..d5c137a5c3 100644 --- a/runtime/hsa-ext-aql-profile/CMakeLists.txt +++ b/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 \{\} \;" ) diff --git a/runtime/hsa-ext-aql-profile/inc/hsa_ext_amd_aql_profile.h b/runtime/hsa-ext-aql-profile/inc/hsa_ext_amd_aql_profile.h index 478b3cf1fb..bd3e072a66 100644 --- a/runtime/hsa-ext-aql-profile/inc/hsa_ext_amd_aql_profile.h +++ b/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: diff --git a/runtime/hsa-ext-aql-profile/src/aqlprofile/CMakeLists.txt b/runtime/hsa-ext-aql-profile/src/aqlprofile/CMakeLists.txt index 7bb981c8ae..f70b3cb169 100644 --- a/runtime/hsa-ext-aql-profile/src/aqlprofile/CMakeLists.txt +++ b/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} ) # diff --git a/runtime/hsa-ext-aql-profile/inc/amd_aql_pm4_ib_packet.h b/runtime/hsa-ext-aql-profile/src/aqlprofile/amd_aql_pm4_ib_packet.h similarity index 100% rename from runtime/hsa-ext-aql-profile/inc/amd_aql_pm4_ib_packet.h rename to runtime/hsa-ext-aql-profile/src/aqlprofile/amd_aql_pm4_ib_packet.h diff --git a/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile.cpp b/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile.cpp index f06fb161b3..7857760591 100644 --- a/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile.cpp +++ b/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 diff --git a/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile.h b/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile.h index b51d754001..f25b8e055b 100644 --- a/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile.h +++ b/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_ diff --git a/runtime/hsa-ext-aql-profile/src/aqlprofile/populate_aql.cpp b/runtime/hsa-ext-aql-profile/src/aqlprofile/populate_aql.cpp index aa3da9f463..c00eff2d92 100644 --- a/runtime/hsa-ext-aql-profile/src/aqlprofile/populate_aql.cpp +++ b/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(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(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(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); } } diff --git a/runtime/hsa-ext-aql-profile/test/CMakeLists.txt b/runtime/hsa-ext-aql-profile/test/CMakeLists.txt index f310202398..9714907504 100644 --- a/runtime/hsa-ext-aql-profile/test/CMakeLists.txt +++ b/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 ) diff --git a/runtime/hsa-ext-aql-profile/test/ctrl/test_hsa.cpp b/runtime/hsa-ext-aql-profile/test/ctrl/test_hsa.cpp index 06c589b70d..c45c81c468 100644 --- a/runtime/hsa-ext-aql-profile/test/ctrl/test_hsa.cpp +++ b/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; diff --git a/runtime/hsa-ext-aql-profile/test/ctrl/test_pmgr.cpp b/runtime/hsa-ext-aql-profile/test/ctrl/test_pmgr.cpp index 33c071bc17..e51dd9f31e 100644 --- a/runtime/hsa-ext-aql-profile/test/ctrl/test_pmgr.cpp +++ b/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(); diff --git a/runtime/hsa-ext-aql-profile/test/ctrl/test_pmgr.h b/runtime/hsa-ext-aql-profile/test/ctrl/test_pmgr.h index 274fa9ab47..62a67e7d6a 100644 --- a/runtime/hsa-ext-aql-profile/test/ctrl/test_pmgr.h +++ b/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 + #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_t; + + bool addPacket(const packet_t* packet); + bool addPacketGfx8(const packet_t* packet); + bool addPacketGfx9(const packet_t* packet); }; #endif // _TEST_SMGR_H_