Add cu masking test.
Change-Id: I8b62ebd60f2edde3ea0b298f0353381855163fea
This commit is contained in:
Sean Keely
committed by
David Yat Sin
David Yat Sin
parent
965df6eef7
commit
d27d4545e2
@@ -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);
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -51,6 +51,17 @@
|
||||
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
#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 <typename lambda>
|
||||
class ScopeGuard {
|
||||
|
||||
@@ -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 <Name of Development Group, Name of Institution>,
|
||||
* 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<Device::Memory>& pools = *reinterpret_cast<std::vector<Device::Memory>*>(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; i<dev.pools.size(); i++) {
|
||||
if(dev.pools[i].fine && dev.pools[i].kernarg && dev.fine==-1u)
|
||||
dev.fine = i;
|
||||
if(dev.pools[i].fine && !dev.pools[i].kernarg)
|
||||
dev.fine = i;
|
||||
if(!dev.pools[i].fine)
|
||||
dev.coarse = i;
|
||||
}
|
||||
|
||||
if(type == HSA_DEVICE_TYPE_CPU)
|
||||
devices->cpu_.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]);
|
||||
}
|
||||
@@ -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 <Name of Development Group, Name of Institution>,
|
||||
* 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 <fcntl.h>
|
||||
#include <assert.h>
|
||||
#include "string.h"
|
||||
|
||||
#include <vector>
|
||||
|
||||
#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<Memory> 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<uint16_t> 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<Device> cpu_, gpu_;
|
||||
std::vector<hsa_agent_t> all_devices_;
|
||||
Device::Memory kernarg_;
|
||||
|
||||
static void Init();
|
||||
static void Shutdown();
|
||||
static std::vector<Device>& cpu() { return sys.cpu_; }
|
||||
static std::vector<Device>& gpu() { return sys.gpu_; }
|
||||
static std::vector<hsa_agent_t>& 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);
|
||||
@@ -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 <Name of Development Group, Name of Institution>,
|
||||
* 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 <string>
|
||||
#include <stdlib.h>
|
||||
#include <algorithm>
|
||||
#include <random>
|
||||
#include <chrono>
|
||||
|
||||
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<uint32_t>& 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<uint32_t> 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<uint32_t> bits;
|
||||
for(uint32_t i=0; i<cu_count; i++)
|
||||
bits.push_back(i);
|
||||
|
||||
std::vector<uint32_t> bitmask, resultmask;
|
||||
uint32_t dwords = (cu_count + 31) / 32;
|
||||
|
||||
bitmask.resize(dwords);
|
||||
resultmask.resize(dwords);
|
||||
|
||||
for(size_t iteration=0; iteration<RealIterationNum(); iteration++) {
|
||||
|
||||
auto setBits = [&](uint32_t start, uint32_t stop, std::vector<uint32_t>& array) {
|
||||
assert(array.size() == dwords && "Bitmask array has incorrect size.");
|
||||
for(uint32_t i=0; i<dwords; i++)
|
||||
array[i] = 0;
|
||||
for(uint32_t i=start; i<stop; i++) {
|
||||
int dword = bits[i] / 32;
|
||||
int offset = bits[i] % 32;
|
||||
array[dword] |= (1 << offset);
|
||||
}
|
||||
};
|
||||
|
||||
auto getMasks = [&](uint32_t start, uint32_t stop, std::vector<uint32_t>& 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<uint32_t>& mask) {
|
||||
printf("0x");
|
||||
for(size_t i=1; i<mask.size()+1; i++)
|
||||
printf("%08X", mask[mask.size()-i]);
|
||||
};
|
||||
|
||||
auto printMasks = [&]() {
|
||||
printf("Set mask: ");
|
||||
printMask(bitmask);
|
||||
printf("\n");
|
||||
printf("Get mask: ");
|
||||
printMask(resultmask);
|
||||
printf("\n");
|
||||
};
|
||||
|
||||
// CU set API check, no overlap
|
||||
std::shuffle(bits.begin(), bits.end(), rand);
|
||||
uint32_t split_index = (rand() % (cu_count - 2)) + 1;
|
||||
|
||||
init();
|
||||
|
||||
getMasks(0, split_index, left);
|
||||
printMasks();
|
||||
printf("Observed %lu CUs.\n", left.size());
|
||||
for(uint32_t i=0; i<dwords; i++)
|
||||
ASSERT_EQ(bitmask[i], resultmask[i]);
|
||||
ASSERT_EQ(split_index, left.size());
|
||||
|
||||
getMasks(split_index, cu_count, right);
|
||||
printMasks();
|
||||
printf("Observed %lu CUs.\n", right.size());
|
||||
for(uint32_t i=0; i<dwords; i++)
|
||||
ASSERT_EQ(bitmask[i], resultmask[i]);
|
||||
ASSERT_EQ(cu_count-split_index, right.size());
|
||||
|
||||
getIsect();
|
||||
printf("Overlap of %lu CUs.\n", isect.size());
|
||||
ASSERT_EQ(0u, isect.size());
|
||||
|
||||
// CU set API check, overlap possible
|
||||
uint32_t high_split_index = (rand() % (cu_count - 2)) + 1;
|
||||
|
||||
if(high_split_index < split_index)
|
||||
std::swap(high_split_index, split_index);
|
||||
|
||||
getMasks(0, high_split_index, left);
|
||||
printMasks();
|
||||
printf("Observed %lu CUs.\n", left.size());
|
||||
for(uint32_t i=0; i<dwords; i++)
|
||||
ASSERT_EQ(bitmask[i], resultmask[i]);
|
||||
ASSERT_EQ(high_split_index, left.size());
|
||||
|
||||
getMasks(split_index, cu_count, right);
|
||||
printMasks();
|
||||
printf("Observed %lu CUs.\n", right.size());
|
||||
for(uint32_t i=0; i<dwords; i++)
|
||||
ASSERT_EQ(bitmask[i], resultmask[i]);
|
||||
ASSERT_EQ(cu_count-split_index, right.size());
|
||||
|
||||
getIsect();
|
||||
printf("Overlap of %lu CUs.\n", isect.size());
|
||||
ASSERT_EQ(high_split_index - split_index, isect.size());
|
||||
|
||||
// HSA_CU_MASK check, default
|
||||
fini();
|
||||
|
||||
// Pick masking bits for env var
|
||||
std::shuffle(bits.begin(), bits.end(), rand);
|
||||
uint32_t mask_index = (rand() % (cu_count - 2)) + 1;
|
||||
std::vector<uint32_t> 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<std::string> ranges;
|
||||
// Append invalid bit so that final loop will emit the last range.
|
||||
env_mask.push_back(-1);
|
||||
for(size_t j=1; j<env_mask.size(); j++) {
|
||||
uint32_t index = env_mask[j];
|
||||
if(index != stop+1) {
|
||||
if(start==stop)
|
||||
ranges.push_back(std::to_string(start));
|
||||
else
|
||||
ranges.push_back(std::to_string(start)+"-"+std::to_string(stop));
|
||||
start=stop=index;
|
||||
} else {
|
||||
stop = index;
|
||||
}
|
||||
}
|
||||
env_mask.pop_back();
|
||||
// Shuffle ranges
|
||||
std::shuffle(ranges.begin(), ranges.end(), rand);
|
||||
// Assemble final env var string.
|
||||
std::string env_var = std::to_string(idx) + ":";
|
||||
env_var += ranges[0];
|
||||
for(uint32_t i=1; i<ranges.size(); i++)
|
||||
env_var += ", " + ranges[i];
|
||||
|
||||
// Set env var and check that default queues are masked.
|
||||
//env_var = "0:41-44, 104-107, 47-50, 67-68, 77-100, 61, 102, 19-24, 109, 70-75, 52-59, 63-65, 0-17, 27-39";
|
||||
setenv("HSA_CU_MASK", env_var.c_str(), 1);
|
||||
printf("HSA_CU_MASK = %s\n", env_var.c_str());
|
||||
env_mask.clear();
|
||||
env_mask.resize(dwords);
|
||||
setBits(0, mask_index, env_mask);
|
||||
printf(" HSA_CU_MASK => ");
|
||||
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<dwords; i++) {
|
||||
if((bitmask[i] & env_mask[i]) != 0)
|
||||
return;
|
||||
}
|
||||
}
|
||||
}();
|
||||
|
||||
getMasks(0, split_index, left);
|
||||
printMasks();
|
||||
printf("Observed %lu CUs.\n", left.size());
|
||||
uint32_t enabledCus = 0;
|
||||
for(uint32_t i=0; i<dwords; i++) {
|
||||
bitmask[i] &= env_mask[i];
|
||||
enabledCus += rocrtst::popcount(bitmask[i]);
|
||||
ASSERT_EQ(bitmask[i], resultmask[i]);
|
||||
}
|
||||
ASSERT_EQ(enabledCus, left.size());
|
||||
ASSERT_LE(enabledCus, mask_index);
|
||||
|
||||
fini();
|
||||
unsetenv("HSA_CU_MASK");
|
||||
|
||||
// Todo: Hex syntax. Syntax errors. Above hw limit bits.
|
||||
|
||||
}
|
||||
idx++;
|
||||
}
|
||||
|
||||
if(!mask_var.empty())
|
||||
setenv("HSA_CU_MASK", mask_var.c_str(), 1);
|
||||
if(!mask_init_var.empty())
|
||||
setenv("HSA_CU_MASK_SKIP_INIT", mask_var.c_str(), 1);
|
||||
}
|
||||
@@ -0,0 +1,80 @@
|
||||
/*
|
||||
* =============================================================================
|
||||
* 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 <Name of Development Group, Name of Institution>,
|
||||
* 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 <vector>
|
||||
|
||||
#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_
|
||||
@@ -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
|
||||
|
||||
@@ -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));
|
||||
}
|
||||
@@ -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;
|
||||
|
||||
@@ -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();
|
||||
}
|
||||
|
||||
@@ -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_;
|
||||
};
|
||||
|
||||
Reference in New Issue
Block a user