From 2e1f363d2f41f696d9a6e3e4a6eba6db50d2737e Mon Sep 17 00:00:00 2001 From: Shweta Khatri Date: Fri, 31 May 2024 17:25:29 -0400 Subject: [PATCH] Revert "Added new ROCr Trap Handler Test" This reverts commit b156e906d9c192bd487d10a8900e3eb6090ef547. Reason for revert: Memory violation test causing a timeout in subsequent test. Change-Id: If3a217575af545a47d6d67bebba4a2c640a43b81 Signed-off-by: Chris Freehill --- 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 deletions(-) delete mode 100644 rocrtst/suites/functional/trap_handler.cc delete mode 100644 rocrtst/suites/functional/trap_handler.h delete mode 100644 rocrtst/suites/test_common/kernels/trap_handler_kernels.cl diff --git a/rocrtst/suites/functional/trap_handler.cc b/rocrtst/suites/functional/trap_handler.cc deleted file mode 100644 index 587e489433..0000000000 --- a/rocrtst/suites/functional/trap_handler.cc +++ /dev/null @@ -1,381 +0,0 @@ -/* - * ============================================================================= - * 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/rocrtst/suites/functional/trap_handler.h b/rocrtst/suites/functional/trap_handler.h deleted file mode 100644 index cbe67fdbcb..0000000000 --- a/rocrtst/suites/functional/trap_handler.h +++ /dev/null @@ -1,75 +0,0 @@ -/* - * ============================================================================= - * 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/rocrtst/suites/test_common/CMakeLists.txt b/rocrtst/suites/test_common/CMakeLists.txt index c8dd3b91e2..f9a9c3af9f 100755 --- a/rocrtst/suites/test_common/CMakeLists.txt +++ b/rocrtst/suites/test_common/CMakeLists.txt @@ -380,11 +380,6 @@ 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/rocrtst/suites/test_common/kernels/trap_handler_kernels.cl b/rocrtst/suites/test_common/kernels/trap_handler_kernels.cl deleted file mode 100644 index 44943db30a..0000000000 --- a/rocrtst/suites/test_common/kernels/trap_handler_kernels.cl +++ /dev/null @@ -1,76 +0,0 @@ -/* - * ============================================================================= - * 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/rocrtst/suites/test_common/main.cc b/rocrtst/suites/test_common/main.cc index 1c61ba33f3..c9406a69f6 100644 --- a/rocrtst/suites/test_common/main.cc +++ b/rocrtst/suites/test_common/main.cc @@ -77,7 +77,6 @@ #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; @@ -221,18 +220,6 @@ 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;