SWDEV-392791 - [catch2][dtest] Migrating locally available test cases of hipStreamCreateWithPriority() api to Catch2 (#269)
Change-Id: I232f9d2e4d469f8b46639d020b58bbbbcf09a6d4
[ROCm/hip-tests commit: 573d14a3bf]
Цей коміт міститься в:
зафіксовано
GitHub
джерело
24c863712a
коміт
4ba3a31215
@@ -42,3 +42,10 @@ THE SOFTWARE.
|
||||
* This section describes the warp shuffle types & functions of HIP runtime API.
|
||||
* @}
|
||||
*/
|
||||
|
||||
/**
|
||||
* @defgroup StreamTest Stream Management
|
||||
* @{
|
||||
* This section describes the stream management types & functions of HIP runtime API.
|
||||
* @}
|
||||
*/
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2022-2023 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
|
||||
@@ -18,38 +18,20 @@ THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/**
|
||||
Testcase Scenarios :
|
||||
* @addtogroup hipStreamCreateWithPriority hipStreamCreateWithPriority
|
||||
* @{
|
||||
* @ingroup StreamTest
|
||||
* `hipStreamCreateWithPriority (hipStream_t *stream, unsigned int flags, int priority)` -
|
||||
* begins graph capture on a stream
|
||||
*/
|
||||
|
||||
1)Create streams with default flag for all available priority levels and
|
||||
queue tasks in each of these streams, perform device synchronize and validate behavior.
|
||||
|
||||
2)Create streams with non-blocking flag for all available priority levels and
|
||||
queue tasks in each of these streams, perform stream synchronize and validate behavior.
|
||||
|
||||
3)Create streams with default flag for all available priority levels and
|
||||
queue tasks in each of these streams, perform stream synchronize and validate behavior.
|
||||
|
||||
4)Create streams with non-blocking flag for all available priority levels and
|
||||
queue tasks in each of these streams, perform device synchronize and validate behavior.
|
||||
|
||||
5)Create a stream for each priority level with default flag, Launch memcpy and kernel
|
||||
tasks on these streams from multiple threads. Validate all the results.
|
||||
|
||||
6)Create a stream for each priority level with non-blocking flag, Launch memcpy and
|
||||
kernel tasks on these streams from multiple threads. Validate all the results.
|
||||
|
||||
7) Validate negative scenarios for hipStreamCreateWithPriority api.
|
||||
|
||||
8) Validate stream priorities with event after classifying them as low, medium, high.
|
||||
*/
|
||||
|
||||
#include "streamCommon.hh"
|
||||
#include <hip_test_kernels.hh>
|
||||
#include <atomic>
|
||||
#include <vector>
|
||||
#include "streamCommon.hh" // NOLINT
|
||||
|
||||
#define MEMCPYSIZE 64*1024*1024
|
||||
#define MEMCPYSIZE2 1024*1024
|
||||
#define MEMCPYSIZE1 (64*1024*1024)
|
||||
#define MEMCPYSIZE2 (1024*1024)
|
||||
#define NUMITERS 2
|
||||
#define GRIDSIZE 1024
|
||||
#define BLOCKSIZE 256
|
||||
@@ -87,7 +69,7 @@ void funcTestsForAllPriorityLevelsWrtNullStrm(unsigned int flags,
|
||||
int priority;
|
||||
int priority_low{};
|
||||
int priority_high{};
|
||||
size_t size = MEMCPYSIZE2*sizeof(int);
|
||||
size_t size = MEMCPYSIZE2 * sizeof(int);
|
||||
// Test is to get the Stream Priority Range
|
||||
HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
|
||||
|
||||
@@ -194,7 +176,7 @@ void funcTestsForAllPriorityLevelsWrtNullStrm(unsigned int flags,
|
||||
*/
|
||||
void queueTasksInStreams(std::vector<hipStream_t> *stream,
|
||||
const int arrsize) {
|
||||
size_t size = MEMCPYSIZE2*sizeof(int);
|
||||
size_t size = MEMCPYSIZE2 * sizeof(int);
|
||||
// Allocate memory
|
||||
int **A_d = reinterpret_cast<int**>(malloc(arrsize*sizeof(int *)));
|
||||
int **C_d = reinterpret_cast<int**>(malloc(arrsize*sizeof(int *)));
|
||||
@@ -211,8 +193,8 @@ void queueTasksInStreams(std::vector<hipStream_t> *stream,
|
||||
HIPASSERT(A_h[idx] != nullptr);
|
||||
C_h[idx] = reinterpret_cast<int*>(malloc(size));
|
||||
HIPASSERT(C_h[idx] != nullptr);
|
||||
HIPCHECK(hipMalloc(&A_d[idx], size));
|
||||
HIPCHECK(hipMalloc(&C_d[idx], size));
|
||||
HIP_CHECK(hipMalloc(&A_d[idx], size));
|
||||
HIP_CHECK(hipMalloc(&C_d[idx], size));
|
||||
}
|
||||
// Initialize host memory
|
||||
constexpr int initVal = 2;
|
||||
@@ -223,20 +205,20 @@ void queueTasksInStreams(std::vector<hipStream_t> *stream,
|
||||
}
|
||||
// Launch task on each stream
|
||||
for (int idx = 0; idx < arrsize; idx++) {
|
||||
HIPCHECK(hipMemcpyAsync(A_d[idx], A_h[idx], size,
|
||||
HIP_CHECK(hipMemcpyAsync(A_d[idx], A_h[idx], size,
|
||||
hipMemcpyHostToDevice, (*stream)[idx]));
|
||||
hipLaunchKernelGGL((HipTest::vector_square), dim3(GRIDSIZE),
|
||||
dim3(BLOCKSIZE), 0, (*stream)[idx], A_d[idx],
|
||||
C_d[idx], MEMCPYSIZE2);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
HIPCHECK(hipMemcpyAsync(C_h[idx], C_d[idx], size,
|
||||
HIP_CHECK(hipMemcpyAsync(C_h[idx], C_d[idx], size,
|
||||
hipMemcpyDeviceToHost, (*stream)[idx]));
|
||||
}
|
||||
|
||||
bool isPassed = true;
|
||||
// Validate the output of each queue
|
||||
for (int idx = 0; idx < arrsize; idx++) {
|
||||
HIPCHECK(hipStreamSynchronize((*stream)[idx]));
|
||||
HIP_CHECK(hipStreamSynchronize((*stream)[idx]));
|
||||
for (int idy = 0; idy < MEMCPYSIZE2; idy++) {
|
||||
if (C_h[idx][idy] != A_h[idx][idy] * A_h[idx][idy]) {
|
||||
UNSCOPED_INFO("Data mismatch at idx:" << idx << " idy:" << idy);
|
||||
@@ -248,8 +230,8 @@ void queueTasksInStreams(std::vector<hipStream_t> *stream,
|
||||
}
|
||||
// Deallocate memory
|
||||
for (int idx = 0; idx < arrsize; idx++) {
|
||||
HIPCHECK(hipFree(reinterpret_cast<void*>(C_d[idx])));
|
||||
HIPCHECK(hipFree(reinterpret_cast<void*>(A_d[idx])));
|
||||
HIP_CHECK(hipFree(reinterpret_cast<void*>(C_d[idx])));
|
||||
HIP_CHECK(hipFree(reinterpret_cast<void*>(A_d[idx])));
|
||||
free(C_h[idx]);
|
||||
free(A_h[idx]);
|
||||
}
|
||||
@@ -322,7 +304,7 @@ bool runFuncTestsForAllPriorityLevelsMultThread(unsigned int flags) {
|
||||
|
||||
template <typename T>
|
||||
bool validateStreamPrioritiesWithEvents() {
|
||||
size_t size = NUMITERS*MEMCPYSIZE;
|
||||
size_t size = NUMITERS * MEMCPYSIZE1;
|
||||
|
||||
// get the range of priorities available
|
||||
#define OP(x) \
|
||||
@@ -421,13 +403,13 @@ bool validateStreamPrioritiesWithEvents() {
|
||||
#undef OP
|
||||
|
||||
// launch kernels repeatedly on each of the prioritiy streams
|
||||
for (int i = 0; i < static_cast<int>(size); i += MEMCPYSIZE) {
|
||||
for (int i = 0; i < static_cast<int>(size); i += MEMCPYSIZE1) {
|
||||
int j = i / sizeof(T);
|
||||
#define OP(x) \
|
||||
if (enable_priority_##x) { \
|
||||
hipLaunchKernelGGL((memcpy_kernel<T>), dim3(GRIDSIZE), \
|
||||
dim3(BLOCKSIZE), 0, stream_##x, dst_d_##x + j, src_d_##x + j, \
|
||||
(MEMCPYSIZE / sizeof(T))); \
|
||||
(MEMCPYSIZE1 / sizeof(T))); \
|
||||
HIP_CHECK(hipGetLastError()); \
|
||||
}
|
||||
OP(low)
|
||||
@@ -522,19 +504,330 @@ bool validateStreamPrioritiesWithEvents() {
|
||||
return true;
|
||||
}
|
||||
|
||||
#define LOW_PRIORITY_STREAMCOUNT 2
|
||||
#define HIGH_PRIORITY_STREAMCOUNT 2
|
||||
#define NORMAL_PRIORITY_STREAMCOUNT 2
|
||||
template <typename T>
|
||||
void TestForMultipleStreamWithPriority(void) {
|
||||
size_t size = NUMITERS * MEMCPYSIZE1;
|
||||
// get the range of priorities available
|
||||
#define OP(x) \
|
||||
int priority_##x; \
|
||||
bool enable_priority_##x = false;
|
||||
OP(low)
|
||||
OP(normal)
|
||||
OP(high)
|
||||
#undef OP
|
||||
HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
|
||||
INFO("HIP stream priority range - low: " << priority_low << ",high: "
|
||||
<< priority_high << ",normal: "
|
||||
<< (priority_low + priority_high)/2);
|
||||
|
||||
// Check if priorities are indeed supported
|
||||
// Enable/disable priorities based on number of available priority levels
|
||||
enable_priority_low = true;
|
||||
enable_priority_high = true;
|
||||
if ((priority_low - priority_high) > 1) {
|
||||
enable_priority_normal = true;
|
||||
}
|
||||
if (enable_priority_normal) {
|
||||
priority_normal = ((priority_low + priority_high) / 2);
|
||||
}
|
||||
// create streams with low priority
|
||||
hipStream_t stream_low[LOW_PRIORITY_STREAMCOUNT];
|
||||
for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_low) {
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream_low[i],
|
||||
hipStreamDefault, priority_low));
|
||||
}
|
||||
}
|
||||
// create streams with normal priority
|
||||
hipStream_t stream_normal[NORMAL_PRIORITY_STREAMCOUNT];
|
||||
for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_normal) {
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream_normal[i],
|
||||
hipStreamDefault, priority_normal));
|
||||
}
|
||||
}
|
||||
// create streams with high priority
|
||||
hipStream_t stream_high[HIGH_PRIORITY_STREAMCOUNT];
|
||||
for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_high) {
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream_high[i],
|
||||
hipStreamDefault, priority_high));
|
||||
}
|
||||
}
|
||||
// allocate and initialise host source and destination buffers for
|
||||
// low streams
|
||||
T* src_h_low[LOW_PRIORITY_STREAMCOUNT];
|
||||
T* dst_h_low[LOW_PRIORITY_STREAMCOUNT];
|
||||
for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_low) {
|
||||
src_h_low[i] = reinterpret_cast<T*>(malloc(size));
|
||||
REQUIRE(src_h_low[i] != nullptr);
|
||||
mem_init<T>(src_h_low[i], (size / sizeof(T)));
|
||||
dst_h_low[i] = reinterpret_cast<T*>(malloc(size));
|
||||
REQUIRE(dst_h_low[i] != nullptr);
|
||||
memset(dst_h_low[i], 0, size);
|
||||
}
|
||||
}
|
||||
// allocate and initialise host source and destination buffers for
|
||||
// normal streams
|
||||
T* src_h_normal[NORMAL_PRIORITY_STREAMCOUNT];
|
||||
T* dst_h_normal[NORMAL_PRIORITY_STREAMCOUNT];
|
||||
for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_normal) {
|
||||
src_h_normal[i] = reinterpret_cast<T*>(malloc(size));
|
||||
REQUIRE(src_h_normal[i] != nullptr);
|
||||
mem_init<T>(src_h_normal[i], (size / sizeof(T)));
|
||||
dst_h_normal[i] = reinterpret_cast<T*>(malloc(size));
|
||||
REQUIRE(dst_h_normal[i] != nullptr);
|
||||
memset(dst_h_normal[i], 0, size);
|
||||
}
|
||||
}
|
||||
// allocate and initialise host source and destination buffers for
|
||||
// high streams
|
||||
T* src_h_high[HIGH_PRIORITY_STREAMCOUNT];
|
||||
T* dst_h_high[HIGH_PRIORITY_STREAMCOUNT];
|
||||
for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_high) {
|
||||
src_h_high[i] = reinterpret_cast<T*>(malloc(size));
|
||||
REQUIRE(src_h_high[i] != nullptr);
|
||||
mem_init<T>(src_h_high[i], (size / sizeof(T)));
|
||||
dst_h_high[i] = reinterpret_cast<T*>(malloc(size));
|
||||
REQUIRE(dst_h_high[i] != nullptr);
|
||||
memset(dst_h_high[i], 0, size);
|
||||
}
|
||||
}
|
||||
// allocate and initialize device source and destination buffers for
|
||||
// low streams
|
||||
T* src_d_low[LOW_PRIORITY_STREAMCOUNT];
|
||||
T* dst_d_low[LOW_PRIORITY_STREAMCOUNT];
|
||||
for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_low) {
|
||||
HIP_CHECK(hipMalloc(&src_d_low[i], size));
|
||||
HIP_CHECK(hipMemcpy(src_d_low[i], src_h_low[i], size,
|
||||
hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMalloc(&dst_d_low[i], size));
|
||||
}
|
||||
}
|
||||
// allocate and initialize device source and destination buffers for
|
||||
// normal streams
|
||||
T* src_d_normal[NORMAL_PRIORITY_STREAMCOUNT];
|
||||
T* dst_d_normal[NORMAL_PRIORITY_STREAMCOUNT];
|
||||
for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_normal) {
|
||||
HIP_CHECK(hipMalloc(&src_d_normal[i], size));
|
||||
HIP_CHECK(hipMemcpy(src_d_normal[i], src_h_normal[i], size,
|
||||
hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMalloc(&dst_d_normal[i], size));
|
||||
}
|
||||
}
|
||||
// allocate and initialize device source and destination buffers for
|
||||
// high streams
|
||||
T* src_d_high[HIGH_PRIORITY_STREAMCOUNT];
|
||||
T* dst_d_high[HIGH_PRIORITY_STREAMCOUNT];
|
||||
for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_high) {
|
||||
HIP_CHECK(hipMalloc(&src_d_high[i], size));
|
||||
HIP_CHECK(hipMemcpy(src_d_high[i], src_h_high[i], size,
|
||||
hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMalloc(&dst_d_high[i], size));
|
||||
}
|
||||
}
|
||||
|
||||
hipEvent_t event_start_low[LOW_PRIORITY_STREAMCOUNT];
|
||||
hipEvent_t event_end_low[LOW_PRIORITY_STREAMCOUNT];
|
||||
// create events for measuring time spent in kernel execution for low
|
||||
// priority streams
|
||||
for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_low) {
|
||||
HIP_CHECK(hipEventCreate(&event_start_low[i]));
|
||||
HIP_CHECK(hipEventCreate(&event_end_low[i]));
|
||||
}
|
||||
}
|
||||
|
||||
hipEvent_t event_start_normal[NORMAL_PRIORITY_STREAMCOUNT];
|
||||
hipEvent_t event_end_normal[NORMAL_PRIORITY_STREAMCOUNT];
|
||||
// create events for measuring time spent in kernel execution for
|
||||
// normal priority streams
|
||||
for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_normal) {
|
||||
HIP_CHECK(hipEventCreate(&event_start_normal[i]));
|
||||
HIP_CHECK(hipEventCreate(&event_end_normal[i]));
|
||||
}
|
||||
}
|
||||
|
||||
hipEvent_t event_start_high[HIGH_PRIORITY_STREAMCOUNT];
|
||||
hipEvent_t event_end_high[HIGH_PRIORITY_STREAMCOUNT];
|
||||
// create events for measuring time spent in kernel execution for
|
||||
// high priority streams
|
||||
for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_high) {
|
||||
HIP_CHECK(hipEventCreate(&event_start_high[i]));
|
||||
HIP_CHECK(hipEventCreate(&event_end_high[i]));
|
||||
}
|
||||
}
|
||||
// record start events for each of the low priority streams
|
||||
for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_low) {
|
||||
HIP_CHECK(hipEventRecord(event_start_low[i], stream_low[i]));
|
||||
}
|
||||
}
|
||||
// record start events for each of the normal priority streams
|
||||
for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_normal) {
|
||||
HIP_CHECK(hipEventRecord(event_start_normal[i], stream_normal[i]));
|
||||
}
|
||||
}
|
||||
// record start events for each of the high priority streams
|
||||
for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_high) {
|
||||
HIP_CHECK(hipEventRecord(event_start_high[i], stream_high[i]));
|
||||
}
|
||||
}
|
||||
// launch kernels repeatedly on each of the low prioritiy stream
|
||||
for (int k = 0; k < LOW_PRIORITY_STREAMCOUNT; ++k) {
|
||||
for (int i = 0; i < size; i += MEMCPYSIZE1) {
|
||||
int j = i / sizeof(T);
|
||||
if (enable_priority_low) {
|
||||
hipLaunchKernelGGL((memcpy_kernel<T>), dim3(GRIDSIZE), dim3(BLOCKSIZE),
|
||||
0, stream_low[k], dst_d_low[k] + j, src_d_low[k] + j,
|
||||
(MEMCPYSIZE1 / sizeof(T)));
|
||||
}
|
||||
}
|
||||
}
|
||||
// launch kernels repeatedly on each of the normal prioritiy stream
|
||||
for (int k = 0; k < NORMAL_PRIORITY_STREAMCOUNT; ++k) {
|
||||
for (int i = 0; i < size; i += MEMCPYSIZE1) {
|
||||
int j = i / sizeof(T);
|
||||
if (enable_priority_normal) {
|
||||
hipLaunchKernelGGL((memcpy_kernel<T>), dim3(GRIDSIZE), dim3(BLOCKSIZE),
|
||||
0, stream_normal[k], dst_d_normal[k] + j, src_d_normal[k] + j,
|
||||
(MEMCPYSIZE1 / sizeof(T)));
|
||||
}
|
||||
}
|
||||
}
|
||||
// launch kernels repeatedly on each of the high prioritiy stream
|
||||
for (int k = 0; k < HIGH_PRIORITY_STREAMCOUNT; ++k) {
|
||||
for (int i = 0; i < size; i += MEMCPYSIZE1) {
|
||||
int j = i / sizeof(T);
|
||||
if (enable_priority_high) {
|
||||
hipLaunchKernelGGL((memcpy_kernel<T>), dim3(GRIDSIZE), dim3(BLOCKSIZE),
|
||||
0, stream_high[k], dst_d_high[k] + j, src_d_high[k] + j,
|
||||
(MEMCPYSIZE1 / sizeof(T)));
|
||||
}
|
||||
}
|
||||
}
|
||||
// record end events for each of the low priority streams
|
||||
for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_low) {
|
||||
HIP_CHECK(hipEventRecord(event_end_low[i], stream_low[i]));
|
||||
}
|
||||
}
|
||||
// record end events for each of the normal priority streams
|
||||
for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_normal) {
|
||||
HIP_CHECK(hipEventRecord(event_end_normal[i], stream_normal[i]));
|
||||
}
|
||||
}
|
||||
// record end events for each of the high priority streams
|
||||
for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_high) {
|
||||
HIP_CHECK(hipEventRecord(event_end_high[i], stream_high[i]));
|
||||
}
|
||||
}
|
||||
// synchronize events for each of the low priority streams
|
||||
for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_low) {
|
||||
HIP_CHECK(hipEventSynchronize(event_end_low[i]));
|
||||
}
|
||||
}
|
||||
// synchronize events for each of the normal priority streams
|
||||
for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_normal) {
|
||||
HIP_CHECK(hipEventSynchronize(event_end_normal[i]));
|
||||
}
|
||||
}
|
||||
// synchronize events for each of the high priority streams
|
||||
for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_high) {
|
||||
HIP_CHECK(hipEventSynchronize(event_end_high[i]));
|
||||
}
|
||||
}
|
||||
// compute time spent for memcpy in each low stream
|
||||
float time_spent_low[LOW_PRIORITY_STREAMCOUNT];
|
||||
for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_low) {
|
||||
HIP_CHECK(hipEventElapsedTime(&time_spent_low[i], event_start_low[i],
|
||||
event_end_low[i]));
|
||||
}
|
||||
}
|
||||
// compute time spent for memcpy in each normal stream
|
||||
float time_spent_normal[NORMAL_PRIORITY_STREAMCOUNT];
|
||||
for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_normal) {
|
||||
HIP_CHECK(hipEventElapsedTime(&time_spent_normal[i],
|
||||
event_start_normal[i], event_end_normal[i]));
|
||||
}
|
||||
}
|
||||
// compute time spent for memcpy in each high stream
|
||||
float time_spent_high[HIGH_PRIORITY_STREAMCOUNT];
|
||||
for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_high) {
|
||||
HIP_CHECK(hipEventElapsedTime(&time_spent_high[i], event_start_high[i],
|
||||
event_end_high[i]));
|
||||
}
|
||||
}
|
||||
// sanity check for low priority streams
|
||||
for (int i = 0; i < LOW_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_low) {
|
||||
HIP_CHECK(hipMemcpy(dst_h_low[i], dst_d_low[i], size,
|
||||
hipMemcpyDeviceToHost));
|
||||
REQUIRE(memcmp(dst_h_low[i], src_h_low[i], size) == 0);
|
||||
}
|
||||
}
|
||||
// sanity check for normal priority streams
|
||||
for (int i = 0; i < NORMAL_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_normal) {
|
||||
HIP_CHECK(hipMemcpy(dst_h_normal[i], dst_d_normal[i], size,
|
||||
hipMemcpyDeviceToHost));
|
||||
REQUIRE(memcmp(dst_h_normal[i], src_h_normal[i], size) == 0);
|
||||
}
|
||||
}
|
||||
// sanity check for high priority streams
|
||||
for (int i = 0; i < HIGH_PRIORITY_STREAMCOUNT; ++i) {
|
||||
if (enable_priority_high) {
|
||||
HIP_CHECK(hipMemcpy(dst_h_high[i], dst_d_high[i], size,
|
||||
hipMemcpyDeviceToHost));
|
||||
REQUIRE(memcmp(dst_h_high[i], src_h_high[i], size) == 0);
|
||||
}
|
||||
}
|
||||
}
|
||||
} // namespace hipStreamCreateWithPriorityTest
|
||||
|
||||
/**
|
||||
Tests following scenarios.
|
||||
1)Create streams with default flag for all available priority levels and
|
||||
queue tasks in each of these streams, perform device synchronize and validate behavior.
|
||||
2)Create streams with non-blocking flag for all available priority levels and
|
||||
queue tasks in each of these streams, perform stream synchronize and validate behavior.
|
||||
3)Create streams with default flag for all available priority levels and
|
||||
queue tasks in each of these streams, perform stream synchronize and validate behavior.
|
||||
4)Create streams with non-blocking flag for all available priority levels and
|
||||
queue tasks in each of these streams, perform device synchronize and validate behavior.
|
||||
*/
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Create streams with default flag for all available priority levels and
|
||||
* queue tasks in each of these streams, perform device synchronize and validate
|
||||
* behavior.
|
||||
* - Create streams with non-blocking flag for all available priority levels
|
||||
* and queue tasks in each of these streams, perform stream synchronize and
|
||||
* validate behavior.
|
||||
* - Create streams with default flag for all available priority levels and
|
||||
* queue tasks in each of these streams, perform stream synchronize and validate
|
||||
* behavior.
|
||||
* - Create streams with non-blocking flag for all available priority levels
|
||||
* and queue tasks in each of these streams, perform device synchronize and validate
|
||||
* behavior.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - catch\unit\stream\hipStreamCreateWithPriority.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamCreateWithPriority_FunctionalForAllPriorities") {
|
||||
SECTION("Default flag and device synchronize") {
|
||||
hipStreamCreateWithPriorityTest::
|
||||
@@ -558,8 +851,16 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_FunctionalForAllPriorities") {
|
||||
}
|
||||
|
||||
/**
|
||||
* Create a stream for each priority level with default flag, Launch memcpy and kernel
|
||||
* tasks on these streams from multiple threads. Validate all the results.
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Create a stream for each priority level with default flag, Launch
|
||||
* memcpy and kernel tasks on these streams from multiple threads. Validate
|
||||
* all the results.
|
||||
* ------------------------
|
||||
* - catch\unit\stream\hipStreamCreateWithPriority.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadDefaultflag") {
|
||||
bool TestPassed = true;
|
||||
@@ -569,8 +870,16 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadDefaultflag") {
|
||||
}
|
||||
|
||||
/**
|
||||
* Create a stream for each priority level with non-blocking flag, Launch memcpy and
|
||||
* kernel tasks on these streams from multiple threads. Validate all the results.
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Create a stream for each priority level with non-blocking flag, Launch
|
||||
* memcpy and kernel tasks on these streams from multiple threads. Validate all
|
||||
* the results.
|
||||
* ------------------------
|
||||
* - catch\unit\stream\hipStreamCreateWithPriority.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadNonblockingflag") {
|
||||
bool TestPassed = true;
|
||||
@@ -579,13 +888,17 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadNonblockingflag") {
|
||||
REQUIRE(TestPassed);
|
||||
}
|
||||
|
||||
|
||||
/**
|
||||
* Scenario1: Validates functionality of hipStreamCreateWithPriority when
|
||||
stream = nullptr.
|
||||
* Scenario2: Validates functionality of hipStreamCreateWithPriority when
|
||||
flag = 0xffffffff.
|
||||
*/
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Validates functionality of hipStreamCreateWithPriority when stream = nullptr
|
||||
* - Validates functionality of hipStreamCreateWithPriority when flag = 0xffffffff
|
||||
* ------------------------
|
||||
* - catch\unit\stream\hipStreamCreateWithPriority.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamCreateWithPriority_NegTst") {
|
||||
hipStream_t stream{nullptr};
|
||||
int priority_low{0};
|
||||
@@ -601,15 +914,26 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_NegTst") {
|
||||
|
||||
SECTION("stream = nullptr test") {
|
||||
REQUIRE(hipErrorInvalidValue ==
|
||||
hipStreamCreateWithPriority(nullptr, hipStreamDefault, priority_low));
|
||||
hipStreamCreateWithPriority(nullptr, hipStreamDefault, priority_low));
|
||||
}
|
||||
|
||||
SECTION("flag value invalid test") {
|
||||
REQUIRE(hipErrorInvalidValue == hipStreamCreateWithPriority(&stream, 0xffffffff, priority_low));
|
||||
REQUIRE(hipErrorInvalidValue ==
|
||||
hipStreamCreateWithPriority(&stream, 0xffffffff, priority_low));
|
||||
}
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipStreamCreateWithPriority") {
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Set and Get Priority Value
|
||||
* ------------------------
|
||||
* - catch\unit\stream\hipStreamCreateWithPriority.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamCreateWithPriority_CheckPriorityVal") {
|
||||
int id = GENERATE(range(0, HipTest::getDeviceCount()));
|
||||
|
||||
HIP_CHECK(hipSetDevice(id));
|
||||
@@ -619,45 +943,73 @@ TEST_CASE("Unit_hipStreamCreateWithPriority") {
|
||||
hipStream_t stream{nullptr};
|
||||
|
||||
SECTION("Setting high priority") {
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_high));
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault,
|
||||
priority_high));
|
||||
REQUIRE(stream != nullptr);
|
||||
REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_high));
|
||||
}
|
||||
|
||||
SECTION("Setting low priority") {
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_low));
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault,
|
||||
priority_low));
|
||||
REQUIRE(stream != nullptr);
|
||||
REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_low));
|
||||
}
|
||||
|
||||
SECTION("Setting lowest possible priority") {
|
||||
HIP_CHECK(
|
||||
hipStreamCreateWithPriority(&stream, hipStreamDefault, std::numeric_limits<int>::max()));
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault,
|
||||
std::numeric_limits<int>::max()));
|
||||
REQUIRE(stream != nullptr);
|
||||
REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_low));
|
||||
}
|
||||
|
||||
SECTION("Setting highest possible priority") {
|
||||
HIP_CHECK(
|
||||
hipStreamCreateWithPriority(&stream, hipStreamDefault, std::numeric_limits<int>::min()));
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault,
|
||||
std::numeric_limits<int>::min()));
|
||||
REQUIRE(stream != nullptr);
|
||||
REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_high));
|
||||
}
|
||||
|
||||
SECTION("Setting flags to hipStreamNonBlocking") {
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_high));
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking,
|
||||
priority_high));
|
||||
REQUIRE(stream != nullptr);
|
||||
REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_high, hipStreamNonBlocking));
|
||||
REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_high,
|
||||
hipStreamNonBlocking));
|
||||
}
|
||||
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
/**
|
||||
* Validate stream priorities with event after classifying them as low, medium and high.
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Validate stream priorities with event after classifying them as low,
|
||||
* medium and high.
|
||||
* ------------------------
|
||||
* - catch\unit\stream\hipStreamCreateWithPriority.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamCreateWithPriority_ValidateWithEvents") {
|
||||
bool TestPassed = true;
|
||||
TestPassed = hipStreamCreateWithPriorityTest::validateStreamPrioritiesWithEvents<int>();
|
||||
TestPassed =
|
||||
hipStreamCreateWithPriorityTest::validateStreamPrioritiesWithEvents<int>();
|
||||
REQUIRE(TestPassed);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Test that create Multiple streams with priority low, normal and high
|
||||
* then these streams used to launch the kernel in random seq high,normal,low.
|
||||
* ------------------------
|
||||
* - catch\unit\stream\hipStreamCreateWithPriority.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 5.2
|
||||
*/
|
||||
TEST_CASE("Unit_hipStreamCreateWithPriority_TestMultipleStreamWithPriority") {
|
||||
hipStreamCreateWithPriorityTest::TestForMultipleStreamWithPriority<int>();
|
||||
}
|
||||
|
||||
Посилання в новій задачі
Заблокувати користувача