AI/rocm-systems
2
0
Vork 0
Code Kwesties Pull-aanvragen Actions Paketten Projecten Publicaties Wiki Activiteit

SWDEV-389399 - [catch2][dtest] Adding hipStreamSynchronize scenarios when streams are in capturing mode (#270)

Bladeren bron 
Change-Id: Ib2f919d84f6acc136db582ca0ac717634e8e5be9

[ROCm/hip-tests commit: b65b09234d]
This commit is contained in:
ROCm CI Service Account
2023-06-21 15:59:45 +05:30
gecommit door GitHub
bovenliggende 4ba3a31215
commit 331a2815a7
1 gewijzigde bestanden met toevoegingen van 361 en 88 verwijderingen
Download Patch-bestand Download Diff-bestand Expand all files Collapse all files
projects/hip-tests/catch/unit/graph/hipStreamBeginCapture.cc
+361 -88
Bestand weergeven
@@ -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
@@ -20,8 +20,7 @@ THE SOFTWARE.
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include <hip_test_defgroups.hh>
#include "stream_capture_common.hh"
#include "stream_capture_common.hh" // NOLINT
/**
* @addtogroup hipStreamBeginCapture hipStreamBeginCapture
@@ -56,7 +55,8 @@ static void hostNodeCallback(void* data) {
}
template <typename T, typename F>
void captureStreamAndLaunchGraph(F graphFunc, hipStreamCaptureMode mode, hipStream_t stream) {
void captureStreamAndLaunchGraph(F graphFunc, hipStreamCaptureMode mode,
hipStream_t stream) {
constexpr size_t N = 1000000;
size_t Nbytes = N * sizeof(T);
@@ -88,7 +88,8 @@ void captureStreamAndLaunchGraph(F graphFunc, hipStreamCaptureMode mode, hipStre
std::fill_n(A_h.host_ptr(), N, static_cast<float>(i));
HIP_CHECK(hipGraphLaunch(graphExec, stream));
HIP_CHECK(hipStreamSynchronize(stream));
ArrayFindIfNot(B_h.host_ptr(), static_cast<float>(i) * static_cast<float>(i), N);
ArrayFindIfNot(B_h.host_ptr(),
static_cast<float>(i) * static_cast<float>(i), N);
}
HIP_CHECK(hipGraphExecDestroy(graphExec))
@@ -115,15 +116,16 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_Functional") {
StreamGuard stream_guard(stream_type);
hipStream_t stream = stream_guard.stream();
const hipStreamCaptureMode captureMode = GENERATE(
hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed);
const hipStreamCaptureMode captureMode = GENERATE(hipStreamCaptureModeGlobal,
hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed);
EventsGuard events_guard(3);
StreamsGuard streams_guard(2);
SECTION("Linear graph capture") {
captureStreamAndLaunchGraph<float>(
[](float* A_h, float* A_d, float* B_h, float* B_d, size_t N, hipStream_t stream) {
[](float* A_h, float* A_d, float* B_h, float* B_d, size_t N,
hipStream_t stream) {
return captureSequenceLinear(A_h, A_d, B_h, B_d, N, stream);
},
captureMode, stream);
@@ -131,10 +133,10 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_Functional") {
SECTION("Branched graph capture") {
captureStreamAndLaunchGraph<float>(
[&streams_guard, &events_guard](float* A_h, float* A_d, float* B_h, float* B_d, size_t N,
hipStream_t stream) {
captureSequenceBranched(A_h, A_d, B_h, B_d, N, stream, streams_guard.stream_list(),
events_guard.event_list());
[&streams_guard, &events_guard](float* A_h, float* A_d, float* B_h,
float* B_d, size_t N, hipStream_t stream) {
captureSequenceBranched(A_h, A_d, B_h, B_d, N, stream,
streams_guard.stream_list(), events_guard.event_list());
},
captureMode, stream);
}
@@ -170,7 +172,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_Parameters") {
hipErrorIllegalState);
}
SECTION("Creating hipStream with invalid mode") {
HIP_CHECK_ERROR(hipStreamBeginCapture(stream, hipStreamCaptureMode(-1)), hipErrorInvalidValue);
HIP_CHECK_ERROR(hipStreamBeginCapture(stream, hipStreamCaptureMode(-1)),
hipErrorInvalidValue);
}
#if HT_NVIDIA // EXSWHTEC-216
SECTION("Stream capture on uninitialized stream returns error code.") {
@@ -178,7 +181,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_Parameters") {
StreamGuard sg(Streams::created);
return sg.stream();
};
HIP_CHECK_ERROR(hipStreamBeginCapture(InvalidStream(), hipStreamCaptureModeGlobal),
HIP_CHECK_ERROR(hipStreamBeginCapture(InvalidStream(),
hipStreamCaptureModeGlobal),
hipErrorContextIsDestroyed);
}
#endif
@@ -202,8 +206,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_Basic") {
StreamGuard stream_guard(stream_type);
hipStream_t s = stream_guard.stream();
const hipStreamCaptureMode captureMode = GENERATE(
hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed);
const hipStreamCaptureMode captureMode = GENERATE(hipStreamCaptureModeGlobal,
hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed);
HIP_CHECK(hipStreamBeginCapture(s, captureMode));
@@ -213,7 +217,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_Basic") {
/* Local function for inter stream event synchronization
*/
static void interStrmEventSyncCapture(const hipStream_t& stream1, const hipStream_t& stream2) {
static void interStrmEventSyncCapture(const hipStream_t& stream1,
const hipStream_t& stream2) {
hipGraph_t graph1{nullptr}, graph2{nullptr};
hipGraphExec_t graphExec1{nullptr}, graphExec2{nullptr};
@@ -260,7 +265,8 @@ static void interStrmEventSyncCapture(const hipStream_t& stream1, const hipStrea
/* Local function for colligated stream capture
*/
static void colligatedStrmCapture(const hipStream_t& stream1, const hipStream_t& stream2) {
static void colligatedStrmCapture(const hipStream_t& stream1,
const hipStream_t& stream2) {
hipGraph_t graph1{nullptr}, graph2{nullptr};
hipGraphExec_t graphExec1{nullptr}, graphExec2{nullptr};
@@ -303,7 +309,8 @@ static void colligatedStrmCapture(const hipStream_t& stream1, const hipStream_t&
/* Local function for colligated stream capture functionality
*/
static void colligatedStrmCaptureFunc(const hipStream_t& stream1, const hipStream_t& stream2) {
static void colligatedStrmCaptureFunc(const hipStream_t& stream1,
const hipStream_t& stream2) {
constexpr size_t N = 1000000;
size_t Nbytes = N * sizeof(int);
@@ -323,8 +330,10 @@ static void colligatedStrmCaptureFunc(const hipStream_t& stream1, const hipStrea
// Capture 2 streams
HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal));
HIP_CHECK(hipStreamBeginCapture(stream2, hipStreamCaptureModeGlobal));
captureSequenceLinear(A_h.host_ptr(), A_d.ptr(), B_h.host_ptr(), B_d.ptr(), N, stream1);
captureSequenceLinear(C_h.host_ptr(), C_d.ptr(), D_h.host_ptr(), D_d.ptr(), N, stream2);
captureSequenceLinear(A_h.host_ptr(), A_d.ptr(), B_h.host_ptr(), B_d.ptr(),
N, stream1);
captureSequenceLinear(C_h.host_ptr(), C_d.ptr(), D_h.host_ptr(), D_d.ptr(),
N, stream2);
captureSequenceCompute(A_d.ptr(), B_h.host_ptr(), B_d.ptr(), N, stream1);
captureSequenceCompute(C_d.ptr(), D_h.host_ptr(), D_d.ptr(), N, stream2);
HIP_CHECK(hipStreamEndCapture(stream1, &graph1));
@@ -360,8 +369,9 @@ static void colligatedStrmCaptureFunc(const hipStream_t& stream1, const hipStrea
/* Stream Capture thread function
*/
static void threadStrmCaptureFunc(hipStream_t stream, int* A_h, int* A_d, int* B_h, int* B_d,
hipGraph_t* graph, size_t N, hipStreamCaptureMode mode) {
static void threadStrmCaptureFunc(hipStream_t stream, int* A_h, int* A_d,
int* B_h, int* B_d, hipGraph_t* graph,
size_t N, hipStreamCaptureMode mode) {
// Capture stream
HIP_CHECK(hipStreamBeginCapture(stream, mode));
captureSequenceLinear(A_h, A_d, B_h, B_d, N, stream);
@@ -393,10 +403,10 @@ static void multithreadedTest(hipStreamCaptureMode mode) {
LinearAllocGuard<int> D_d(LinearAllocs::hipMalloc, Nbytes);
// Launch 2 threads to capture the 2 streams into graphs
std::thread t1(threadStrmCaptureFunc, stream1, A_h.host_ptr(), A_d.ptr(), B_h.host_ptr(),
B_d.ptr(), &graph1, N, mode);
std::thread t2(threadStrmCaptureFunc, stream2, C_h.host_ptr(), C_d.ptr(), D_h.host_ptr(),
D_d.ptr(), &graph2, N, mode);
std::thread t1(threadStrmCaptureFunc, stream1, A_h.host_ptr(), A_d.ptr(),
B_h.host_ptr(), B_d.ptr(), &graph1, N, mode);
std::thread t2(threadStrmCaptureFunc, stream2, C_h.host_ptr(), C_d.ptr(),
D_h.host_ptr(), D_d.ptr(), &graph2, N, mode);
t1.join();
t2.join();
@@ -469,9 +479,11 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_InterStrmEventSync_Flags") {
TEST_CASE("Unit_hipStreamBeginCapture_Positive_InterStrmEventSync_Priority") {
int minPriority = 0, maxPriority = 0;
HIP_CHECK(hipDeviceGetStreamPriorityRange(&minPriority, &maxPriority));
StreamGuard stream_guard1(Streams::withPriority, hipStreamDefault, minPriority);
StreamGuard stream_guard1(Streams::withPriority, hipStreamDefault,
minPriority);
hipStream_t stream1 = stream_guard1.stream();
StreamGuard stream_guard2(Streams::withPriority, hipStreamDefault, maxPriority);
StreamGuard stream_guard2(Streams::withPriority, hipStreamDefault,
maxPriority);
hipStream_t stream2 = stream_guard2.stream();
interStrmEventSyncCapture(stream1, stream2);
}
@@ -517,12 +529,14 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_ColligatedStrmCapture_Flags") {
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_hipStreamBeginCapture_Positive_ColligatedStrmCapture_Priority") {
TEST_CASE("Unit_hipStreamBeginCapture_Positive_ColligatedStrmCapture_Prio") {
int minPriority = 0, maxPriority = 0;
HIP_CHECK(hipDeviceGetStreamPriorityRange(&minPriority, &maxPriority));
StreamGuard stream_guard1(Streams::withPriority, hipStreamDefault, minPriority);
StreamGuard stream_guard1(Streams::withPriority, hipStreamDefault,
minPriority);
hipStream_t stream1 = stream_guard1.stream();
StreamGuard stream_guard2(Streams::withPriority, hipStreamDefault, maxPriority);
StreamGuard stream_guard2(Streams::withPriority, hipStreamDefault,
maxPriority);
hipStream_t stream2 = stream_guard2.stream();
colligatedStrmCapture(stream1, stream2);
}
@@ -563,8 +577,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_ColligatedStrmCaptureFunc") {
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_hipStreamBeginCapture_Positive_Multithreaded") {
const hipStreamCaptureMode captureMode = GENERATE(
hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed);
const hipStreamCaptureMode captureMode = GENERATE(hipStreamCaptureModeGlobal,
hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed);
multithreadedTest(captureMode);
}
@@ -693,7 +707,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_CapturingFromWithinStrms") {
HIP_CHECK(hipEventRecord(events[2], streams[2]));
HIP_CHECK(hipStreamWaitEvent(streams[0], events[1], 0));
HIP_CHECK(hipStreamWaitEvent(streams[0], events[2], 0));
HIP_CHECK(hipMemcpyAsync(hostMem, devMem, sizeof(int), hipMemcpyDefault, streams[0]));
HIP_CHECK(hipMemcpyAsync(hostMem, devMem, sizeof(int), hipMemcpyDefault,
streams[0]));
HIP_CHECK(hipStreamEndCapture(streams[0], &graph)); // End Capture
// Reset device memory
HIP_CHECK(hipMemset(devMem, 0, sizeof(int)));
@@ -735,7 +750,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_DetectingInvalidCapture") {
dummyKernel<<<1, 1, 0, streams[0]>>>();
// Since stream[1] is already in capture mode due to event wait
// hipStreamBeginCapture on stream[1] is expected to return error.
HIP_CHECK_ERROR(hipStreamBeginCapture(streams[1], hipStreamCaptureModeGlobal),
HIP_CHECK_ERROR(hipStreamBeginCapture(streams[1],
hipStreamCaptureModeGlobal),
hipErrorIllegalState);
}
@@ -768,7 +784,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_CapturingMultGraphsFrom1Strm") {
for (int i = 0; i < 3; i++) {
HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal));
for (int j = 0; j <= i; j++) incrementKernel<<<1, 1, 0, stream1>>>(devMem);
HIP_CHECK(hipMemcpyAsync(hostMem, devMem, sizeof(int), hipMemcpyDefault, stream1));
HIP_CHECK(hipMemcpyAsync(hostMem, devMem, sizeof(int), hipMemcpyDefault,
stream1));
HIP_CHECK(hipStreamEndCapture(stream1, &graphs[i]));
}
// Instantiate and execute all graphs
@@ -784,7 +801,6 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_CapturingMultGraphsFrom1Strm") {
}
}
#if HT_NVIDIA
/**
* Test Description
* ------------------------
@@ -808,29 +824,35 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_CheckingSyncDuringCapture") {
EventsGuard events_guard(1);
hipEvent_t e = events_guard[0];
const hipStreamCaptureMode captureMode = GENERATE(
hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed);
const hipStreamCaptureMode captureMode = GENERATE(hipStreamCaptureModeGlobal,
hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed);
HIP_CHECK(hipStreamBeginCapture(stream, captureMode));
SECTION("Synchronize stream during capture") {
HIP_CHECK_ERROR(hipStreamSynchronize(stream), hipErrorStreamCaptureUnsupported);
HIP_CHECK_ERROR(hipStreamSynchronize(stream),
hipErrorStreamCaptureUnsupported);
}
SECTION("Query stream during capture") {
HIP_CHECK_ERROR(hipStreamQuery(stream),
hipErrorStreamCaptureUnsupported);
}
#if HT_NVIDIA
SECTION("Synchronize device during capture") {
HIP_CHECK_ERROR(hipDeviceSynchronize(), hipErrorStreamCaptureUnsupported);
HIP_CHECK_ERROR(hipDeviceSynchronize(),
hipErrorStreamCaptureUnsupported);
}
SECTION("Synchronize event during capture") {
HIP_CHECK(hipEventRecord(e, stream));
HIP_CHECK_ERROR(hipEventSynchronize(e), hipErrorCapturedEvent);
}
SECTION("Query stream during capture") {
HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorStreamCaptureUnsupported);
}
SECTION("Query for an event during capture") {
HIP_CHECK(hipEventRecord(e, stream));
HIP_CHECK_ERROR(hipEventQuery(e), hipErrorCapturedEvent);
}
#endif
}
#if HT_NVIDIA
/**
* Test Description
* ------------------------
@@ -861,14 +883,17 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_UnsafeCallsDuringCapture") {
HIP_CHECK(hipStreamBeginCapture(stream, captureMode));
SECTION("hipMalloc during capture") {
HIP_CHECK_ERROR(hipMalloc(&devMem2, sizeof(int)), hipErrorStreamCaptureUnsupported);
HIP_CHECK_ERROR(hipMalloc(&devMem2, sizeof(int)),
hipErrorStreamCaptureUnsupported);
}
SECTION("hipMemcpy during capture") {
HIP_CHECK_ERROR(hipMemcpy(devMem.ptr(), hostMem.host_ptr(), sizeof(int), hipMemcpyHostToDevice),
HIP_CHECK_ERROR(hipMemcpy(devMem.ptr(), hostMem.host_ptr(), sizeof(int),
hipMemcpyHostToDevice),
hipErrorStreamCaptureImplicit);
}
SECTION("hipMemset during capture") {
HIP_CHECK_ERROR(hipMemset(devMem.ptr(), 0, sizeof(int)), hipErrorStreamCaptureImplicit);
HIP_CHECK_ERROR(hipMemset(devMem.ptr(), 0, sizeof(int)),
hipErrorStreamCaptureImplicit);
}
}
#endif
@@ -889,7 +914,7 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_UnsafeCallsDuringCapture") {
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_hipStreamBeginCapture_Negative_EndingCapturewhenCaptureInProgress") {
TEST_CASE("Unit_hipStreamBeginCapture_Negative_EndingCapwhenCapInProg") {
hipGraph_t graph{nullptr};
StreamsGuard streams_guard(2);
@@ -905,7 +930,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_EndingCapturewhenCaptureInProgres
HIP_CHECK(hipEventRecord(e, stream1));
HIP_CHECK(hipStreamWaitEvent(stream2, e, 0));
dummyKernel<<<1, 1, 0, stream2>>>();
HIP_CHECK_ERROR(hipStreamEndCapture(stream1, &graph), hipErrorStreamCaptureUnjoined);
HIP_CHECK_ERROR(hipStreamEndCapture(stream1, &graph),
hipErrorStreamCaptureUnjoined);
}
SECTION("End strm capture when forked strm still has operations") {
EventsGuard events_guard(2);
@@ -919,7 +945,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_EndingCapturewhenCaptureInProgres
HIP_CHECK(hipEventRecord(e2, stream2));
HIP_CHECK(hipStreamWaitEvent(stream1, e2, 0));
dummyKernel<<<1, 1, 0, stream2>>>();
HIP_CHECK_ERROR(hipStreamEndCapture(stream1, &graph), hipErrorStreamCaptureUnjoined);
HIP_CHECK_ERROR(hipStreamEndCapture(stream1, &graph),
hipErrorStreamCaptureUnjoined);
}
}
/**
@@ -942,17 +969,19 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_MultiGPU") {
SUCCEED("skipping the testcases as numDevices < 2");
return;
}
hipStream_t* stream = reinterpret_cast<hipStream_t*>(malloc(devcount * sizeof(hipStream_t)));
hipStream_t* stream = reinterpret_cast<hipStream_t*>
(malloc(devcount * sizeof(hipStream_t)));
REQUIRE(stream != nullptr);
hipGraph_t* graph = reinterpret_cast<hipGraph_t*>(malloc(devcount * sizeof(hipGraph_t)));
hipGraph_t* graph = reinterpret_cast<hipGraph_t*>
(malloc(devcount * sizeof(hipGraph_t)));
REQUIRE(graph != nullptr);
int **devMem{nullptr}, **hostMem{nullptr};
hostMem = reinterpret_cast<int**>(malloc(sizeof(int*) * devcount));
REQUIRE(hostMem != nullptr);
devMem = reinterpret_cast<int**>(malloc(sizeof(int*) * devcount));
REQUIRE(devMem != nullptr);
hipGraphExec_t* graphExec =
reinterpret_cast<hipGraphExec_t*>(malloc(devcount * sizeof(hipGraphExec_t)));
hipGraphExec_t* graphExec = reinterpret_cast<hipGraphExec_t*>
(malloc(devcount * sizeof(hipGraphExec_t)));
// Capture stream in each device
for (int dev = 0; dev < devcount; dev++) {
HIP_CHECK(hipSetDevice(dev));
@@ -964,14 +993,15 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_MultiGPU") {
for (int i = 0; i < (dev + 1); i++) {
incrementKernel<<<1, 1, 0, stream[dev]>>>(devMem[dev]);
}
HIP_CHECK(
hipMemcpyAsync(hostMem[dev], devMem[dev], sizeof(int), hipMemcpyDefault, stream[dev]));
HIP_CHECK(hipMemcpyAsync(hostMem[dev], devMem[dev], sizeof(int),
hipMemcpyDefault, stream[dev]));
HIP_CHECK(hipStreamEndCapture(stream[dev], &graph[dev]));
}
// Launch the captured graphs in the respective device
for (int dev = 0; dev < devcount; dev++) {
HIP_CHECK(hipSetDevice(dev));
HIP_CHECK(hipGraphInstantiate(&graphExec[dev], graph[dev], nullptr, nullptr, 0));
HIP_CHECK(hipGraphInstantiate(&graphExec[dev], graph[dev], nullptr,
nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec[dev], stream[dev]));
}
// Validate output
@@ -1038,8 +1068,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_nestedStreamCapture") {
HIP_CHECK(hipEventRecord(events[3], streams[2]));
HIP_CHECK(hipStreamWaitEvent(streams[0], events[3], 0));
HIP_CHECK(hipStreamWaitEvent(streams[0], events[2], 0));
HIP_CHECK(hipMemcpyAsync(hostMem_g.host_ptr(), devMem_g.ptr(), sizeof(int), hipMemcpyDefault,
streams[0]));
HIP_CHECK(hipMemcpyAsync(hostMem_g.host_ptr(), devMem_g.ptr(), sizeof(int),
hipMemcpyDefault, streams[0]));
HIP_CHECK(hipStreamEndCapture(streams[0], &graph)); // End Capture
// Reset device memory
HIP_CHECK(hipMemset(devMem_g.ptr(), 0, sizeof(int)));
@@ -1077,15 +1107,23 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_streamReuse") {
hipGraph_t graphs[3];
StreamsGuard streams(3);
EventsGuard events(4);
LinearAllocGuard<int> hostMem_g1 = LinearAllocGuard<int>(LinearAllocs::malloc, sizeof(int));
LinearAllocGuard<int> hostMem_g2 = LinearAllocGuard<int>(LinearAllocs::malloc, sizeof(int));
LinearAllocGuard<int> hostMem_g3 = LinearAllocGuard<int>(LinearAllocs::malloc, sizeof(int));
LinearAllocGuard<int> devMem_g1 = LinearAllocGuard<int>(LinearAllocs::hipMalloc, sizeof(int));
LinearAllocGuard<int> devMem_g2 = LinearAllocGuard<int>(LinearAllocs::hipMalloc, sizeof(int));
LinearAllocGuard<int> devMem_g3 = LinearAllocGuard<int>(LinearAllocs::hipMalloc, sizeof(int));
LinearAllocGuard<int> hostMem_g1 = LinearAllocGuard<int>
(LinearAllocs::malloc, sizeof(int));
LinearAllocGuard<int> hostMem_g2 = LinearAllocGuard<int>
(LinearAllocs::malloc, sizeof(int));
LinearAllocGuard<int> hostMem_g3 = LinearAllocGuard<int>
(LinearAllocs::malloc, sizeof(int));
LinearAllocGuard<int> devMem_g1 = LinearAllocGuard<int>
(LinearAllocs::hipMalloc, sizeof(int));
LinearAllocGuard<int> devMem_g2 = LinearAllocGuard<int>
(LinearAllocs::hipMalloc, sizeof(int));
LinearAllocGuard<int> devMem_g3 = LinearAllocGuard<int>
(LinearAllocs::hipMalloc, sizeof(int));
std::vector<int*> hostMem = {hostMem_g1.host_ptr(), hostMem_g2.host_ptr(), hostMem_g3.host_ptr()};
std::vector<int*> devMem = {devMem_g1.ptr(), devMem_g2.ptr(), devMem_g3.ptr()};
std::vector<int*> hostMem = {hostMem_g1.host_ptr(), hostMem_g2.host_ptr(),
hostMem_g3.host_ptr()};
std::vector<int*> devMem = {devMem_g1.ptr(), devMem_g2.ptr(),
devMem_g3.ptr()};
// Create a device memory of size int and initialize it to 0
for (int i = 0; i < 3; i++) {
memset(hostMem[i], 0, sizeof(int));
@@ -1109,14 +1147,16 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_streamReuse") {
HIP_CHECK(hipEventRecord(events[3], streams[2]));
HIP_CHECK(hipStreamWaitEvent(streams[0], events[3], 0));
HIP_CHECK(hipStreamWaitEvent(streams[0], events[2], 0));
HIP_CHECK(hipMemcpyAsync(hostMem[0], devMem[0], sizeof(int), hipMemcpyDefault, streams[0]));
HIP_CHECK(hipMemcpyAsync(hostMem[0], devMem[0], sizeof(int),
hipMemcpyDefault, streams[0]));
HIP_CHECK(hipStreamEndCapture(streams[0], &graphs[0])); // End Capture
// Start capturing graph2 from stream 2
HIP_CHECK(hipStreamBeginCapture(streams[1], hipStreamCaptureModeGlobal));
incrementKernel<<<1, 1, 0, streams[1]>>>(devMem[1]);
incrementKernel<<<1, 1, 0, streams[1]>>>(devMem[1]);
incrementKernel<<<1, 1, 0, streams[1]>>>(devMem[1]);
HIP_CHECK(hipMemcpyAsync(hostMem[1], devMem[1], sizeof(int), hipMemcpyDefault, streams[1]));
HIP_CHECK(hipMemcpyAsync(hostMem[1], devMem[1], sizeof(int),
hipMemcpyDefault, streams[1]));
HIP_CHECK(hipStreamEndCapture(streams[1], &graphs[1])); // End Capture
// Start capturing graph3 from stream 3
HIP_CHECK(hipStreamBeginCapture(streams[2], hipStreamCaptureModeGlobal));
@@ -1125,7 +1165,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_streamReuse") {
incrementKernel<<<1, 1, 0, streams[2]>>>(devMem[2]);
incrementKernel<<<1, 1, 0, streams[2]>>>(devMem[2]);
incrementKernel<<<1, 1, 0, streams[2]>>>(devMem[2]);
HIP_CHECK(hipMemcpyAsync(hostMem[2], devMem[2], sizeof(int), hipMemcpyDefault, streams[2]));
HIP_CHECK(hipMemcpyAsync(hostMem[2], devMem[2], sizeof(int),
hipMemcpyDefault, streams[2]));
HIP_CHECK(hipStreamEndCapture(streams[2], &graphs[2])); // End Capture
// Reset device memory
HIP_CHECK(hipMemset(devMem[0], 0, sizeof(int)));
@@ -1169,32 +1210,40 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_captureComplexGraph") {
EventsGuard events(7);
// Allocate Device memory and Host memory
size_t N = GRIDSIZE * BLOCKSIZE;
LinearAllocGuard<int> Ah = LinearAllocGuard<int>(LinearAllocs::malloc, N * sizeof(int));
LinearAllocGuard<int> Bh = LinearAllocGuard<int>(LinearAllocs::malloc, N * sizeof(int));
LinearAllocGuard<int> Ch = LinearAllocGuard<int>(LinearAllocs::malloc, N * sizeof(int));
LinearAllocGuard<int> Ad = LinearAllocGuard<int>(LinearAllocs::hipMalloc, N * sizeof(int));
LinearAllocGuard<int> Bd = LinearAllocGuard<int>(LinearAllocs::hipMalloc, N * sizeof(int));
LinearAllocGuard<int> Ah = LinearAllocGuard<int>
(LinearAllocs::malloc, N * sizeof(int));
LinearAllocGuard<int> Bh = LinearAllocGuard<int>
(LinearAllocs::malloc, N * sizeof(int));
LinearAllocGuard<int> Ch = LinearAllocGuard<int>
(LinearAllocs::malloc, N * sizeof(int));
LinearAllocGuard<int> Ad = LinearAllocGuard<int>
(LinearAllocs::hipMalloc, N * sizeof(int));
LinearAllocGuard<int> Bd = LinearAllocGuard<int>
(LinearAllocs::hipMalloc, N * sizeof(int));
// Capture streams into graph
HIP_CHECK(hipStreamBeginCapture(streams[0], hipStreamCaptureModeGlobal));
HIP_CHECK(hipEventRecord(events[0], streams[0]));
HIP_CHECK(hipStreamWaitEvent(streams[3], events[0], 0));
HIP_CHECK(hipStreamWaitEvent(streams[4], events[0], 0));
HIP_CHECK(
hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), (N * sizeof(int)), hipMemcpyDefault, streams[0]));
HIP_CHECK(
hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), (N * sizeof(int)), hipMemcpyDefault, streams[4]));
HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), (N * sizeof(int)),
hipMemcpyDefault, streams[0]));
HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), (N * sizeof(int)),
hipMemcpyDefault, streams[4]));
hipHostFn_t fn = hostNodeCallback;
HIPCHECK(hipLaunchHostFunc(streams[3], fn, nullptr));
HIP_CHECK(hipEventRecord(events[1], streams[0]));
HIP_CHECK(hipStreamWaitEvent(streams[1], events[1], 0));
int* Ad_2nd_half = Ad.ptr() + N / 2;
int* Ad_1st_half = Ad.ptr();
mymul<<<GRIDSIZE / 2, BLOCKSIZE, 0, streams[0]>>>(Ad_2nd_half, CONST_KER2_VAL);
mymul<<<GRIDSIZE / 2, BLOCKSIZE, 0, streams[1]>>>(Ad_1st_half, CONST_KER1_VAL);
mymul<<<GRIDSIZE / 2, BLOCKSIZE, 0, streams[0]>>>(Ad_2nd_half,
CONST_KER2_VAL);
mymul<<<GRIDSIZE / 2, BLOCKSIZE, 0, streams[1]>>>(Ad_1st_half,
CONST_KER1_VAL);
HIP_CHECK(hipEventRecord(events[2], streams[1]));
HIP_CHECK(hipStreamWaitEvent(streams[2], events[2], 0));
mymul<<<GRIDSIZE / 2, BLOCKSIZE, 0, streams[1]>>>(Ad_1st_half, CONST_KER3_VAL);
mymul<<<GRIDSIZE / 2, BLOCKSIZE, 0, streams[1]>>>(Ad_1st_half,
CONST_KER3_VAL);
HIPCHECK(hipLaunchHostFunc(streams[2], fn, nullptr));
HIP_CHECK(hipEventRecord(events[6], streams[1]));
HIP_CHECK(hipStreamWaitEvent(streams[0], events[6], 0));
@@ -1205,8 +1254,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_captureComplexGraph") {
HIP_CHECK(hipStreamWaitEvent(streams[0], events[3], 0));
HIP_CHECK(hipEventRecord(events[4], streams[3]));
HIP_CHECK(hipStreamWaitEvent(streams[0], events[4], 0));
HIP_CHECK(
hipMemcpyAsync(Ch.host_ptr(), Ad.ptr(), (N * sizeof(int)), hipMemcpyDefault, streams[0]));
HIP_CHECK(hipMemcpyAsync(Ch.host_ptr(), Ad.ptr(), (N * sizeof(int)),
hipMemcpyDefault, streams[0]));
HIP_CHECK(hipStreamEndCapture(streams[0], &graph)); // End Capture
// Execute and test the graph
hipGraphExec_t graphExec{nullptr};
@@ -1219,10 +1268,11 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_captureComplexGraph") {
HIP_CHECK(hipStreamSynchronize(streams[0]));
for (size_t i = 0; i < N; i++) {
if (i > (N / 2 - 1)) {
REQUIRE(Ch.host_ptr()[i] == (Bh.host_ptr()[i] + Ah.host_ptr()[i] * CONST_KER2_VAL));
REQUIRE(Ch.host_ptr()[i] == (Bh.host_ptr()[i] +
Ah.host_ptr()[i] * CONST_KER2_VAL));
} else {
REQUIRE(Ch.host_ptr()[i] ==
(Bh.host_ptr()[i] + Ah.host_ptr()[i] * CONST_KER1_VAL * CONST_KER3_VAL));
REQUIRE(Ch.host_ptr()[i] == (Bh.host_ptr()[i] +
Ah.host_ptr()[i] * CONST_KER1_VAL * CONST_KER3_VAL));
}
}
}
@@ -1267,3 +1317,226 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_captureEmptyStreams") {
HIP_CHECK(hipGraphDestroy(graph));
}
/**
* Test Description
* ------------------------
* - Test to verify hipStreamSynchronize on a stream works when stream capture
* on another stream is ongoing.
* Test source
* ------------------------
* - catch\unit\graph\hipStreamBeginCapture.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
TEST_CASE("Unit_hipStreamBeginCapture_StreamSync_OngoingCapture") {
hipStreamCaptureMode flag = hipStreamCaptureModeRelaxed;
constexpr int GRIDSIZE = 1;
constexpr int BLOCKSIZE = 512;
constexpr int VALUE1 = 7, VALUE2 = 11;
hipGraph_t graph{nullptr};
hipGraphExec_t graphExec{nullptr};
// Allocate device memory
LinearAllocGuard<int> Ah = LinearAllocGuard<int>(LinearAllocs::malloc,
BLOCKSIZE * sizeof(int));
LinearAllocGuard<int> Ad = LinearAllocGuard<int>(LinearAllocs::hipMalloc,
BLOCKSIZE * sizeof(int));
LinearAllocGuard<int> Bh = LinearAllocGuard<int>(LinearAllocs::malloc,
BLOCKSIZE * sizeof(int));
LinearAllocGuard<int> Bd = LinearAllocGuard<int>(LinearAllocs::hipMalloc,
BLOCKSIZE * sizeof(int));
// Fill input data
std::fill_n(Ah.host_ptr(), BLOCKSIZE, VALUE1);
std::fill_n(Bh.host_ptr(), BLOCKSIZE, VALUE2);
// Stream create
StreamsGuard stream0(1);
// Capture streams into graph
SECTION("Stream Creation Before Capture") {
StreamsGuard stream1(1);
HIP_CHECK(hipStreamBeginCapture(stream0[0], flag));
HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream1[0]));
HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream1[0]));
HIP_CHECK(hipStreamSynchronize(stream1[0]));
myadd<<<GRIDSIZE, BLOCKSIZE, 0, stream0[0]>>>(Ad.ptr(), Bd.ptr());
HIP_CHECK(hipStreamEndCapture(stream0[0], &graph)); // End Capture
}
SECTION("Synchronizing multiple streams during Capture") {
StreamsGuard stream1(1), stream2(1);
HIP_CHECK(hipStreamBeginCapture(stream0[0], flag));
HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream1[0]));
HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream2[0]));
HIP_CHECK(hipStreamSynchronize(stream1[0]));
HIP_CHECK(hipStreamSynchronize(stream2[0]));
myadd<<<GRIDSIZE, BLOCKSIZE, 0, stream0[0]>>>(Ad.ptr(), Bd.ptr());
HIP_CHECK(hipStreamEndCapture(stream0[0], &graph)); // End Capture
}
SECTION("Stream Creation After Capture") {
HIP_CHECK(hipStreamBeginCapture(stream0[0], flag));
StreamsGuard stream1(1);
HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream1[0]));
HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream1[0]));
HIP_CHECK(hipStreamSynchronize(stream1[0]));
myadd<<<GRIDSIZE, BLOCKSIZE, 0, stream0[0]>>>(Ad.ptr(), Bd.ptr());
HIP_CHECK(hipStreamEndCapture(stream0[0], &graph)); // End Capture
}
SECTION("Stream Synchronize Before Capture") {
StreamsGuard stream1(1);
HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream1[0]));
HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream1[0]));
HIP_CHECK(hipStreamSynchronize(stream1[0]));
HIP_CHECK(hipStreamBeginCapture(stream0[0], flag));
myadd<<<GRIDSIZE, BLOCKSIZE, 0, stream0[0]>>>(Ad.ptr(), Bd.ptr());
HIP_CHECK(hipStreamEndCapture(stream0[0], &graph)); // End Capture
}
SECTION("Stream Synchronize After Capture") {
HIP_CHECK(hipStreamBeginCapture(stream0[0], flag));
myadd<<<GRIDSIZE, BLOCKSIZE, 0, stream0[0]>>>(Ad.ptr(), Bd.ptr());
HIP_CHECK(hipStreamEndCapture(stream0[0], &graph)); // End Capture
StreamsGuard stream1(1);
HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream1[0]));
HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream1[0]));
HIP_CHECK(hipStreamSynchronize(stream1[0]));
}
// Execute and test the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, stream0[0]));
HIP_CHECK(hipStreamSynchronize(stream0[0]));
// Check output
HIP_CHECK(hipMemcpy(Ah.host_ptr(), Ad.ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDeviceToHost));
for (int idx = 0; idx < BLOCKSIZE; idx++) {
REQUIRE(Ah.host_ptr()[idx] == (VALUE1 + VALUE2));
}
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
}
/**
* Test Description
* ------------------------
* - Test to verify hipStreamSynchronize on a stream behavior when stream capture
* on another stream is ongoing in another thread.
* Test source
* ------------------------
* - catch\unit\graph\hipStreamBeginCapture.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.6
*/
// Local function executed as thread
static void strmSyncThread(int *Ah, int *Ad, int *Bh, int *Bd,
int BLOCKSIZE, hipError_t *error) {
StreamsGuard stream(1);
HIP_CHECK(hipMemcpyAsync(Ad, Ah, BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream[0]));
HIP_CHECK(hipMemcpyAsync(Bd, Bh, BLOCKSIZE * sizeof(int),
hipMemcpyDefault, stream[0]));
*error = hipStreamSynchronize(stream[0]);
}
// Local function executed as thread
static void captureStrmThread(hipGraph_t *graph, int *Ah, int *Ad,
int *Bh, int *Bd, int BLOCKSIZE, int GRIDSIZE,
hipStreamCaptureMode flag, hipError_t *error) {
StreamsGuard stream(1);
// Capture streams into graph
HIP_CHECK(hipStreamBeginCapture(stream[0], flag));
std::thread t1(strmSyncThread, Ah, Ad, Bh, Bd, BLOCKSIZE, error);
t1.join();
myadd<<<GRIDSIZE, BLOCKSIZE, 0, stream[0]>>>(Ad, Bd);
HIP_CHECK(hipStreamEndCapture(stream[0], graph)); // End Capture
}
TEST_CASE("Unit_hipStreamBeginCapture_StreamSync_OngoingCapture_MThread") {
constexpr int GRIDSIZE = 1;
constexpr int BLOCKSIZE = 512;
constexpr int VALUE1 = 7, VALUE2 = 11;
hipGraph_t graph{nullptr};
// Allocate device memory
LinearAllocGuard<int> Ah = LinearAllocGuard<int>(LinearAllocs::malloc,
BLOCKSIZE * sizeof(int));
LinearAllocGuard<int> Ad = LinearAllocGuard<int>(LinearAllocs::hipMalloc,
BLOCKSIZE * sizeof(int));
LinearAllocGuard<int> Bh = LinearAllocGuard<int>(LinearAllocs::malloc,
BLOCKSIZE * sizeof(int));
LinearAllocGuard<int> Bd = LinearAllocGuard<int>(LinearAllocs::hipMalloc,
BLOCKSIZE * sizeof(int));
// Fill input data
std::fill_n(Ah.host_ptr(), BLOCKSIZE, VALUE1);
std::fill_n(Bh.host_ptr(), BLOCKSIZE, VALUE2);
// Stream create
hipError_t error = hipSuccess;
SECTION("Capture Flag = hipStreamCaptureModeGlobal Single Threaded") {
StreamsGuard stream(2);
// Capture streams into graph
HIP_CHECK(hipStreamBeginCapture(stream[0], hipStreamCaptureModeGlobal));
HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(),
BLOCKSIZE * sizeof(int), hipMemcpyDefault, stream[1]));
HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(),
BLOCKSIZE * sizeof(int), hipMemcpyDefault, stream[1]));
error = hipStreamSynchronize(stream[1]);
REQUIRE(error == hipErrorStreamCaptureUnsupported);
}
#if HT_NVIDIA
SECTION("Capture Flag = hipStreamCaptureModeThreadLocal Single Threaded") {
StreamsGuard stream(2);
// Capture streams into graph
HIP_CHECK(hipStreamBeginCapture(stream[0],
hipStreamCaptureModeThreadLocal));
HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(),
BLOCKSIZE * sizeof(int), hipMemcpyDefault, stream[1]));
HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(),
BLOCKSIZE * sizeof(int), hipMemcpyDefault, stream[1]));
error = hipStreamSynchronize(stream[1]);
REQUIRE(error == hipErrorStreamCaptureUnsupported);
}
#endif
#if HT_AMD
SECTION("Capture Flag = hipStreamCaptureModeGlobal Multithreaded") {
captureStrmThread(&graph, Ah.host_ptr(), Ad.ptr(),
Bh.host_ptr(), Bd.ptr(), BLOCKSIZE, GRIDSIZE,
hipStreamCaptureModeGlobal, &error);
REQUIRE(error == hipErrorStreamCaptureUnsupported);
}
#endif
SECTION("Capture Flag = hipStreamCaptureModeThreadLocal Multithreaded") {
captureStrmThread(&graph, Ah.host_ptr(), Ad.ptr(),
Bh.host_ptr(), Bd.ptr(), BLOCKSIZE, GRIDSIZE,
hipStreamCaptureModeThreadLocal, &error);
REQUIRE(error == hipSuccess);
}
SECTION("Capture Flag = hipStreamCaptureModeRelaxed Multithreaded") {
captureStrmThread(&graph, Ah.host_ptr(), Ad.ptr(),
Bh.host_ptr(), Bd.ptr(), BLOCKSIZE, GRIDSIZE,
hipStreamCaptureModeRelaxed, &error);
REQUIRE(error == hipSuccess);
}
if (graph != nullptr) {
hipGraphExec_t graphExec{nullptr};
StreamsGuard stream(1);
// Execute and test the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, stream[0]));
HIP_CHECK(hipStreamSynchronize(stream[0]));
// Check output
HIP_CHECK(hipMemcpy(Ah.host_ptr(), Ad.ptr(), BLOCKSIZE * sizeof(int),
hipMemcpyDeviceToHost));
for (int idx = 0; idx < BLOCKSIZE; idx++) {
REQUIRE(Ah.host_ptr()[idx] == (VALUE1 + VALUE2));
}
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
}
}