From e7af8fb99ac5dd0335dc17f1b5b580ca682f689d Mon Sep 17 00:00:00 2001 From: "Shweta.Khatri" Date: Thu, 18 Apr 2024 15:39:45 -0400 Subject: [PATCH] Added new ROCr Trap Handler Test Intentionally trigger sw exceptions to verify GPU can handle abnormal scenarios Change-Id: Ie80aa21883a912834ce6d917ecbb21ff6b3145f5 Signed-off-by: Chris Freehill [ROCm/ROCR-Runtime commit: 6f6f02f3522da9801d8515e7ed8edb8579d0252a] --- .../rocrtst/suites/functional/trap_handler.cc | 381 ++++++++++++++++++ .../rocrtst/suites/functional/trap_handler.h | 75 ++++ .../rocrtst/suites/test_common/CMakeLists.txt | 5 + .../kernels/trap_handler_kernels.cl | 76 ++++ .../rocrtst/suites/test_common/main.cc | 13 + 5 files changed, 550 insertions(+) create mode 100644 projects/rocr-runtime/rocrtst/suites/functional/trap_handler.cc create mode 100644 projects/rocr-runtime/rocrtst/suites/functional/trap_handler.h create mode 100644 projects/rocr-runtime/rocrtst/suites/test_common/kernels/trap_handler_kernels.cl diff --git a/projects/rocr-runtime/rocrtst/suites/functional/trap_handler.cc b/projects/rocr-runtime/rocrtst/suites/functional/trap_handler.cc new file mode 100644 index 0000000000..587e489433 --- /dev/null +++ b/projects/rocr-runtime/rocrtst/suites/functional/trap_handler.cc @@ -0,0 +1,381 @@ +/* + * ============================================================================= + * ROC Runtime Conformance Release License + * ============================================================================= + * The University of Illinois/NCSA + * Open Source License (NCSA) + * + * Copyright (c) 2024, Advanced Micro Devices, Inc. + * All rights reserved. + * + * Developed by: + * + * AMD Research and AMD ROC Software Development + * + * Advanced Micro Devices, Inc. + * + * www.amd.com + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal with 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: + * + * - Redistributions of source code must retain the above copyright notice, + * this list of conditions and the following disclaimers. + * - Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimers in + * the documentation and/or other materials provided with the distribution. + * - Neither the names of , + * nor the names of its contributors may be used to endorse or promote + * products derived from this Software without specific prior written + * permission. + * + * 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 CONTRIBUTORS 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 WITH THE SOFTWARE. + * + */ + +#include "suites/functional/trap_handler.h" +#include "common/base_rocr_utils.h" +#include "common/common.h" +#include "common/os.h" +#include "common/helper_funcs.h" +#include "gtest/gtest.h" +#include "hsa/hsa.h" +#include "hsa/hsa_ext_amd.h" + +#include + +#define RETRY_LIMIT 5 +#define DELAY_IN_MILLISECONDS 1 +#define TIMEOUT_LIMIT 5000 // milliseconds + +static const char kSubTestSeparator[] = " **************************"; +static void PrintDebugSubtestHeader(const char* header) { + std::cout << " *** TrapHandler Subtest: " << header << " ***" << std::endl; +} + +static const uint32_t kNumBufferElements = 256; + +typedef struct q_callback_data_t { + hsa_queue_t** qptr; + bool trap; +} q_callback_data; + +static void CallbackQueueErrorHandler(hsa_status_t status, hsa_queue_t* source, void* data); +static hsa_status_t CallbackEventHandler(const hsa_amd_event_t* event, void* data); + +TrapHandler::TrapHandler(bool trigger_s_trap, bool trigger_memory_violation) : TestBase() { + std::string name; + std::string desc; + + name = "ROCr Trap Handler Test"; + desc = + "This set of tests intentionally trigger software exceptions and verify " + "that the GPU can handle abnormal situations."; + + if (trigger_s_trap) { + name += ": Trigger a software trap"; + desc += + "\n\nCurrent sub-test intentionally triggers a software exception using" + " the 's_trap' instruction, to validate if the queue's error handling" + " callback is triggered."; + } else if (trigger_memory_violation) { + name += ": Trigger illegal memory access"; + desc += + "\n\nCurrent sub-test intentionally triggers a memory violation error" + " to attempt accessing an invalid memory address. It verifies if the " + " GPU Memory protection exception is triggered."; + } + + set_title(name); + set_description(desc); + set_kernel_file_name("trap_handler_kernels.hsaco"); + kernel_names_ = {"trigger_s_trap", "trigger_memory_violation"}; +} + +void TrapHandler::SetUp() { + hsa_status_t err; + + TestBase::SetUp(); + + err = rocrtst::SetDefaultAgents(this); + CHECK(err); + + err = rocrtst::SetPoolsTypical(this); + CHECK(err) + + return; +} + +static inline void AtomicSetPacketHeader(uint16_t header, uint16_t setup, + hsa_kernel_dispatch_packet_t* queue_packet) { + __atomic_store_n(reinterpret_cast(queue_packet), header | (setup << 16), + __ATOMIC_RELEASE); +} + +void TrapHandler::Run(void) { + if (!rocrtst::CheckProfile(this)) { + return; + } + TestBase::Run(); +} + +void TrapHandler::TriggerSoftwareTrap(void) { + std::vector cpus, gpus; + hsa_status_t err = hsa_iterate_agents(rocrtst::IterateCPUAgents, &cpus); + CHECK(err) + + err = hsa_iterate_agents(rocrtst::IterateGPUAgents, &gpus); + CHECK(err) + + // Check if there are any GPUs available + if (gpus.empty()) { + std::cerr << "No GPUs found." << std::endl; + return; + } + + // Select the first GPU in the vector + hsa_agent_t selectedGpuAgent = gpus[0]; + uint32_t node_id; + err = hsa_agent_get_info(selectedGpuAgent, HSA_AGENT_INFO_NODE, &node_id); + CHECK(err); + + std::cout << "*** Running test on GPU node ID: 0x" << std::hex << node_id << "***\n" << std::endl; + execute_kernel("trigger_s_trap", cpus[0], selectedGpuAgent); + return; +} + +void TrapHandler::TriggerMemoryViolation(void) { + std::vector cpus, gpus; + hsa_status_t err = hsa_iterate_agents(rocrtst::IterateCPUAgents, &cpus); + CHECK(err) + + err = hsa_iterate_agents(rocrtst::IterateGPUAgents, &gpus); + CHECK(err) + + // Check if there are any GPUs available + if (gpus.empty()) { + std::cerr << "No GPUs found." << std::endl; + return; + } + + // Select the first GPU from the list + hsa_agent_t selectedGpuAgent = gpus[0]; + uint32_t node_id; + err = hsa_agent_get_info(selectedGpuAgent, HSA_AGENT_INFO_NODE, &node_id); + CHECK(err); + + std::cout << "*** Running test on GPU node ID: 0x" << std::hex << node_id << "***\n" << std::endl; + execute_kernel("trigger_memory_violation", cpus[0], selectedGpuAgent); + return; +} + +void TrapHandler::execute_kernel(const char* kernel_name, hsa_agent_t cpuAgent, + hsa_agent_t gpuAgent) { + hsa_signal_t signal; + uint32_t queue_size = 0; + + set_kernel_name(kernel_name); + + hsa_status_t err = hsa_agent_get_info(gpuAgent, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_size); + CHECK(err) + q_callback_data data; + data.qptr = &queue; + data.trap = false; + + // Create queue and register queue error handler callback function + err = hsa_queue_create(gpuAgent, queue_size, HSA_QUEUE_TYPE_MULTI, CallbackQueueErrorHandler, + &data, 0, 0, &queue); + CHECK(err) + ASSERT_NE(queue, nullptr); + set_main_queue(queue); + + // Register system event handler callback function + hsa_status_t status = hsa_amd_register_system_event_handler(&CallbackEventHandler, this); + ASSERT_EQ(status, HSA_STATUS_SUCCESS); + + hsa_amd_memory_pool_t kernarg_pool; + err = hsa_amd_agent_iterate_memory_pools(cpuAgent, rocrtst::GetKernArgMemoryPool, &kernarg_pool); + CHECK(err) + + hsa_amd_memory_pool_t global_pool; + err = hsa_amd_agent_iterate_memory_pools(cpuAgent, rocrtst::GetGlobalMemoryPool, &global_pool); + CHECK(err) + + err = hsa_amd_memory_pool_allocate(global_pool, kNumBufferElements * sizeof(uint32_t), 0, + reinterpret_cast(&src_buffer_)); + + CHECK(err) + + err = hsa_amd_agents_allow_access(1, &gpuAgent, NULL, src_buffer_); + CHECK(err) + + for (uint32_t i = 0; i < kNumBufferElements; ++i) { + reinterpret_cast(src_buffer_)[i] = i; + } + + err = hsa_amd_memory_pool_allocate(global_pool, kNumBufferElements * sizeof(uint32_t), 0, + reinterpret_cast(&dst_buffer_)); + CHECK(err) + err = hsa_amd_agents_allow_access(1, &gpuAgent, NULL, dst_buffer_); + CHECK(err) + + typedef struct __attribute__((aligned(16))) local_args_t { + uint32_t* dstArray; + uint32_t* srcArray; + uint32_t size; + } local_args; + + local_args* KernArgs = NULL; + err = hsa_amd_memory_pool_allocate(kernarg_pool, sizeof(local_args), 0, + reinterpret_cast(&KernArgs)); + CHECK(err) + + err = hsa_amd_agents_allow_access(1, &gpuAgent, NULL, KernArgs); + CHECK(err) + + KernArgs->dstArray = reinterpret_cast(dst_buffer_); + KernArgs->srcArray = reinterpret_cast(src_buffer_); + KernArgs->size = kNumBufferElements; + + err = rocrtst::LoadKernelFromObjFile(this, &gpuAgent); + CHECK(err) + + err = hsa_signal_create(1, 0, NULL, &signal); + CHECK(err) + + hsa_kernel_dispatch_packet_t aql; + memset(&aql, 0, sizeof(aql)); + + aql.header = 0; + aql.setup = 1; + aql.workgroup_size_x = kNumBufferElements; + aql.workgroup_size_y = 1; + aql.workgroup_size_z = 1; + aql.grid_size_x = kNumBufferElements; + aql.grid_size_y = 1; + aql.grid_size_z = 1; + aql.private_segment_size = 0; + aql.group_segment_size = 0; + aql.kernel_object = kernel_object(); + aql.kernarg_address = KernArgs; + aql.completion_signal = signal; + + const uint32_t queue_mask = queue->size - 1; + + uint64_t index = hsa_queue_load_write_index_relaxed(queue); + + hsa_queue_store_write_index_relaxed(queue, index + 1); + + rocrtst::WriteAQLToQueueLoc(queue, index, &aql); + + uint32_t 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; + + void* q_base = queue->base_address; + rocrtst::AtomicSetPacketHeader( + aql_header, aql.setup, + &(reinterpret_cast(q_base))[index & queue_mask]); + + hsa_signal_store_relaxed(queue->doorbell_signal, index); + + hsa_signal_value_t completion = 0; + int retry_count = 0; + + if (strcmp(kernel_name, "trigger_s_trap") == 0) { + completion = hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, 1, TIMEOUT_LIMIT, + HSA_WAIT_STATE_BLOCKED); + ASSERT_EQ(completion, 1); + + while (data.trap != true && retry_count < RETRY_LIMIT) { + ++retry_count; + sleep(DELAY_IN_MILLISECONDS); + } + ASSERT_EQ(data.trap, true); + } + + else if (strcmp(kernel_name, "trigger_memory_violation") == 0) { + while (event_occured != 1 && retry_count < RETRY_LIMIT) { + ++retry_count; + sleep(DELAY_IN_MILLISECONDS); + } + ASSERT_EQ(event_occured, 1); + } + + // Cleanup + if (KernArgs) hsa_memory_free(KernArgs); + if (src_buffer_) hsa_memory_free(src_buffer_); + if (dst_buffer_) hsa_memory_free(dst_buffer_); + if (signal.handle) hsa_signal_destroy(signal); + + if (!event_occured && queue) { + hsa_queue_destroy(queue); + queue = nullptr; + } + + return; +} + +void TrapHandler::DisplayTestInfo(void) { TestBase::DisplayTestInfo(); } + +void TrapHandler::DisplayResults(void) const { + if (!rocrtst::CheckProfile(this)) { + return; + } + + TestBase::DisplayResults(); + return; +} + +void TrapHandler::Close() { TestBase::Close(); } + +hsa_status_t CallbackEventHandler(const hsa_amd_event_t* event, void* data) { + TrapHandler* ptr = reinterpret_cast(data); + + switch (event->event_type) { + case HSA_AMD_GPU_MEMORY_FAULT_EVENT: + ptr->event_occured = true; + printf("Subtest Passed: Runtime caught GPU Memory Fault Event successfully!\n"); + break; + + case HSA_AMD_GPU_HW_EXCEPTION_EVENT: + ptr->event_occured = true; + printf("Subtest Passed: Runtime caught GPU HW Exception Event successfully!\n"); + break; + default: + // Unknown event type + printf("Subtest Failed: Unknown event type occurred\n"); + break; + }; + + std::cout << kSubTestSeparator << std::endl; + // Returning success to indicate that event was handled + return HSA_STATUS_SUCCESS; +} + +void CallbackQueueErrorHandler(hsa_status_t status, hsa_queue_t* source, void* data) { + std::cout << "Subtest Passed: Runtime caught trap instruction successfully!" << std::endl; + + ASSERT_NE(source, nullptr); + ASSERT_NE(data, nullptr); + + q_callback_data* debug_data = reinterpret_cast(data); + hsa_queue_t* queue = *(debug_data->qptr); + // check the queue id and user data + ASSERT_EQ(source->id, queue->id); + debug_data->trap = true; + + std::cout << kSubTestSeparator << std::endl; + return; +} diff --git a/projects/rocr-runtime/rocrtst/suites/functional/trap_handler.h b/projects/rocr-runtime/rocrtst/suites/functional/trap_handler.h new file mode 100644 index 0000000000..cbe67fdbcb --- /dev/null +++ b/projects/rocr-runtime/rocrtst/suites/functional/trap_handler.h @@ -0,0 +1,75 @@ +/* + * ============================================================================= + * ROC Runtime Conformance Release License + * ============================================================================= + * The University of Illinois/NCSA + * Open Source License (NCSA) + * + * Copyright (c) 2024, Advanced Micro Devices, Inc. + * All rights reserved. + * + * Developed by: + * + * AMD Research and AMD ROC Software Development + * + * Advanced Micro Devices, Inc. + * + * www.amd.com + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal with 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: + * + * - Redistributions of source code must retain the above copyright notice, + * this list of conditions and the following disclaimers. + * - Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimers in + * the documentation and/or other materials provided with the distribution. + * - Neither the names of , + * nor the names of its contributors may be used to endorse or promote + * products derived from this Software without specific prior written + * permission. + * + * 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 CONTRIBUTORS 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 WITH THE SOFTWARE. + * + */ + +#ifndef ROCRTST_SUITES_FUNCTIONAL_TRAP_HANDLER_H_ +#define ROCRTST_SUITES_FUNCTIONAL_TRAP_HANDLER_H_ + +#include "suites/test_common/test_base.h" +#include "common/base_rocr.h" +#include "common/common.h" + +class TrapHandler : public TestBase { + public: + bool event_occured = false; + hsa_queue_t* queue = nullptr; + std::vector kernel_names_; + + TrapHandler(bool trigger_s_trap, bool trigger_memory_violation); + virtual ~TrapHandler() {} + virtual void Run(); + virtual void Close(); + virtual void DisplayResults() const; + virtual void DisplayTestInfo(void); + void SetUp(); + void TriggerSoftwareTrap(void); + void TriggerMemoryViolation(void); + + private: + void* src_buffer_; + void* dst_buffer_; + void execute_kernel(const char* kernel_name, hsa_agent_t cpuAgent, hsa_agent_t gpuAgent); +}; + +#endif // ROCRTST_SUITES_FUNCTIONAL_TRAP_HANDLER_H_ diff --git a/projects/rocr-runtime/rocrtst/suites/test_common/CMakeLists.txt b/projects/rocr-runtime/rocrtst/suites/test_common/CMakeLists.txt index f9a9c3af9f..c8dd3b91e2 100755 --- a/projects/rocr-runtime/rocrtst/suites/test_common/CMakeLists.txt +++ b/projects/rocr-runtime/rocrtst/suites/test_common/CMakeLists.txt @@ -380,6 +380,11 @@ set(BITCODE_LIBS "${COMMON_BITCODE_LIBS}") set(CL_FILE_LIST "${KERNELS_DIR}/cu_mask_kernels.cl") build_sample_for_devices("cu_mask") +#Trap Handler Test case +set(BITCODE_LIBS "${COMMON_BITCODE_LIBS}") +set(CL_FILE_LIST "${KERNELS_DIR}/trap_handler_kernels.cl") +build_sample_for_devices("trap_handler") + set(CMAKE_BUILD_WITH_INSTALL_RPATH ON) # Build rules diff --git a/projects/rocr-runtime/rocrtst/suites/test_common/kernels/trap_handler_kernels.cl b/projects/rocr-runtime/rocrtst/suites/test_common/kernels/trap_handler_kernels.cl new file mode 100644 index 0000000000..44943db30a --- /dev/null +++ b/projects/rocr-runtime/rocrtst/suites/test_common/kernels/trap_handler_kernels.cl @@ -0,0 +1,76 @@ +/* + * ============================================================================= + * ROC Runtime Conformance Release License + * ============================================================================= + * The University of Illinois/NCSA + * Open Source License (NCSA) + * + * Copyright (c) 2024, Advanced Micro Devices, Inc. + * All rights reserved. + * + * Developed by: + * + * AMD Research and AMD ROC Software Development + * + * Advanced Micro Devices, Inc. + * + * www.amd.com + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal with 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: + * + * - Redistributions of source code must retain the above copyright notice, + * this list of conditions and the following disclaimers. + * - Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimers in + * the documentation and/or other materials provided with the distribution. + * - Neither the names of , + * nor the names of its contributors may be used to endorse or promote + * products derived from this Software without specific prior written + * permission. + * + * 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 CONTRIBUTORS 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 WITH THE SOFTWARE. + * + */ + +__kernel void trigger_s_trap(__global int* dstArray, __global const int* srcArray, const int sz) { + // Obtain the unique ID of the work-item + size_t id = get_global_id(0); + + // Adding this line _just_ to avoid any compiler optimization of this kernel + dstArray[id] = srcArray[id] * srcArray[id]; + + // Explicitly trigger a software trap on the first thread using AMDGCN assembly + if (id == 0) { + __asm__ volatile("s_trap 1"); + } +} + + +__kernel void trigger_memory_violation(__global int* dstArray, __global const int* srcArray, + const int sz) { + // Obtain the unique ID of the work-item + size_t id = get_global_id(0); + + // Adding this line _just_ to avoid any compiler optimization of this kernel + dstArray[id] = srcArray[id] * srcArray[id]; + + if (id == 0) { + // Point to a garbage memory location + int* invalidPtr = (int*)0xdeadbeef; + + // Try to read from and also write to the garbage memory location to trigger a memory violation + int invalidData = *invalidPtr; + *invalidPtr = 42; + } +} diff --git a/projects/rocr-runtime/rocrtst/suites/test_common/main.cc b/projects/rocr-runtime/rocrtst/suites/test_common/main.cc index c9406a69f6..1c61ba33f3 100644 --- a/projects/rocr-runtime/rocrtst/suites/test_common/main.cc +++ b/projects/rocr-runtime/rocrtst/suites/test_common/main.cc @@ -77,6 +77,7 @@ #include "suites/functional/aql_barrier_bit.h" #include "suites/functional/signal_kernel.h" #include "suites/functional/cu_masking.h" +#include "suites/functional/trap_handler.h" #include "rocm_smi/rocm_smi.h" static RocrTstGlobals *sRocrtstGlvalues = nullptr; @@ -220,6 +221,18 @@ TEST(rocrtstFunc, DISABLED_CU_Masking) { RunGenericTest(&sd); } +TEST(rocrtstFunc, TriggerSoftwareTrap) { + TrapHandler trap(true, false); + RunCustomTestProlog(&trap); + trap.TriggerSoftwareTrap(); +} + +TEST(rocrtstFunc, TriggerMemoryViolation) { + TrapHandler trap(false, true); + RunCustomTestProlog(&trap); + trap.TriggerMemoryViolation(); +} + #ifndef ROCRTST_EMULATOR_BUILD TEST(rocrtstFunc, IPC) { IPCTest ipc;