Making v2 tests CI ready
Change-Id: Ia32c9b2a1b5f833d1c7b078b678895b736d1e2a1
[ROCm/rocprofiler commit: 13c12bc7e7]
This commit is contained in:
gobhardw
Ammar ELWazir
@@ -0,0 +1,387 @@
|
||||
/*
|
||||
* =============================================================================
|
||||
* ROC Runtime Conformance Release License
|
||||
* =============================================================================
|
||||
* The University of Illinois/NCSA
|
||||
* Open Source License (NCSA)
|
||||
*
|
||||
* Copyright (c) 2017, 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 <cassert>
|
||||
#include <iostream>
|
||||
|
||||
#include "hsa/hsa.h"
|
||||
#include "hsa/hsa_ext_amd.h"
|
||||
|
||||
#define RET_IF_HSA_ERR(err) \
|
||||
{ \
|
||||
if ((err) != HSA_STATUS_SUCCESS) { \
|
||||
const char *msg = 0; \
|
||||
hsa_status_string(err, &msg); \
|
||||
std::cout << "hsa api call failure at line " << __LINE__ \
|
||||
<< ", file: " << __FILE__ << ". Call returned " << err \
|
||||
<< std::endl; \
|
||||
std::cout << msg << std::endl; \
|
||||
return (err); \
|
||||
} \
|
||||
}
|
||||
|
||||
static const uint32_t kTestFillValue1 = 0xabcdef12;
|
||||
static const uint32_t kTestFillValue2 = 0xba5eba11;
|
||||
static const uint32_t kTestFillValue3 = 0xfeed5a1e;
|
||||
static const uint32_t kTestInitValue = 0xbaadf00d;
|
||||
|
||||
// This structure holds an agent pointer and associated memory pool to be used
|
||||
// for this test program.
|
||||
struct async_mem_cpy_agent {
|
||||
hsa_agent_t dev;
|
||||
hsa_amd_memory_pool_t pool;
|
||||
size_t granule;
|
||||
void *ptr;
|
||||
};
|
||||
struct async_mem_cpy_pool_query {
|
||||
async_mem_cpy_agent *pool_info;
|
||||
hsa_agent_t peer_device;
|
||||
};
|
||||
struct callback_args {
|
||||
struct async_mem_cpy_agent cpu;
|
||||
struct async_mem_cpy_agent gpu1;
|
||||
struct async_mem_cpy_agent gpu2;
|
||||
};
|
||||
// Find the least common multiple of 2 numbers
|
||||
static uint32_t lcm(uint32_t a, uint32_t b) {
|
||||
int tmp_a;
|
||||
int tmp_b;
|
||||
tmp_a = a;
|
||||
tmp_b = b;
|
||||
while (tmp_a != tmp_b) {
|
||||
if (tmp_a < tmp_b) {
|
||||
tmp_a = tmp_a + a;
|
||||
} else {
|
||||
tmp_b = tmp_b + b;
|
||||
}
|
||||
}
|
||||
return tmp_a;
|
||||
}
|
||||
// This function is a callback for hsa_amd_agent_iterate_memory_pools()
|
||||
// and will test whether the provided memory pool is 1) in the GLOBAL
|
||||
// segment, 2) allows allocation and 3) is accessible by the provided
|
||||
// agent. The "data" input parameter is assumed to be pointing to a
|
||||
// struct async_mem_cpy_agent. If the provided pool meets these criteria,
|
||||
// HSA_STATUS_INFO_BREAK is returned.
|
||||
static hsa_status_t FindPool(hsa_amd_memory_pool_t in_pool, void *data) {
|
||||
hsa_amd_segment_t segment;
|
||||
hsa_status_t err;
|
||||
if (nullptr == data) {
|
||||
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
|
||||
}
|
||||
struct async_mem_cpy_pool_query *args =
|
||||
(struct async_mem_cpy_pool_query *)data;
|
||||
err = hsa_amd_memory_pool_get_info(in_pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT,
|
||||
&segment);
|
||||
RET_IF_HSA_ERR(err);
|
||||
if (segment != HSA_AMD_SEGMENT_GLOBAL) {
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
bool canAlloc;
|
||||
err = hsa_amd_memory_pool_get_info(
|
||||
in_pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED, &canAlloc);
|
||||
RET_IF_HSA_ERR(err);
|
||||
if (!canAlloc) {
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
if (args->peer_device.handle != 0) {
|
||||
hsa_amd_memory_pool_access_t access =
|
||||
HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED;
|
||||
err = hsa_amd_agent_memory_pool_get_info(
|
||||
args->peer_device, in_pool, HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS,
|
||||
&access);
|
||||
RET_IF_HSA_ERR(err);
|
||||
if (access == HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED) {
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
}
|
||||
err = hsa_amd_memory_pool_get_info(
|
||||
in_pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE,
|
||||
&args->pool_info->granule);
|
||||
RET_IF_HSA_ERR(err);
|
||||
args->pool_info->pool = in_pool;
|
||||
return HSA_STATUS_INFO_BREAK;
|
||||
}
|
||||
// This function is meant to be a callback to hsa_iterate_agents. For each
|
||||
// input agent the iterator provides as input, this function will check to
|
||||
// see if the input agent is a CPU agent. If so, it will update the
|
||||
// async_mem_cpy_agent structure pointed to by the input parameter "data".
|
||||
// Return values:
|
||||
// HSA_STATUS_INFO_BREAK -- CPU agent has been found and stored. Iterator
|
||||
// should stop iterating
|
||||
// HSA_STATUS_SUCCESS -- CPU agent has not yet been found; iterator
|
||||
// should keep iterating
|
||||
// Other -- Some error occurred
|
||||
static hsa_status_t FindCPUDevice(hsa_agent_t agent, void *data) {
|
||||
if (data == NULL) {
|
||||
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
|
||||
}
|
||||
hsa_device_type_t hsa_device_type;
|
||||
hsa_status_t err =
|
||||
hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &hsa_device_type);
|
||||
RET_IF_HSA_ERR(err);
|
||||
if (hsa_device_type == HSA_DEVICE_TYPE_CPU) {
|
||||
struct async_mem_cpy_agent *args = (struct async_mem_cpy_agent *)data;
|
||||
args->dev = agent;
|
||||
async_mem_cpy_pool_query pool_query;
|
||||
pool_query.peer_device.handle = 0;
|
||||
pool_query.pool_info = args;
|
||||
err = hsa_amd_agent_iterate_memory_pools(agent, FindPool, &pool_query);
|
||||
if (err == HSA_STATUS_INFO_BREAK) { // we found what we were looking for
|
||||
return HSA_STATUS_INFO_BREAK;
|
||||
} else {
|
||||
args->dev = {0};
|
||||
return err;
|
||||
}
|
||||
}
|
||||
// Returning HSA_STATUS_SUCCESS tells the calling iterator to keep iterating
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
// This function is meant to be a callback to hsa_iterate_agents. It will
|
||||
// attempt to find 2, or at least 1 GPU agent suitable for our test. The data
|
||||
// input parameter should point to a callback_args struct. The 2 GPU fields
|
||||
// will be updated as GPUs are discovered.
|
||||
// Return values:
|
||||
// HSA_STATUS_INFO_BREAK -- 2 GPU agents have been found and stored. Iterator
|
||||
// should stop iterating
|
||||
// HSA_STATUS_SUCCESS -- 2 GPU agents have not yet been found; 0 or 1 may
|
||||
// have been found; iterator function should keep iterating
|
||||
// Other -- Some error occurred
|
||||
static hsa_status_t FindGPUs(hsa_agent_t agent, void *data) {
|
||||
if (data == NULL) {
|
||||
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
|
||||
}
|
||||
hsa_device_type_t hsa_device_type;
|
||||
hsa_status_t err =
|
||||
hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &hsa_device_type);
|
||||
RET_IF_HSA_ERR(err);
|
||||
if (hsa_device_type != HSA_DEVICE_TYPE_GPU) {
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
struct callback_args *args = (struct callback_args *)data;
|
||||
struct async_mem_cpy_agent *gpu;
|
||||
async_mem_cpy_pool_query pool_query = {0, 0};
|
||||
if (args->gpu1.dev.handle == 0) {
|
||||
gpu = &args->gpu1;
|
||||
} else {
|
||||
gpu = &args->gpu2;
|
||||
// Check that gpu1 has peer access into the selected pool.
|
||||
pool_query.peer_device = args->gpu1.dev;
|
||||
}
|
||||
// Make sure GPU device has pool host can access
|
||||
gpu->dev = agent;
|
||||
pool_query.pool_info = gpu;
|
||||
err = hsa_amd_agent_iterate_memory_pools(agent, FindPool, &pool_query);
|
||||
if (err == HSA_STATUS_INFO_BREAK) {
|
||||
if (gpu == &args->gpu2) {
|
||||
// We found 2 gpu's
|
||||
return HSA_STATUS_INFO_BREAK;
|
||||
} else {
|
||||
// Keep looking for another gpu
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
} else {
|
||||
gpu->dev = {0};
|
||||
}
|
||||
RET_IF_HSA_ERR(err);
|
||||
// Returning HSA_STATUS_SUCCESS tells the calling iterator to keep iterating
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
// This is the main test, showing various paths of async. copy. Source and
|
||||
// destination agents and their respective pools should already be discovered.
|
||||
// Additionally, buffer from the pools should already be allocated and availble
|
||||
// from the input parameters.
|
||||
static hsa_status_t AsyncCpyTest(async_mem_cpy_agent *dst,
|
||||
async_mem_cpy_agent *src, callback_args *args,
|
||||
size_t sz, uint32_t val) {
|
||||
hsa_status_t err;
|
||||
hsa_signal_t copy_signal;
|
||||
// Initialize the system and destination buffers with a value so we can later
|
||||
// validate it has been overwritten
|
||||
void *sysPtr = args->cpu.ptr;
|
||||
err = hsa_amd_memory_fill(sysPtr, kTestInitValue, sz / sizeof(uint32_t));
|
||||
RET_IF_HSA_ERR(err);
|
||||
if (dst->ptr != sysPtr) {
|
||||
err = hsa_amd_memory_fill(dst->ptr, kTestInitValue, sz / sizeof(uint32_t));
|
||||
RET_IF_HSA_ERR(err);
|
||||
}
|
||||
// Fill the source buffer with the provided uint32_t value
|
||||
err = hsa_amd_memory_fill(src->ptr, val, sz / sizeof(uint32_t));
|
||||
RET_IF_HSA_ERR(err);
|
||||
// Make sure the target and destination agents have access to the buffer.
|
||||
hsa_agent_t ag_list[2] = {dst->dev, src->dev};
|
||||
err = hsa_amd_agents_allow_access(2, ag_list, NULL, dst->ptr);
|
||||
RET_IF_HSA_ERR(err);
|
||||
// Create a signal that will be used to inform us when the copy is done
|
||||
err = hsa_signal_create(1, 0, NULL, ©_signal);
|
||||
RET_IF_HSA_ERR(err);
|
||||
// Do the copy...
|
||||
err = hsa_amd_memory_async_copy(dst->ptr, dst->dev, src->ptr, src->dev, sz, 0,
|
||||
NULL, copy_signal);
|
||||
RET_IF_HSA_ERR(err);
|
||||
// Here we do a blocking wait. Alternatively, we could also use a
|
||||
// non-blocking wait in a loop, and do other work while waiting.
|
||||
if (hsa_signal_wait_relaxed(copy_signal, HSA_SIGNAL_CONDITION_LT, 1, -1,
|
||||
HSA_WAIT_STATE_BLOCKED) != 0) {
|
||||
printf("Async copy returned error value.\n");
|
||||
return HSA_STATUS_ERROR;
|
||||
}
|
||||
// Verify the copy was successful; copy from the dst buffer to the sysBuf,
|
||||
// (if the result is not already in sys. mem.) and check the sysBuf values
|
||||
if (dst->ptr != sysPtr) {
|
||||
if (src->ptr != sysPtr) {
|
||||
// In this case, we need to give the gpu dev that owns dst->ptr access
|
||||
// to the system memory we are going to copy to.
|
||||
hsa_agent_t ag_list_ck[2] = {dst->dev, args->cpu.dev};
|
||||
err = hsa_amd_agents_allow_access(2, ag_list_ck, NULL, sysPtr);
|
||||
RET_IF_HSA_ERR(err);
|
||||
}
|
||||
// Reset signal to 1
|
||||
hsa_signal_store_screlease(copy_signal, 1);
|
||||
err = hsa_amd_memory_async_copy(sysPtr, args->cpu.dev, dst->ptr, dst->dev,
|
||||
sz, 0, NULL, copy_signal);
|
||||
RET_IF_HSA_ERR(err);
|
||||
if (hsa_signal_wait_relaxed(copy_signal, HSA_SIGNAL_CONDITION_LT, 1, -1,
|
||||
HSA_WAIT_STATE_BLOCKED) != 0) {
|
||||
printf("Async copy returned error value.\n");
|
||||
return HSA_STATUS_ERROR;
|
||||
}
|
||||
}
|
||||
// Check that the contents of the buffer are what is expected.
|
||||
for (uint32_t i = 0; i < sz / sizeof(uint32_t); ++i) {
|
||||
if (reinterpret_cast<uint32_t *>(sysPtr)[i] != val) {
|
||||
fprintf(stdout, "Expected 0x%x but got 0x%x in buffer at index %d.\n",
|
||||
val, reinterpret_cast<uint32_t *>(sysPtr)[i], i);
|
||||
return HSA_STATUS_ERROR;
|
||||
}
|
||||
}
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
// This program illustrates the usage of the asynchronous copy capability of
|
||||
// the RocR runtime library. The program will create a system memory buffer and
|
||||
// a local buffer for each GPU, up to 2 GPUs, if the system has at least 2
|
||||
// GPUs. The program will copy data to/from the host from/to the GPU. If 2
|
||||
// GPUs are available, the program will also copy data from one to the other.
|
||||
int main() {
|
||||
hsa_status_t err;
|
||||
struct callback_args args;
|
||||
bool twoGPUs = false;
|
||||
err = hsa_init();
|
||||
RET_IF_HSA_ERR(err);
|
||||
// First, find the cpu agent and associated pool
|
||||
args.cpu = {0, 0, 0};
|
||||
err = hsa_iterate_agents(FindCPUDevice, reinterpret_cast<void *>(&args.cpu));
|
||||
assert(err == HSA_STATUS_INFO_BREAK);
|
||||
if (err != HSA_STATUS_INFO_BREAK) {
|
||||
return -1;
|
||||
}
|
||||
// Now, find 1 or 2 (if possible) GPUs and associated pool(s) for our test
|
||||
args.gpu1 = {0, 0, 0};
|
||||
args.gpu2 = {0, 0, 0};
|
||||
err = hsa_iterate_agents(FindGPUs, &args);
|
||||
if (err == HSA_STATUS_INFO_BREAK) {
|
||||
twoGPUs = true;
|
||||
} else {
|
||||
// See if we at least have 1 GPU
|
||||
if (args.gpu1.dev.handle == 0) {
|
||||
fprintf(
|
||||
stdout,
|
||||
"GPU with accessible VRAM not found; at least 1 required. Exiting\n");
|
||||
return -1;
|
||||
}
|
||||
fprintf(stdout, "Only 1 GPU found with required VRAM. "
|
||||
"Peer-to-Peer copy will be skipped.\n");
|
||||
}
|
||||
// We will use the smallest amount of allocatable memory that works for all
|
||||
// potential sources and destinations of the copy
|
||||
size_t sz = lcm(args.cpu.granule, args.gpu1.granule);
|
||||
// Allocate memory on each source/destination
|
||||
if (twoGPUs) {
|
||||
sz = lcm(sz, args.gpu2.granule);
|
||||
err = hsa_amd_memory_pool_allocate(
|
||||
args.gpu2.pool, sz, 0, reinterpret_cast<void **>(&args.gpu2.ptr));
|
||||
RET_IF_HSA_ERR(err);
|
||||
}
|
||||
err = hsa_amd_memory_pool_allocate(args.cpu.pool, sz, 0,
|
||||
reinterpret_cast<void **>(&args.cpu.ptr));
|
||||
RET_IF_HSA_ERR(err);
|
||||
err = hsa_amd_memory_pool_allocate(args.gpu1.pool, sz, 0,
|
||||
reinterpret_cast<void **>(&args.gpu1.ptr));
|
||||
RET_IF_HSA_ERR(err);
|
||||
char name[64];
|
||||
err = hsa_agent_get_info(args.cpu.dev, HSA_AGENT_INFO_NAME, &name);
|
||||
fprintf(stdout, "CPU is \"%s\"\n", name);
|
||||
err = hsa_agent_get_info(args.gpu1.dev, HSA_AGENT_INFO_NAME, &name);
|
||||
fprintf(stdout, "GPU1 is \"%s\"\n", name);
|
||||
if (twoGPUs) {
|
||||
err = hsa_agent_get_info(args.gpu2.dev, HSA_AGENT_INFO_NAME, &name);
|
||||
fprintf(stdout, "GPU2 is \"%s\"\n", name);
|
||||
}
|
||||
fprintf(stdout, "Copying %lu bytes from gpu1 memory to system memory...\n",
|
||||
sz);
|
||||
err = AsyncCpyTest(&args.cpu, &args.gpu1, &args, sz, kTestFillValue1);
|
||||
RET_IF_HSA_ERR(err);
|
||||
fprintf(stdout, "Success!\n");
|
||||
fprintf(stdout, "Copying %lu bytes from system memory to gpu1 memory...\n",
|
||||
sz);
|
||||
err = AsyncCpyTest(&args.gpu1, &args.cpu, &args, sz, kTestFillValue2);
|
||||
RET_IF_HSA_ERR(err);
|
||||
fprintf(stdout, "Success!\n");
|
||||
if (twoGPUs) {
|
||||
fprintf(stdout, "Copying %lu bytes from gpu1 memory to gpu2 memory...\n",
|
||||
sz);
|
||||
err = AsyncCpyTest(&args.gpu2, &args.gpu1, &args, sz, kTestFillValue3);
|
||||
RET_IF_HSA_ERR(err);
|
||||
fprintf(stdout, "Success!\n");
|
||||
}
|
||||
// Clean up
|
||||
err = hsa_amd_memory_pool_free(args.cpu.ptr);
|
||||
RET_IF_HSA_ERR(err);
|
||||
err = hsa_amd_memory_pool_free(args.gpu1.ptr);
|
||||
RET_IF_HSA_ERR(err);
|
||||
if (twoGPUs) {
|
||||
err = hsa_amd_memory_pool_free(args.gpu2.ptr);
|
||||
RET_IF_HSA_ERR(err);
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,32 @@
|
||||
/* Copyright (c) 2022 Advanced Micro Devices, Inc.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in 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:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
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
|
||||
AUTHORS 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 IN
|
||||
THE SOFTWARE. */
|
||||
|
||||
__kernel void copyA(__global unsigned int* a, __global unsigned int* b) {
|
||||
uint tid = get_global_id(0);
|
||||
a[tid] = b[tid];
|
||||
}
|
||||
__kernel void copyB(__global unsigned int* a, __global unsigned int* b) {
|
||||
uint tid = get_global_id(0);
|
||||
a[tid] = b[tid];
|
||||
}
|
||||
__kernel void copyC(__global unsigned int* a, __global unsigned int* b) {
|
||||
uint tid = get_global_id(0);
|
||||
a[tid] = b[tid];
|
||||
}
|
||||
@@ -0,0 +1 @@
|
||||
pmc: SQ_WAVES GRBM_COUNT GRBM_GUI_ACTIVE
|
||||
Executable
+20
@@ -0,0 +1,20 @@
|
||||
0x5faaa0 agent cpu
|
||||
0x5fbb30 agent gpu
|
||||
844346791235313
|
||||
|
||||
|
||||
ROCMTools: Collecting the following counters:
|
||||
- GRBM_COUNT
|
||||
|
||||
Enabling Counter Collection
|
||||
System minor 0
|
||||
System major 9
|
||||
agent prop name AMD Radeon VII
|
||||
input string:
|
||||
GdkknVnqkc
|
||||
|
||||
output string:
|
||||
HelloWorld
|
||||
Passed!
|
||||
dispatch[2], gpu_id(0), queue_id(1), queue_index(0), pid(1531373), tid(0), grd(0), wgr(0), lds(0), scr(0), arch_vgpr(0), accum_vgpr(0), sgpr(0), wave_size(0), sig(0), obj(140646043297024), kernel-name("helloworld"), start_time(844346969235689), end_time(844346969239689)
|
||||
, GRBM_COUNT (20292)
|
||||
Executable
+16
@@ -0,0 +1,16 @@
|
||||
0x2147aa0 agent cpu
|
||||
0x2148b30 agent gpu
|
||||
844395826761899
|
||||
|
||||
|
||||
ROCMTools: Collecting the following counters:
|
||||
- GRBM_COUNT
|
||||
|
||||
Enabling Counter Collection
|
||||
System minor 0
|
||||
System major 9
|
||||
agent prop name AMD Radeon VII
|
||||
hip Device prop succeeded
|
||||
PASSED!
|
||||
dispatch[2], gpu_id(0), queue_id(1), queue_index(0), pid(1531435), tid(0), grd(0), wgr(0), lds(0), scr(0), arch_vgpr(0), accum_vgpr(0), sgpr(0), wave_size(0), sig(0), obj(140553153857024), kernel-name("vectoradd_float"), start_time(844396006072252), end_time(844396006104732)
|
||||
, GRBM_COUNT (67002)
|
||||
+16
@@ -0,0 +1,16 @@
|
||||
0xd1eeb0 agent cpu
|
||||
0xd4b380 agent gpu
|
||||
844434431085362
|
||||
|
||||
|
||||
ROCMTools: Collecting the following counters:
|
||||
- GRBM_COUNT
|
||||
|
||||
Enabling Counter Collection
|
||||
Only 1 GPU found with required VRAM. Peer-to-Peer copy will be skipped.
|
||||
CPU is "AMD Ryzen 9 5950X 16-Core Processor"
|
||||
GPU1 is "gfx906"
|
||||
Copying 4096 bytes from gpu1 memory to system memory...
|
||||
Success!
|
||||
Copying 4096 bytes from system memory to gpu1 memory...
|
||||
Success!
|
||||
@@ -0,0 +1 @@
|
||||
pmc: GRBM_COUNT
|
||||
Executable
+37
@@ -0,0 +1,37 @@
|
||||
0x55e2aaab1540 agent cpu
|
||||
0x55e2aab9f700 agent gpu
|
||||
844463523587280
|
||||
|
||||
|
||||
ROCMTools: Collecting the following counters:
|
||||
- GRBM_COUNT
|
||||
|
||||
Enabling Counter Collection
|
||||
0x12f9580 agent cpu
|
||||
0x1340580 agent gpu
|
||||
844463824808245
|
||||
|
||||
|
||||
ROCMTools: Collecting the following counters:
|
||||
- GRBM_COUNT
|
||||
|
||||
0xdcb320 agent cpu
|
||||
0xe122e0 agent gpu
|
||||
844463824808355
|
||||
|
||||
|
||||
ROCMTools: Collecting the following counters:
|
||||
- GRBM_COUNT
|
||||
|
||||
Enabling Counter Collection
|
||||
Enabling Counter Collection
|
||||
device count and rank is1: 2
|
||||
Rank Id: 0 | Device Id : 0 | Num Devices: 1
|
||||
device count and rank is1: 2
|
||||
Rank Id: 1 | Device Id : 0 | Num Devices: 1
|
||||
Max error: 0.000000
|
||||
Max error: 0.000000
|
||||
dispatch[2], gpu_id(0), queue_id(1), queue_index(0), pid(1531660), tid(0), grd(0), wgr(0), lds(0), scr(0), arch_vgpr(0), accum_vgpr(0), sgpr(0), wave_size(0), sig(0), obj(140604145903232), kernel-name("add"), start_time(844464004374381), end_time(844464006775011)
|
||||
, GRBM_COUNT (3724176)
|
||||
dispatch[2], gpu_id(0), queue_id(1), queue_index(0), pid(1531661), tid(0), grd(0), wgr(0), lds(0), scr(0), arch_vgpr(0), accum_vgpr(0), sgpr(0), wave_size(0), sig(0), obj(140242024941184), kernel-name("add"), start_time(844464004374753), end_time(844464006776661)
|
||||
, GRBM_COUNT (3724418)
|
||||
+5
@@ -0,0 +1,5 @@
|
||||
ROCMTools: Collecting the following counters:
|
||||
- GRBM_COUNT
|
||||
PASSED!
|
||||
dispatch[2], gpu-id(0), kernel-name("hip_helloworld"), time(7853273641921013,7853273641924568)
|
||||
GRBM_COUNT (21840)
|
||||
@@ -0,0 +1,84 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in 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:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
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
|
||||
AUTHORS 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 IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <string.h>
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
#include <fstream>
|
||||
#include <iostream>
|
||||
#include <string>
|
||||
|
||||
#define SUCCESS 0
|
||||
#define FAILURE 1
|
||||
|
||||
__global__ void helloworld(char *in, char *out) {
|
||||
int num = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x;
|
||||
out[num] = in[num] + 1;
|
||||
}
|
||||
|
||||
int main(int argc, char *argv[]) {
|
||||
hipDeviceProp_t devProp;
|
||||
hipGetDeviceProperties(&devProp, 0);
|
||||
std::cout << " System minor " << devProp.minor << std::endl;
|
||||
std::cout << " System major " << devProp.major << std::endl;
|
||||
std::cout << " agent prop name " << devProp.name << std::endl;
|
||||
|
||||
/* Initial input,output for the host and create memory objects for the
|
||||
* kernel*/
|
||||
const char *input = "GdkknVnqkc";
|
||||
size_t strlength = strlen(input);
|
||||
std::cout << "input string:" << std::endl;
|
||||
std::cout << input << std::endl;
|
||||
char *output = reinterpret_cast<char *>(malloc(strlength + 1));
|
||||
|
||||
char *inputBuffer;
|
||||
char *outputBuffer;
|
||||
hipMalloc(reinterpret_cast<void **>(&inputBuffer),
|
||||
(strlength + 1) * sizeof(char));
|
||||
hipMalloc(reinterpret_cast<void **>(&outputBuffer),
|
||||
(strlength + 1) * sizeof(char));
|
||||
|
||||
hipMemcpy(inputBuffer, input, (strlength + 1) * sizeof(char),
|
||||
hipMemcpyHostToDevice);
|
||||
|
||||
hipLaunchKernelGGL(helloworld, dim3(1), dim3(strlength), 0, 0, inputBuffer,
|
||||
outputBuffer);
|
||||
|
||||
hipMemcpy(output, outputBuffer, (strlength + 1) * sizeof(char),
|
||||
hipMemcpyDeviceToHost);
|
||||
|
||||
hipFree(inputBuffer);
|
||||
hipFree(outputBuffer);
|
||||
|
||||
output[strlength] = '\0'; // Add the terminal character to the end of output.
|
||||
std::cout << "\noutput string:" << std::endl;
|
||||
std::cout << output << std::endl;
|
||||
|
||||
free(output);
|
||||
|
||||
std::cout << "Passed!\n";
|
||||
|
||||
return SUCCESS;
|
||||
}
|
||||
@@ -0,0 +1,91 @@
|
||||
/*
|
||||
Copyright (c) 2020-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in 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:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
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
|
||||
AUTHORS 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 IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
// OpenMP program to print Hello World
|
||||
// using C language is supported by HIP
|
||||
|
||||
// OpenMP header
|
||||
#include <omp.h>
|
||||
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
// HIP header
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
#define NUM_THREADS 16
|
||||
#define CHECK(cmd) \
|
||||
{ \
|
||||
hipError_t error = cmd; \
|
||||
if (error != hipSuccess) { \
|
||||
fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), \
|
||||
error, __FILE__, __LINE__); \
|
||||
exit(EXIT_FAILURE); \
|
||||
} \
|
||||
}
|
||||
|
||||
__global__ void hip_helloworld(unsigned omp_id, int *A_d) {
|
||||
// Note: the printf command will only work if printf is enabled in your build.
|
||||
// printf("Hello World... from HIP thread = %u\n", omp_id);
|
||||
|
||||
A_d[omp_id] = omp_id;
|
||||
}
|
||||
|
||||
int main(int argc, char *argv[]) {
|
||||
int *A_h, *A_d;
|
||||
size_t Nbytes = NUM_THREADS * sizeof(int);
|
||||
|
||||
hipDeviceProp_t props;
|
||||
CHECK(hipGetDeviceProperties(&props, 0 /*deviceID*/));
|
||||
// printf("info: running on device %s\n", props.name);
|
||||
|
||||
A_h = reinterpret_cast<int *>(malloc(Nbytes));
|
||||
CHECK(hipMalloc(&A_d, Nbytes));
|
||||
for (int i = 0; i < NUM_THREADS; i++) {
|
||||
A_h[i] = 0;
|
||||
}
|
||||
CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
|
||||
// Beginning of parallel region
|
||||
#pragma omp parallel num_threads(NUM_THREADS)
|
||||
{
|
||||
// fprintf(stderr, "Hello World... from OMP thread = %d\n",
|
||||
// omp_get_thread_num());
|
||||
|
||||
hipLaunchKernelGGL(hip_helloworld, dim3(1), dim3(1), 0, 0,
|
||||
omp_get_thread_num(), A_d);
|
||||
}
|
||||
// Ending of parallel region
|
||||
|
||||
hipStreamSynchronize(0);
|
||||
CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
// printf("Device Results:\n");
|
||||
for (int i = 0; i < NUM_THREADS; i++) {
|
||||
// printf(" A_d[%d] = %d\n", i, A_h[i]);
|
||||
}
|
||||
|
||||
printf("PASSED!\n");
|
||||
|
||||
free(A_h);
|
||||
CHECK(hipFree(A_d));
|
||||
return 0;
|
||||
}
|
||||
@@ -0,0 +1,37 @@
|
||||
'
|
||||
--------------------------------------------------------------------------
|
||||
Running as root is *strongly* discouraged as any mistake (e.g., in
|
||||
defining TMPDIR) or bug can result in catastrophic damage to the OS
|
||||
file system, leaving your system in an unusable state.
|
||||
|
||||
We strongly suggest that you run mpirun as a non-root user.
|
||||
|
||||
You can override this protection by adding the --allow-run-as-root option
|
||||
to the cmd line or by setting two environment variables in the following way:
|
||||
the variable OMPI_ALLOW_RUN_AS_ROOT=1 to indicate the desire to override this
|
||||
protection, and OMPI_ALLOW_RUN_AS_ROOT_CONFIRM=1 to confirm the choice and
|
||||
add one more layer of certainty that you want to do so.
|
||||
We reiterate our advice against doing so - please proceed at your own risk.
|
||||
--------------------------------------------------------------------------
|
||||
'
|
||||
|
||||
MPIRUN=mpirun
|
||||
if ! command -v $MPIRUN &> /dev/null
|
||||
then
|
||||
echo "$MPIRUN could not be found. checking libs"
|
||||
if [ -f "/usr/lib64/openmpi/bin/mpirun" ]
|
||||
then
|
||||
MPIRUN=/usr/lib64/openmpi/bin/mpirun
|
||||
else
|
||||
if [ -f "/usr/lib64/mpi/gcc/openmpi2/bin/mpirun" ]
|
||||
then
|
||||
MPIRUN=/usr/lib64/mpi/gcc/openmpi2/bin/mpirun
|
||||
else
|
||||
echo "$MPIRUN could not be found. exiting"
|
||||
exit
|
||||
fi
|
||||
fi
|
||||
fi
|
||||
SCRIPT=$(realpath "$0")
|
||||
SCRIPTPATH=$(dirname "$SCRIPT")
|
||||
$MPIRUN --allow-run-as-root -np 2 $SCRIPTPATH/mpi_vectoradd mdrun -pin on -nsteps 10 -resetstep 9 -ntomp 64 -noconfout -nb gpu -bonded gpu -pme gpu -v -gpu_id 0
|
||||
@@ -0,0 +1,307 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in 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:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
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
|
||||
AUTHORS 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 IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/** \mainpage ROC Profiler Multi Queue Dependency Test
|
||||
*
|
||||
* \section introduction Introduction
|
||||
*
|
||||
* The goal of this test is to ensure ROC profiler does not go to deadlock
|
||||
* when multiple queue are created and they are dependent on each other
|
||||
*
|
||||
*/
|
||||
|
||||
#include "multiqueue_testapp.h"
|
||||
#include "src/utils/exception.h"
|
||||
|
||||
namespace fs = std::experimental::filesystem;
|
||||
std::vector<hsa_agent_t> Device::all_devices;
|
||||
|
||||
std::string GetRunningPath(std::string string_to_erase);
|
||||
|
||||
int main() {
|
||||
hsa_status_t status;
|
||||
MQDependencyTest obj;
|
||||
|
||||
// Get Agent info
|
||||
obj.DeviceDiscovery();
|
||||
|
||||
char agent_name[64];
|
||||
status = hsa_agent_get_info(gpu[0].agent, HSA_AGENT_INFO_NAME, agent_name);
|
||||
ASSERT_EQ(status, HSA_STATUS_SUCCESS);
|
||||
|
||||
// Getting Current Path
|
||||
std::string app_path = GetRunningPath("tests/featuretests/profiler/apps/multiqueue_testapp");
|
||||
// Getting hasco Path
|
||||
std::string ko_path =
|
||||
app_path + "tests/featuretests/profiler/" + std::string(agent_name) + "_copy.hsaco";
|
||||
|
||||
MQDependencyTest::CodeObject code_object;
|
||||
if (!obj.LoadCodeObject(ko_path, gpu[0].agent, code_object)) {
|
||||
printf("Kernel file not found or not usable with given agent.\n");
|
||||
abort();
|
||||
}
|
||||
|
||||
MQDependencyTest::Kernel copyA;
|
||||
if (!obj.GetKernel(code_object, "copyA", gpu[0].agent, copyA)) {
|
||||
printf("Test kernel A not found.\n");
|
||||
abort();
|
||||
}
|
||||
MQDependencyTest::Kernel copyB;
|
||||
if (!obj.GetKernel(code_object, "copyB", gpu[0].agent, copyB)) {
|
||||
printf("Test kernel B not found.\n");
|
||||
abort();
|
||||
}
|
||||
MQDependencyTest::Kernel copyC;
|
||||
if (!obj.GetKernel(code_object, "copyC", gpu[0].agent, copyC)) {
|
||||
printf("Test kernel C not found.\n");
|
||||
abort();
|
||||
}
|
||||
|
||||
struct args_t {
|
||||
uint32_t* a;
|
||||
uint32_t* b;
|
||||
MQDependencyTest::OCLHiddenArgs hidden;
|
||||
};
|
||||
|
||||
args_t* args;
|
||||
args = static_cast<args_t*>(obj.hsaMalloc(sizeof(args_t), kernarg));
|
||||
memset(args, 0, sizeof(args_t));
|
||||
|
||||
uint32_t* a = static_cast<uint32_t*>(obj.hsaMalloc(64 * sizeof(uint32_t), kernarg));
|
||||
uint32_t* b = static_cast<uint32_t*>(obj.hsaMalloc(64 * sizeof(uint32_t), kernarg));
|
||||
|
||||
memset(a, 0, 64 * sizeof(uint32_t));
|
||||
memset(b, 1, 64 * sizeof(uint32_t));
|
||||
|
||||
// Create queue in gpu agent and prepare a kernel dispatch packet
|
||||
hsa_queue_t* queue1;
|
||||
status = hsa_queue_create(gpu[0].agent, 1024, HSA_QUEUE_TYPE_SINGLE, NULL, NULL, UINT32_MAX,
|
||||
UINT32_MAX, &queue1);
|
||||
ASSERT_EQ(status, HSA_STATUS_SUCCESS);
|
||||
|
||||
// Create a signal with a value of 1 and attach it to the first kernel
|
||||
// dispatch packet
|
||||
hsa_signal_t completion_signal_1;
|
||||
status = hsa_signal_create(1, 0, NULL, &completion_signal_1);
|
||||
ASSERT_EQ(status, HSA_STATUS_SUCCESS);
|
||||
|
||||
// First dispath packet on queue 1, Kernel A
|
||||
{
|
||||
MQDependencyTest::Aql packet{};
|
||||
packet.header.type = HSA_PACKET_TYPE_KERNEL_DISPATCH;
|
||||
packet.header.barrier = 1;
|
||||
packet.header.acquire = HSA_FENCE_SCOPE_SYSTEM;
|
||||
packet.header.release = HSA_FENCE_SCOPE_SYSTEM;
|
||||
|
||||
packet.dispatch.setup = 1;
|
||||
packet.dispatch.workgroup_size_x = 64;
|
||||
packet.dispatch.workgroup_size_y = 1;
|
||||
packet.dispatch.workgroup_size_z = 1;
|
||||
packet.dispatch.grid_size_x = 64;
|
||||
packet.dispatch.grid_size_y = 1;
|
||||
packet.dispatch.grid_size_z = 1;
|
||||
|
||||
packet.dispatch.group_segment_size = copyA.group;
|
||||
packet.dispatch.private_segment_size = copyA.scratch;
|
||||
packet.dispatch.kernel_object = copyA.handle;
|
||||
|
||||
packet.dispatch.kernarg_address = args;
|
||||
packet.dispatch.completion_signal = completion_signal_1;
|
||||
|
||||
args->a = a;
|
||||
args->b = b;
|
||||
// Tell packet processor of A to launch the first kernel dispatch packet
|
||||
obj.SubmitPacket(queue1, packet);
|
||||
}
|
||||
|
||||
// Create a signal with a value of 1 and attach it to the second kernel
|
||||
// dispatch packet
|
||||
hsa_signal_t completion_signal_2;
|
||||
status = hsa_signal_create(1, 0, NULL, &completion_signal_2);
|
||||
ASSERT_EQ(status, HSA_STATUS_SUCCESS);
|
||||
|
||||
hsa_signal_t completion_signal_3;
|
||||
status = hsa_signal_create(1, 0, NULL, &completion_signal_3);
|
||||
ASSERT_EQ(status, HSA_STATUS_SUCCESS);
|
||||
|
||||
// Create barrier-AND packet that is enqueued in queue 1
|
||||
{
|
||||
MQDependencyTest::Aql packet{};
|
||||
packet.header.type = HSA_PACKET_TYPE_BARRIER_AND;
|
||||
packet.header.barrier = 1;
|
||||
packet.header.acquire = HSA_FENCE_SCOPE_SYSTEM;
|
||||
packet.header.release = HSA_FENCE_SCOPE_SYSTEM;
|
||||
|
||||
packet.barrier_and.dep_signal[0] = completion_signal_2;
|
||||
obj.SubmitPacket(queue1, packet);
|
||||
}
|
||||
|
||||
// Second dispath packet on queue 1, Kernel C
|
||||
{
|
||||
MQDependencyTest::Aql packet{};
|
||||
packet.header.type = HSA_PACKET_TYPE_KERNEL_DISPATCH;
|
||||
packet.header.barrier = 1;
|
||||
packet.header.acquire = HSA_FENCE_SCOPE_SYSTEM;
|
||||
packet.header.release = HSA_FENCE_SCOPE_SYSTEM;
|
||||
|
||||
packet.dispatch.setup = 1;
|
||||
packet.dispatch.workgroup_size_x = 64;
|
||||
packet.dispatch.workgroup_size_y = 1;
|
||||
packet.dispatch.workgroup_size_z = 1;
|
||||
packet.dispatch.grid_size_x = 64;
|
||||
packet.dispatch.grid_size_y = 1;
|
||||
packet.dispatch.grid_size_z = 1;
|
||||
|
||||
packet.dispatch.group_segment_size = copyC.group;
|
||||
packet.dispatch.private_segment_size = copyC.scratch;
|
||||
packet.dispatch.kernel_object = copyC.handle;
|
||||
packet.dispatch.completion_signal = completion_signal_3;
|
||||
packet.dispatch.kernarg_address = args;
|
||||
|
||||
args->a = a;
|
||||
args->b = b;
|
||||
// Tell packet processor to launch the second kernel dispatch packet
|
||||
obj.SubmitPacket(queue1, packet);
|
||||
}
|
||||
|
||||
// Create queue 2
|
||||
hsa_queue_t* queue2;
|
||||
status = hsa_queue_create(gpu[0].agent, 1024, HSA_QUEUE_TYPE_SINGLE, NULL, NULL, UINT32_MAX,
|
||||
UINT32_MAX, &queue2);
|
||||
ASSERT_EQ(status, HSA_STATUS_SUCCESS);
|
||||
|
||||
// Create barrier-AND packet that is enqueued in queue 2
|
||||
{
|
||||
MQDependencyTest::Aql packet{};
|
||||
packet.header.type = HSA_PACKET_TYPE_BARRIER_AND;
|
||||
packet.header.barrier = 1;
|
||||
packet.header.acquire = HSA_FENCE_SCOPE_SYSTEM;
|
||||
packet.header.release = HSA_FENCE_SCOPE_SYSTEM;
|
||||
|
||||
packet.barrier_and.dep_signal[0] = completion_signal_1;
|
||||
obj.SubmitPacket(queue2, packet);
|
||||
}
|
||||
|
||||
// Third dispath packet on queue 2, Kernel B
|
||||
{
|
||||
MQDependencyTest::Aql packet{};
|
||||
packet.header.type = HSA_PACKET_TYPE_KERNEL_DISPATCH;
|
||||
packet.header.barrier = 1;
|
||||
packet.header.acquire = HSA_FENCE_SCOPE_SYSTEM;
|
||||
packet.header.release = HSA_FENCE_SCOPE_SYSTEM;
|
||||
|
||||
packet.dispatch.setup = 1;
|
||||
packet.dispatch.workgroup_size_x = 64;
|
||||
packet.dispatch.workgroup_size_y = 1;
|
||||
packet.dispatch.workgroup_size_z = 1;
|
||||
packet.dispatch.grid_size_x = 64;
|
||||
packet.dispatch.grid_size_y = 1;
|
||||
packet.dispatch.grid_size_z = 1;
|
||||
|
||||
packet.dispatch.group_segment_size = copyB.group;
|
||||
packet.dispatch.private_segment_size = copyB.scratch;
|
||||
packet.dispatch.kernel_object = copyB.handle;
|
||||
|
||||
packet.dispatch.kernarg_address = args;
|
||||
packet.dispatch.completion_signal = completion_signal_2;
|
||||
|
||||
args->a = a;
|
||||
args->b = b;
|
||||
// Tell packet processor to launch the third kernel dispatch packet
|
||||
obj.SubmitPacket(queue2, packet);
|
||||
}
|
||||
|
||||
// Wait on the completion signal
|
||||
hsa_signal_wait_relaxed(completion_signal_1, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX,
|
||||
HSA_WAIT_STATE_BLOCKED);
|
||||
|
||||
// Wait on the completion signal
|
||||
hsa_signal_wait_relaxed(completion_signal_2, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX,
|
||||
HSA_WAIT_STATE_BLOCKED);
|
||||
|
||||
// Wait on the completion signal
|
||||
hsa_signal_wait_relaxed(completion_signal_3, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX,
|
||||
HSA_WAIT_STATE_BLOCKED);
|
||||
|
||||
for (int i = 0; i < 64; i++) {
|
||||
if (a[i] != b[i]) {
|
||||
printf("error at %d: expected %d, got %d\n", i, b[i], a[i]);
|
||||
abort();
|
||||
}
|
||||
}
|
||||
|
||||
// Clearing data structures and memory
|
||||
status = hsa_signal_destroy(completion_signal_1);
|
||||
ASSERT_EQ(status, HSA_STATUS_SUCCESS);
|
||||
|
||||
status = hsa_signal_destroy(completion_signal_2);
|
||||
ASSERT_EQ(status, HSA_STATUS_SUCCESS);
|
||||
|
||||
status = hsa_signal_destroy(completion_signal_3);
|
||||
ASSERT_EQ(status, HSA_STATUS_SUCCESS);
|
||||
|
||||
if (queue1 != nullptr) {
|
||||
status = hsa_queue_destroy(queue1);
|
||||
ASSERT_EQ(status, HSA_STATUS_SUCCESS);
|
||||
}
|
||||
|
||||
if (queue2 != nullptr) {
|
||||
status = hsa_queue_destroy(queue2);
|
||||
ASSERT_EQ(status, HSA_STATUS_SUCCESS);
|
||||
}
|
||||
|
||||
status = hsa_memory_free(a);
|
||||
ASSERT_EQ(status, HSA_STATUS_SUCCESS);
|
||||
status = hsa_memory_free(b);
|
||||
ASSERT_EQ(status, HSA_STATUS_SUCCESS);
|
||||
|
||||
status = hsa_executable_destroy(code_object.executable);
|
||||
ASSERT_EQ(status, HSA_STATUS_SUCCESS);
|
||||
|
||||
status = hsa_code_object_reader_destroy(code_object.code_obj_rdr);
|
||||
ASSERT_EQ(status, HSA_STATUS_SUCCESS);
|
||||
close(code_object.file);
|
||||
}
|
||||
|
||||
// This function returns the running path of executable
|
||||
std::string GetRunningPath(std::string string_to_erase) {
|
||||
std::string path;
|
||||
char* real_path;
|
||||
Dl_info dl_info;
|
||||
|
||||
if (0 != dladdr(reinterpret_cast<void*>(main), &dl_info)) {
|
||||
std::string to_erase = string_to_erase;
|
||||
path = dl_info.dli_fname;
|
||||
real_path = realpath(path.c_str(), NULL);
|
||||
if (real_path == nullptr) {
|
||||
throw(std::string("Error! in extracting real path"));
|
||||
}
|
||||
path.clear(); // reset path
|
||||
path.append(real_path);
|
||||
|
||||
size_t pos = path.find(to_erase);
|
||||
if (pos != std::string::npos) path.erase(pos, to_erase.length());
|
||||
} else {
|
||||
throw(std::string("Error! in extracting real path"));
|
||||
}
|
||||
return path;
|
||||
}
|
||||
@@ -0,0 +1,330 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in 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:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
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
|
||||
AUTHORS 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 IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
#ifndef TESTS_FEATURETESTS_PROFILER_DISCRETETESTS_BINARY_MULTIQUEUE_TESTAPP_H_
|
||||
#define TESTS_FEATURETESTS_PROFILER_DISCRETETESTS_BINARY_MULTIQUEUE_TESTAPP_H_
|
||||
|
||||
#include <assert.h>
|
||||
#include <dlfcn.h>
|
||||
#include <fcntl.h>
|
||||
#include <hsa/hsa.h>
|
||||
#include <hsa/hsa_api_trace.h>
|
||||
#include <stdio.h>
|
||||
#include <string.h>
|
||||
#include <unistd.h>
|
||||
#include <cstdlib>
|
||||
|
||||
#include <experimental/filesystem>
|
||||
#include <iostream>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
#include "src/utils/exception.h"
|
||||
|
||||
#define ASSERT_EQ(val1, val2) \
|
||||
do { \
|
||||
if ((val1) != val2) { \
|
||||
assert(false); \
|
||||
abort(); \
|
||||
} \
|
||||
} while (false)
|
||||
|
||||
struct Device {
|
||||
struct Memory {
|
||||
hsa_amd_memory_pool_t pool;
|
||||
bool fine;
|
||||
bool kernarg;
|
||||
size_t size;
|
||||
size_t granule;
|
||||
};
|
||||
|
||||
hsa_agent_t agent;
|
||||
char name[64];
|
||||
std::vector<Memory> pools;
|
||||
uint32_t fine;
|
||||
uint32_t coarse;
|
||||
static std::vector<hsa_agent_t> all_devices;
|
||||
};
|
||||
|
||||
std::vector<Device> cpu, gpu;
|
||||
Device::Memory kernarg;
|
||||
|
||||
class MQDependencyTest {
|
||||
public:
|
||||
MQDependencyTest() { hsa_init(); }
|
||||
~MQDependencyTest() { hsa_shut_down(); }
|
||||
|
||||
struct CodeObject {
|
||||
hsa_file_t file;
|
||||
hsa_code_object_reader_t code_obj_rdr;
|
||||
hsa_executable_t executable;
|
||||
};
|
||||
|
||||
struct Kernel {
|
||||
uint64_t handle;
|
||||
uint32_t scratch;
|
||||
uint32_t group;
|
||||
uint32_t kernarg_size;
|
||||
uint32_t kernarg_align;
|
||||
};
|
||||
|
||||
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;
|
||||
};
|
||||
|
||||
bool LoadCodeObject(std::string filename, hsa_agent_t agent, CodeObject& code_object) {
|
||||
hsa_status_t err;
|
||||
printf("%s", filename.c_str());
|
||||
code_object.file = open(filename.c_str(), O_RDONLY);
|
||||
if (code_object.file == -1) {
|
||||
abort();
|
||||
return false;
|
||||
}
|
||||
|
||||
err = hsa_code_object_reader_create_from_file(code_object.file, &code_object.code_obj_rdr);
|
||||
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
|
||||
|
||||
err = hsa_executable_create_alt(HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT,
|
||||
nullptr, &code_object.executable);
|
||||
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
|
||||
|
||||
err = hsa_executable_load_agent_code_object(code_object.executable, agent,
|
||||
code_object.code_obj_rdr, nullptr, nullptr);
|
||||
if (err != HSA_STATUS_SUCCESS) return false;
|
||||
|
||||
err = hsa_executable_freeze(code_object.executable, nullptr);
|
||||
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
bool GetKernel(const CodeObject& code_object, std::string kernel, hsa_agent_t agent,
|
||||
Kernel& kern) {
|
||||
hsa_executable_symbol_t symbol;
|
||||
hsa_status_t err =
|
||||
hsa_executable_get_symbol_by_name(code_object.executable, kernel.c_str(), &agent, &symbol);
|
||||
if (err != HSA_STATUS_SUCCESS) {
|
||||
err = hsa_executable_get_symbol_by_name(code_object.executable, (kernel + ".kd").c_str(),
|
||||
&agent, &symbol);
|
||||
if (err != HSA_STATUS_SUCCESS) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
// printf("\nkernel-name: %s\n", kernel.c_str());
|
||||
err = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT,
|
||||
&kern.handle);
|
||||
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
|
||||
|
||||
err = hsa_executable_symbol_get_info(
|
||||
symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, &kern.scratch);
|
||||
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
|
||||
// printf("Scratch: %d\n", kern.scratch);
|
||||
|
||||
err = hsa_executable_symbol_get_info(
|
||||
symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, &kern.group);
|
||||
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
|
||||
// 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);
|
||||
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
|
||||
// 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);
|
||||
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
|
||||
// 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 = static_cast<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);
|
||||
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
|
||||
|
||||
err = hsa_amd_agents_allow_access(Device::all_devices.size(), &Device::all_devices[0], nullptr,
|
||||
ret);
|
||||
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
|
||||
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]);
|
||||
}
|
||||
|
||||
bool DeviceDiscovery() {
|
||||
hsa_status_t err;
|
||||
err = hsa_iterate_agents(
|
||||
[](hsa_agent_t agent, void*) {
|
||||
hsa_status_t err;
|
||||
|
||||
Device dev;
|
||||
dev.agent = agent;
|
||||
|
||||
dev.fine = -1u;
|
||||
dev.coarse = -1u;
|
||||
|
||||
err = hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, dev.name);
|
||||
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
|
||||
|
||||
hsa_device_type_t type;
|
||||
err = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type);
|
||||
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
|
||||
|
||||
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);
|
||||
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
|
||||
|
||||
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);
|
||||
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
|
||||
|
||||
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);
|
||||
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
|
||||
|
||||
err = hsa_amd_memory_pool_get_info(
|
||||
pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &mem.granule);
|
||||
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
|
||||
|
||||
pools.push_back(mem);
|
||||
return HSA_STATUS_SUCCESS;
|
||||
},
|
||||
static_cast<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)
|
||||
cpu.push_back(dev);
|
||||
else
|
||||
gpu.push_back(dev);
|
||||
|
||||
Device::all_devices.push_back(dev.agent);
|
||||
}
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
},
|
||||
nullptr);
|
||||
|
||||
[]() {
|
||||
for (auto& dev : cpu) {
|
||||
for (auto& mem : dev.pools) {
|
||||
if (mem.fine && mem.kernarg) {
|
||||
kernarg = mem;
|
||||
return;
|
||||
}
|
||||
}
|
||||
}
|
||||
}();
|
||||
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
|
||||
|
||||
if (cpu.empty() || gpu.empty() || kernarg.pool.handle == 0) return false;
|
||||
return true;
|
||||
}
|
||||
};
|
||||
#endif // TESTS_FEATURETESTS_PROFILER_DISCRETETESTS_BINARY_MULTIQUEUE_TESTAPP_H_
|
||||
@@ -0,0 +1,83 @@
|
||||
/******************************************************************************
|
||||
Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in 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:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
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
|
||||
AUTHORS 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 IN
|
||||
THE SOFTWARE.
|
||||
*******************************************************************************/
|
||||
|
||||
/** \mainpage ROC Profiler Multi-Threaded Test Application
|
||||
*
|
||||
* \section introduction Introduction
|
||||
*
|
||||
* Test application launches an empty kernel on multiple threads.
|
||||
*
|
||||
* In subsequent tests, ROC profiler is run against this applicaiton
|
||||
* to confirm if collected contexts are valid.
|
||||
*
|
||||
*/
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
#include <functional>
|
||||
#include <thread>
|
||||
#include <vector>
|
||||
|
||||
#include "utils/test_utils.h"
|
||||
|
||||
/** \mainpage ROC Profiler Test APplication
|
||||
*
|
||||
* \section introduction Introduction
|
||||
*
|
||||
* The goal of this test application is to launch an empty kernel
|
||||
* on multiple threads and multiple gpu's.
|
||||
*
|
||||
* Number of threads are caluculated based on the cores in the system
|
||||
* Number of gpus's are calculated based on the gpu's in the system
|
||||
*/
|
||||
|
||||
// empty kernel
|
||||
__global__ void kernel() {}
|
||||
|
||||
// launches kernel on multiple gpu's
|
||||
void KernelLaunch() {
|
||||
// Multi-GPU
|
||||
int gpu_count = 0;
|
||||
hipGetDeviceCount(&gpu_count);
|
||||
|
||||
for (uint32_t gpu_id = 0; gpu_id < gpu_count; gpu_id++) {
|
||||
// run empty kernel
|
||||
kernel<<<1, 1>>>();
|
||||
}
|
||||
}
|
||||
|
||||
int main(int argc, char** argv) {
|
||||
// create as many threads as number of cores in system
|
||||
int threads_count = GetNumberOfCores();
|
||||
|
||||
// create a pool of thrads
|
||||
std::vector<std::thread> threads(threads_count);
|
||||
|
||||
// launch kernel on each thread
|
||||
for (int n = 0; n < threads_count; ++n) {
|
||||
threads[n] = std::thread(KernelLaunch);
|
||||
}
|
||||
// wait for all kernel launches to complete
|
||||
for (int n = 0; n < threads_count; ++n) {
|
||||
threads[n].join();
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,130 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in 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:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
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
|
||||
AUTHORS 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 IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
#include <assert.h>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
#include <algorithm>
|
||||
#include <iostream>
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
#define HIP_ASSERT(x) (assert((x) == hipSuccess))
|
||||
|
||||
#define WIDTH 1024
|
||||
#define HEIGHT 1024
|
||||
|
||||
#define NUM (WIDTH * HEIGHT)
|
||||
|
||||
#define THREADS_PER_BLOCK_X 16
|
||||
#define THREADS_PER_BLOCK_Y 16
|
||||
#define THREADS_PER_BLOCK_Z 1
|
||||
|
||||
__global__ void vectoradd_float(float *__restrict__ a,
|
||||
const float *__restrict__ b,
|
||||
const float *__restrict__ c, int width,
|
||||
int height) {
|
||||
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
|
||||
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
|
||||
|
||||
int i = y * width + x;
|
||||
if (i < (width * height)) {
|
||||
a[i] = b[i] + c[i];
|
||||
}
|
||||
}
|
||||
|
||||
int main() {
|
||||
float *hostA;
|
||||
float *hostB;
|
||||
float *hostC;
|
||||
|
||||
float *deviceA;
|
||||
float *deviceB;
|
||||
float *deviceC;
|
||||
|
||||
hipDeviceProp_t devProp;
|
||||
hipGetDeviceProperties(&devProp, 0);
|
||||
std::cout << " System minor " << devProp.minor << std::endl;
|
||||
std::cout << " System major " << devProp.major << std::endl;
|
||||
std::cout << " agent prop name " << devProp.name << std::endl;
|
||||
|
||||
std::cout << "hip Device prop succeeded " << std::endl;
|
||||
|
||||
int i;
|
||||
int errors;
|
||||
|
||||
hostA = reinterpret_cast<float *>(malloc(NUM * sizeof(float)));
|
||||
hostB = reinterpret_cast<float *>(malloc(NUM * sizeof(float)));
|
||||
hostC = reinterpret_cast<float *>(malloc(NUM * sizeof(float)));
|
||||
|
||||
// initialize the input data
|
||||
for (i = 0; i < NUM; i++) {
|
||||
hostB[i] = static_cast<float>(i);
|
||||
hostC[i] = static_cast<float>(i) * 100.0f;
|
||||
}
|
||||
|
||||
HIP_ASSERT(
|
||||
hipMalloc(reinterpret_cast<void **>(&deviceA), NUM * sizeof(float)));
|
||||
HIP_ASSERT(
|
||||
hipMalloc(reinterpret_cast<void **>(&deviceB), NUM * sizeof(float)));
|
||||
HIP_ASSERT(
|
||||
hipMalloc(reinterpret_cast<void **>(&deviceC), NUM * sizeof(float)));
|
||||
|
||||
HIP_ASSERT(
|
||||
hipMemcpy(deviceB, hostB, NUM * sizeof(float), hipMemcpyHostToDevice));
|
||||
HIP_ASSERT(
|
||||
hipMemcpy(deviceC, hostC, NUM * sizeof(float), hipMemcpyHostToDevice));
|
||||
|
||||
hipLaunchKernelGGL(
|
||||
vectoradd_float,
|
||||
dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT / THREADS_PER_BLOCK_Y),
|
||||
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, deviceA, deviceB,
|
||||
deviceC, WIDTH, HEIGHT);
|
||||
|
||||
HIP_ASSERT(
|
||||
hipMemcpy(hostA, deviceA, NUM * sizeof(float), hipMemcpyDeviceToHost));
|
||||
|
||||
// verify the results
|
||||
errors = 0;
|
||||
for (i = 0; i < NUM; i++) {
|
||||
if (hostA[i] != (hostB[i] + hostC[i])) {
|
||||
errors++;
|
||||
}
|
||||
}
|
||||
if (errors != 0) {
|
||||
printf("FAILED: %d errors\n", errors);
|
||||
} else {
|
||||
printf("PASSED!\n");
|
||||
}
|
||||
|
||||
HIP_ASSERT(hipFree(deviceA));
|
||||
HIP_ASSERT(hipFree(deviceB));
|
||||
HIP_ASSERT(hipFree(deviceC));
|
||||
|
||||
free(hostA);
|
||||
free(hostB);
|
||||
free(hostC);
|
||||
|
||||
// hipResetDefaultAccelerator();
|
||||
|
||||
return errors;
|
||||
}
|
||||
@@ -0,0 +1,132 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in 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:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
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
|
||||
AUTHORS 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 IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
#include <mpi.h>
|
||||
|
||||
#include <math.h>
|
||||
#include <stdio.h>
|
||||
|
||||
#include <iostream>
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
#define HIP_RC(call) \
|
||||
do { \
|
||||
hipError_t err = call; \
|
||||
if (hipSuccess != err) { \
|
||||
printf("HIP ERROR (code = %d, %s) at %s:%d\n", err, \
|
||||
hipGetErrorString(err), __FILE__, __LINE__); \
|
||||
assert(0); \
|
||||
exit(1); \
|
||||
} \
|
||||
} while (0)
|
||||
|
||||
#define HIP_KL(call) \
|
||||
do { \
|
||||
call; \
|
||||
hipError_t err = hipGetLastError(); \
|
||||
if (hipSuccess != err) { \
|
||||
printf("HIP ERROR (code = %d, %s) at %s:%d\n", err, \
|
||||
hipGetErrorString(err), __FILE__, __LINE__); \
|
||||
assert(0); \
|
||||
exit(1); \
|
||||
} \
|
||||
} while (0)
|
||||
|
||||
// CUDA kernel to add elements of two arrays
|
||||
__global__ void add(int n, float *x, float *y) {
|
||||
int index = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int stride = blockDim.x * gridDim.x;
|
||||
for (int i = index; i < n; i += stride)
|
||||
y[i] = x[i] + y[i];
|
||||
}
|
||||
|
||||
int main(int argc, char *argv[]) {
|
||||
int N = 1 << 20;
|
||||
float *x = new float[N];
|
||||
float *y = new float[N];
|
||||
float *d_x;
|
||||
float *d_y;
|
||||
|
||||
int myId;
|
||||
int devId;
|
||||
int numRank;
|
||||
int deviceCount;
|
||||
|
||||
// init MPI
|
||||
MPI_Init(&argc, &argv);
|
||||
MPI_Comm_rank(MPI_COMM_WORLD, &myId);
|
||||
MPI_Comm_size(MPI_COMM_WORLD, &numRank);
|
||||
hipGetDeviceCount(&deviceCount);
|
||||
|
||||
std::cout << "device count and rank is" << deviceCount << ": " << numRank
|
||||
<< std::endl;
|
||||
|
||||
// set the device ID to the rank ID mod deviceCount (in this case 4 since
|
||||
// there are 4 devices on a node)
|
||||
devId = myId % deviceCount;
|
||||
// set the device ID
|
||||
hipSetDevice(devId);
|
||||
|
||||
printf("Rank Id: %d | Device Id : %d | Num Devices: %d\n", myId, devId,
|
||||
deviceCount);
|
||||
fflush(stdout);
|
||||
|
||||
// Allocate Unified Memory -- accessible from CPU or GPU
|
||||
hipMallocManaged(&d_x, N * sizeof(float));
|
||||
hipMallocManaged(&d_y, N * sizeof(float));
|
||||
|
||||
// initialize x and y arrays on the host
|
||||
for (int i = 0; i < N; i++) {
|
||||
x[i] = 1.0f;
|
||||
y[i] = 2.0f;
|
||||
}
|
||||
|
||||
HIP_RC(hipMemcpy(d_x, x, N * sizeof(float), hipMemcpyHostToDevice));
|
||||
HIP_RC(hipMemcpy(d_y, y, N * sizeof(float), hipMemcpyHostToDevice));
|
||||
|
||||
// Launch kernel on 1M elements on the GPU
|
||||
int blockSize = 256;
|
||||
int numBlocks = (N + blockSize - 1) / blockSize;
|
||||
HIP_KL(hipLaunchKernelGGL(add, numBlocks, blockSize, 0, 0, N, d_x, d_y));
|
||||
|
||||
// Wait for GPU to finish before accessing on host
|
||||
HIP_RC(hipDeviceSynchronize());
|
||||
|
||||
HIP_RC(hipMemcpy(x, d_x, N * sizeof(float), hipMemcpyDeviceToHost));
|
||||
HIP_RC(hipMemcpy(y, d_y, N * sizeof(float), hipMemcpyDeviceToHost));
|
||||
|
||||
// Check for errors (all values should be 3.0f)
|
||||
float maxError = 0.0f;
|
||||
for (int i = 0; i < N; i++)
|
||||
maxError = fmax(maxError, fabs(y[i] - 3.0f));
|
||||
printf("Max error: %f\n", maxError);
|
||||
|
||||
// Free memory
|
||||
HIP_RC(hipFree(d_x));
|
||||
HIP_RC(hipFree(d_y));
|
||||
|
||||
delete[] x;
|
||||
delete[] y;
|
||||
|
||||
MPI_Finalize();
|
||||
return 0;
|
||||
}
|
||||
Reference in New Issue
Block a user