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

EXSWHTEC-281 - Implement Unit Tests for assert functions #210

Browse Source 
Change-Id: I6c05915c957d9b67951b3e97cc35cb1ca72a945f
This commit is contained in:
Mirza Halilcevic
2023-12-28 17:41:54 +00:00
committed by Rakesh Roy
parent cf5ebc5990
commit 044a59496c
31 changed files with 674 additions and 463 deletions
Download Patch File Download Diff File Expand all files Collapse all files
catch/unit/graph/hipGraphAddMemsetNode.cc
+40 -56
View File
@@ -22,7 +22,6 @@ THE SOFTWARE.
#include <functional>
#include <vector>
#include <hip_test_defgroups.hh>
#include <hip_test_common.hh>
#include <resource_guards.hh>
#include <utils.hh>
@@ -129,7 +128,7 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_Negative_Parameters") {
* Allocate a 2D array using hipMallocPitch. Initialize the allocated memory
* using hipGraphAddMemsetNode. Copy the values in device memory to host using
* hipGraphAddMemcpyNode. Verify the results.
*/
*/
TEST_CASE("Unit_hipGraphAddMemsetNode_hipMallocPitch_2D") {
CHECK_IMAGE_SUPPORT
@@ -147,22 +146,20 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_hipMallocPitch_2D") {
}
}
// 2D Memory allocation hipMallocPitch
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A, width,
numH));
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A, width, numH));
// Create Graph
HIP_CHECK(hipGraphCreate(&graph, 0));
hipGraphNode_t memsetNode, memcpyNode;
// Add MemSet Node
hipMemsetParams memsetParams{};
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void *>(A_d);
memsetParams.dst = reinterpret_cast<void*>(A_d);
memsetParams.value = memSetVal;
memsetParams.pitch = pitch_A;
memsetParams.elementSize = sizeof(char);
memsetParams.width = numW;
memsetParams.height = numH;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams));
nodeDependencies.push_back(memsetNode);
// Add MemCpy Node
hipMemcpy3DParms myparms{};
@@ -173,21 +170,20 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_hipMallocPitch_2D") {
myparms.extent = make_hipExtent(width, numH, 1);
myparms.kind = hipMemcpyDeviceToHost;
HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(),
nodeDependencies.size(), &myparms));
nodeDependencies.size(), &myparms));
nodeDependencies.clear();
// Create executable graph
hipStream_t streamForGraph;
hipGraphExec_t graphExec;
HIP_CHECK(hipStreamCreate(&streamForGraph));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr,
nullptr, 0));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Verfication
for (size_t i = 0; i < numW; i++) {
for (size_t j = 0; j < numH; j++) {
REQUIRE(*(A_h + i*numH + j) == memSetVal);
REQUIRE(*(A_h + i * numH + j) == memSetVal);
}
}
HIP_CHECK(hipGraphExecDestroy(graphExec));
@@ -200,12 +196,12 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_hipMallocPitch_2D") {
* Allocate a 1D array using hipMallocPitch. Initialize the allocated memory using
* hipGraphAddMemsetNode. Copy the values in device memory to host using
* hipGraphAddMemcpyNode. Verify the results.
*/
*/
TEST_CASE("Unit_hipGraphAddMemsetNode_hipMallocPitch_1D") {
CHECK_IMAGE_SUPPORT
size_t width = SIZE * sizeof(char), numW{SIZE}, pitch_A;
char *A_d;
char* A_d;
// Initialize the host memory
std::vector<char> A_h(numW, ' ');
@@ -213,22 +209,20 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_hipMallocPitch_1D") {
hipGraph_t graph;
std::vector<hipGraphNode_t> nodeDependencies;
// 1D Memory allocation hipMallocPitch
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A, width,
1));
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A, width, 1));
// Create Graph
HIP_CHECK(hipGraphCreate(&graph, 0));
hipGraphNode_t memsetNode, memcpyNode;
// Add MemSet Node
hipMemsetParams memsetParams{};
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void *>(A_d);
memsetParams.dst = reinterpret_cast<void*>(A_d);
memsetParams.value = memSetVal;
memsetParams.pitch = pitch_A;
memsetParams.elementSize = sizeof(char);
memsetParams.width = numW;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams));
nodeDependencies.push_back(memsetNode);
// Add MemCpy Node
hipMemcpy3DParms myparms{};
@@ -239,15 +233,14 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_hipMallocPitch_1D") {
myparms.extent = make_hipExtent(width, 1, 1);
myparms.kind = hipMemcpyDeviceToHost;
HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(),
nodeDependencies.size(), &myparms));
nodeDependencies.size(), &myparms));
nodeDependencies.clear();
// Create executable graph
hipStream_t streamForGraph;
hipGraphExec_t graphExec;
HIP_CHECK(hipStreamCreate(&streamForGraph));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr,
nullptr, 0));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
@@ -264,7 +257,7 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_hipMallocPitch_1D") {
* Allocate a 2D array using hipMalloc3D. Initialize the allocated memory using
* hipGraphAddMemsetNode. Copy the values in device memory to host using
* hipGraphAddMemcpyNode. Verify the results.
*/
*/
TEST_CASE("Unit_hipGraphAddMemsetNode_hipMalloc3D_2D") {
CHECK_IMAGE_SUPPORT
@@ -300,8 +293,7 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_hipMalloc3D_2D") {
memsetParams.elementSize = sizeof(char);
memsetParams.width = numW;
memsetParams.height = numH;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams));
nodeDependencies.push_back(memsetNode);
// MemCpy params
@@ -315,22 +307,21 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_hipMalloc3D_2D") {
// Add MemCpy Node
HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(),
nodeDependencies.size(), &myparms));
nodeDependencies.size(), &myparms));
nodeDependencies.clear();
// Create executable graph
hipStream_t streamForGraph;
hipGraphExec_t graphExec;
HIP_CHECK(hipStreamCreate(&streamForGraph));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr,
nullptr, 0));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Verfication
for (size_t i = 0; i < numW; i++) {
for (size_t j = 0; j < numH; j++) {
REQUIRE(*(A_h + i*numH + j) == memSetVal);
REQUIRE(*(A_h + i * numH + j) == memSetVal);
}
}
HIP_CHECK(hipGraphExecDestroy(graphExec));
@@ -343,7 +334,7 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_hipMalloc3D_2D") {
* Allocate a 1D array using hipMalloc3D. Initialize the allocated
* memory using hipGraphAddMemsetNode. Copy the values in device
* memory to host using hipGraphAddMemcpyNode. Verify the results.
*/
*/
TEST_CASE("Unit_hipGraphAddMemsetNode_hipMalloc3D_1D") {
CHECK_IMAGE_SUPPORT
@@ -375,8 +366,7 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_hipMalloc3D_1D") {
memsetParams.elementSize = sizeof(char);
memsetParams.width = numW;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams));
nodeDependencies.push_back(memsetNode);
// MemCpy params
@@ -390,21 +380,20 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_hipMalloc3D_1D") {
// Add MemCpy Node
HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(),
nodeDependencies.size(), &myparms));
nodeDependencies.size(), &myparms));
nodeDependencies.clear();
// Create executable graph
hipStream_t streamForGraph;
hipGraphExec_t graphExec;
HIP_CHECK(hipStreamCreate(&streamForGraph));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr,
nullptr, 0));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Verfication
for (size_t i = 0; i < numW; i++) {
REQUIRE(A_h[i] == memSetVal);
REQUIRE(A_h[i] == memSetVal);
}
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
@@ -415,9 +404,9 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_hipMalloc3D_1D") {
* Allocate a 1D array using hipMalloc. Initialize the allocated memory using
* hipGraphAddMemsetNode. Copy the values in device memory to host using
* hipGraphAddMemcpyNode. Verify the results.
*/
*/
TEST_CASE("Unit_hipGraphAddMemsetNode_hipMalloc_1D") {
char *A_d;
char* A_d;
size_t NumW = SIZE;
size_t Nbytes1D = SIZE * sizeof(char);
@@ -436,14 +425,13 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_hipMalloc_1D") {
// Add Memset node
hipMemsetParams memsetParams{};
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void *>(A_d);
memsetParams.dst = reinterpret_cast<void*>(A_d);
memsetParams.value = memSetVal;
memsetParams.pitch = Nbytes1D;
memsetParams.elementSize = sizeof(char);
memsetParams.width = NumW;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams));
nodeDependencies.push_back(memsetNode);
// Add MemCpy Node
hipPitchedPtr devPitchedPtr{A_d, Nbytes1D, NumW, 0};
@@ -456,20 +444,19 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_hipMalloc_1D") {
myparms.extent = make_hipExtent(Nbytes1D, 1, 1);
myparms.kind = hipMemcpyDeviceToHost;
HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(),
nodeDependencies.size(), &myparms));
nodeDependencies.size(), &myparms));
nodeDependencies.clear();
// Create executable graph
hipStream_t streamForGraph;
hipGraphExec_t graphExec;
HIP_CHECK(hipStreamCreate(&streamForGraph));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr,
nullptr, 0));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Verfication
for (size_t i = 0; i < NumW; i++) {
REQUIRE(A_h[i] == memSetVal);
REQUIRE(A_h[i] == memSetVal);
}
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
@@ -479,16 +466,15 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_hipMalloc_1D") {
TEST_CASE("Unit_hipGraphAddMemsetNode_hipMallocManaged") {
int managed = 0;
HIP_CHECK(hipDeviceGetAttribute(&managed,
hipDeviceAttributeManagedMemory, 0));
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, 0));
INFO("hipDeviceAttributeManagedMemory: " << managed);
if (managed != 1) {
WARN(
"GPU 0 doesn't support hipDeviceAttributeManagedMemory attribute"
"so defaulting to system memory.");
"GPU 0 doesn't support hipDeviceAttributeManagedMemory attribute"
"so defaulting to system memory.");
}
size_t Nbytes1D = SIZE * sizeof(char);
char *A_d;
char* A_d;
// Initialize the host memory
std::vector<char> A_h(SIZE, ' ');
// Device Memory
@@ -502,14 +488,13 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_hipMallocManaged") {
// Add Memset node
hipMemsetParams memsetParams{};
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void *>(A_d);
memsetParams.dst = reinterpret_cast<void*>(A_d);
memsetParams.value = memSetVal;
memsetParams.pitch = Nbytes1D;
memsetParams.elementSize = sizeof(char);
memsetParams.width = SIZE;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams));
nodeDependencies.push_back(memsetNode);
// Add MemCpy Node
@@ -524,21 +509,20 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_hipMallocManaged") {
myparms.extent = make_hipExtent(Nbytes1D, 1, 1);
myparms.kind = hipMemcpyDeviceToHost;
HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(),
nodeDependencies.size(), &myparms));
nodeDependencies.size(), &myparms));
nodeDependencies.clear();
// Create executable graph
hipStream_t streamForGraph;
hipGraphExec_t graphExec;
HIP_CHECK(hipStreamCreate(&streamForGraph));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr,
nullptr, 0));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Verfication
for (size_t i = 0; i < SIZE; i++) {
REQUIRE(A_h[i] == memSetVal);
REQUIRE(A_h[i] == memSetVal);
}
HIP_CHECK(hipGraphExecDestroy(graphExec));