AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge 'develop' into 'amd-staging'

Browse Source 
Change-Id: I1db8477632ddecdbdb9a963c7fec6a72308bed03


[ROCm/hip-tests commit: 682ad05404]
This commit is contained in:
Jenkins
2023-01-12 00:09:59 +00:00
parent 609fdd43b3 4f18c1fc48
commit 27d2764a66
7 changed files with 1077 additions and 7 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/multiproc/hipMemGetInfo.cc
+95 -2
View File
@@ -1,5 +1,5 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 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
@@ -296,5 +296,98 @@ TEST_CASE("Unit_hipMemGetInfo_Functional_MultiDevice_Scenario5") {
wait(NULL);
}
}
#endif
#if HT_AMD
static bool testHiddenFreeMemFromChild() {
bool result = true;
int testResult = 0, result_dummy = 0;
int fd_c2p[2], fd_p2c[2];
pipe(fd_c2p);
pipe(fd_p2c);
pid_t cPid;
cPid = fork();
if (cPid == 0) { // child
size_t free = 0, total = 0, min_size = 0;
close(fd_c2p[ReadEnd]);
close(fd_p2c[WriteEnd]);
int64_t size_tohide = (FREE_MEM_TO_HIDE/(1024*1024)); // in MB
// set environment variable from shell
unsetenv("HIP_HIDDEN_FREE_MEM");
setenv("HIP_HIDDEN_FREE_MEM", std::to_string(size_tohide).c_str(), 1);
// allocate memory in device
char* d_ptr{nullptr};
HIP_CHECK(hipMalloc(&d_ptr, SIZE_TO_ALLOCATE));
HIP_CHECK(hipMemGetInfo(&free, &total));
min_size = (FREE_MEM_TO_HIDE + SIZE_TO_ALLOCATE);
if ((total - free) >= min_size) {
testResult = 1;
}
// Write to and signal parent
write(fd_c2p[WriteEnd], &testResult, sizeof(testResult));
close(fd_c2p[WriteEnd]);
// Wait for signal from parent
read(fd_p2c[ReadEnd], &result_dummy, sizeof(result_dummy));
close(fd_p2c[ReadEnd]);
exit(0);
} else if (cPid > 0) { // parent
close(fd_c2p[WriteEnd]);
close(fd_p2c[ReadEnd]);
// wait for result from child
read(fd_c2p[ReadEnd], &testResult, sizeof(testResult));
close(fd_c2p[ReadEnd]);
if (testResult) {
result &= true;
} else {
result &= false;
}
size_t free = 0, total = 0, min_size = SIZE_TO_ALLOCATE;
HIP_CHECK(hipMemGetInfo(&free, &total));
if ((total - free) >= min_size) {
result &= true;
} else {
result &= false;
}
// Write to and signal child
write(fd_p2c[WriteEnd], &result_dummy, sizeof(result_dummy));
close(fd_p2c[WriteEnd]);
wait(NULL);
} else {
WARN("fork() failed");
HIP_ASSERT(false);
}
return result;
}
/**
* Scenario: Fork() a child process. In child, get free and total memory.
* Set the HIP_HIDDEN_FREE_MEM to 4GB. Allocate 2 GB of device memory.
* Get the free and total memory. Free memory available should be
* (actual free - 6 GB). Signal parent process. Wait for signal from child
* in parent. Get free and total memory. Free memory available should be
* actual (actual free - 4 GB).
*/
TEST_CASE("Unit_hipMemGetInfo_SetHiddenFreeMemFromChild") {
REQUIRE(true == testHiddenFreeMemFromChild());
}
/**
* Scenario: Set the HIP_HIDDEN_FREE_MEM to 4GB. Invoke hipMemGetInfo to
* verify that 4GB free memory is hidden for all available GPUs.
*/
TEST_CASE("Unit_hipMemGetInfo_VerifyHiddenFreeMemForAllGpu") {
int numDevices = 0;
int64_t size_tohide = (FREE_MEM_TO_HIDE/(1024*1024)); // in MB
// set environment variable from shell
unsetenv("HIP_HIDDEN_FREE_MEM");
setenv("HIP_HIDDEN_FREE_MEM", std::to_string(size_tohide).c_str(), 1);
HIP_CHECK(hipGetDeviceCount(&numDevices));
for (int dev = 0; dev < numDevices; dev++) {
HIP_CHECK(hipSetDevice(dev));
size_t free = 0, total = 0;
HIP_CHECK(hipMemGetInfo(&free, &total));
REQUIRE((total - free) >= FREE_MEM_TO_HIDE);
}
}
#endif
#endif
projects/hip-tests/catch/unit/device/hipDeviceEnableDisablePeerAccess.cc
+12
View File
@@ -74,6 +74,12 @@ TEST_CASE("Unit_hipDeviceEnablePeerAccess_negative") {
}
SECTION("Peer Access already enabled") {
HIP_CHECK(hipSetDevice(0));
int canAccessPeer = 0;
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 1, 0));
if (canAccessPeer == 0) {
HipTest::HIP_SKIP_TEST("Skipping because no P2P support");
return;
}
HIP_CHECK(hipDeviceEnablePeerAccess(1, 0));
HIP_CHECK_ERROR(hipDeviceEnablePeerAccess(1, 0), hipErrorPeerAccessAlreadyEnabled);
HIP_CHECK(hipDeviceDisablePeerAccess(1));
@@ -97,6 +103,12 @@ TEST_CASE("Unit_hipDeviceDisablePeerAccess_negative") {
}
SECTION("Peer Access disabled twice") {
HIP_CHECK(hipSetDevice(0));
int canAccessPeer = 0;
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 1, 0));
if (canAccessPeer == 0) {
HipTest::HIP_SKIP_TEST("Skipping because no P2P support");
return;
}
HIP_CHECK(hipDeviceEnablePeerAccess(1, 0));
HIP_CHECK(hipDeviceDisablePeerAccess(1));
HIP_CHECK_ERROR(hipDeviceDisablePeerAccess(1), hipErrorPeerAccessNotEnabled);
projects/hip-tests/catch/unit/graph/hipStreamGetCaptureInfo.cc
+394 -4
View File
@@ -36,6 +36,30 @@ Argument Validation/Negative:
1) Pass pId as nullptr and verify api doesn’t crash and returns success.
2) Pass pCaptureStatus as nullptr and verify api doesn’t crash and returns error code.
Extended Scenarios
------------------
1.Create 2 streams s1 and s2. Start capturing s1. Record event e1 on s1 and wait for event e1 on s2. Queue some operations
in s1 and s2. Invoke hipStreamGetCaptureInfo on both s1 and s2. Verify that the capture info (status and id) of both s1 and s2
are identical. Record event e2 on s2 and wait for event e2 on s1. End the capture of stream s1. Verify that the capture info
(status and id) of both s1 and s2 are identical.
2.Create a stream s1. Start capturing s1. Get the capture info of s1. Launch a thread. In the thread get the capture info of s1
using hipStreamGetCaptureInfo. Verify that it is in state hipStreamCaptureStatusActive and capture id inside thread is same as
capture id in main function. Exit the thread and end the capture
3.Verify that the id remains same througout the capture. Create a stream s1. Start capturing s1. Get the capture info of s1.
Queue some oprations in s1. Again get the capture info. Queue different operations in s1. Again get the capture info.
Verify that all the capture info are identical.
4.Create a stream with default flag (hipStreamDefault). Start capturing the stream. Invoke hipStreamGetCaptureInfo() on the null
stream. Verify hipErrorStreamCaptureImplicit is returned by hipStreamGetCaptureInfo(). Verify capture status of created stream.
Do some operatoins. End the capture on the created stream.Verify the capture status. Execute the graph and verify the output
from the operations.
5. Test scenario 1 using hipStreamGetCaptureInfo_v2.
6. Test scenario 2 using hipStreamGetCaptureInfo_v2.
7. Test scenario 3 using hipStreamGetCaptureInfo_v2.
8. Test scenario 4 using hipStreamGetCaptureInfo_v2.
*/
#include <hip_test_common.hh>
@@ -43,6 +67,9 @@ Argument Validation/Negative:
#include <hip_test_kernels.hh>
constexpr size_t N = 1000000;
constexpr unsigned blocks = 512;
constexpr unsigned threadsPerBlock = 256;
size_t Nbytes = N * sizeof(float);
constexpr int LAUNCH_ITERS = 1;
/**
@@ -53,9 +80,6 @@ void validateStreamCaptureInfo(hipStream_t mstream) {
hipEvent_t memsetEvent1, memsetEvent2, forkStreamEvent;
hipGraph_t graph{nullptr};
hipGraphExec_t graphExec{nullptr};
constexpr unsigned blocks = 512;
constexpr unsigned threadsPerBlock = 256;
size_t Nbytes = N * sizeof(float);
float *A_d, *C_d;
float *A_h, *C_h;
A_h = reinterpret_cast<float*>(malloc(Nbytes));
@@ -176,7 +200,7 @@ TEST_CASE("Unit_hipStreamGetCaptureInfo_UniqueID") {
hipStream_t streams[numStreams]{};
hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone};
std::vector<int> idlist;
unsigned long long capSequenceID{}; // NOLINT
unsigned long long capSequenceID{}; //NOLINT
hipGraph_t graph{nullptr};
for (int i = 0; i < numStreams; i++) {
@@ -229,3 +253,369 @@ TEST_CASE("Unit_hipStreamGetCaptureInfo_ArgValidation") {
HIP_CHECK(hipStreamDestroy(stream));
}
/*
* Create 2 streams s1 and s2. Start capturing s1. Record event e1 on s1 and
* wait for event e1 on s2. Queue some operations in s1 and s2. Invoke
* hipStreamGetCaptureInfo on both s1 and s2. Verify that the capture info
* (status and id) of both s1 and s2 are identical. Record event e2 on s2
* and wait for event e2 on s1. End the capture of stream s1. Verify that the
* capture info (status and id) of both s1 and s2 are identical.
* The above scenario using hipStreamGetCaptureInfo_v2 API
*/
TEST_CASE("Unit_hipStreamGetCaptureInfo_ParentAndForkedStrm_CaptureStatus") {
hipStream_t stream1{nullptr}, stream2{nullptr};
hipEvent_t event2{nullptr}, forkStreamEvent{nullptr};
hipGraph_t graph{nullptr};
float *A_d, *B_d, *C_d, *D_d;
float *A_h, *B_h, *C_h, *D_h;
// Memory allocation to Host pointers
A_h = reinterpret_cast<float*>(malloc(Nbytes));
B_h = reinterpret_cast<float*>(malloc(Nbytes));
C_h = reinterpret_cast<float*>(malloc(Nbytes));
D_h = reinterpret_cast<float*>(malloc(Nbytes));
REQUIRE(A_h != nullptr);
REQUIRE(B_h != nullptr);
REQUIRE(C_h != nullptr);
REQUIRE(D_h != nullptr);
// Memory allocation to Device pointers
HIP_CHECK(hipMalloc(&A_d, Nbytes));
HIP_CHECK(hipMalloc(&B_d, Nbytes));
HIP_CHECK(hipMalloc(&C_d, Nbytes));
HIP_CHECK(hipMalloc(&D_d, Nbytes));
REQUIRE(A_d != nullptr);
REQUIRE(B_d != nullptr);
REQUIRE(C_d != nullptr);
REQUIRE(D_d != nullptr);
HIP_CHECK(hipStreamCreate(&stream1));
HIP_CHECK(hipStreamCreate(&stream2));
HIP_CHECK(hipEventCreate(&event2));
HIP_CHECK(hipEventCreate(&forkStreamEvent));
// Start capture on stream1
HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal));
HIP_CHECK(hipEventRecord(forkStreamEvent, stream1));
HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0));
// Copy data to Device
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream1));
HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream2));
// Kernal Operations
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, stream1, A_d, C_d, N);
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, stream2, B_d, D_d, N);
// Copy data back to the Host
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream1));
HIP_CHECK(hipMemcpyAsync(D_h, D_d, Nbytes, hipMemcpyDeviceToHost, stream2));
hipStreamCaptureStatus captureStatus1{hipStreamCaptureStatusNone},
captureStatus2{hipStreamCaptureStatusNone},
captureStatus3{hipStreamCaptureStatusNone},
captureStatus4{hipStreamCaptureStatusNone};
unsigned long long capSequenceID1, capSequenceID2, capSequenceID3, //NOLINT
capSequenceID4;
SECTION("hipStreamGetCaptureInfo verification before End capture") {
// Capture info
HIP_CHECK(hipStreamGetCaptureInfo(stream1, &captureStatus1,
&capSequenceID1));
HIP_CHECK(hipStreamGetCaptureInfo(stream2, &captureStatus2,
&capSequenceID2));
// Verfication of results
REQUIRE(capSequenceID1 == capSequenceID2);
REQUIRE(captureStatus1 == hipStreamCaptureStatusActive);
REQUIRE(captureStatus2 == hipStreamCaptureStatusActive);
}
SECTION("hipStreamGetCaptureInfo_v2 verification before End capture") {
// Capture info
HIP_CHECK(hipStreamGetCaptureInfo_v2(stream1, &captureStatus1,
&capSequenceID1, nullptr, nullptr, nullptr));
HIP_CHECK(hipStreamGetCaptureInfo_v2(stream2, &captureStatus2,
&capSequenceID2, nullptr, nullptr, nullptr));
// Verfication of results
REQUIRE(capSequenceID1 == capSequenceID2);
REQUIRE(captureStatus1 == hipStreamCaptureStatusActive);
REQUIRE(captureStatus2 == hipStreamCaptureStatusActive);
}
HIP_CHECK(hipEventRecord(event2, stream2));
HIP_CHECK(hipStreamWaitEvent(stream1, event2, 0));
// End the capture
HIP_CHECK(hipStreamEndCapture(stream1, &graph));
REQUIRE(graph != nullptr);
SECTION("hipStreamGetCaptureInfo verification after End capture") {
// Capture Info
HIP_CHECK(hipStreamGetCaptureInfo(stream1, &captureStatus3,
&capSequenceID3));
HIP_CHECK(hipStreamGetCaptureInfo(stream2, &captureStatus4,
&capSequenceID4));
// Verification of results
REQUIRE(captureStatus3 == hipStreamCaptureStatusNone);
REQUIRE(captureStatus4 == hipStreamCaptureStatusNone);
}
SECTION("hipStreamGetCaptureInfo_v2 verification after End capture") {
// Capture Info
HIP_CHECK(hipStreamGetCaptureInfo_v2(stream1, &captureStatus3,
&capSequenceID3, nullptr, nullptr, nullptr));
HIP_CHECK(hipStreamGetCaptureInfo_v2(stream2, &captureStatus4,
&capSequenceID4, nullptr, nullptr, nullptr));
// Verification of results
REQUIRE(captureStatus3 == hipStreamCaptureStatusNone);
REQUIRE(captureStatus4 == hipStreamCaptureStatusNone);
}
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(stream1));
HIP_CHECK(hipStreamDestroy(stream2));
HIP_CHECK(hipEventDestroy(forkStreamEvent));
HIP_CHECK(hipEventDestroy(event2));
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(B_d));
HIP_CHECK(hipFree(C_d));
HIP_CHECK(hipFree(D_d));
free(A_h);
free(B_h);
free(C_h);
free(D_h);
}
// Thread Function
static void thread_func(hipStream_t stream, unsigned long long capSequenceID1, //NOLINT
unsigned long long capSequenceID2) { //NOLINT
hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone};
unsigned long long capSequenceID3, capSequenceID4; //NOLINT
SECTION("hipStreamGetCaptureInfo CaptureStatus in Thread") {
HIP_CHECK(hipStreamGetCaptureInfo(stream, &captureStatus, &capSequenceID3));
REQUIRE(capSequenceID1 == capSequenceID3);
REQUIRE(captureStatus == hipStreamCaptureStatusActive);
}
SECTION("hipStreamGetCaptureInfo_v2 CaptureStatus in Thread") {
HIP_CHECK(hipStreamGetCaptureInfo_v2(stream, &captureStatus,
&capSequenceID4, nullptr, nullptr, nullptr));
REQUIRE(capSequenceID2 == capSequenceID4);
REQUIRE(captureStatus == hipStreamCaptureStatusActive);
}
}
/*
* Create a stream s1. Start capturing s1. Get the capture info of s1. Launch
* a thread. In the thread get the capture info of s1 using hipStreamGetCaptureInfo.
* Verify that it is in state hipStreamCaptureStatusActive and capture id inside
* thread is same as capture id in main function. Exit the thread and end the capture
* The above scenario using hipStreamGetCaptureInfo_v2 API
*/
TEST_CASE("Unit_hipStreamGetCaptureInfo_CaptureStatus_InThread") {
hipStream_t stream{nullptr};
hipGraph_t graph{nullptr};
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
// Capture info
hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone};
unsigned long long capSequenceID1, capSequenceID2; //NOLINT
// hipStreamGetCaptureInfo Capture status
HIP_CHECK(hipStreamGetCaptureInfo(stream, &captureStatus, &capSequenceID1));
// hipStreamGetCaptureInfo_v2 Capture status
HIP_CHECK(hipStreamGetCaptureInfo_v2(stream, &captureStatus,
&capSequenceID2, nullptr, nullptr, nullptr));
// Thread launch
std::thread t(thread_func, stream, capSequenceID1, capSequenceID2);
t.join();
HIP_CHECK(hipStreamEndCapture(stream, &graph));
REQUIRE(graph != nullptr);
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(stream));
}
/*
* Verify that the id remains same througout the capture. Create a stream s1.
* Start capturing s1. Get the capture info of s1. Queue some oprations in s1.
* Again get the capture info. Queue different operations in s1. Again get the
* capture info. Verify that all the capture info are identical.
* The above scenario using hipStreamGetCaptureInfo_v2 API
*/
TEST_CASE("Unit_hipStreamGetCaptureInfo_CaptureStatus_Througout_Capture") {
hipStream_t stream{nullptr};
hipGraph_t graph{nullptr};
float *A_d, *B_d, *C_d, *D_d;
float *A_h, *B_h, *C_h, *D_h;
// Memory allocation to Host pointers
A_h = reinterpret_cast<float*>(malloc(Nbytes));
B_h = reinterpret_cast<float*>(malloc(Nbytes));
C_h = reinterpret_cast<float*>(malloc(Nbytes));
D_h = reinterpret_cast<float*>(malloc(Nbytes));
REQUIRE(A_h != nullptr);
REQUIRE(B_h != nullptr);
REQUIRE(C_h != nullptr);
REQUIRE(D_h != nullptr);
// Memory allocation to Device pointers
HIP_CHECK(hipMalloc(&A_d, Nbytes));
HIP_CHECK(hipMalloc(&B_d, Nbytes));
HIP_CHECK(hipMalloc(&C_d, Nbytes));
HIP_CHECK(hipMalloc(&D_d, Nbytes));
REQUIRE(A_d != nullptr);
REQUIRE(B_d != nullptr);
REQUIRE(C_d != nullptr);
REQUIRE(D_d != nullptr);
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
// Capture Info
hipStreamCaptureStatus captureStatus1{hipStreamCaptureStatusNone},
captureStatus2{hipStreamCaptureStatusNone},
captureStatus3{hipStreamCaptureStatusNone},
captureStatus4{hipStreamCaptureStatusNone},
captureStatus5{hipStreamCaptureStatusNone},
captureStatus6{hipStreamCaptureStatusNone};
unsigned long long capSequenceID1, capSequenceID2, capSequenceID3, //NOLINT
capSequenceID4, capSequenceID5, capSequenceID6;
// hipStreamGetCaptureInfo Capture status
HIP_CHECK(hipStreamGetCaptureInfo(stream, &captureStatus1, &capSequenceID1));
// hipStreamGetCaptureInfo_v2 Capture status
HIP_CHECK(hipStreamGetCaptureInfo_v2(stream, &captureStatus2,
&capSequenceID2, nullptr, nullptr, nullptr));
// Copy data to Device
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
// Kernal Operations
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, stream, A_d, C_d, N);
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));
// hipStreamGetCaptureInfo Capture status
HIP_CHECK(hipStreamGetCaptureInfo(stream, &captureStatus3, &capSequenceID3));
REQUIRE(captureStatus1 == captureStatus3);
REQUIRE(capSequenceID1 == capSequenceID3);
// hipStreamGetCaptureInfo_v2 Capture status
HIP_CHECK(hipStreamGetCaptureInfo_v2(stream, &captureStatus4,
&capSequenceID4, nullptr, nullptr, nullptr));
REQUIRE(captureStatus2 == captureStatus4);
REQUIRE(capSequenceID2 == capSequenceID4);
// Kernal Operations
HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks),
dim3(threadsPerBlock), 0, stream, A_d, B_d, D_d, N);
HIP_CHECK(hipMemcpyAsync(D_h, D_d, Nbytes, hipMemcpyDeviceToHost, stream));
// hipStreamGetCaptureInfo Capture status
HIP_CHECK(hipStreamGetCaptureInfo(stream, &captureStatus5, &capSequenceID5));
REQUIRE(captureStatus3 == captureStatus5);
REQUIRE(capSequenceID3 == capSequenceID5);
// hipStreamGetCaptureInfo_v2 Capture status
HIP_CHECK(hipStreamGetCaptureInfo_v2(stream, &captureStatus6,
&capSequenceID6, nullptr, nullptr, nullptr));
REQUIRE(captureStatus4 == captureStatus6);
REQUIRE(capSequenceID4 == capSequenceID6);
HIP_CHECK(hipStreamEndCapture(stream, &graph));
REQUIRE(graph != nullptr);
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(B_d));
HIP_CHECK(hipFree(C_d));
HIP_CHECK(hipFree(D_d));
free(A_h);
free(B_h);
free(C_h);
free(D_h);
}
/*
* Create a stream with default flag (hipStreamDefault). Start capturing the stream.
* Invoke hipStreamGetCaptureInfo() on the null stream. Verify hipErrorStreamCaptureImplicit
* is returned by hipStreamGetCaptureInfo(). Verify capture status of created stream. Do some
* operatoins. End the capture on the created stream.Verify the capture status. Execute the
* graph and verify the output from the operations.
* The above scenario using hipStreamGetCaptureInfo_v2 API
*/
TEST_CASE("Unit_hipStreamGetCaptureInfo_Nullstream_CaptureInfo") {
hipStream_t stream{nullptr}, streamForGraph{nullptr};
hipGraph_t graph{nullptr};
hipError_t ret;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipStreamCreate(&streamForGraph));
float *A_d, *C_d;
float *A_h, *C_h, *D_h;
// Memory allocation to Host pointers
A_h = reinterpret_cast<float*>(malloc(Nbytes));
C_h = reinterpret_cast<float*>(malloc(Nbytes));
D_h = reinterpret_cast<float*>(malloc(Nbytes));
REQUIRE(A_h != nullptr);
REQUIRE(C_h != nullptr);
REQUIRE(D_h != nullptr);
// Memory allocation to Device pointers
HIP_CHECK(hipMalloc(&A_d, Nbytes));
HIP_CHECK(hipMalloc(&C_d, Nbytes));
REQUIRE(A_d != nullptr);
REQUIRE(C_d != nullptr);
// Initialize input buffer
for (size_t i = 0; i < N; ++i) {
A_h[i] = 1.0f + i;
D_h[i] = 0.0f;
}
HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone},
captureStatus1{hipStreamCaptureStatusNone},
captureStatus2{hipStreamCaptureStatusNone};
unsigned long long capSequenceID = 0, // NOLINT
capSequenceID1 = 0;
// Verify the Error returned with null stream.
SECTION("hipStreamGetCaptureInfo with null stream") {
ret = hipStreamGetCaptureInfo(0, &captureStatus, &capSequenceID);
REQUIRE(ret == hipErrorStreamCaptureImplicit);
}
SECTION("hipStreamGetCaptureInfo_v2 with null stream") {
ret = hipStreamGetCaptureInfo_v2(0, &captureStatus, &capSequenceID,
nullptr, nullptr, nullptr);
REQUIRE(ret == hipErrorStreamCaptureImplicit);
}
// Check the capture status of the stream
HIP_CHECK(hipStreamIsCapturing(stream, &captureStatus1));
REQUIRE(captureStatus1 == hipStreamCaptureStatusActive);
// Copy data to Device
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
// Kernal Operation
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, stream, A_d, C_d, N);
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));
// End the capture
HIP_CHECK(hipStreamEndCapture(stream, &graph));
REQUIRE(graph != nullptr);
// Capture Status
SECTION("hipStreamGetCaptureInfo with null stream after End capture") {
ret = hipStreamGetCaptureInfo(0, &captureStatus2, &capSequenceID1);
REQUIRE(ret == hipSuccess);
}
SECTION("hipStreamGetCaptureInfo_v2 with null stream after End capture") {
ret = hipStreamGetCaptureInfo_v2(0, &captureStatus2, &capSequenceID1,
nullptr, nullptr, nullptr);
REQUIRE(ret == hipSuccess);
}
// Launch graph
hipGraphExec_t graphExec;
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Verify Output
for (size_t i = 0; i < N; i++) {
D_h[i] = A_h[i] * A_h[i];
REQUIRE(C_h[i] == D_h[i]);
}
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipStreamDestroy(streamForGraph));
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(C_d));
free(A_h);
free(C_h);
free(D_h);
}
projects/hip-tests/catch/unit/graph/hipStreamIsCapturing.cc
+222 -1
View File
@@ -45,9 +45,22 @@ Functional Testcase Scenarios :
capture status returned as hipStreamCaptureStatusActive.
8) Functional : Stop capturing using hipStreamPerThread and check
status is returned as hipStreamCaptureStatusNone.
9) Functional : Create 2 streams s1 and s2. Start capturing s1. Record event e1
on s1 and wait for event e1 on s2. Queue some operations in s1 and s2. Invoke
hipStreamIsCapturing on both s1 and s2. Verify that the capture info (status)
of both s1 and s2 are identical. Record event e2 on s2 and wait for event e2
on s1. End the capture of stream s1. Invoke hipStreamIsCapturing on both streams.
Verify that the capture info(status)of both s1 and s2 are identical
10)Functional : Create a stream s1. Start capturing s1. Get the capture info using
hipStreamIsCapturing of s1. Launch a thread. In the thread get the capture info
of s1 using hipStreamIsCapturing. Verify that it is in state hipStreamCaptureStatusActive
in thread. Exit the thread and end the capture.
11)Functional : Create a stream with default flag (hipStreamDefault). Start capturing
the stream. Invoke hipStreamIsCapturing() on the null stream. Verify hipErrorStreamCaptureImplicit
is returned by hipStreamIsCapturing(). Verify capture status of created stream. Do some operatoins.
End the capture on the created stream. Execute the graph and verify the output from the operations.
*/
TEST_CASE("Unit_hipStreamIsCapturing_Negative") {
hipError_t ret;
hipStream_t stream{};
@@ -213,3 +226,211 @@ TEST_CASE("Unit_hipStreamIsCapturing_hipStreamPerThread") {
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(C_d));
}
/*
* Create 2 streams s1 and s2. Start capturing s1. Record event e1 on s1 and wait
* for event e1 on s2. Queue some operations in s1 and s2. Invoke hipStreamIsCapturing
* on both s1 and s2. Verify that the capture info (status) of both s1 and s2 are identical.
* Record event e2 on s2 and wait for event e2 on s1. End the capture of stream s1.
* Invoke hipStreamIsCapturing on both streams. Verify that the capture info(status)
* of both s1 and s2 are identical.
*/
TEST_CASE("Unit_hipStreamIsCapturing_ParentAndForkedStream") {
hipStream_t stream1{nullptr}, stream2{nullptr};
hipEvent_t event2{nullptr}, forkStreamEvent{nullptr};
hipGraph_t graph{nullptr};
constexpr unsigned blocks = 512;
constexpr unsigned threadsPerBlock = 256;
size_t Nbytes = N * sizeof(float);
float *A_d, *B_d, *C_d, *D_d;
float *A_h, *B_h, *C_h, *D_h;
// Memory allocation to Host pointers
A_h = reinterpret_cast<float*>(malloc(Nbytes));
B_h = reinterpret_cast<float*>(malloc(Nbytes));
C_h = reinterpret_cast<float*>(malloc(Nbytes));
D_h = reinterpret_cast<float*>(malloc(Nbytes));
REQUIRE(A_h != nullptr);
REQUIRE(B_h != nullptr);
REQUIRE(C_h != nullptr);
REQUIRE(D_h != nullptr);
// Memory allocation to Device pointers
HIP_CHECK(hipMalloc(&A_d, Nbytes));
HIP_CHECK(hipMalloc(&B_d, Nbytes));
HIP_CHECK(hipMalloc(&C_d, Nbytes));
HIP_CHECK(hipMalloc(&D_d, Nbytes));
REQUIRE(A_d != nullptr);
REQUIRE(B_d != nullptr);
REQUIRE(C_d != nullptr);
REQUIRE(D_d != nullptr);
// Initialize input buffer
for (size_t i = 0; i < N; ++i) {
A_h[i] = 3.146f + i; // Pi
B_h[i] = A_h[i];
}
HIP_CHECK(hipStreamCreate(&stream1));
HIP_CHECK(hipStreamCreate(&stream2));
HIP_CHECK(hipEventCreate(&event2));
HIP_CHECK(hipEventCreate(&forkStreamEvent));
// Start capture on stream1
HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal));
HIP_CHECK(hipEventRecord(forkStreamEvent, stream1));
HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0));
// Copy data to Device
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream1));
HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream2));
// Kernal Operations
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, stream1, A_d, C_d, N);
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, stream2, B_d, D_d, N);
// Copy data back to the Host
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream1));
HIP_CHECK(hipMemcpyAsync(D_h, D_d, Nbytes, hipMemcpyDeviceToHost, stream2));
hipStreamCaptureStatus captureStatus1{hipStreamCaptureStatusNone},
captureStatus2{hipStreamCaptureStatusNone},
captureStatus3{hipStreamCaptureStatusNone},
captureStatus4{hipStreamCaptureStatusNone};
// Capturing info
HIP_CHECK(hipStreamIsCapturing(stream1, &captureStatus1));
HIP_CHECK(hipStreamIsCapturing(stream2, &captureStatus2));
// Verfication of results
REQUIRE(captureStatus1 == hipStreamCaptureStatusActive);
REQUIRE(captureStatus2 == hipStreamCaptureStatusActive);
HIP_CHECK(hipEventRecord(event2, stream2));
HIP_CHECK(hipStreamWaitEvent(stream1, event2, 0));
// End the capture
HIP_CHECK(hipStreamEndCapture(stream1, &graph));
REQUIRE(graph != nullptr);
// Capture Info
HIP_CHECK(hipStreamIsCapturing(stream1, &captureStatus3));
HIP_CHECK(hipStreamIsCapturing(stream2, &captureStatus4));
// Verification of results
REQUIRE(captureStatus3 == hipStreamCaptureStatusNone);
REQUIRE(captureStatus4 == hipStreamCaptureStatusNone);
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(stream1));
HIP_CHECK(hipStreamDestroy(stream2));
HIP_CHECK(hipEventDestroy(forkStreamEvent));
HIP_CHECK(hipEventDestroy(event2));
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(B_d));
HIP_CHECK(hipFree(C_d));
HIP_CHECK(hipFree(D_d));
free(A_h);
free(B_h);
free(C_h);
free(D_h);
}
/*
* Create a stream s1. Start capturing s1. Get the capture info using hipStreamIsCapturing
* of s1. Launch a thread. In the thread get the capture info of s1 using hipStreamIsCapturing.
* Verify that it is in state hipStreamCaptureStatusActive in thread. Exit the thread and end
* the capture.
*/
// Thread Function
static void thread_func(hipStream_t stream) {
hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone};
HIP_CHECK(hipStreamIsCapturing(stream, &captureStatus));
REQUIRE(captureStatus == hipStreamCaptureStatusActive);
}
TEST_CASE("Unit_hipStreamIsCapturing_CheckCaptureStatus_FromThread") {
hipStream_t stream{nullptr};
hipGraph_t graph{nullptr};
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
// Capture info
hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone};
HIP_CHECK(hipStreamIsCapturing(stream, &captureStatus));
REQUIRE(captureStatus == hipStreamCaptureStatusActive);
// Thread launch
std::thread t(thread_func, stream);
t.join();
HIP_CHECK(hipStreamEndCapture(stream, &graph));
REQUIRE(graph != nullptr);
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(stream));
}
/*
* Create a stream with default flag (hipStreamDefault). Start capturing the stream.
* Invoke hipStreamIsCapturing() on the null stream. Verify hipErrorStreamCaptureImplicit
* is returned by hipStreamIsCapturing(). Verify capture status of created stream. Do some operatoins.
* End the capture on the created stream. Execute the graph and verify the output from the operations.
*/
TEST_CASE("Unit_hipStreamIsCapturing_ChkNullStrmStatus") {
hipStream_t stream{nullptr}, streamForGraph{nullptr};
hipGraph_t graph{nullptr};
hipError_t ret;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipStreamCreate(&streamForGraph));
float *A_d, *C_d;
float *A_h, *C_h, *D_h;
// Memory allocation to Host pointers
A_h = reinterpret_cast<float*>(malloc(Nbytes));
C_h = reinterpret_cast<float*>(malloc(Nbytes));
D_h = reinterpret_cast<float*>(malloc(Nbytes));
REQUIRE(A_h != nullptr);
REQUIRE(C_h != nullptr);
REQUIRE(D_h != nullptr);
// Memory allocation to Device pointers
HIP_CHECK(hipMalloc(&A_d, Nbytes));
HIP_CHECK(hipMalloc(&C_d, Nbytes));
REQUIRE(A_d != nullptr);
REQUIRE(C_d != nullptr);
// Initialize input buffer
for (size_t i = 0; i < N; ++i) {
A_h[i] = 1.0f + i;
D_h[i] = 0.0f;
}
HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone},
captureStatus1{hipStreamCaptureStatusNone},
captureStatus2{hipStreamCaptureStatusNone};
// Verify the Error returned if null stream is passed.
ret = hipStreamIsCapturing(0, &captureStatus);
REQUIRE(ret == hipErrorStreamCaptureImplicit);
// Check the capture status of the stream
HIP_CHECK(hipStreamIsCapturing(stream, &captureStatus1));
REQUIRE(captureStatus1 == hipStreamCaptureStatusActive);
// Copy data to Device
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
// Kernal Operations
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0, stream, A_d, C_d, N);
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));
// End the capture
HIP_CHECK(hipStreamEndCapture(stream, &graph));
REQUIRE(graph != nullptr);
ret = hipStreamIsCapturing(0, &captureStatus2);
REQUIRE(ret == hipSuccess);
// Launch graph
hipGraphExec_t graphExec;
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Verify Output
for (size_t i = 0; i < N; i++) {
D_h[i] = A_h[i] * A_h[i];
REQUIRE(C_h[i] == D_h[i]);
}
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipStreamDestroy(streamForGraph));
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(C_d));
free(A_h);
free(C_h);
free(D_h);
}
projects/hip-tests/catch/unit/memory/hipMemAdvise.cc
+2
View File
@@ -224,6 +224,8 @@ TEST_CASE("Unit_hipMemAdvise_TstFlags") {
}
TEST_CASE("Unit_hipMemAdvise_NegtveTsts") {
HipTest::HIP_SKIP_TEST("Fixed few issues to match with Nvidia, Skip now to avoid CI failures");
return;
int MangdMem = HmmAttrPrint();
if (MangdMem == 1) {
bool IfTestPassed = true;
projects/hip-tests/catch/unit/occupancy/CMakeLists.txt
+1
View File
@@ -2,6 +2,7 @@
set(TEST_SRC
hipOccupancyMaxActiveBlocksPerMultiprocessor.cc
hipOccupancyMaxPotentialBlockSize.cc
hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags.cc
)
hip_add_exe_to_target(NAME OccupancyTest
projects/hip-tests/catch/unit/occupancy/hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags.cc
+351
View File
@@ -0,0 +1,351 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#define SHARED_MEM_CONST 256
#define UNUSED(expr) do { (void)(expr); } while (0)
// global variables
static int gArrSize = 0;
// sample global functions
static __global__ void f1(float *a) { *a = 1.0; }
// Dynamic shared
static __global__ void copyKerDyn(int* out, int* in) {
extern __shared__ int sharedMem[];
int tid = blockDim.x * blockIdx.x + threadIdx.x;
sharedMem[tid] = in[tid];
__syncthreads();
out[tid] = sharedMem[tid];
}
// Without Dynamic shared
static __global__ void copyKer(int* out, int* in) {
int tid = blockDim.x * blockIdx.x + threadIdx.x;
out[tid] = in[tid];
}
// sample function
static size_t blockSizeToDynamicSMemSize(int blocksize) {
return (static_cast<size_t>(blocksize*SHARED_MEM_CONST));
}
// sample functor
class functorBlockSizeToDynamicSMemSize {
int myconst;
public:
explicit functorBlockSizeToDynamicSMemSize(int n):myconst(n) {
}
int operator () (int blocksize) const {
return (static_cast<size_t>(blocksize*myconst));
}
};
/**
Local function to check hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags
functionality for different block_size_limit.
*/
void hipOccupancyMaxPotBlkSizeVariableSMemWithFlags_chkRange(
int block_size_limit, int maxThreadsPerBlock) {
int minGridSize = 0, blockSize = 0;
hipError_t ret;
// Get potential blocksize
ret = hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags(&minGridSize,
&blockSize, f1, blockSizeToDynamicSMemSize, block_size_limit, 0);
REQUIRE(ret == hipSuccess);
REQUIRE(minGridSize > 0);
REQUIRE(blockSize > 0);
REQUIRE(blockSize <= maxThreadsPerBlock);
}
/**
Check the basic functionality of hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags
- for block_size_limit = 0
- for 0 < block_size_limit < attr.maxThreadsPerBlock
- for block_size_limit > attr.maxThreadsPerBlock
*/
TEST_CASE("Unit_hipOccupancyMaxPotBlkSizeVariableSMemWithFlags_chkRange") {
hipDeviceProp_t devProp;
// Get current device property
HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
SECTION("block_size_limit = 0") {
hipOccupancyMaxPotBlkSizeVariableSMemWithFlags_chkRange(0,
devProp.maxThreadsPerBlock);
}
SECTION("block_size_limit < maxThreadsPerBlock") {
hipOccupancyMaxPotBlkSizeVariableSMemWithFlags_chkRange(
(devProp.maxThreadsPerBlock - 1), devProp.maxThreadsPerBlock);
}
SECTION("block_size_limit = maxThreadsPerBlock") {
hipOccupancyMaxPotBlkSizeVariableSMemWithFlags_chkRange(
devProp.maxThreadsPerBlock, devProp.maxThreadsPerBlock);
}
SECTION("block_size_limit > maxThreadsPerBlock") {
hipOccupancyMaxPotBlkSizeVariableSMemWithFlags_chkRange(
(devProp.maxThreadsPerBlock + 1), devProp.maxThreadsPerBlock);
}
}
/**
Check range of minGridSize and blockSize for multiple GPU
- for block_size_limit = 0
- for 0 < block_size_limit < attr.maxThreadsPerBlock
- for block_size_limit > attr.maxThreadsPerBlock
*/
TEST_CASE("Unit_hipOccupancyMaxPotBlkSizeVariableSMemWithFlags_mgpu") {
int devcount = 0;
HIP_CHECK(hipGetDeviceCount(&devcount));
// If only single GPU is detected then return
if (devcount < 2) {
SUCCEED("Skipping the test as number of Devices found less than 2");
return;
}
// Get current device property
for (int dev = 0; dev < devcount; dev++) {
hipDeviceProp_t devProp;
HIP_CHECK(hipGetDeviceProperties(&devProp, dev));
HIP_CHECK(hipSetDevice(dev));
hipOccupancyMaxPotBlkSizeVariableSMemWithFlags_chkRange(0,
devProp.maxThreadsPerBlock);
hipOccupancyMaxPotBlkSizeVariableSMemWithFlags_chkRange(
(devProp.maxThreadsPerBlock - 1), devProp.maxThreadsPerBlock);
hipOccupancyMaxPotBlkSizeVariableSMemWithFlags_chkRange(
devProp.maxThreadsPerBlock, devProp.maxThreadsPerBlock);
HIP_CHECK(hipSetDevice(0));
}
}
/**
Check the basic functionality of hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags
by passing a functor as 4th parameter.
*/
TEST_CASE("Unit_hipOccupancyMaxPotBlkSizeVariableSMemWithFlags_Functor") {
hipDeviceProp_t devProp;
HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
functorBlockSizeToDynamicSMemSize testFunc(SHARED_MEM_CONST);
// Get current device property
int minGridSize = 0, blockSize = 0;
hipError_t ret;
// Get potential blocksize
ret = hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags(&minGridSize,
&blockSize, f1, testFunc, 0, 0);
REQUIRE(ret == hipSuccess);
REQUIRE(minGridSize > 0);
REQUIRE(blockSize > 0);
REQUIRE(blockSize <= devProp.maxThreadsPerBlock);
}
/**
Check the basic functionality of hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags
by passing a lambda function as 4th parameter.
*/
TEST_CASE("Unit_hipOccupancyMaxPotBlkSizeVariableSMemWithFlags_Lambda") {
hipDeviceProp_t devProp;
HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
auto testFunc = [](const int blockSize){
return (static_cast<size_t>(blockSize*SHARED_MEM_CONST));
};
// Get current device property
int minGridSize = 0, blockSize = 0;
hipError_t ret;
// Get potential blocksize
ret = hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags(&minGridSize,
&blockSize, f1, testFunc, 0, 0);
REQUIRE(ret == hipSuccess);
REQUIRE(minGridSize > 0);
REQUIRE(blockSize > 0);
REQUIRE(blockSize <= devProp.maxThreadsPerBlock);
// Test again by passing the lamda function directly
ret = hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags(&minGridSize,
&blockSize, f1,
[](const int blockSize){
return (static_cast<size_t>(blockSize*SHARED_MEM_CONST));
}, 0, 0);
REQUIRE(ret == hipSuccess);
REQUIRE(minGridSize > 0);
REQUIRE(blockSize > 0);
REQUIRE(blockSize <= devProp.maxThreadsPerBlock);
}
/**
Negative tests hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags
- null min_grid_size
- null block_size
- null func
- Invalid flag
*/
TEST_CASE("Unit_hipOccupancyMaxPotBlkSizeVariableSMemWithFlags_NegTst") {
hipError_t ret;
int minGridSize = 0, blockSize = 0;
SECTION("null min_grid_size") {
ret = hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags(nullptr,
&blockSize, f1, blockSizeToDynamicSMemSize, 0, 0);
REQUIRE(ret == hipErrorInvalidValue);
}
SECTION("null block_size") {
ret = hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags(&minGridSize,
nullptr, f1, blockSizeToDynamicSMemSize, 0, 0);
REQUIRE(ret == hipErrorInvalidValue);
}
SECTION("null func") {
ret = hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags
<size_t(*)(int), void(*)(float*)>(&minGridSize, &blockSize, nullptr,
blockSizeToDynamicSMemSize, 0, 0);
REQUIRE(ret == hipErrorInvalidValue);
}
#if HT_NVIDIA
SECTION("invalid flag") {
ret = hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags(&minGridSize,
&blockSize, f1, blockSizeToDynamicSMemSize, 0, 0xffff);
REQUIRE(ret == hipErrorInvalidValue);
}
#endif
}
/**
Local function to launch kernel with gridsize and blocksize derived from
hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags.
*/
static void checkFunc(void(*kerFn)(int*, int*), int num,
int sharedMemBytes, int blockSize) {
int SIZE = num * sizeof(int);
int *inpArr_h, *outArr_h;
int *inpArr_d, *outArr_d;
// allocate host matrix
inpArr_h = reinterpret_cast<int*>(malloc(SIZE));
REQUIRE(inpArr_h != nullptr);
outArr_h = reinterpret_cast<int*>(malloc(SIZE));
REQUIRE(outArr_h != nullptr);
// initialize the input data
for (int i = 0; i < num; i++) {
inpArr_h[i] = i;
}
// allocate the memory on the device side
HIP_CHECK(hipMalloc(&inpArr_d, SIZE));
HIP_CHECK(hipMalloc(&outArr_d, SIZE));
// Memory transfer from host to device
HIP_CHECK(hipMemcpy(inpArr_d, inpArr_h, SIZE, hipMemcpyHostToDevice));
// Lauching kernel from host
dim3 gridsize = dim3(num / blockSize);
dim3 blocksize = dim3(blockSize);
hipLaunchKernelGGL(kerFn, gridsize, blocksize, sharedMemBytes, 0,
outArr_d, inpArr_d);
// Memory transfer from device to host
HIP_CHECK(hipMemcpy(outArr_h, outArr_d, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipDeviceSynchronize());
// verify the results
for (int i = 0; i < num; i++) {
REQUIRE(outArr_h[i] == inpArr_h[i]);
}
// free the resources on device side
HIP_CHECK(hipFree(inpArr_d));
HIP_CHECK(hipFree(outArr_d));
// free the resources on host side
free(inpArr_h);
free(outArr_h);
}
/**
Local function to return appropriate array size which consumes
memory less than the maximum allowed shared memory per block.
*/
static int getAppropriateDynShMemSize(int sharedMemPerBlock) {
int size = 1;
while (static_cast<int>(size*size*sizeof(int)) < sharedMemPerBlock) {
size = size * 2;
}
return (size/2);
}
// functor to return 0 dynamic shared memory
static size_t getZeroDynShMem(int blocksize) {
UNUSED(blocksize);
return 0;
}
// functor to return maximum possible dynamic shared memory.
static size_t getMaxDynShMem(int blocksize) {
UNUSED(blocksize);
return static_cast<size_t>(gArrSize*gArrSize*sizeof(int));
}
/**
Functional tests for hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags.
Scenario1:
Calculate the gridsize and blocksize that give theoretical maximum potential
occupancy for a kernel function that does not use dynamic shared memory.
Using the derived gridsize and blocksize launch the kernel and validate its
output.
Scenario2:
Calculate the gridsize and blocksize that give theoretical maximum potential
occupancy for a kernel function that uses dynamic shared memory. Ensure that
allocated dynamic shared memory is less than the maximum allowed by system.
Using the derived gridsize and blocksize launch the kernel and validate its
output.
*/
TEST_CASE("Unit_hipOccupancyMaxPotBlkSizeVariableSMemWithFlags_Functional") {
hipDeviceProp_t devProp;
HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
SECTION("Non Dynamic Shared Kernel") {
int arrSize;
int minGridSize = 0, blockSize = 0;
hipError_t ret;
// Get potential blocksize
ret = hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags(&minGridSize,
&blockSize, copyKer, getZeroDynShMem, 0, 0);
REQUIRE(ret == hipSuccess);
REQUIRE(minGridSize > 0);
REQUIRE(blockSize > 0);
REQUIRE(blockSize <= devProp.maxThreadsPerBlock);
arrSize = minGridSize*blockSize;
checkFunc(copyKer, arrSize, 0, blockSize);
}
SECTION("Dynamic Shared Kernel") {
int arrSize = getAppropriateDynShMemSize(devProp.sharedMemPerBlock);
gArrSize = arrSize;
int minGridSize = 0, blockSize = 0;
hipError_t ret;
// Get potential blocksize
ret = hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags(&minGridSize,
&blockSize, copyKerDyn, getMaxDynShMem, 0, 0);
REQUIRE(ret == hipSuccess);
REQUIRE(minGridSize > 0);
REQUIRE(blockSize > 0);
REQUIRE(blockSize <= devProp.maxThreadsPerBlock);
int totalThreads;
totalThreads = minGridSize*blockSize;
// allow launching kernel with occupancy derived blocksize and gridsize
// only if allocated dynamic memory is less than system limit.
if ((totalThreads*sizeof(int)) < devProp.sharedMemPerBlock) {
checkFunc(copyKerDyn, totalThreads, (totalThreads*sizeof(int)),
blockSize);
} else {
totalThreads = arrSize*arrSize;
// allow launching kernel only if blockSize is a multiple of
// totalThreads
if (((totalThreads % blockSize) == 0) &&
((totalThreads / blockSize) > 0)) {
checkFunc(copyKerDyn, totalThreads, (totalThreads*sizeof(int)),
blockSize);
}
}
}
}