diff --git a/projects/hip-tests/catch/unit/graph/hipGraphAddMemsetNode.cc b/projects/hip-tests/catch/unit/graph/hipGraphAddMemsetNode.cc index 0984e194c8..c89d1a6f5c 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphAddMemsetNode.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphAddMemsetNode.cc @@ -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 @@ -31,14 +31,27 @@ Negative Testcase Scenarios for api hipGraphAddMemsetNode : 10) Pass hipMemsetParams::dst as nullptr should return error code. 11) Pass hipMemsetParams::element size other than 1, 2, or 4 and check api should return error code. 12) Pass hipMemsetParams::height as zero and check api should return error code. +Functional Scenarios for api hipGraphAddMemsetNode : +1. 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 +2. 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.. +3. 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. +4. 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. +5. 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. +6. Allocate memory using hipMallocManaged. Initialize the allocated memory using hipGraphAddMemsetNode. + Copy the values in device memory to host using hipGraphAddMemcpyNode. Verify the results. */ #include - /** * Negative Test for API hipGraphAddMemsetNode */ - +#define SIZE 1024 +static char memSetVal = 'a'; TEST_CASE("Unit_hipGraphAddMemsetNode_Negative") { hipError_t ret; hipGraph_t graph; @@ -121,3 +134,416 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_Negative") { HIP_CHECK(hipFree(devData)); HIP_CHECK(hipGraphDestroy(graph)); } +/* + * 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") { + size_t width = SIZE * sizeof(char), numW{SIZE}, + numH{SIZE}, pitch_A; + char *A_d; + + hipGraph_t graph; + std::vector nodeDependencies; + // Host memory. + char* A_h = new char[numW * numH]; + for (size_t i = 0; i < numW; i++) { + for (size_t j = 0; j < numH; j++) { + *(A_h + i * numH + j) = ' '; + } + } + // 2D Memory allocation hipMallocPitch + HIP_CHECK(hipMallocPitch(reinterpret_cast(&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(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)); + nodeDependencies.push_back(memsetNode); + // Add MemCpy Node + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.srcPtr = make_hipPitchedPtr(A_d, pitch_A, numW, numH); + myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH); + myparms.extent = make_hipExtent(width, numH, 1); + myparms.kind = hipMemcpyDeviceToHost; + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(), + 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(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); + } + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + delete[] A_h; + HIP_CHECK(hipFree(A_d)); +} +/* + * 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") { + size_t width = SIZE * sizeof(char), numW{SIZE}, pitch_A; + char *A_d; + + // Initialize the host memory + std::vector A_h(numW, ' '); + + hipGraph_t graph; + std::vector nodeDependencies; + // 1D Memory allocation hipMallocPitch + HIP_CHECK(hipMallocPitch(reinterpret_cast(&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(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)); + nodeDependencies.push_back(memsetNode); + // Add MemCpy Node + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.srcPtr = make_hipPitchedPtr(A_d, pitch_A, numW, 1); + myparms.dstPtr = make_hipPitchedPtr(A_h.data(), width, numW, 1); + myparms.extent = make_hipExtent(width, 1, 1); + myparms.kind = hipMemcpyDeviceToHost; + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(), + 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(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verfication + for (size_t i = 0; i < numW; i++) { + REQUIRE(A_h[i] == memSetVal); + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipFree(A_d)); +} +/* + * 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") { + size_t width = SIZE * sizeof(char); + size_t numW = SIZE, numH = SIZE; + + // Host Memory + char* A_h = new char[numW * numH]; + for (size_t i = 0; i < numW; i++) { + for (size_t j = 0; j < numH; j++) { + *(A_h + i * numH + j) = ' '; + } + } + hipGraph_t graph; + std::vector nodeDependencies; + + hipPitchedPtr A_d; + hipExtent extent3D = make_hipExtent(width, numH, 1); + + // Allocate 3D memory. + HIPCHECK(hipMalloc3D(&A_d, extent3D)); + + // Create Graph + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memsetNode, memcpyNode; + + // Add MemSet Node + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = A_d.ptr; + memsetParams.value = memSetVal; + memsetParams.pitch = A_d.pitch; + memsetParams.elementSize = sizeof(char); + memsetParams.width = numW; + memsetParams.height = numH; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + nodeDependencies.push_back(memsetNode); + + // MemCpy params + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.srcPtr = A_d; + myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH); + myparms.extent = make_hipExtent(width, numH, 1); + myparms.kind = hipMemcpyDeviceToHost; + + // Add MemCpy Node + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(), + 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(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); + } + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + delete[] A_h; + HIP_CHECK(hipFree(A_d.ptr)); +} +/* + * 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") { + size_t width = SIZE * sizeof(char); + size_t numW = SIZE; + + // Initialize the host memory + std::vector A_h(numW, ' '); + + hipGraph_t graph; + std::vector nodeDependencies; + + hipPitchedPtr A_d; + hipExtent extent1D = make_hipExtent(width, 1, 1); + + // Allocate 3D memory. + HIPCHECK(hipMalloc3D(&A_d, extent1D)); + + // Create Graph + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memsetNode, memcpyNode; + + // Add MemSet Node + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = A_d.ptr; + memsetParams.value = memSetVal; + memsetParams.pitch = A_d.pitch; + memsetParams.elementSize = sizeof(char); + memsetParams.width = numW; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + nodeDependencies.push_back(memsetNode); + + // MemCpy params + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.srcPtr = A_d; + myparms.dstPtr = make_hipPitchedPtr(A_h.data(), width, numW, 1); + myparms.extent = make_hipExtent(width, 1, 1); + myparms.kind = hipMemcpyDeviceToHost; + + // Add MemCpy Node + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(), + 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(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verfication + for (size_t i = 0; i < numW; i++) { + REQUIRE(A_h[i] == memSetVal); + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)) + HIP_CHECK(hipFree(A_d.ptr)); +} +/* + * 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; + size_t NumW = SIZE; + size_t Nbytes1D = SIZE * sizeof(char); + + // Initialize the host memory + std::vector A_h(NumW, ' '); + + // Allocate memory to Device pointer + HIP_CHECK(hipMalloc(reinterpret_cast(&A_d), Nbytes1D)); + + // Create the graph + hipGraph_t graph; + std::vector nodeDependencies; + hipGraphNode_t memsetNode, memcpyNode; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Add Memset node + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(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)); + nodeDependencies.push_back(memsetNode); + // Add MemCpy Node + hipPitchedPtr devPitchedPtr{A_d, Nbytes1D, NumW, 0}; + hipPitchedPtr hostPitchedPtr{A_h.data(), Nbytes1D, NumW, 0}; + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.srcPtr = devPitchedPtr; + myparms.dstPtr = hostPitchedPtr; + myparms.extent = make_hipExtent(Nbytes1D, 1, 1); + myparms.kind = hipMemcpyDeviceToHost; + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(), + 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(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verfication + for (size_t i = 0; i < NumW; i++) { + REQUIRE(A_h[i] == memSetVal); + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipFree(A_d)); +} + +TEST_CASE("Unit_hipGraphAddMemsetNode_hipMallocManaged") { + int managed = 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."); + } + size_t Nbytes1D = SIZE * sizeof(char); + char *A_d; + // Initialize the host memory + std::vector A_h(SIZE, ' '); + // Device Memory + HIP_CHECK(hipMallocManaged(&A_d, SIZE * sizeof(char))); + // Create the graph + hipGraph_t graph; + std::vector nodeDependencies; + hipGraphNode_t memsetNode, memcpyNode; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Add Memset node + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(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)); + nodeDependencies.push_back(memsetNode); + + // Add MemCpy Node + hipPitchedPtr devPitchedPtr{A_d, Nbytes1D, SIZE, 1}; + hipPitchedPtr hostPitchedPtr{A_h.data(), Nbytes1D, SIZE, 1}; + + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.srcPtr = devPitchedPtr; + myparms.dstPtr = hostPitchedPtr; + myparms.extent = make_hipExtent(Nbytes1D, 1, 1); + myparms.kind = hipMemcpyDeviceToHost; + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(), + 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(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verfication + for (size_t i = 0; i < SIZE; i++) { + REQUIRE(A_h[i] == memSetVal); + } + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipFree(A_d)); +}