AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

Public C++ header files and samples updates (#819)

Przeglądaj źródła 
* Public C++ header files (source/include/rocprofiler-sdk/cxx)

* Update samples/api_buffered_tracing

- scratch memory and page migration
- README

* Update samples/api_buffered_tracing

- page migration component in sample

* Update tests/page-migration/validate.py

- fix checks for page migration operation names

* Update tests/page-migration/validate.py

- fix get_allocated_pages

* Update scratch memory and page migration validations

* Fix include/rocprofiler-sdk/cxx installation

* Rework include/rocprofiler-sdk/cxx

- Improve name_info to support const char*, string_view, string

* Update samples/api_{buffered,callback}_tracing

* External correlation ID request sample

- includes correlation ID retirement demo

* Update samples/api_buffered_tracing/README.md

* Update lib/rocprofiler-sdk/hsa/queue.cpp

- generate correlation ID for kernel launch if one doesn't exist

* Remove priority check from tool libraries (samples/tests)

- if(priority > 0) return nullptr check in rocprofiler_configure has proliferated beyond its intended use

* Apply suggestions from code review

[ROCm/rocprofiler-sdk commit: de13d2ac5d]
This commit is contained in:
Jonathan R. Madsen
2024-04-25 20:09:11 -05:00
zatwierdzone przez GitHub
rodzic 9b747eb63c
commit f167317524
37 zmienionych plików z 2232 dodań i 512 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
projects/rocprofiler-sdk/samples/api_buffered_tracing/README.md
+18
Wyświetl plik
@@ -0,0 +1,18 @@
# API Buffer Tracing Sample
## Services
- Code object callback tracing for mapping kernel IDs to kernel names
- HSA API (Core, AMD Ext)
- HIP API (Runtime)
- Kernel dispatch
- Memory copy
- Page Migration
- Scratch Memory
## Properties
- Buffer size of 4096 bytes which is automatically flushed once >= 87.5% of buffer is filled (3584 bytes)
- Creation of dedicated thread for buffer callback delivery
- Push external correlation IDs once per thread (value is thread ID)
- Receives notifications for internal thread creation
projects/rocprofiler-sdk/samples/api_buffered_tracing/client.cpp
+160 -49
Wyświetl plik
@@ -50,6 +50,7 @@
#include <atomic>
#include <cassert>
#include <chrono>
#include <cmath>
#include <cstddef>
#include <cstdint>
#include <cstdio>
@@ -148,21 +149,29 @@ tool_tracing_callback(rocprofiler_context_id_t context,
{
auto* header = headers[i];
if(header == nullptr)
auto kind_name = std::string{};
if(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING)
{
throw std::runtime_error{
"rocprofiler provided a null pointer to header. this should never happen"};
const char* _name = nullptr;
auto _kind = static_cast<rocprofiler_buffer_tracing_kind_t>(header->kind);
ROCPROFILER_CALL(rocprofiler_query_buffer_tracing_kind_name(_kind, &_name, nullptr),
"query buffer tracing kind name");
if(_name)
{
static size_t len = 15;
kind_name = std::string{_name};
len = std::max(len, kind_name.length());
kind_name.resize(len, ' ');
kind_name += " :: ";
}
}
else if(header->hash !=
rocprofiler_record_header_compute_hash(header->category, header->kind))
{
throw std::runtime_error{"rocprofiler_record_header_t (category | kind) != hash"};
}
else if(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING &&
(header->kind == ROCPROFILER_BUFFER_TRACING_HSA_CORE_API ||
header->kind == ROCPROFILER_BUFFER_TRACING_HSA_AMD_EXT_API ||
header->kind == ROCPROFILER_BUFFER_TRACING_HSA_IMAGE_EXT_API ||
header->kind == ROCPROFILER_BUFFER_TRACING_HSA_FINALIZE_EXT_API))
if(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING &&
(header->kind == ROCPROFILER_BUFFER_TRACING_HSA_CORE_API ||
header->kind == ROCPROFILER_BUFFER_TRACING_HSA_AMD_EXT_API ||
header->kind == ROCPROFILER_BUFFER_TRACING_HSA_IMAGE_EXT_API ||
header->kind == ROCPROFILER_BUFFER_TRACING_HSA_FINALIZE_EXT_API))
{
auto* record =
static_cast<rocprofiler_buffer_tracing_hsa_api_record_t*>(header->payload);
@@ -173,7 +182,7 @@ tool_tracing_callback(rocprofiler_context_id_t context,
<< ", extern_cid=" << record->correlation_id.external.value
<< ", kind=" << record->kind << ", operation=" << record->operation
<< ", start=" << record->start_timestamp << ", stop=" << record->end_timestamp
<< ", name=" << client_name_info.operation_names[record->kind][record->operation];
<< ", name=" << client_name_info.at(record->kind, record->operation);
if(record->start_timestamp > record->end_timestamp)
{
@@ -186,7 +195,7 @@ tool_tracing_callback(rocprofiler_context_id_t context,
}
static_cast<call_stack_t*>(user_data)->emplace_back(
source_location{__FUNCTION__, __FILE__, __LINE__, info.str()});
source_location{__FUNCTION__, __FILE__, __LINE__, kind_name + info.str()});
}
else if(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING &&
header->kind == ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API)
@@ -200,7 +209,7 @@ tool_tracing_callback(rocprofiler_context_id_t context,
<< ", extern_cid=" << record->correlation_id.external.value
<< ", kind=" << record->kind << ", operation=" << record->operation
<< ", start=" << record->start_timestamp << ", stop=" << record->end_timestamp
<< ", name=" << client_name_info.operation_names[record->kind][record->operation];
<< ", name=" << client_name_info[record->kind][record->operation];
if(record->start_timestamp > record->end_timestamp)
{
@@ -213,7 +222,7 @@ tool_tracing_callback(rocprofiler_context_id_t context,
}
static_cast<call_stack_t*>(user_data)->emplace_back(
source_location{__FUNCTION__, __FILE__, __LINE__, info.str()});
source_location{__FUNCTION__, __FILE__, __LINE__, kind_name + info.str()});
}
else if(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING &&
header->kind == ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH)
@@ -223,15 +232,16 @@ tool_tracing_callback(rocprofiler_context_id_t context,
auto info = std::stringstream{};
info << "agent_id=" << record->dispatch_info.agent_id.handle
info << "tid=" << record->thread_id << ", context=" << context.handle
<< ", buffer_id=" << buffer_id.handle
<< ", cid=" << record->correlation_id.internal
<< ", extern_cid=" << record->correlation_id.external.value
<< ", kind=" << record->kind << ", operation=" << record->operation
<< ", agent_id=" << record->dispatch_info.agent_id.handle
<< ", queue_id=" << record->dispatch_info.queue_id.handle
<< ", kernel_id=" << record->dispatch_info.kernel_id
<< ", kernel=" << client_kernels.at(record->dispatch_info.kernel_id).kernel_name
<< ", context=" << context.handle << ", buffer_id=" << buffer_id.handle
<< ", cid=" << record->correlation_id.internal
<< ", extern_cid=" << record->correlation_id.external.value
<< ", kind=" << record->kind << ", start=" << record->start_timestamp
<< ", stop=" << record->end_timestamp
<< ", start=" << record->start_timestamp << ", stop=" << record->end_timestamp
<< ", private_segment_size=" << record->dispatch_info.private_segment_size
<< ", group_segment_size=" << record->dispatch_info.group_segment_size
<< ", workgroup_size=(" << record->dispatch_info.workgroup_size.x << ","
@@ -244,7 +254,7 @@ tool_tracing_callback(rocprofiler_context_id_t context,
throw std::runtime_error("kernel dispatch: start > end");
static_cast<call_stack_t*>(user_data)->emplace_back(
source_location{__FUNCTION__, __FILE__, __LINE__, info.str()});
source_location{__FUNCTION__, __FILE__, __LINE__, kind_name + info.str()});
}
else if(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING &&
header->kind == ROCPROFILER_BUFFER_TRACING_MEMORY_COPY)
@@ -254,20 +264,111 @@ tool_tracing_callback(rocprofiler_context_id_t context,
auto info = std::stringstream{};
info << "src_agent_id=" << record->src_agent_id.handle
<< ", dst_agent_id=" << record->dst_agent_id.handle
<< ", direction=" << record->operation << ", context=" << context.handle
info << "tid=" << record->thread_id << ", context=" << context.handle
<< ", buffer_id=" << buffer_id.handle
<< ", cid=" << record->correlation_id.internal
<< ", extern_cid=" << record->correlation_id.external.value
<< ", kind=" << record->kind << ", start=" << record->start_timestamp
<< ", stop=" << record->end_timestamp;
<< ", kind=" << record->kind << ", operation=" << record->operation
<< ", src_agent_id=" << record->src_agent_id.handle
<< ", dst_agent_id=" << record->dst_agent_id.handle
<< ", direction=" << record->operation << ", start=" << record->start_timestamp
<< ", stop=" << record->end_timestamp
<< ", name=" << client_name_info.at(record->kind, record->operation);
if(record->start_timestamp > record->end_timestamp)
throw std::runtime_error("memory copy: start > end");
static_cast<call_stack_t*>(user_data)->emplace_back(
source_location{__FUNCTION__, __FILE__, __LINE__, info.str()});
source_location{__FUNCTION__, __FILE__, __LINE__, kind_name + info.str()});
}
else if(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING &&
header->kind == ROCPROFILER_BUFFER_TRACING_PAGE_MIGRATION)
{
auto* record =
static_cast<rocprofiler_buffer_tracing_page_migration_record_t*>(header->payload);
auto info = std::stringstream{};
info << "kind=" << record->kind << ", operation=" << record->operation
<< ", pid=" << record->pid << ", start=" << record->start_timestamp
<< ", stop=" << record->end_timestamp
<< ", name=" << client_name_info.at(record->kind, record->operation);
switch(record->operation)
{
case ROCPROFILER_PAGE_MIGRATION_PAGE_MIGRATE:
{
info << ", page_fault=(" << record->page_fault.read_fault << ", "
<< record->page_fault.migrated << ", " << record->page_fault.node_id
<< ", " << std::hex << "0x" << record->page_fault.address << ")";
break;
}
case ROCPROFILER_PAGE_MIGRATION_PAGE_FAULT:
{
info << ", page_migrate=(" << std::hex << "0x"
<< record->page_migrate.start_addr << ", 0x"
<< record->page_migrate.end_addr << ", " << std::dec
<< record->page_migrate.from_node << ", " << record->page_migrate.to_node
<< ", " << record->page_migrate.prefetch_node << ", "
<< record->page_migrate.preferred_node << ", "
<< record->page_migrate.trigger << ")";
break;
}
case ROCPROFILER_PAGE_MIGRATION_QUEUE_SUSPEND:
{
info << ", queue_suspend=(" << record->queue_suspend.rescheduled << ", "
<< record->queue_suspend.node_id << ", " << record->queue_suspend.trigger
<< ")";
break;
}
case ROCPROFILER_PAGE_MIGRATION_UNMAP_FROM_GPU:
{
info << ", unmap_from_gpu=(" << record->unmap_from_gpu.node_id << std::hex
<< ", 0x" << record->unmap_from_gpu.start_addr << ", 0x"
<< record->unmap_from_gpu.end_addr << ", " << std::dec
<< record->unmap_from_gpu.trigger << ")";
break;
}
case ROCPROFILER_PAGE_MIGRATION_NONE:
case ROCPROFILER_PAGE_MIGRATION_LAST:
{
throw std::runtime_error{"unexpected page migration value"};
break;
}
}
if(record->start_timestamp > record->end_timestamp)
throw std::runtime_error("page migration: start > end");
static_cast<call_stack_t*>(user_data)->emplace_back(
source_location{__FUNCTION__, __FILE__, __LINE__, kind_name + info.str()});
}
else if(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING &&
header->kind == ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY)
{
auto* record =
static_cast<rocprofiler_buffer_tracing_scratch_memory_record_t*>(header->payload);
auto info = std::stringstream{};
auto _elapsed =
std::chrono::duration_cast<std::chrono::duration<double, std::micro>>(
std::chrono::nanoseconds{record->end_timestamp - record->start_timestamp})
.count();
info << "tid=" << record->thread_id << ", context=" << context.handle
<< ", buffer_id=" << buffer_id.handle
<< ", cid=" << record->correlation_id.internal
<< ", extern_cid=" << record->correlation_id.external.value
<< ", kind=" << record->kind << ", operation=" << record->operation
<< ", agent_id=" << record->agent_id.handle
<< ", queue_id=" << record->queue_id.handle << ", thread_id=" << record->thread_id
<< ", elapsed=" << std::setprecision(3) << std::fixed << _elapsed
<< " usec, flags=" << record->flags
<< ", name=" << client_name_info.at(record->kind, record->operation);
static_cast<call_stack_t*>(user_data)->emplace_back(
source_location{__FUNCTION__, __FILE__, __LINE__, kind_name + info.str()});
}
else
{
@@ -312,25 +413,25 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
client_name_info = common::get_buffer_tracing_names();
for(const auto& itr : client_name_info.operation_names)
for(const auto& itr : client_name_info)
{
auto name_idx = std::stringstream{};
name_idx << " [" << std::setw(3) << static_cast<int32_t>(itr.first) << "]";
name_idx << " [" << std::setw(3) << itr.value << "]";
call_stack_v->emplace_back(
source_location{"rocprofiler_buffer_tracing_kind_names " + name_idx.str(),
__FILE__,
__LINE__,
client_name_info.kind_names.at(itr.first)});
std::string{itr.name}});
for(const auto& ditr : itr.second)
for(auto [didx, ditr] : itr.items())
{
auto operation_idx = std::stringstream{};
operation_idx << " [" << std::setw(3) << static_cast<int32_t>(ditr.first) << "]";
operation_idx << " [" << std::setw(3) << didx << "]";
call_stack_v->emplace_back(source_location{
"rocprofiler_buffer_tracing_kind_operation_names" + operation_idx.str(),
__FILE__,
__LINE__,
std::string{"- "} + std::string{ditr.second}});
std::string{"- "} + std::string{*ditr}});
}
}
@@ -338,28 +439,32 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
ROCPROFILER_CALL(rocprofiler_create_context(&client_ctx), "context creation");
auto code_object_ops = std::vector<rocprofiler_tracing_operation_t>{
ROCPROFILER_CODE_OBJECT_DEVICE_KERNEL_SYMBOL_REGISTER};
ROCPROFILER_CALL(
rocprofiler_configure_callback_tracing_service(client_ctx,
ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT,
nullptr,
0,
code_object_ops.data(),
code_object_ops.size(),
tool_code_object_callback,
nullptr),
"code object tracing service configure");
constexpr auto buffer_size_bytes = 4096;
constexpr auto buffer_watermark_bytes = buffer_size_bytes - (buffer_size_bytes / 8);
ROCPROFILER_CALL(rocprofiler_create_buffer(client_ctx,
4096,
2048,
buffer_size_bytes,
buffer_watermark_bytes,
ROCPROFILER_BUFFER_POLICY_LOSSLESS,
tool_tracing_callback,
tool_data,
&client_buffer),
"buffer creation");
for(auto itr : {ROCPROFILER_BUFFER_TRACING_HSA_CORE_API,
ROCPROFILER_BUFFER_TRACING_HSA_AMD_EXT_API,
ROCPROFILER_BUFFER_TRACING_HSA_IMAGE_EXT_API,
ROCPROFILER_BUFFER_TRACING_HSA_FINALIZE_EXT_API})
for(auto itr :
{ROCPROFILER_BUFFER_TRACING_HSA_CORE_API, ROCPROFILER_BUFFER_TRACING_HSA_AMD_EXT_API})
{
ROCPROFILER_CALL(rocprofiler_configure_buffer_tracing_service(
client_ctx, itr, nullptr, 0, client_buffer),
@@ -381,6 +486,15 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
client_ctx, ROCPROFILER_BUFFER_TRACING_MEMORY_COPY, nullptr, 0, client_buffer),
"buffer tracing service for memory copy configure");
// May have incompatible kernel so only emit a warning here
ROCPROFILER_WARN(rocprofiler_configure_buffer_tracing_service(
client_ctx, ROCPROFILER_BUFFER_TRACING_PAGE_MIGRATION, nullptr, 0, client_buffer));
ROCPROFILER_CALL(
rocprofiler_configure_buffer_tracing_service(
client_ctx, ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY, nullptr, 0, client_buffer),
"buffer tracing service for page migration configure");
auto client_thread = rocprofiler_callback_thread_t{};
ROCPROFILER_CALL(rocprofiler_create_callback_thread(&client_thread),
"creating callback thread");
@@ -469,9 +583,6 @@ rocprofiler_configure(uint32_t version,
uint32_t priority,
rocprofiler_client_id_t* id)
{
// only activate if main tool
if(priority > 0) return nullptr;
// set the client name
id->name = "ExampleTool";
@@ -485,8 +596,8 @@ rocprofiler_configure(uint32_t version,
// generate info string
auto info = std::stringstream{};
info << id->name << " is using rocprofiler-sdk v" << major << "." << minor << "." << patch
<< " (" << runtime_version << ")";
info << id->name << " (priority=" << priority << ") is using rocprofiler-sdk v" << major << "."
<< minor << "." << patch << " (" << runtime_version << ")";
std::clog << info.str() << std::endl;
projects/rocprofiler-sdk/samples/api_buffered_tracing/main.cpp
+201 -37
Wyświetl plik
@@ -30,6 +30,7 @@
#include <iostream>
#include <mutex>
#include <random>
#include <sstream>
#include <stdexcept>
#define HIP_API_CALL(CALL) \
@@ -51,7 +52,7 @@ namespace
{
using auto_lock_t = std::unique_lock<std::mutex>;
auto print_lock = std::mutex{};
size_t nthreads = 2;
size_t nthread_per_device = 2;
size_t nitr = 500;
size_t nsync = 10;
constexpr unsigned shared_mem_tile_dim = 32;
@@ -64,10 +65,19 @@ verify(int* in, int* out, int M, int N);
} // namespace
__global__ void
transpose_a(const int* in, int* out, int M, int N);
transpose(const int* in, int* out, int M, int N);
void
run(int rank, int tid, hipStream_t stream, int argc, char** argv);
run(int rank, int tid, int devid, int argc, char** argv);
void
run_transpose(int rank, int tid, hipStream_t stream, int argc, char** argv);
void
run_migrate(int rank, int tid, hipStream_t stream, int, char** argv);
void
run_scratch(int rank, int tid, hipStream_t stream, int argc, char** argv);
int
main(int argc, char** argv)
@@ -76,6 +86,8 @@ main(int argc, char** argv)
client::start(); // starts context before any API tables are available
client::identify(1);
auto* exe_name = ::basename(argv[0]);
int rank = 0;
for(int i = 1; i < argc; ++i)
{
@@ -83,47 +95,38 @@ main(int argc, char** argv)
if(_arg == "?" || _arg == "-h" || _arg == "--help")
{
fprintf(stderr,
"usage: transpose [NUM_THREADS (%zu)] [NUM_ITERATION (%zu)] "
"usage: %s [NUM_THREADS_PER_DEVICE (%zu)] [NUM_ITERATION (%zu)] "
"[SYNC_EVERY_N_ITERATIONS (%zu)]\n",
nthreads,
exe_name,
nthread_per_device,
nitr,
nsync);
exit(EXIT_SUCCESS);
}
}
if(argc > 1) nthreads = atoll(argv[1]);
if(argc > 1) nthread_per_device = atoll(argv[1]);
if(argc > 2) nitr = atoll(argv[2]);
if(argc > 3) nsync = atoll(argv[3]);
printf("[transpose] Number of threads: %zu\n", nthreads);
printf("[transpose] Number of iterations: %zu\n", nitr);
printf("[transpose] Syncing every %zu iterations\n", nsync);
// this is a temporary workaround in omnitrace when HIP + MPI is enabled
int ndevice = 0;
int devid = rank;
HIP_API_CALL(hipGetDeviceCount(&ndevice));
printf("[transpose] Number of devices found: %i\n", ndevice);
if(ndevice > 0)
auto nthreads = (ndevice * nthread_per_device);
printf("[%s] Number of devices found: %i\n", exe_name, ndevice);
printf("[%s] Number of threads (per device): %zu\n", exe_name, nthread_per_device);
printf("[%s] Number of threads (total): %zu\n", exe_name, nthreads);
printf("[%s] Number of iterations: %zu\n", exe_name, nitr);
printf("[%s] Syncing every %zu iterations\n", exe_name, nsync);
{
devid = rank % ndevice;
HIP_API_CALL(hipSetDevice(devid));
printf("[transpose] Rank %i assigned to device %i\n", rank, devid);
}
if(rank == devid && rank < ndevice)
{
std::vector<std::thread> _threads{};
std::vector<hipStream_t> _streams(nthreads);
auto _threads = std::vector<std::thread>{};
for(size_t i = 0; i < nthreads; ++i)
HIP_API_CALL(hipStreamCreate(&_streams.at(i)));
for(size_t i = 1; i < nthreads; ++i)
_threads.emplace_back(run, rank, i, _streams.at(i), argc, argv);
run(rank, 0, _streams.at(0), argc, argv);
_threads.emplace_back(run, rank, i, i % ndevice, argc, argv);
for(auto& itr : _threads)
itr.join();
for(size_t i = 0; i < nthreads; ++i)
HIP_API_CALL(hipStreamDestroy(_streams.at(i)));
}
HIP_API_CALL(hipDeviceSynchronize());
HIP_API_CALL(hipDeviceReset());
@@ -134,7 +137,7 @@ main(int argc, char** argv)
}
__global__ void
transpose_a(const int* in, int* out, int M, int N)
transpose(const int* in, int* out, int M, int N)
{
__shared__ int tile[shared_mem_tile_dim][shared_mem_tile_dim];
@@ -145,17 +148,91 @@ transpose_a(const int* in, int* out, int M, int N)
out[idx] = tile[threadIdx.x][threadIdx.y];
}
template <typename Tp>
__global__ void
test_page_migrate(Tp* data, Tp val)
{
int idx = (blockIdx.x * blockDim.x) + threadIdx.x;
data[idx] += val;
}
__global__ void
test_kern_large(uint64_t* output)
{
uint64_t result = 0;
int test[4000];
memset(test, 5, 4000);
for(int& i : test)
{
i = i + 7;
*output += i;
result += i;
}
*output ^= result;
*output ^= result;
}
__global__ void
test_kern_medium(uint64_t* output)
{
uint64_t result = 0;
int test[175];
memset(test, 5, 175);
for(int& i : test)
{
i = i + 7;
*output += i;
result += i;
}
*output ^= result;
*output ^= result;
}
__global__ void
test_kern_small(uint64_t* output)
{
uint64_t result = 0;
int test[2];
for(int& i : test)
{
i = i + 7;
*output += i;
result += i;
}
*output ^= result;
*output ^= result;
}
void
run(int rank, int tid, hipStream_t stream, int argc, char** argv)
run(int rank, int tid, int devid, int argc, char** argv)
{
client::identify(tid + 1);
auto* stream = hipStream_t{};
HIP_API_CALL(hipSetDevice(devid));
HIP_API_CALL(hipStreamCreate(&stream));
run_migrate(rank, tid, stream, argc, argv);
run_scratch(rank, tid, stream, argc, argv);
run_transpose(rank, tid, stream, argc, argv);
HIP_API_CALL(hipStreamSynchronize(stream));
HIP_API_CALL(hipStreamDestroy(stream));
}
void
run_transpose(int rank, int tid, hipStream_t stream, int argc, char** argv)
{
auto* exe_name = ::basename(argv[0]);
unsigned int M = 4960 * 2;
unsigned int N = 4960 * 2;
if(argc > 2) nitr = atoll(argv[2]);
if(argc > 3) nsync = atoll(argv[3]);
auto_lock_t _lk{print_lock};
std::cout << "[transpose][" << rank << "][" << tid << "] M: " << M << " N: " << N << std::endl;
std::cout << "[" << exe_name << "][transpose][" << rank << "][" << tid << "] M: " << M
<< " N: " << N << std::endl;
_lk.unlock();
std::default_random_engine _engine{std::random_device{}() * (rank + 1) * (tid + 1)};
@@ -180,10 +257,11 @@ run(int rank, int tid, hipStream_t stream, int argc, char** argv)
HIP_API_CALL(hipStreamSynchronize(stream));
dim3 grid(M / 32, N / 32, 1);
dim3 block(32, 32, 1); // transpose_a
dim3 block(32, 32, 1); // transpose
print_lock.lock();
printf("[transpose][%i][%i] grid=(%i,%i,%i), block=(%i,%i,%i)\n",
printf("[%s][transpose][%i][%i] grid=(%i,%i,%i), block=(%i,%i,%i)\n",
exe_name,
rank,
tid,
grid.x,
@@ -197,7 +275,7 @@ run(int rank, int tid, hipStream_t stream, int argc, char** argv)
auto t1 = std::chrono::high_resolution_clock::now();
for(size_t i = 0; i < nitr; ++i)
{
transpose_a<<<grid, block, 0, stream>>>(in, out, M, N);
transpose<<<grid, block, 0, stream>>>(in, out, M, N);
check_hip_error();
if(i % nsync == (nsync - 1)) HIP_API_CALL(hipStreamSynchronize(stream));
}
@@ -208,9 +286,9 @@ run(int rank, int tid, hipStream_t stream, int argc, char** argv)
float GB = (float) size * nitr * 2 / (1 << 30);
print_lock.lock();
std::cout << "[transpose][" << rank << "][" << tid << "] Runtime of transpose is " << time
<< " sec\n";
std::cout << "[transpose][" << rank << "][" << tid
std::cout << "[" << exe_name << "][transpose][" << rank << "][" << tid
<< "] Runtime of transpose is " << time << " sec\n";
std::cout << "[" << exe_name << "][transpose][" << rank << "][" << tid
<< "] The average performance of transpose is " << GB / time << " GBytes/sec"
<< std::endl;
print_lock.unlock();
@@ -227,6 +305,92 @@ run(int rank, int tid, hipStream_t stream, int argc, char** argv)
delete[] out_matrix;
}
void
run_scratch(int rank, int tid, hipStream_t stream, int, char** argv)
{
auto t1 = std::chrono::high_resolution_clock::now();
HIP_API_CALL(hipStreamSynchronize(stream));
const auto* exe_name = ::basename(argv[0]);
uint64_t* data_ptr = nullptr;
HIP_API_CALL(hipHostMalloc(&data_ptr, sizeof(uint64_t), 0));
*data_ptr = 0;
test_kern_small<<<1000, 1, 0, stream>>>(data_ptr);
test_kern_medium<<<1000, 1, 0, stream>>>(data_ptr);
test_kern_small<<<1000, 1, 0, stream>>>(data_ptr);
test_kern_large<<<1100, 1, 0, stream>>>(data_ptr);
HIP_API_CALL(hipStreamSynchronize(stream));
test_kern_small<<<1000, 1, 0, stream>>>(data_ptr);
HIP_API_CALL(hipStreamSynchronize(stream));
test_kern_medium<<<1000, 1, 0, stream>>>(data_ptr);
HIP_API_CALL(hipStreamSynchronize(stream));
test_kern_small<<<1000, 1, 0, stream>>>(data_ptr);
HIP_API_CALL(hipStreamSynchronize(stream));
test_kern_large<<<1100, 1, 0, stream>>>(data_ptr);
HIP_API_CALL(hipStreamSynchronize(stream));
auto t2 = std::chrono::high_resolution_clock::now();
double time = std::chrono::duration_cast<std::chrono::duration<double>>(t2 - t1).count();
print_lock.lock();
std::cout << "[" << exe_name << "][scratch][" << rank << "][" << tid
<< "] Runtime of scratch is " << time << " sec\n";
print_lock.unlock();
}
void
run_migrate(int rank, int tid, hipStream_t stream, int, char** argv)
{
using data_type = uint64_t;
constexpr data_type init_v = 1;
constexpr data_type incr_v = 1;
auto t1 = std::chrono::high_resolution_clock::now();
HIP_API_CALL(hipStreamSynchronize(stream));
const auto* exe_name = ::basename(argv[0]);
auto page_data = std::vector<data_type>(1024, 0);
HIP_API_CALL(hipHostRegister(
page_data.data(), page_data.size() * sizeof(data_type), hipHostRegisterDefault));
for(auto& itr : page_data)
itr = init_v;
test_page_migrate<<<1, 1024, 0, stream>>>(page_data.data(), incr_v);
HIP_API_CALL(hipStreamSynchronize(stream));
for(auto& itr : page_data)
{
auto diff = (itr - incr_v);
if(diff != init_v)
{
auto msg = std::stringstream{};
msg << "invalid diff: " << diff << ". expected: " << init_v;
throw std::runtime_error{msg.str()};
}
}
HIP_API_CALL(hipHostUnregister(page_data.data()));
auto t2 = std::chrono::high_resolution_clock::now();
double time = std::chrono::duration_cast<std::chrono::duration<double>>(t2 - t1).count();
print_lock.lock();
std::cout << "[" << exe_name << "][migrate][" << rank << "][" << tid
<< "] Runtime of migrate is " << time << " sec\n";
print_lock.unlock();
}
namespace
{
void