From d27d4545e254105f83dbf2064bf036a31a29822e Mon Sep 17 00:00:00 2001 From: Sean Keely Date: Thu, 7 Oct 2021 19:36:35 -0500 Subject: [PATCH] Add cu masking test. Change-Id: I8b62ebd60f2edde3ea0b298f0353381855163fea --- rocrtst/common/base_rocr_utils.cc | 53 +- rocrtst/common/base_rocr_utils.h | 3 + rocrtst/common/helper_funcs.h | 19 +- rocrtst/common/rocr.cc | 254 ++++++++++ rocrtst/common/rocr.h | 177 +++++++ rocrtst/suites/functional/cu_masking.cc | 459 ++++++++++++++++++ rocrtst/suites/functional/cu_masking.h | 80 +++ rocrtst/suites/test_common/CMakeLists.txt | 8 +- .../test_common/kernels/cu_mask_kernels.cl | 20 + rocrtst/suites/test_common/main.cc | 6 + rocrtst/suites/test_common/test_base.cc | 15 +- rocrtst/suites/test_common/test_base.h | 6 + 12 files changed, 1073 insertions(+), 27 deletions(-) create mode 100644 rocrtst/common/rocr.cc create mode 100644 rocrtst/common/rocr.h create mode 100644 rocrtst/suites/functional/cu_masking.cc create mode 100644 rocrtst/suites/functional/cu_masking.h create mode 100644 rocrtst/suites/test_common/kernels/cu_mask_kernels.cl diff --git a/rocrtst/common/base_rocr_utils.cc b/rocrtst/common/base_rocr_utils.cc index b33679908b..e16e3b4d75 100755 --- a/rocrtst/common/base_rocr_utils.cc +++ b/rocrtst/common/base_rocr_utils.cc @@ -72,6 +72,16 @@ namespace rocrtst { } \ } +#define RET_IF_HSA_UTILS_ERR_RET(err, ret) \ + { \ + if ((err) != HSA_STATUS_SUCCESS) { \ + const char* msg = 0; \ + hsa_status_string(err, &msg); \ + EXPECT_EQ(HSA_STATUS_SUCCESS, err) << msg; \ + return (ret); \ + } \ + } + // Clean up some of the common handles and memory used by BaseRocR code, then // shut down hsa. Restore HSA_ENABLE_INTERRUPT to original value, if necessary hsa_status_t CommonCleanUp(BaseRocR* test) { @@ -271,6 +281,28 @@ bool CheckProfile(BaseRocR const* test) { return (test->requires_profile() == test->profile()); } } + +/// Locate file using local and device named file paths. +std::string LocateKernelFile(std::string filename, hsa_agent_t agent) { + char agent_name[64]; + std::string obj_file; + hsa_status_t err = hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, agent_name); + RET_IF_HSA_UTILS_ERR_RET(err, obj_file); + + obj_file = "./" + filename; + int file_handle = open(obj_file.c_str(), O_RDONLY); + if (file_handle == -1) { + obj_file = "./" + std::string(agent_name) + "/" + filename; + file_handle = open(obj_file.c_str(), O_RDONLY); + } + + if(file_handle == -1) + obj_file.clear(); + else + close(file_handle); + return obj_file; +} + // Load the specified kernel code from the specified file, inspect and fill // in BaseRocR member variables related to the kernel and executable. // Required Input BaseRocR member variables: @@ -294,25 +326,11 @@ hsa_status_t LoadKernelFromObjFile(BaseRocR* test, hsa_agent_t* agent) { agent = test->gpu_device1(); // Assume GPU agent for now } - // if agent name is not set, then set the agent name - if (!test->get_agent_name().size()) { - char agent_name[64]; - err = hsa_agent_get_info(*agent, HSA_AGENT_INFO_NAME, agent_name); - RET_IF_HSA_UTILS_ERR(err); - test->set_agent_name(agent_name); - } - std::string obj_file = "./" + test->kernel_file_name(); - std::string kern_name = test->kernel_name(); - - hsa_file_t file_handle = open(obj_file.c_str(), O_RDONLY); + std::string filename = LocateKernelFile(test->kernel_file_name(), *agent); + hsa_file_t file_handle = open(filename.c_str(), O_RDONLY); if (file_handle == -1) { - obj_file = "./" + test->get_agent_name() + "/" + test->kernel_file_name(); - file_handle = open(obj_file.c_str(), O_RDONLY); - } - - if (file_handle == -1) { - std::cout << "failed to open " << obj_file.c_str() << " at line " + std::cout << "failed to open " << filename.c_str() << " at line " << __LINE__ << ", file: " << __FILE__ << std::endl; return (hsa_status_t) errno; @@ -332,6 +350,7 @@ hsa_status_t LoadKernelFromObjFile(BaseRocR* test, hsa_agent_t* agent) { err = hsa_executable_freeze(executable, NULL); RET_IF_HSA_UTILS_ERR(err); + std::string kern_name = test->kernel_name(); hsa_executable_symbol_t kern_sym; err = hsa_executable_get_symbol(executable, NULL, (kern_name + ".kd").c_str(), *agent, 0, &kern_sym); diff --git a/rocrtst/common/base_rocr_utils.h b/rocrtst/common/base_rocr_utils.h index c86d6f42d2..f86b82bcf8 100755 --- a/rocrtst/common/base_rocr_utils.h +++ b/rocrtst/common/base_rocr_utils.h @@ -54,6 +54,9 @@ namespace rocrtst { +/// Locate kernel code object file and return path suitable for use with open(). +std::string LocateKernelFile(std::string filename, hsa_agent_t agent); + /// Open binary kernel object file and set all member data related to the /// kernel. Assumes that input test already has the kernel file name, /// agent name and kernel function specifed diff --git a/rocrtst/common/helper_funcs.h b/rocrtst/common/helper_funcs.h index 3777783586..5e05836668 100755 --- a/rocrtst/common/helper_funcs.h +++ b/rocrtst/common/helper_funcs.h @@ -51,6 +51,17 @@ #include #include + +#if defined(__GNUC__) +#define __forceinline __inline__ __attribute__((always_inline)) +#endif + +#define STRING2(x) #x +#define STRING(x) STRING2(x) + +#define PASTE2(x, y) x##y +#define PASTE(x, y) PASTE2(x, y) + namespace rocrtst { bool Compare(const float* refData, const float* data, @@ -100,10 +111,10 @@ uint64_t RoundToPowerOf2(uint64_t val); /// Checks if a value is a power of 2 bool IsPowerOf2(uint64_t val); -#define PASTE2(x, y) x##y -#define PASTE(x, y) PASTE2(x, y) - -#define __forceinline __inline__ __attribute__((always_inline)) +// Count set bits. +static __forceinline uint32_t popcount(uint32_t value) { + return __builtin_popcount(value); +} template class ScopeGuard { diff --git a/rocrtst/common/rocr.cc b/rocrtst/common/rocr.cc new file mode 100644 index 0000000000..4be44734db --- /dev/null +++ b/rocrtst/common/rocr.cc @@ -0,0 +1,254 @@ +/* + * ============================================================================= + * ROC Runtime Conformance Release License + * ============================================================================= + * The University of Illinois/NCSA + * Open Source License (NCSA) + * + * Copyright (c) 2021-2021, 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 "common/rocr.h" +System System::sys; + +bool DeviceDiscovery(System& devices) { + hsa_status_t err; + + err = hsa_iterate_agents([](hsa_agent_t agent, void* data) { + hsa_status_t err; + + System* devices = (System*)data; + + Device dev; + dev.agent = agent; + + dev.fine = -1u; + dev.coarse = -1u; + + hsa_device_type_t type; + err = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type); + CHECK(err); + + err = hsa_amd_agent_iterate_memory_pools(agent, [](hsa_amd_memory_pool_t pool, void* data) { + std::vector& pools = *reinterpret_cast*>(data); + hsa_status_t err; + + hsa_amd_segment_t segment; + err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment); + CHECK(err); + + if(segment != HSA_AMD_SEGMENT_GLOBAL) + return HSA_STATUS_SUCCESS; + + uint32_t flags; + err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flags); + CHECK(err); + + Device::Memory mem; + mem.pool=pool; + mem.fine = (flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED); + mem.kernarg = (flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT); + + err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SIZE, &mem.size); + CHECK(err); + + err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &mem.granule); + CHECK(err); + + pools.push_back(mem); + return HSA_STATUS_SUCCESS; + }, (void*)&dev.pools); + + if(!dev.pools.empty()) { + for(size_t i=0; icpu_.push_back(dev); + else + devices->gpu_.push_back(dev); + + devices->all_devices_.push_back(dev.agent); + } + + return HSA_STATUS_SUCCESS; + }, &devices); + + [&]() { + for(auto& dev : devices.cpu_) { + for(auto& mem : dev.pools) { + if(mem.fine && mem.kernarg) { + devices.kernarg_ = mem; + return; + } + } + } + }(); + + if(devices.cpu_.empty() || devices.gpu_.empty() || devices.kernarg_.pool.handle == 0) + return false; + return true; +} + +void System::Init() { + hsa_status_t err = hsa_init(); + CHECK(err); + + DeviceDiscovery(sys); +} + +void System::Shutdown() { + sys.~System(); + new (&sys) System(); + hsa_status_t err = hsa_shut_down(); + ASSERT_EQ(HSA_STATUS_SUCCESS, err); + err = hsa_shut_down(); + EXPECT_EQ(HSA_STATUS_ERROR_NOT_INITIALIZED, err); +} + +CodeObject::CodeObject(std::string filename, Device& agent) : agent(agent.agent) { + hsa_status_t err; + + file = open(filename.c_str(), O_RDONLY); + if(file == -1) { + throw std::runtime_error("Could not open file.\n"); + } + MAKE_NAMED_SCOPE_GUARD(fileGuard, [&](){ close(file); }); + + err = hsa_code_object_reader_create_from_file(file, &code_obj_rdr); + CHECK(err); + MAKE_NAMED_SCOPE_GUARD(readerGuard, [&](){ hsa_code_object_reader_destroy(code_obj_rdr); }); + + err = hsa_executable_create_alt(HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, nullptr, &executable); + CHECK(err); + MAKE_NAMED_SCOPE_GUARD(exeGuard, [&](){ hsa_executable_destroy(executable); }); + + err = hsa_executable_load_agent_code_object(executable, agent.agent, code_obj_rdr, nullptr, nullptr); + CHECK(err); + + err = hsa_executable_freeze(executable, nullptr); + CHECK(err); + + exeGuard.Dismiss(); + readerGuard.Dismiss(); + fileGuard.Dismiss(); +} + +CodeObject::~CodeObject() { + hsa_executable_destroy(executable); + hsa_code_object_reader_destroy(code_obj_rdr); + close(file); +} + +bool CodeObject::GetKernel(std::string name, Kernel& kern) { + hsa_executable_symbol_t symbol; + hsa_status_t err = hsa_executable_get_symbol_by_name(executable, name.c_str(), &agent, &symbol); + if(err != HSA_STATUS_SUCCESS) { + err = hsa_executable_get_symbol_by_name(executable, (name+".kd").c_str(), &agent, &symbol); + if(err != HSA_STATUS_SUCCESS) { + return false; + } + } + + err = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kern.handle); + CHECK(err); + + err = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, &kern.scratch); + CHECK(err); + //printf("Scratch: %d\n", kern.scratch); + + err = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, &kern.group); + CHECK(err); + //printf("LDS: %d\n", kern.group); + + // Remaining needs code object v2 or comgr. + err = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, &kern.kernarg_size); + CHECK(err); + //printf("Kernarg Size: %d\n", kern.kernarg_size); + + err = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT, &kern.kernarg_align); + CHECK(err); + //printf("Kernarg Align: %d\n", kern.kernarg_align); + + return true; +} + +// Not for parallel insertion. +bool SubmitPacket(hsa_queue_t* queue, Aql& pkt) { + size_t mask = queue->size - 1; + Aql* ring = (Aql*)queue->base_address; + + uint64_t write = hsa_queue_load_write_index_relaxed(queue); + uint64_t read = hsa_queue_load_read_index_relaxed(queue); + //if(write - read + 1 > queue->size) + // return false; + + Aql& dst = ring[write & mask]; + + uint16_t header = pkt.header.raw; + pkt.header.raw = dst.header.raw; + dst = pkt; + __atomic_store_n(&dst.header.raw, header, __ATOMIC_RELEASE); + pkt.header.raw = header; + + hsa_queue_store_write_index_release(queue, write+1); + hsa_signal_store_screlease(queue->doorbell_signal, write); + + return true; +} + +void* hsaMalloc(size_t size, const Device::Memory& mem) { + void* ret; + hsa_status_t err = hsa_amd_memory_pool_allocate(mem.pool, size, 0, &ret); + CHECK(err); + err = hsa_amd_agents_allow_access(System::all_devices().size(), &System::all_devices()[0], nullptr, ret); + CHECK(err); + return ret; +} + +void* hsaMalloc(size_t size, const Device& dev, bool fine) { + uint32_t index = fine ? dev.fine : dev.coarse; + assert(index != -1u && "Memory type unavailable."); + return hsaMalloc(size, dev.pools[index]); +} diff --git a/rocrtst/common/rocr.h b/rocrtst/common/rocr.h new file mode 100644 index 0000000000..71c8625c75 --- /dev/null +++ b/rocrtst/common/rocr.h @@ -0,0 +1,177 @@ +/* + * ============================================================================= + * ROC Runtime Conformance Release License + * ============================================================================= + * The University of Illinois/NCSA + * Open Source License (NCSA) + * + * Copyright (c) 2021-2021, 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 "hsa/hsa.h" +#include "hsa/hsa_ext_amd.h" +#include "hsa/hsa_ext_image.h" + +#include "common/helper_funcs.h" + +#include "gtest/gtest.h" + +#include +#include +#include "string.h" + +#include + +#define CHECK(err) [&](){ \ + if(err != HSA_STATUS_SUCCESS) { \ + EXPECT_EQ(HSA_STATUS_SUCCESS, err); \ + throw std::runtime_error("CHECK failure."); \ + } \ + }(); + +struct Device { + struct Memory { + hsa_amd_memory_pool_t pool; + bool fine; + bool kernarg; + size_t size; + size_t granule; + }; + + hsa_agent_t agent; + std::vector pools; + uint32_t fine; + uint32_t coarse; +}; + +struct Kernel { + uint64_t handle; + uint32_t scratch; + uint32_t group; + uint32_t kernarg_size; + uint32_t kernarg_align; +}; + +// Assumes bitfield layout is little endian. +// Assumes std::atomic is binary compatible with uint16_t and uses HW atomics. +union AqlHeader { + struct { + uint16_t type : 8; + uint16_t barrier : 1; + uint16_t acquire : 2; + uint16_t release : 2; + uint16_t reserved : 3; + }; + uint16_t raw; +}; + +struct BarrierValue { + AqlHeader header; + uint8_t AmdFormat; + uint8_t reserved; + uint32_t reserved1; + hsa_signal_t signal; + hsa_signal_value_t value; + hsa_signal_value_t mask; + uint32_t cond; + uint32_t reserved2; + uint64_t reserved3; + uint64_t reserved4; + hsa_signal_t completion_signal; +}; + +union Aql { + AqlHeader header; + hsa_kernel_dispatch_packet_t dispatch; + hsa_barrier_and_packet_t barrier_and; + hsa_barrier_or_packet_t barrier_or; + BarrierValue barrier_value; +}; + +struct OCLHiddenArgs { + uint64_t offset_x; + uint64_t offset_y; + uint64_t offset_z; + void* printf_buffer; + void* enqueue; + void* enqueue2; + void* multi_grid; +}; + +struct hip_hiddens { + uint64_t offset_x; + uint64_t offset_y; + uint64_t offset_z; + uint64_t _; + uint64_t _2; + uint64_t _3; + uint64_t multi_grid_sync; +}; + +class System { +public: + std::vector cpu_, gpu_; + std::vector all_devices_; + Device::Memory kernarg_; + + static void Init(); + static void Shutdown(); + static std::vector& cpu() { return sys.cpu_; } + static std::vector& gpu() { return sys.gpu_; } + static std::vector& all_devices() { return sys.all_devices_; } + static Device::Memory& kernarg() { return sys.kernarg_; } + static System sys; +}; + +class CodeObject { +public: + CodeObject(std::string filename, Device& agent); + ~CodeObject(); + bool GetKernel(std::string name, Kernel& kernel); +private: + hsa_file_t file; + hsa_code_object_reader_t code_obj_rdr; + hsa_executable_t executable; + hsa_agent_t agent; +}; + +// Not for parallel insertion. +bool SubmitPacket(hsa_queue_t* queue, Aql& pkt); + +void* hsaMalloc(size_t size, const Device::Memory& mem); +void* hsaMalloc(size_t size, const Device& dev, bool fine); diff --git a/rocrtst/suites/functional/cu_masking.cc b/rocrtst/suites/functional/cu_masking.cc new file mode 100644 index 0000000000..5a70614fee --- /dev/null +++ b/rocrtst/suites/functional/cu_masking.cc @@ -0,0 +1,459 @@ +/* + * ============================================================================= + * ROC Runtime Conformance Release License + * ============================================================================= + * The University of Illinois/NCSA + * Open Source License (NCSA) + * + * Copyright (c) 2021-2021, 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/cu_masking.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 +#include +#include +#include +#include + +CU_Masking::CU_Masking() : TestBase() { + std::string name; + std::string desc; + + name = "CU Masking"; + desc = "This test checks CU masking functionality via hsa_amd_queue_cu_get(set)_mask and HSA_CU_MASK."; + + set_title(name); + set_description(desc); + + set_kernel_file_name("cu_mask_kernels.hsaco"); +} + +void CU_Masking::Run() { + hsa_status_t err; + TestBase::Run(); + + printf("Running %lu iterations\n", RealIterationNum()); + + // Random source + std::mt19937 rand(std::chrono::system_clock::now().time_since_epoch().count()); + + // Store cu masking variable + std::string mask_var; + char* temp = getenv("HSA_CU_MASK"); + if(temp!=nullptr) + mask_var = temp; + unsetenv("HSA_CU_MASK"); + + std::string mask_init_var; + temp = getenv("HSA_CU_MASK_SKIP_INIT"); + if(temp!=nullptr) + mask_init_var = temp; + unsetenv("HSA_CU_MASK_SKIP_INIT"); + + // Loop over and test all GPUs + uint32_t idx = 0; + while(true) { + Device* gpu; + CodeObject* obj; + Kernel kern; + + struct args_t { + uint32_t* hw_ids; + OCLHiddenArgs _; + }; + args_t* args; + + hsa_signal_t signal; + hsa_queue_t* q; + + uint32_t cu_count; + uint32_t group_size; + uint32_t max_grid_size; + uint32_t threads; + + auto init = [&]() { + System::Init(); + if(idx == System::gpu().size()) + return false; + + gpu = &System::gpu()[idx]; + std::string filename = rocrtst::LocateKernelFile(kernel_file_name(), gpu->agent); + + obj = new CodeObject(filename, *gpu); + + err = hsa_agent_get_info(gpu->agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT, &cu_count); + CHECK(err); + + err = hsa_agent_get_info(gpu->agent, (hsa_agent_info_t)HSA_AGENT_INFO_WORKGROUP_MAX_SIZE, &group_size); + CHECK(err); + + err = hsa_agent_get_info(gpu->agent, (hsa_agent_info_t)HSA_AGENT_INFO_GRID_MAX_SIZE, &max_grid_size); + CHECK(err); + + uint64_t max_threads = uint64_t(cu_count)*group_size*10; + threads = max_threads < max_grid_size ? max_threads : max_grid_size; + threads = (threads / group_size) * group_size; + + // All CU enabled check + if(!obj->GetKernel("get_hw_id", kern)) { + ADD_FAILURE(); + return false; + } + + args = (args_t*)hsaMalloc(sizeof(args_t), System::kernarg()); + memset(args, 0, sizeof(args_t)); + + args->hw_ids = (uint32_t*)hsaMalloc(sizeof(uint32_t)*threads, System::kernarg()); + + err = hsa_signal_create(1, 0, nullptr, &signal); + CHECK(err); + + err = hsa_queue_create(gpu->agent, 4096, HSA_QUEUE_TYPE_SINGLE, nullptr, nullptr, 0, 0, &q); + CHECK(err); + + return true; + }; + + auto fini = [&]() { + err = hsa_queue_destroy(q); + CHECK(err); + err = hsa_signal_destroy(signal); + CHECK(err); + err = hsa_memory_free(args->hw_ids); + CHECK(err); + err = hsa_memory_free(args); + CHECK(err); + delete obj; + gpu = nullptr; + System::Shutdown(); + }; + + auto dispatch = [&]() { + memset(args->hw_ids, 0, sizeof(uint32_t)*threads); + + Aql pkt = {0}; + pkt.header.type = HSA_PACKET_TYPE_KERNEL_DISPATCH; + pkt.header.acquire = HSA_FENCE_SCOPE_SYSTEM; + pkt.header.release = HSA_FENCE_SCOPE_SYSTEM; + pkt.dispatch.kernel_object = kern.handle; + pkt.dispatch.private_segment_size = kern.scratch; + pkt.dispatch.group_segment_size = kern.group; + pkt.dispatch.setup = 1; + pkt.dispatch.workgroup_size_x = group_size; + pkt.dispatch.workgroup_size_y = 1; + pkt.dispatch.workgroup_size_z = 1; + pkt.dispatch.grid_size_x = threads; + pkt.dispatch.grid_size_y = 1; + pkt.dispatch.grid_size_z = 1; + pkt.dispatch.kernarg_address = args; + pkt.dispatch.completion_signal = signal; + + SubmitPacket(q, pkt); + + hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_EQ, 0, -1ull, HSA_WAIT_STATE_BLOCKED); + hsa_signal_store_relaxed(signal, 1); + }; + + auto getHwIds = [&](std::vector& ids){ + dispatch(); + std::sort(&args->hw_ids[0], &args->hw_ids[threads]); + uint32_t* end = std::unique(&args->hw_ids[0], &args->hw_ids[threads]); + ids.clear(); + ids.insert(ids.begin(), &args->hw_ids[0], end); + }; + + // Check fully unconstrained. + unsetenv("HSA_CU_MASK_SKIP_INIT"); + setenv("HSA_CU_MASK_SKIP_INIT", "1", 1); + + if(!init()) + break; + + { + char name[64]; + hsa_agent_get_info(gpu->agent, HSA_AGENT_INFO_NAME, name); + name[63]='\0'; + printf("Testing gpu index %u, %s\n", idx, name); + } + + std::vector left, right, isect; + + // Check unconstrained cu set. + getHwIds(left); + printf("Expecting %u CUs, found %lu with HSA_CU_MASK_SKIP_INIT.\n", cu_count, left.size()); + ASSERT_EQ(cu_count, left.size()); + fini(); + unsetenv("HSA_CU_MASK_SKIP_INIT"); + + // Check fully enabled, but mask used, set. + setenv("HSA_CU_MASK", (std::to_string(idx)+":0-"+std::to_string(cu_count-1)).c_str(), 1); + init(); + getHwIds(right); + printf("Expecting %u CUs, found %lu with HSA_CU_MASK.\n", cu_count, right.size()); + if(cu_count != right.size()) { + isect.resize(left.size()); + auto isect_end = std::set_difference(left.begin(), left.end(), right.begin(), right.end(), isect.begin()); + isect.resize(isect_end - isect.begin()); + printf("Missing CUs: "); + for(auto cu : isect) + printf("%u ", cu); + printf("\n"); + } + ASSERT_EQ(cu_count, right.size()); + fini(); + unsetenv("HSA_CU_MASK"); + + // Check rocr default mask. + init(); + getHwIds(right); + printf("Expecting %u CUs, found %lu.\n", cu_count, right.size()); + if(cu_count != right.size()) { + isect.resize(left.size()); + auto isect_end = std::set_difference(left.begin(), left.end(), right.begin(), right.end(), isect.begin()); + isect.resize(isect_end - isect.begin()); + printf("Missing CUs: "); + for(auto cu : isect) + printf("%u ", cu); + printf("\n"); + } + ASSERT_EQ(cu_count, right.size()); + fini(); + + std::vector bits; + for(uint32_t i=0; i bitmask, resultmask; + uint32_t dwords = (cu_count + 31) / 32; + + bitmask.resize(dwords); + resultmask.resize(dwords); + + for(size_t iteration=0; iteration& array) { + assert(array.size() == dwords && "Bitmask array has incorrect size."); + for(uint32_t i=0; i& hw_ids) { + setBits(start, stop, bitmask); + err = hsa_amd_queue_cu_set_mask(q, dwords*32, &bitmask[0]); + if((err!=HSA_STATUS_SUCCESS) && (err!=(hsa_status_t)HSA_STATUS_CU_MASK_REDUCED)) + CHECK(err); + err = hsa_amd_queue_cu_get_mask(q, dwords*32, &resultmask[0]); + CHECK(err); + getHwIds(hw_ids); + }; + + auto getIsect = [&]() { + isect.resize(left.size()); + auto isect_end = std::set_intersection(left.begin(), left.end(), right.begin(), right.end(), isect.begin()); + isect.resize(isect_end - isect.begin()); + }; + + auto printMask = [](std::vector& mask) { + printf("0x"); + for(size_t i=1; i env_mask(&bits[0], &bits[mask_index]); + + // Convert to string range syntax + std::sort(env_mask.begin(), env_mask.end()); + uint32_t start, stop; + start=stop=env_mask[0]; + std::vector ranges; + // Append invalid bit so that final loop will emit the last range. + env_mask.push_back(-1); + for(size_t j=1; j "); + printMask(env_mask); + printf("\n"); + + init(); + + getHwIds(left); + printf("Expecting %u CUs, found %lu\n", mask_index, left.size()); + ASSERT_EQ(left.size(), mask_index); + + // Check that HSA_CU_MASK constrains the API + // Find at least partially enabled CU mask. + [&]() { + while(true) { + std::shuffle(bits.begin(), bits.end(), rand); + split_index = (rand() % (cu_count - 2)) + 1; + setBits(0, split_index, bitmask); + for(uint32_t i=0; i, + * 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_CU_MASKING_H_ +#define ROCRTST_SUITES_FUNCTIONAL_CU_MASKING_H_ +#include + +#include "suites/test_common/test_base.h" +#include "common/base_rocr.h" +#include "common/common.h" +#include "common/rocr.h" + +// @Brief: This class is defined to measure the mean latency of enqueuing +// the packets to an empty kernel + +class CU_Masking : public TestBase { + public: + // @Brief: Constructor + explicit CU_Masking(); + + // @Brief: Destructor + virtual ~CU_Masking() {} + + // @Brief: Set up the environment for the test + virtual void SetUp() { TestBase::SetupPrint(); } + + // @Brief: Run the test case + virtual void Run(); + + // @Brief: Clean up and close the runtime + virtual void Close() { TestBase::ClosePrint(); } + + private: + // @Brief: Get actual iteration number + virtual size_t RealIterationNum() { return num_iteration(); } +}; + +#endif // ROCRTST_SUITES_FUNCTIONAL_CU_MASKING_H_ diff --git a/rocrtst/suites/test_common/CMakeLists.txt b/rocrtst/suites/test_common/CMakeLists.txt index 84c3522cff..d71d75a57a 100755 --- a/rocrtst/suites/test_common/CMakeLists.txt +++ b/rocrtst/suites/test_common/CMakeLists.txt @@ -254,8 +254,9 @@ function(build_kernel S_NAME TARG_DEV) set(TARG_NAME "${S_NAME}_hsaco.${TARG_DEV}") set(HSACO_TARG_LIST ${HSACO_TARG_LIST} "${KERNEL_DIR}/${SNAME_KERNEL}" PARENT_SCOPE) + string(SUBSTRING ${TARG_DEV} 3 -1 gfxNum) separate_arguments(CLANG_ARG_LIST UNIX_COMMAND - "-x cl -target amdgcn-amd-amdhsa -include ${OPENCL_INC_DIR}opencl-c.h -mcpu=${TARG_DEV} ${BITCODE_ARGS} -cl-std=CL${OPENCL_VER} ${CL_FILE_LIST} -o ${KERNEL_DIR}/${SNAME_KERNEL}") + "-D ROCRTST_GPU=0x${gfxNum} -x cl -target amdgcn-amd-amdhsa -include ${OPENCL_INC_DIR}opencl-c.h -mcpu=${TARG_DEV} ${BITCODE_ARGS} -cl-std=CL${OPENCL_VER} ${CL_FILE_LIST} -o ${KERNEL_DIR}/${SNAME_KERNEL}") add_custom_command(OUTPUT "${KERNEL_DIR}/${SNAME_KERNEL}" COMMAND ${CLANG} ${CLANG_ARG_LIST} DEPENDS ${CL_FILE_LIST} ${CLANG} COMMENT "BUILDING ${KERNEL_DIR}/${SNAME_KERNEL}" VERBATIM) endfunction(build_kernel) @@ -337,6 +338,11 @@ set(BITCODE_LIBS "${COMMON_BITCODE_LIBS}") set(CL_FILE_LIST "${KERNELS_DIR}/groupMemoryDynamic_kernels.cl") build_sample_for_devices("groupMemoryDynamic") +# groupMemoryDynamic +set(BITCODE_LIBS "${COMMON_BITCODE_LIBS}") +set(CL_FILE_LIST "${KERNELS_DIR}/cu_mask_kernels.cl") +build_sample_for_devices("cu_mask") + set(CMAKE_BUILD_WITH_INSTALL_RPATH ON) # Build rules diff --git a/rocrtst/suites/test_common/kernels/cu_mask_kernels.cl b/rocrtst/suites/test_common/kernels/cu_mask_kernels.cl new file mode 100644 index 0000000000..9527c04adf --- /dev/null +++ b/rocrtst/suites/test_common/kernels/cu_mask_kernels.cl @@ -0,0 +1,20 @@ +#define GETREG_IMMED(SIZE, OFFSET, REG) ((SIZE-1)<<11)|(OFFSET<<6)|REG + +#if ROCRTST_GPU < 0x1000 + #define HW_ID_CU_ID_OFFSET 8 + #define HW_ID 4 + #if (ROCRTST_GPU == 0x908) || (ROCRTST_GPU == 0x90a) || (ROCRTST_GPU == 0x940) + #define HW_ID_CU_ID_SIZE 8 + #else + #define HW_ID_CU_ID_SIZE 7 + #endif +#else + #define HW_ID_CU_ID_OFFSET 9 //Skips first bit of SIMD ID, could be wrong. + #define HW_ID 23 + #define HW_ID_CU_ID_SIZE 10 +#endif + +__kernel void get_hw_id(__global uint* hw_ids) { + uint idx = get_global_id(0); + hw_ids[idx] = __builtin_amdgcn_s_getreg(GETREG_IMMED(HW_ID_CU_ID_SIZE, HW_ID_CU_ID_OFFSET, HW_ID)); +} diff --git a/rocrtst/suites/test_common/main.cc b/rocrtst/suites/test_common/main.cc index 3a4ece4b84..f9e1735969 100755 --- a/rocrtst/suites/test_common/main.cc +++ b/rocrtst/suites/test_common/main.cc @@ -75,6 +75,7 @@ #include "suites/functional/signal_concurrent.h" #include "suites/functional/aql_barrier_bit.h" #include "suites/functional/signal_kernel.h" +#include "suites/functional/cu_masking.h" #include "rocm_smi/rocm_smi.h" static RocrTstGlobals *sRocrtstGlvalues = nullptr; @@ -205,6 +206,11 @@ TEST(rocrtstFunc, Signal_Create_Concurrently) { RunCustomTestEpilog(&sd); } +TEST(rocrtstFunc, CU_Masking) { + CU_Masking sd; + RunGenericTest(&sd); +} + #ifndef ROCRTST_EMULATOR_BUILD TEST(rocrtstFunc, IPC) { IPCTest ipc; diff --git a/rocrtst/suites/test_common/test_base.cc b/rocrtst/suites/test_common/test_base.cc index ef3b5f6b18..47cc48eb35 100755 --- a/rocrtst/suites/test_common/test_base.cc +++ b/rocrtst/suites/test_common/test_base.cc @@ -77,12 +77,15 @@ static void MakeHeaderStr(const char *inStr, std::string *outStr) { *outStr += kLabelDelimiter; } -void TestBase::SetUp(void) { - hsa_status_t err; +void TestBase::SetupPrint() { std::string label; MakeHeaderStr(kSetupLabel, &label); printf("\n\t%s\n", label.c_str()); +} +void TestBase::SetUp(void) { + hsa_status_t err; + SetupPrint(); err = rocrtst::InitAndSetupHSA(this); ASSERT_EQ(HSA_STATUS_SUCCESS, err); @@ -95,13 +98,15 @@ void TestBase::Run(void) { printf("\n\t%s\n", label.c_str()); } -void TestBase::Close(void) { - hsa_status_t err; +void TestBase::ClosePrint() { std::string label; MakeHeaderStr(kCloseLabel, &label); printf("\n\t%s\n", label.c_str()); +} - +void TestBase::Close(void) { + hsa_status_t err; + ClosePrint(); if (monitor_verbosity() > 0) { DumpMonitorInfo(); } diff --git a/rocrtst/suites/test_common/test_base.h b/rocrtst/suites/test_common/test_base.h index dcb6677be9..c7b44b2c93 100755 --- a/rocrtst/suites/test_common/test_base.h +++ b/rocrtst/suites/test_common/test_base.h @@ -79,6 +79,12 @@ class TestBase : public rocrtst::BaseRocR { void set_description(std::string d); + // @Brief: Emit setup output string only. For tests with custom setup. + void SetupPrint(void); + + // @Brief: Emit close output string only. For tests with custom close. + void ClosePrint(void); + private: std::string description_; };