Public C++ header files and samples updates (#819)
* 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
Этот коммит содержится в:
Jonathan R. Madsen
коммит произвёл
GitHub
GitHub
родитель
e2bce49655
Коммит
de13d2ac5d
@@ -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
|
||||
|
||||
Ссылка в новой задаче
Block a user