diff --git a/catch/include/memcpy1d_tests_common.hh b/catch/include/memcpy1d_tests_common.hh index fe0d46788b..72c1500fc8 100644 --- a/catch/include/memcpy1d_tests_common.hh +++ b/catch/include/memcpy1d_tests_common.hh @@ -141,7 +141,7 @@ void MemcpyDeviceToDeviceShell(F memcpy_func, const hipStream_t kernel_stream = HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, src_device, dst_device)); if (!can_access_peer) { INFO("Peer access cannot be enabled between devices " << src_device << " " << dst_device); - REQUIRE(can_access_peer); + return; } HIP_CHECK(hipDeviceEnablePeerAccess(dst_device, 0)); } diff --git a/catch/unit/compiler/hipClassKernel.cc b/catch/unit/compiler/hipClassKernel.cc index b205c6bffb..71286b706a 100644 --- a/catch/unit/compiler/hipClassKernel.cc +++ b/catch/unit/compiler/hipClassKernel.cc @@ -87,6 +87,8 @@ TEST_CASE("Unit_hipClassKernel_Friend") { 0, 0, result_ecd); + HIP_CHECK(hipStreamSynchronize(nullptr)); + HIP_CHECK(hipFree(result_ecd)); } // check sizeof empty class is 1 diff --git a/catch/unit/event/hipEventSynchronize.cc b/catch/unit/event/hipEventSynchronize.cc index a347badb22..dc70323b15 100644 --- a/catch/unit/event/hipEventSynchronize.cc +++ b/catch/unit/event/hipEventSynchronize.cc @@ -110,8 +110,6 @@ TEST_CASE("Unit_hipEventSynchronize_NoEventRecord_Positive") { // Record the end_event HIP_CHECK(hipEventRecord(end_event, NULL)); - // End event has not been completed - HIP_CHECK_ERROR(hipEventQuery(end_event), hipErrorNotReady); // When hipEventSynchronized is called on event that has not been recorded, // the function returns immediately diff --git a/catch/unit/graph/CMakeLists.txt b/catch/unit/graph/CMakeLists.txt index 6332680877..beea08f1aa 100644 --- a/catch/unit/graph/CMakeLists.txt +++ b/catch/unit/graph/CMakeLists.txt @@ -83,6 +83,7 @@ set(TEST_SRC hipGraphDebugDotPrint.cc hipGraphCloneComplx.cc hipGraphNodeSetEnabled.cc + hipGraphNodeGetEnabled.cc hipGraphCreate.cc hipGraphDestroy.cc hipGraphExecDestroy.cc diff --git a/catch/unit/graph/hipGraphAddMemsetNode.cc b/catch/unit/graph/hipGraphAddMemsetNode.cc index 0984e194c8..c89d1a6f5c 100644 --- a/catch/unit/graph/hipGraphAddMemsetNode.cc +++ b/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)); +} diff --git a/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc b/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc index e83bfe9cea..4f32551228 100644 --- a/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc +++ b/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc @@ -282,6 +282,7 @@ void hipGraphExecMemcpyNodeSetParamsToSymbol_GlobalMem(bool useConstVar) { } HIP_CHECK(hipGraphLaunch(graphExec, 0)); + HIP_CHECK(hipStreamSynchronize(0)); // Validating the result for (int i = 0; i < SIZE; i++) { diff --git a/catch/unit/graph/hipGraphNodeGetEnabled.cc b/catch/unit/graph/hipGraphNodeGetEnabled.cc new file mode 100644 index 0000000000..9f9d8279a5 --- /dev/null +++ b/catch/unit/graph/hipGraphNodeGetEnabled.cc @@ -0,0 +1,257 @@ +/* +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 +#include +#include + +#define N 1024 + +/** + * Functional Test for API - hipGraphNodeGetEnabled + 1) Add MemCpy node to the graph and verify in graphExec it enabled status + 2) Add MemSet node to the graph and verify in graphExec it enabled status + 3) Add Kernel node to the graph and verify in graphExec it enabled status +- Can check other node type, as only above 3 types are mentioned in the doc + 4) Add HostNode node to the graph and verify in graphExec it enabled status + 5) Add emptyNode node to the graph and verify in graphExec it enabled status + 6) Add ChildNode node to the graph and verify in graphExec it enabled status + 7) Add EventWait node to the graph and verify in graphExec it enabled status + 8) Add EventRecord node to the graph and verify in graphExec it enabled status + */ + +static void callbackfunc(void *A_h) { + int *A = reinterpret_cast(A_h); + for (int i = 0; i < N; i++) { + A[i] = i; + } +} + +TEST_CASE("Unit_hipGraphNodeGetEnabled_Functional_Basic") { + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + hipGraph_t graph, childGraph; + hipGraphNode_t memcpy_A, memcpy_B, memsetNode, kNodeAdd; + hipGraphNode_t hostNode, emptyNode, childGraphNode, eventWait, eventRecord; + int *A_d, *B_d, *C_d, *A_h, *B_h, *C_h; + hipGraphExec_t graphExec; + size_t NElem{N}; + unsigned int isEnabled = 0; + hipError_t ret; + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + hipKernelNodeParams kNodeParams{}; + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.gridDim = dim3(blocks); + kNodeParams.blockDim = dim3(threadsPerBlock); + kNodeParams.sharedMemBytes = 0; + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kNodeAdd, graph, nullptr, 0, &kNodeParams)); + + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(B_d); + memsetParams.value = 7; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + + HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, graph, NULL, 0)); + + hipHostNodeParams hostParams = {0, 0}; + hostParams.fn = callbackfunc; + hostParams.userData = A_h; + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, nullptr, 0, &hostParams)); + + hipEvent_t event; + HIP_CHECK(hipEventCreate(&event)); + HIP_CHECK(hipGraphAddEventRecordNode(&eventRecord, graph, nullptr, + 0, event)); + HIP_CHECK(hipGraphAddEventWaitNode(&eventWait, graph, nullptr, 0, event)); + + HIP_CHECK(hipGraphCreate(&childGraph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, childGraph, nullptr, 0, + B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + // Adding child node to clonedGraph + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph, + nullptr, 0, childGraph)); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + + SECTION("Create graphExec with Kernel node and verify its enabled status") { + HIP_CHECK(hipGraphNodeGetEnabled(graphExec, kNodeAdd, &isEnabled)); + REQUIRE(1 == isEnabled); + } + SECTION("Create graphExec with MemCpy node and verify its enabled status") { + HIP_CHECK(hipGraphNodeGetEnabled(graphExec, memcpy_A, &isEnabled)); + REQUIRE(1 == isEnabled); + } + SECTION("Create graphExec with MemSet node and verify its enabled status") { + HIP_CHECK(hipGraphNodeGetEnabled(graphExec, memsetNode, &isEnabled)); + REQUIRE(1 == isEnabled); + } + SECTION("Create graphExec with hostNode and verify its enabled status") { + ret = hipGraphNodeGetEnabled(graphExec, hostNode, &isEnabled); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Create graphExec with emptyNode and verify its enabled status") { + ret = hipGraphNodeGetEnabled(graphExec, emptyNode, &isEnabled); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Create graphExec with ChildGraphNode & verify its enabled status") { + ret = hipGraphNodeGetEnabled(graphExec, childGraphNode, &isEnabled); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Create graphExec with EventWait and verify its enabled status") { + ret = hipGraphNodeGetEnabled(graphExec, eventWait, &isEnabled); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Create graphExec with EventRecord and verify its enabled status") { + ret = hipGraphNodeGetEnabled(graphExec, eventRecord, &isEnabled); + REQUIRE(hipErrorInvalidValue == ret); + } + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(childGraph)); + HIP_CHECK(hipGraphDestroy(graph)); + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); +} + +/** + Negative Test for API - hipGraphNodeGetEnabled + 1) Pass graphExec as nullptr + 2) Pass graphExec as uninitialized object + 3) Pass Node as nullptr + 4) Pass Node as uninitialized object + 5) Pass isEnabled as nullptr + + Negative Functional Test for API - hipGraphNodeGetEnabled + 6) Pass hNode from different graph and verify + 7) Create graphExec and then add one more new node to the graph verify + 8) Pass hNode a deleted node from same graph where exec was created + 9) Create graphExec and then delete the graph and verify a node + 10) Create graphExec and then delete the graphExec and verify a node + */ + +TEST_CASE("Unit_hipGraphNodeGetEnabled_Negative_Functional") { + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + hipGraph_t graph, graph2; + hipGraphNode_t memcpy_A, memcpy_B, memcpy_A2, memcpy_C, kNodeAdd; + hipKernelNodeParams kNodeParams{}; + int *A_d, *B_d, *C_d, *A_h, *B_h, *C_h; + hipGraphExec_t graphExec, graphExec2; + size_t NElem{N}; + unsigned int isEnabled; + hipError_t ret; + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph, nullptr, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.gridDim = dim3(blocks); + kNodeParams.blockDim = dim3(threadsPerBlock); + kNodeParams.sharedMemBytes = 0; + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kNodeAdd, graph, nullptr, 0, &kNodeParams)); + + HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_A, &kNodeAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_B, &kNodeAdd, 1)); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + HIP_CHECK(hipGraphInstantiate(&graphExec2, graph, NULL, NULL, 0)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddDependencies(graph, &kNodeAdd, &memcpy_C, 1)); + + HIP_CHECK(hipGraphCreate(&graph2, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A2, graph2, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + + SECTION("Pass hGraphExec as nullptr") { + ret = hipGraphNodeGetEnabled(nullptr, memcpy_B, &isEnabled); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass hGraphExec as uninitialized object") { + hipGraphExec_t graphExec_uninit{}; + ret = hipGraphNodeGetEnabled(graphExec_uninit, memcpy_B, &isEnabled); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass Node as nullptr") { + ret = hipGraphNodeGetEnabled(graphExec, nullptr, &isEnabled); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass Node as uninitialized object") { + hipGraphNode_t node_uninit{}; + ret = hipGraphNodeGetEnabled(graphExec, node_uninit, &isEnabled); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass isEnabled as nullptr") { + ret = hipGraphNodeGetEnabled(graphExec, memcpy_B, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass hNode from different graph and verify") { + ret = hipGraphNodeGetEnabled(graphExec, memcpy_A2, &isEnabled); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Create graphExec and add one more new node to the graph & verify") { + ret = hipGraphNodeGetEnabled(graphExec, memcpy_C, &isEnabled); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass hNode a deleted node from same graph where exec was created") { + HIP_CHECK(hipGraphDestroyNode(memcpy_A)); + ret = hipGraphNodeGetEnabled(graphExec, memcpy_A, &isEnabled); + REQUIRE(hipErrorInvalidValue == ret); + } + HIP_CHECK(hipGraphExecDestroy(graphExec2)); + SECTION("Create graphExec and then delete the graphExec and verify a node") { + ret = hipGraphNodeGetEnabled(graphExec2, memcpy_B, &isEnabled); + REQUIRE(hipErrorInvalidValue == ret); + } + HIP_CHECK(hipGraphDestroy(graph)); + SECTION("Create graphExec and then delete the graph and verify a node") { + ret = hipGraphNodeGetEnabled(graphExec, memcpy_B, &isEnabled); + REQUIRE(hipErrorInvalidValue == ret); + } + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph2)); +} diff --git a/catch/unit/graph/hipGraphUpload.cc b/catch/unit/graph/hipGraphUpload.cc index d17fb27236..53b22e7e6f 100644 --- a/catch/unit/graph/hipGraphUpload.cc +++ b/catch/unit/graph/hipGraphUpload.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 @@ -161,8 +161,17 @@ TEST_CASE("Unit_hipGraphUpload_Functional_multidevice_test") { SECTION("Pass a common stream for all device") { hipStream_t stream; HIP_CHECK(hipStreamCreate(&stream)); - + int currDevice = -1; + HIP_CHECK(hipGetDevice(&currDevice)); for (int i = 0; i < numDevices; i++) { + if (i != currDevice) { + int can_access_peer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, currDevice, i)); + if (!can_access_peer) { + INFO("Peer access cannot be enabled between devices " << currDevice << " " << i); + continue; + } + } HIP_CHECK(hipSetDevice(i)); hipGraphUploadFunctional_with_hipStreamBeginCapture(stream); hipGraphUploadFunctional_with_stream(stream); @@ -200,6 +209,53 @@ TEST_CASE("Unit_hipGraphUpload_Functional_multidevice_test") { } } +/** + * Functional Test for API - hipGraphUpload + - Make graph by using hipStreamBeginCapture with a low priority stream. + Upload the graph into high priority stream and execute the graph and verify. + */ + +TEST_CASE("Unit_hipGraphUpload_Functional_With_Priority_Stream") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph; + hipGraphExec_t graphExec; + + int *A_d, *B_d, *C_d, *A_h, *B_h, *C_h; + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + + hipStream_t stream1, stream2; + int minPriority = 0, maxPriority = 0; + HIP_CHECK(hipDeviceGetStreamPriorityRange(&minPriority, &maxPriority)); + HIP_CHECK(hipStreamCreateWithPriority(&stream1, hipStreamDefault, + minPriority)); + HIP_CHECK(hipStreamCreateWithPriority(&stream2, hipStreamDefault, + maxPriority)); + + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream1)); + HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream1)); + HipTest::vectorADD<<<1, 1, 0, stream1>>>(A_d, B_d, C_d, N); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream1)); + HIP_CHECK(hipStreamEndCapture(stream1, &graph)); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + HIP_CHECK(hipGraphUpload(graphExec, stream2)); + + HIP_CHECK(hipGraphLaunch(graphExec, stream2)); + HIP_CHECK(hipStreamSynchronize(stream2)); + + // Verify graph execution result + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream1)); + HIP_CHECK(hipStreamDestroy(stream2)); +} + /** * Negative Test for API - hipGraphUpload Argument Check 1) Pass graphExec node as nullptr. @@ -233,4 +289,3 @@ TEST_CASE("Unit_hipGraphUpload_Negative_Argument_Check") { } HIP_CHECK(hipStreamDestroy(stream)); } - diff --git a/catch/unit/graph/hipStreamEndCapture.cc b/catch/unit/graph/hipStreamEndCapture.cc index dd6db743a7..740764d4b2 100644 --- a/catch/unit/graph/hipStreamEndCapture.cc +++ b/catch/unit/graph/hipStreamEndCapture.cc @@ -30,6 +30,23 @@ Negative Testcase Scenarios : 7) Begin capture on a thread with mode other than hipStreamCaptureModeRelaxed and try to end capture from different thread. Expect to return hipErrorStreamCaptureWrongThread. +8) Start stream capture on stream1 using mode hipStreamCaptureModeRelaxed. + In stream1 queue a memcpy operation, queue a kernel square of a number operation. + Launch a thread. In the thread, queue a memcpy operation. End the capture on + stream1 and return the captured graph. Wait for the thread in main function. + Create an executable graph and launch the graph on input data and validate the + output. +9) Create 2 streams s1 and s2. Begin stream capture in s1, spawn a + captured fork stream on s2. Queue some operations + (like increment kernel) on both s1 and s2. End the stream capture + on s2 and verify the error returned by the End capture. +10)Create 2 streams s1 and s2. Begin stream capture in s1 and spawn a captured + fork stream s2. In main thread, queue a memcpy operation on s1. + Launch a thread, queue a memcpy operation on s2. Perform hipEventRecord on + s2 and wait Event on S1. Wait for the thread to complete. Queue operations + kernel addition(Cd = Ad + Bd) operation and memcpy(Ch <- Cd) in s1. End the + stream capture in s1. Create an executable graph and launch the graph on input + data and validate the output. */ #include @@ -123,14 +140,37 @@ static void thread_func(hipStream_t stream, hipGraph_t graph) { HIP_ASSERT(hipErrorStreamCaptureWrongThread == hipStreamEndCapture(stream, &graph)); } - -TEST_CASE("Unit_hipStreamEndCapture_Thread_Negative") { +static void StreamEndCaptureThreadNegative(float* A_d, float* A_h, + float* C_d, float* C_h, hipStreamCaptureMode mode) { hipStream_t stream{nullptr}; hipGraph_t graph{nullptr}; constexpr unsigned blocks = 512; constexpr unsigned threadsPerBlock = 256; constexpr size_t N = 100000; size_t Nbytes = N * sizeof(float); + + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamBeginCapture(stream, mode)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); + + HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream)); + 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)); + + std::thread t(thread_func, stream, graph); + t.join(); + +#if HT_AMD + HIP_CHECK(hipStreamEndCapture(stream, &graph)); +#endif + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipGraphDestroy(graph)); +} +TEST_CASE("Unit_hipStreamEndCapture_Thread_Negative") { + constexpr size_t N = 100000; + size_t Nbytes = N * sizeof(float); float *A_d, *C_d; float *A_h, *C_h; @@ -149,28 +189,249 @@ TEST_CASE("Unit_hipStreamEndCapture_Thread_Negative") { REQUIRE(A_d != nullptr); REQUIRE(C_d != nullptr); - HIP_CHECK(hipStreamCreate(&stream)); - HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); - HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); - - HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream)); - 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)); - - std::thread t(thread_func, stream, graph); - t.join(); - -#if HT_AMD - HIP_CHECK(hipStreamEndCapture(stream, &graph)); -#endif - + SECTION("Capture Mode:hipStreamCaptureModeGlobal") { + StreamEndCaptureThreadNegative(A_d, A_h, C_d, C_h, + hipStreamCaptureModeGlobal); + } + SECTION("Capture Mode:hipStreamCaptureModeThreadLocal") { + StreamEndCaptureThreadNegative(A_d, A_h, C_d, C_h, + hipStreamCaptureModeThreadLocal); + } free(A_h); free(C_h); HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(C_d)); +} +// Thread function +static void thread_func1(hipStream_t stream, hipGraph_t *graph, + size_t Nbytes, float* A_d, float* B_h) { + HIP_CHECK(hipMemcpyAsync(B_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamEndCapture(stream, graph)); +} +/* + * Start stream capture on stream1 using mode hipStreamCaptureModeRelaxed. + * In stream1 queue a memcpy operation, queue a kernel square of a number operation. + * Launch a thread. In the thread, queue a memcpy operation. End the capture on + * stream1 and return the captured graph. Wait for the thread in main function. + * Create an executable graph and launch the graph on input data and validate the output. + * */ +TEST_CASE("Unit_hipStreamEndCapture_mode_hipStreamCaptureModeRelaxed") { + hipStream_t stream{nullptr}, streamForGraph{nullptr}; + hipGraph_t graph{nullptr}; + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + constexpr size_t N = 10; + size_t Nbytes = N * sizeof(float); + // Device Pointers + float *A_d; + // Host Pointers + float *A_h, *B_h, *C_h; + + // Memory allocation to Host pointers + A_h = reinterpret_cast(malloc(Nbytes)); + B_h = reinterpret_cast(malloc(Nbytes)); + C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A_h != nullptr); + REQUIRE(B_h != nullptr); + REQUIRE(C_h != nullptr); + + // Initialize the Host data + for (size_t i = 0; i < N; i++) { + A_h[i] = 1.0f + i; + C_h[i] = A_h[i]; + } + // Memory allocation to Device pointers + HIP_CHECK(hipMalloc(reinterpret_cast(&A_d), Nbytes)); + REQUIRE(A_d != nullptr); + + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeRelaxed)); + // Copy data from Host to Device + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); + + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, stream, A_d, A_d, N); + // Thread Launch + std::thread t(thread_func1, stream, &graph, Nbytes, A_d, B_h); + t.join(); + + // Launch the graph + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Output verification + for (size_t i = 0; i < N; i++) { + C_h[i] = C_h[i] * C_h[i]; + REQUIRE(B_h[i] == C_h[i]); + } + + free(A_h); + free(B_h); + free(C_h); + HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); } +static __global__ void increment(int* A_d) { + atomicAdd(A_d, 1); +} +/* + * Create 2 streams s1 and s2. Begin stream capture in s1, spawn a + * captured fork stream on s2. Queue some operations + * (like increment kernel) on both s1 and s2. End the stream capture + * on s2 and verify the error returned by the End capture. +*/ +TEST_CASE("Unit_hipStreamEndCapture_chkError_on_wrongStream") { + int *A_d{nullptr}, *A_h{nullptr}; + hipStream_t stream1{nullptr}, stream2{nullptr}; + hipEvent_t forkStreamEvent{nullptr}; + hipGraph_t graph{nullptr}; + hipError_t err; + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + size_t Nbytes = sizeof(int); + + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + HIP_CHECK(hipEventCreate(&forkStreamEvent)); + + A_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A_h != nullptr); + // Initialize the Host data + *A_h = 0; + HIP_CHECK(hipMalloc(reinterpret_cast(&A_d), Nbytes)); + REQUIRE(A_d != nullptr); + + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(forkStreamEvent, stream1)); + HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0)); + + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, + hipMemcpyHostToDevice, stream1)); + + hipLaunchKernelGGL(increment, dim3(blocks), + dim3(threadsPerBlock), 0, stream1, A_d); + hipLaunchKernelGGL(increment, dim3(blocks), + dim3(threadsPerBlock), 0, stream2, A_d); + + err = hipStreamEndCapture(stream2, &graph); + REQUIRE(err == hipErrorStreamCaptureUnmatched); + + HIP_CHECK(hipStreamDestroy(stream1)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipEventDestroy(forkStreamEvent)); + free(A_h); + HIP_CHECK(hipFree(A_d)); +} +static void thread_func4(hipStream_t stream1, hipStream_t stream2, + hipEvent_t event, size_t Nbytes, int* B_d, int* B_h) { + HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream2)); + HIP_CHECK(hipEventRecord(event, stream2)); + HIP_CHECK(hipStreamWaitEvent(stream1, event, 0)); +} +/* + * Create 2 streams s1 and s2. Begin stream capture in s1 and spawn a captured + * fork stream s2. In main thread, queue a memcpy operation on s1. + * Launch a thread, queue a memcpy operation on s2. Perform hipEventRecord on + * s2 and wait Event on S1. Wait for the thread to complete. Queue operations + * kernel addition(Cd = Ad + Bd) operation and memcpy(Ch <- Cd) in s1. End the + * stream capture in s1. Create an executable graph and launch the graph on input + * data and validate the output. + * */ +TEST_CASE("Unit_hipStreamEndCapture_streamMerge_in_thread") { + // Device Pointers + int *A_d, *B_d, *C_d; + // Host Pointers + int *A_h, *B_h, *C_h, *D_h; + hipStream_t stream1{nullptr}, stream2{nullptr}, streamForGraph{nullptr}; + hipEvent_t forkStreamEvent{nullptr}, event{nullptr}; + hipGraph_t graph{nullptr}; + + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + constexpr size_t N = 5; + size_t Nbytes = N * sizeof(int); + + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipEventCreate(&forkStreamEvent)); + HIP_CHECK(hipEventCreate(&event)); + // Memory allocation to Host Pointers + A_h = reinterpret_cast(malloc(Nbytes)); + B_h = reinterpret_cast(malloc(Nbytes)); + C_h = reinterpret_cast(malloc(Nbytes)); + D_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A_h != nullptr); + REQUIRE(B_h != nullptr); + REQUIRE(C_h != nullptr); + REQUIRE(D_h != nullptr); + // Initialize the Host data + for (size_t i = 0; i < N; i++) { + A_h[i] = 1 + i; + B_h[i] = 2 + i; + C_h[i] = 0; + D_h[i] = 0; + } + // Memory allocation to Device Pointers + HIP_CHECK(hipMalloc(reinterpret_cast(&A_d), Nbytes)); + HIP_CHECK(hipMalloc(reinterpret_cast(&B_d), Nbytes)); + HIP_CHECK(hipMalloc(reinterpret_cast(&C_d), Nbytes)); + REQUIRE(A_d != nullptr); + REQUIRE(B_d != nullptr); + REQUIRE(C_d != nullptr); + + // Begin Capture + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + + HIP_CHECK(hipEventRecord(forkStreamEvent, stream1)); + HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0)); + + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, + hipMemcpyHostToDevice, stream1)); + // Thread Launch + std::thread t(thread_func4, stream1, stream2, event, Nbytes, B_d, B_h); + t.join(); + // Launch kernal + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), + dim3(threadsPerBlock), 0, stream1, A_d, + B_d, C_d, N); + + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, + hipMemcpyDeviceToHost, stream1)); + HIP_CHECK(hipStreamEndCapture(stream1, &graph)); + + // 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] + B_h[i]; + REQUIRE(C_h[i] == D_h[i]); + } + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream1)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipEventDestroy(forkStreamEvent)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + + // Release the memory + free(A_h); + free(B_h); + free(C_h); + free(D_h); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(B_d)); + HIP_CHECK(hipFree(C_d)); +} diff --git a/catch/unit/graph/hipUserObjectCreate.cc b/catch/unit/graph/hipUserObjectCreate.cc index d2765d7978..0d67e168b0 100644 --- a/catch/unit/graph/hipUserObjectCreate.cc +++ b/catch/unit/graph/hipUserObjectCreate.cc @@ -427,43 +427,37 @@ TEST_CASE("Unit_hipGraphRetainUserObject_Functional_2") { unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); dependencies.push_back(memcpyNode); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); dependencies.push_back(memcpyNode); - void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; - kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); kNodeParams.gridDim = dim3(blocks); kNodeParams.blockDim = dim3(threadsPerBlock); kNodeParams.sharedMemBytes = 0; kNodeParams.kernelParams = reinterpret_cast(kernelArgs); kNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, dependencies.data(), - dependencies.size(), &kNodeParams)); + HIP_CHECK( + hipGraphAddKernelNode(&kNode, graph, dependencies.data(), dependencies.size(), &kNodeParams)); dependencies.clear(); dependencies.push_back(kNode); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(), - dependencies.size(), C_h, C_d, - Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(), dependencies.size(), + C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); - int refCount = 2; - int refCountRetain = 3; - - float *object = new float(); + float* object = new float(); REQUIRE(object != nullptr); hipUserObject_t hObject; - HIP_CHECK(hipUserObjectCreate(&hObject, object, - destroyFloatObj, - refCount, hipUserObjectNoDestructorSync)); + HIP_CHECK( + hipUserObjectCreate(&hObject, object, destroyFloatObj, 1, hipUserObjectNoDestructorSync)); REQUIRE(hObject != nullptr); - HIP_CHECK(hipUserObjectRetain(hObject, refCountRetain)); - HIP_CHECK(hipGraphRetainUserObject(graph, hObject, refCountRetain, - hipGraphUserObjectMove)); + HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 1, + hipGraphUserObjectMove)); // Pass ownership to hipGraph // Instantiate and launch the graph HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); @@ -473,9 +467,6 @@ TEST_CASE("Unit_hipGraphRetainUserObject_Functional_2") { // Verify result HipTest::checkVectorADD(A_h, B_h, C_h, N); - HIP_CHECK(hipUserObjectRelease(hObject, refCount + refCountRetain)); - HIP_CHECK(hipGraphReleaseUserObject(graph, hObject, refCountRetain)); - HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(graph)); diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index 2fa38dd802..6c528fcc16 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -7,213 +7,127 @@ # 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 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. +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +# SOFTWARE. set(COMMON_SHARED_SRC DriverContext.cc) # Common Tests - Test independent of all platforms +set(TEST_SRC + memset.cc + malloc.cc + hipMemcpy2DToArray.cc + hipMemcpy2DToArray_old.cc + hipMemcpy2DToArrayAsync.cc + hipMemcpy2DToArrayAsync_old.cc + hipMemcpy3D.cc + hipMemcpy3DAsync.cc + hipMemcpyParam2D.cc + hipMemcpyParam2DAsync.cc + hipMemcpy2D.cc + hipMemcpy2DAsync.cc + hipMemcpy2DFromArray.cc + hipMemcpy2DFromArray_old.cc + hipMemcpy2DFromArrayAsync.cc + hipMemcpy2DFromArrayAsync_old.cc + hipMemcpyAtoH.cc + hipMemcpyAtoH_old.cc + hipMemcpyHtoA.cc + hipMemcpyHtoA_old.cc + hipMemcpyAllApiNegative.cc + hipMemcpy_MultiThread.cc + hipHostRegister.cc + hipHostUnregister.cc + hipHostGetFlags.cc + hipHostGetDevicePointer.cc + hipMallocManaged_MultiScenario.cc + hipMemsetNegative.cc + hipMemset.cc + hipMemsetAsyncMultiThread.cc + hipMemset3D.cc + hipMemset2D.cc + hipHostMallocTests.cc + hipMemset3DFunctional.cc + hipMemset3DRegressMultiThread.cc + hipMallocManagedFlagsTst.cc + hipMemPrefetchAsyncExtTsts.cc + hipMemAdviseMmap.cc + hipMallocManaged.cc + hipMemRangeGetAttribute.cc + hipMemRangeGetAttribute_old.cc + hipMemcpyFromSymbol.cc + hipPtrGetAttribute.cc + hipMemPoolApi.cc + hipMemcpyPeer.cc + hipMemcpyPeer_old.cc + hipMemcpyPeerAsync.cc + hipMemcpyPeerAsync_old.cc + hipMemcpyWithStream_old.cc + hipMemcpyWithStream.cc + hipMemcpyWithStreamMultiThread.cc + hipMemsetAsyncAndKernel.cc + hipMemset2DAsyncMultiThreadAndKernel.cc + hipMallocConcurrency.cc + hipMemcpyDtoD.cc + hipMemcpyDtoDAsync.cc + hipHostMalloc.cc + hipMemcpy_old.cc + hipMemcpy_derivatives.cc + hipMemcpyAsync.cc + hipMemsetFunctional.cc + hipMalloc.cc + hipMallocPitch.cc + hipMallocArray.cc + hipMalloc3D.cc + hipMalloc3DArray.cc + hipArrayCreate.cc + hipArray3DCreate.cc + hipDrvMemcpy3D.cc + hipDrvMemcpy3DAsync.cc + hipPointerGetAttribute.cc + hipDrvPtrGetAttributes.cc + hipMemPrefetchAsync.cc + hipMemGetInfo.cc + hipFree.cc + hipMemcpySync.cc + hipMemsetSync.cc + hipMemsetAsync.cc + hipMemAdvise_old.cc + hipMemAdvise.cc + hipMemRangeGetAttributes.cc + hipStreamAttachMemAsync.cc + hipMemRangeGetAttributes_old.cc + hipMemGetAddressRange.cc) + if(HIP_PLATFORM MATCHES "amd") -set(TEST_SRC - memset.cc - malloc.cc - hipMemcpy2DToArray.cc - hipMemcpy2DToArray_old.cc - hipMemcpy2DToArrayAsync.cc - hipMemcpy2DToArrayAsync_old.cc - hipMemcpy3D.cc - hipMemcpy3DAsync.cc - hipMemcpyParam2D.cc - hipMemcpyParam2DAsync.cc - hipMemcpy2D.cc - hipMemcpy2DAsync.cc - hipMemcpy2DFromArray.cc - hipMemcpy2DFromArray_old.cc - hipMemcpy2DFromArrayAsync.cc - hipMemcpy2DFromArrayAsync_old.cc - hipMemcpyAtoH.cc - hipMemcpyAtoH_old.cc - hipMemcpyHtoA.cc - hipMemcpyHtoA_old.cc - hipMemcpyAllApiNegative.cc - hipMemcpy_MultiThread.cc - hipHostRegister.cc - hipHostUnregister.cc - hipMemPtrGetInfo.cc - hipPointerGetAttributes.cc - hipHostGetFlags.cc - hipHostGetDevicePointer.cc - hipMallocManaged_MultiScenario.cc - hipMemsetNegative.cc - hipMemset.cc - hipMemsetAsyncMultiThread.cc - hipMemset3D.cc - hipMemset2D.cc - hipHostMallocTests.cc - hipMemset3DFunctional.cc - hipMemset3DRegressMultiThread.cc - hipMallocManagedFlagsTst.cc - hipMemPrefetchAsyncExtTsts.cc - hipMemAdviseMmap.cc - hipMemCoherencyTst.cc - hipMallocManaged.cc - hipMemRangeGetAttribute.cc - hipMemRangeGetAttribute_old.cc - hipMemcpyFromSymbol.cc - hipPtrGetAttribute.cc - hipMemPoolApi.cc - hipMemcpyPeer.cc - hipMemcpyPeer_old.cc - hipMemcpyPeerAsync.cc - hipMemcpyPeerAsync_old.cc - hipMemcpyWithStream_old.cc - hipMemcpyWithStream.cc - hipMemcpyWithStreamMultiThread.cc - hipMemsetAsyncAndKernel.cc - hipMemset2DAsyncMultiThreadAndKernel.cc - hipMallocConcurrency.cc - hipMemcpyDtoD.cc - hipMemcpyDtoDAsync.cc - hipHostMalloc.cc - hipMemcpy_old.cc - hipMemcpy_derivatives.cc - hipMemcpyAsync.cc - hipMemsetFunctional.cc - hipMalloc.cc - hipExtMallocWithFlags.cc - hipMallocPitch.cc - hipMallocArray.cc - hipMalloc3D.cc - hipMalloc3DArray.cc - hipArrayCreate.cc - hipArray3DCreate.cc - hipDrvMemcpy3D.cc - hipDrvMemcpy3DAsync.cc - hipPointerGetAttribute.cc - hipDrvPtrGetAttributes.cc - hipMallocMngdMultiThread.cc - hipMemPrefetchAsync.cc - hipArray.cc - hipMemVmm.cc - hipMemGetInfo.cc - hipFree.cc - hipMemcpySync.cc - hipMemsetSync.cc - hipMemsetAsync.cc - hipMemAdvise_old.cc - hipMemAdvise.cc - hipMemRangeGetAttributes.cc - hipStreamAttachMemAsync.cc - hipMemRangeGetAttributes_old.cc - hipMemGetAddressRange.cc -) + set(TEST_SRC + ${TEST_SRC} + hipMemPtrGetInfo.cc + hipPointerGetAttributes.cc + hipMemCoherencyTst.cc + hipExtMallocWithFlags.cc + hipMallocMngdMultiThread.cc + hipArray.cc + hipMemVmm.cc) else() -set(TEST_SRC - memset.cc - malloc.cc - hipMemcpy2DToArray.cc - hipMemcpy2DToArray_old.cc - hipMemcpy2DToArrayAsync.cc - hipMemcpy2DToArrayAsync_old.cc - hipMemcpy3D.cc - hipMemcpy3DAsync.cc - hipMemcpyParam2D.cc - hipMemcpyParam2DAsync.cc - hipMemcpy2D.cc - hipMemcpy2DAsync.cc - hipMemcpy2DFromArray.cc - hipMemcpy2DFromArray_old.cc - hipMemcpy2DFromArrayAsync.cc - hipMemcpy2DFromArrayAsync_old.cc - hipMemcpyAtoH.cc - hipMemcpyAtoH_old.cc - hipMemcpyHtoA.cc - hipMemcpyHtoA_old.cc - hipMemcpyAllApiNegative.cc - hipMemcpy_MultiThread.cc - hipHostRegister.cc - hipHostUnregister.cc - hipHostGetFlags.cc - hipHostGetDevicePointer.cc - hipMallocManaged_MultiScenario.cc - hipMemsetNegative.cc - hipMemset.cc - hipMemsetAsyncMultiThread.cc - hipMemset3D.cc - hipMemset2D.cc - hipHostMallocTests.cc - hipMemset3DFunctional.cc - hipMemset3DRegressMultiThread.cc - hipMallocManagedFlagsTst.cc - hipMemPrefetchAsyncExtTsts.cc - hipMemAdviseMmap.cc - hipMallocManaged.cc - hipMemRangeGetAttribute.cc - hipMemRangeGetAttribute_old.cc - hipMemcpyFromSymbol.cc - hipPtrGetAttribute.cc - hipMemPoolApi.cc - hipMemcpyPeer.cc - hipMemcpyPeer_old.cc - hipMemcpyPeerAsync.cc - hipMemcpyPeerAsync_old.cc - hipMemcpyWithStream_old.cc - hipMemcpyWithStream.cc - hipMemcpyWithStreamMultiThread.cc - hipMemsetAsyncAndKernel.cc - hipMemset2DAsyncMultiThreadAndKernel.cc - hipMallocConcurrency.cc - hipMemcpyDtoD.cc - hipMemcpyDtoDAsync.cc - hipHostMalloc.cc - hipMemcpy_old.cc - hipMemcpy_derivatives.cc - hipMemcpyAsync.cc - hipMemsetFunctional.cc - hipMalloc.cc - hipMallocPitch.cc - hipMallocArray.cc - hipMalloc3D.cc - hipMalloc3DArray.cc - hipArrayCreate.cc - hipArray3DCreate.cc - hipDrvMemcpy3D.cc - hipDrvMemcpy3DAsync.cc - hipPointerGetAttribute.cc - hipDrvPtrGetAttributes.cc - hipMemPrefetchAsync.cc - hipMemGetInfo.cc - hipFree.cc - hipMemcpySync.cc - hipMemsetSync.cc - hipMemsetAsync.cc - hipMemAdvise_old.cc - hipMemAdvise.cc - hipMemRangeGetAttributes.cc - hipMemRangeGetAttributes_old.cc - hipGetSymbolSizeAddress.cc - hipStreamAttachMemAsync.cc - hipMemGetAddressRange.cc -) + set(TEST_SRC ${TEST_SRC} hipGetSymbolSizeAddress.cc) endif() -# skipped due to os related code in tests -# need to work on them when all the tests are enabled +# skipped due to os related code in tests need to work on them when all the +# tests are enabled if(UNIX) - set(TEST_SRC ${TEST_SRC} - hipHmmOvrSubscriptionTst.cc - hipMemoryAllocateCoherent.cc) + set(TEST_SRC ${TEST_SRC} hipHmmOvrSubscriptionTst.cc + hipMemoryAllocateCoherent.cc) endif() hip_add_exe_to_target(NAME MemoryTest - TEST_SRC ${TEST_SRC} - TEST_TARGET_NAME build_tests - COMMON_SHARED_SRC ${COMMON_SHARED_SRC}) + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC}) diff --git a/catch/unit/memory/hipPointerGetAttributes.cc b/catch/unit/memory/hipPointerGetAttributes.cc index ceefa1f9ba..42ab07e289 100644 --- a/catch/unit/memory/hipPointerGetAttributes.cc +++ b/catch/unit/memory/hipPointerGetAttributes.cc @@ -122,7 +122,8 @@ void checkPointer(const SuperPointerAttribute& ref, int major, int minor, void* // we do this in the testMultiThreaded_1 test. void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) { Nbytes = N * sizeof(char); - printf("clusterAllocs numAllocs=%d size=%lu..%lu\n", numAllocs, minSize, maxSize); + printf("clusterAllocs numAllocs=%d size=%lu..%lu\n", + numAllocs, static_cast(minSize), static_cast(maxSize)); const int Max_Devices = 256; std::vector reference(numAllocs); @@ -174,7 +175,7 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) { " device#%d: hipMemGetInfo: " "free=%zu (%4.2fMB) totalDevice=%lu (%4.2fMB) total=%zu " "(%4.2fMB)\n", - i, free, (free / 1024.0 / 1024.0), totalDeviceAllocated[i], + i, free, (free / 1024.0 / 1024.0), static_cast(totalDeviceAllocated[i]), (totalDeviceAllocated[i]) / 1024.0 / 1024.0, total, (total / 1024.0 / 1024.0)); REQUIRE(free + totalDeviceAllocated[i] <= total); } @@ -221,7 +222,7 @@ TEST_CASE("Unit_hipPointerGetAttributes_Basic") { size_t free, total; HIP_CHECK(hipMemGetInfo(&free, &total)); printf("hipMemGetInfo: free=%zu (%4.2f) Nbytes=%lu total=%zu (%4.2f)\n", free, - (free / 1024.0 / 1024.0), Nbytes, total, (total / 1024.0 / 1024.0)); + (free / 1024.0 / 1024.0), static_cast(Nbytes), total, (total / 1024.0 / 1024.0)); REQUIRE(free + Nbytes <= total); diff --git a/catch/unit/memory/hipPtrGetAttribute.cc b/catch/unit/memory/hipPtrGetAttribute.cc index eaf4c8b153..c87ea5a5f7 100644 --- a/catch/unit/memory/hipPtrGetAttribute.cc +++ b/catch/unit/memory/hipPtrGetAttribute.cc @@ -50,7 +50,7 @@ TEST_CASE("Unit_hipPtrGetAttribute_Simple") { size_t free, total; HIP_CHECK(hipMemGetInfo(&free, &total)); printf("hipMemGetInfo: free=%zu (%4.2f) Nbytes=%lu total=%zu (%4.2f)\n", free, - (free / 1024.0 / 1024.0), Nbytes, total, + (free / 1024.0 / 1024.0), static_cast(Nbytes), total, (total / 1024.0 / 1024.0)); REQUIRE(free + Nbytes <= total); diff --git a/catch/unit/stream/CMakeLists.txt b/catch/unit/stream/CMakeLists.txt index 74024fce2a..b2228290fe 100644 --- a/catch/unit/stream/CMakeLists.txt +++ b/catch/unit/stream/CMakeLists.txt @@ -1,6 +1,5 @@ set(COMMON_SHARED_SRC streamCommon.cc) -if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC hipStreamCreate.cc hipStreamGetFlags.cc @@ -10,43 +9,30 @@ set(TEST_SRC hipStreamCreateWithFlags.cc hipStreamCreateWithPriority.cc hipStreamDestroy.cc - hipStreamGetCUMask.cc hipAPIStreamDisable.cc hipStreamValue.cc - hipStreamWithCUMask.cc - hipStreamACb_MultiThread.cc - hipStreamSynchronize.cc - hipStreamQuery.cc - hipStreamWaitEvent.cc - hipDeviceGetStreamPriorityRange.cc - hipStreamACb_StrmSyncTiming.cc - ) -else() -set(TEST_SRC - hipStreamCreate.cc - hipStreamGetFlags.cc - hipStreamGetPriority.cc - hipMultiStream.cc - hipStreamACb_MultiThread.cc - hipStreamAddCallback.cc - hipStreamCreateWithFlags.cc - hipStreamCreateWithPriority.cc - hipStreamDestroy.cc - hipAPIStreamDisable.cc - # hipStreamAttachMemAsync_old.cc # Disabling it on nvidia due to issue in function definition of hipStreamAttachMemAsync - # Fixing would break ABI, to be re-enabled when the fix is made. - hipStreamValue.cc hipStreamSynchronize.cc hipStreamQuery.cc hipDeviceGetStreamPriorityRange.cc hipStreamACb_StrmSyncTiming.cc - ) +) - # set_source_files_properties(hipStreamAttachMemAsync_old.cc PROPERTIES COMPILE_FLAGS -std=c++17) +if(HIP_PLATFORM MATCHES "amd") + set(TEST_SRC ${TEST_SRC} hipStreamGetCUMask.cc hipStreamWithCUMask.cc + hipStreamACb_MultiThread.cc hipStreamWaitEvent.cc) +else() + set(TEST_SRC + ${TEST_SRC} + # hipStreamAttachMemAsync_old.cc # Disabling it on nvidia due to issue + # in function definition of hipStreamAttachMemAsync + # Fixing would break ABI, to be re-enabled when the fix is made. + hipStreamACb_MultiThread.cc) + + # set_source_files_properties(hipStreamAttachMemAsync_old.cc PROPERTIES + # COMPILE_FLAGS -std=c++17) endif() -hip_add_exe_to_target(NAME StreamTest - TEST_SRC ${TEST_SRC} - TEST_TARGET_NAME build_tests - COMPILE_OPTIONS -std=c++17 - COMMON_SHARED_SRC ${COMMON_SHARED_SRC}) +hip_add_exe_to_target( + NAME StreamTest + TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests + COMPILE_OPTIONS -std=c++17 COMMON_SHARED_SRC ${COMMON_SHARED_SRC}) diff --git a/catch/unit/streamperthread/CMakeLists.txt b/catch/unit/streamperthread/CMakeLists.txt index dd5b54dabc..e185942712 100644 --- a/catch/unit/streamperthread/CMakeLists.txt +++ b/catch/unit/streamperthread/CMakeLists.txt @@ -5,8 +5,17 @@ set(TEST_SRC hipStreamPerThread_MultiThread.cc hipStreamPerThread_DeviceReset.cc hipStreamPerThrdTsts.cc + hipStreamPerThrdCompilerOptn.cc ) +if(HIP_PLATFORM MATCHES "amd") + set_source_files_properties(hipStreamPerThrdCompilerOptn.cc PROPERTIES COMPILE_OPTIONS "-fgpu-default-stream=per-thread") +endif() + +if(HIP_PLATFORM MATCHES "nvidia") + set_source_files_properties(hipStreamPerThrdCompilerOptn.cc PROPERTIES COMPILE_OPTIONS "--default-stream=per-thread") +endif() + hip_add_exe_to_target(NAME StreamPerThreadTest TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests) diff --git a/catch/unit/streamperthread/hipStreamPerThrdCompilerOptn.cc b/catch/unit/streamperthread/hipStreamPerThrdCompilerOptn.cc new file mode 100644 index 0000000000..15e15a38bf --- /dev/null +++ b/catch/unit/streamperthread/hipStreamPerThrdCompilerOptn.cc @@ -0,0 +1,793 @@ +/* +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* This file will be compiled using the compiler option given below. + The functions will be calculating the time taken to execute under + the influence of the compiler flag: + -fgpu-default-stream=per-thread +*/ + +#include +#include +#include "hip/hip_cooperative_groups.h" + +namespace cg = cooperative_groups; + +namespace DefltStrmPT { + int64_t N = 1024 * 1024 * 100; + int64_t Sz = N * sizeof(int64_t); + int64_t *DevA, *HstA, *HstRes; + int64_t OneMB = 1024 * 1024; + int64_t OneMBSz = OneMB * sizeof(int64_t); + hipStream_t Strm; + int clockrate, CONST = 123; + size_t numH = 1024, numW = 1024; + size_t pitch_A, width = numW * sizeof(int64_t); + size_t sizeElements = width * numH; + size_t elements = numW * numH; +} // namespace DefltStrmPT + +__device__ int64_t globalInDStrmPT[1024 * 1024]; +__device__ int SigComplte = 0; + +// Kernel codes +__global__ void DefltStrmPT_Square(int64_t *C_d, int64_t N) { + int64_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); + int64_t stride = hipBlockDim_x * hipGridDim_x; + + for (int64_t i = offset; i < N; i += stride) { + C_d[i] = C_d[i] * C_d[i]; + } +} + +__global__ void Wait_Kernel3(int clockrate, uint64_t WaitSecs, + int PassSignal = 0) { + uint64_t num_cycles = WaitSecs * clockrate * 1000; + uint64_t start = clock64(), cycles = 0; + while (cycles < num_cycles) { + cycles = clock64() - start; + } + if (PassSignal) { + SigComplte = 1; + } +} + +__global__ void DefltStrmPT_Test_gws(uint* buf, uint bufSize, + int64_t* tmpBuf, int64_t* result) { + extern __shared__ int64_t tmp[]; + uint offset = blockIdx.x * blockDim.x + threadIdx.x; + uint stride = gridDim.x * blockDim.x; + cg::grid_group gg = cg::this_grid(); + + int64_t sum = 0; + for (uint i = offset; i < bufSize; i += stride) { + sum += buf[i]; + } + tmp[threadIdx.x] = sum; + + __syncthreads(); + + if (threadIdx.x == 0) { + sum = 0; + for (uint i = 0; i < blockDim.x; i++) { + sum += tmp[i]; + } + tmpBuf[blockIdx.x] = sum; + } + + gg.sync(); + + if (offset == 0) { + for (uint i = 1; i < gridDim.x; ++i) { + sum += tmpBuf[i]; + } + *result = sum; + } +} + +float DefaultPT2_Memcpy_MemSet(int CpyAsync, int MemSetAsync) { + bool IfTstPassed = true; + DefltStrmPT::HstA = reinterpret_cast (malloc(DefltStrmPT::Sz)); + DefltStrmPT::HstRes = reinterpret_cast (malloc(DefltStrmPT::Sz)); + HIP_CHECK(hipDeviceGetAttribute(&(DefltStrmPT::clockrate), + hipDeviceAttributeMemoryClockRate, 0)); + HIP_CHECK(hipMalloc(&(DefltStrmPT::DevA), DefltStrmPT::Sz)); + for (int64_t i = 0; i < DefltStrmPT::N; ++i) { + DefltStrmPT::HstA[i] = DefltStrmPT::CONST; + } + HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm))); + if (CpyAsync) { + HIP_CHECK(hipMemcpyAsync(DefltStrmPT::DevA, DefltStrmPT::HstA, + DefltStrmPT::Sz, hipMemcpyHostToDevice, + DefltStrmPT::Strm)); + HIP_CHECK(hipStreamSynchronize(DefltStrmPT::Strm)); + } else { + HIP_CHECK(hipMemcpy(DefltStrmPT::DevA, DefltStrmPT::HstA, + DefltStrmPT::Sz, hipMemcpyHostToDevice)); + } + DefltStrmPT_Square<<<(DefltStrmPT::N/256 + 1), 256, 0, DefltStrmPT::Strm>>> + (DefltStrmPT::DevA, DefltStrmPT::N); + HIP_CHECK(hipStreamSynchronize(DefltStrmPT::Strm)); + HIP_CHECK(hipMemcpy(DefltStrmPT::HstRes, DefltStrmPT::DevA, + DefltStrmPT::Sz, hipMemcpyDeviceToHost)); + // Verifying the result + for (int64_t i = 0; i < DefltStrmPT::N; ++i) { + if (DefltStrmPT::HstRes[i] != + (DefltStrmPT::HstA[i] * DefltStrmPT::HstA[i])) { + IfTstPassed = false; + } + } + if (MemSetAsync) { + HIP_CHECK(hipMemsetAsync(DefltStrmPT::DevA, 0, DefltStrmPT::Sz, + DefltStrmPT::Strm)); + HIP_CHECK(hipStreamSynchronize(DefltStrmPT::Strm)); + } else { + HIP_CHECK(hipMemset(DefltStrmPT::DevA, 0, + DefltStrmPT::Sz)); + } + // Copying the device memory to host to check if Memset is successful + HIP_CHECK(hipMemcpy(DefltStrmPT::HstA, DefltStrmPT::DevA, + DefltStrmPT::Sz, hipMemcpyDeviceToHost)); + // verifying if memset was successful + for (int64_t i = 0; i < DefltStrmPT::N; ++i) { + if (DefltStrmPT::HstA[i] != 0) { + IfTstPassed = false; + } + } + HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm)); + HIP_CHECK(hipFree(DefltStrmPT::DevA)); + free(DefltStrmPT::HstA); + free(DefltStrmPT::HstRes); + return IfTstPassed; +} + +void DefaultPT2_Memset2D(int Async) { + constexpr int memsetval = 0x24; + constexpr size_t numH = 256; + constexpr size_t numW = 256; + size_t pitch_A; + size_t width = numW * sizeof(char); + size_t sizeElements = width * numH; + size_t elements = numW * numH; + char *A_d, *A_h; + + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width, + numH)); + A_h = reinterpret_cast(malloc(sizeElements)); + REQUIRE(A_h != nullptr); + + for (size_t i = 0; i < elements; i++) { + A_h[i] = 1; + } + + if (Async) { + hipStream_t Strm; + HIP_CHECK(hipStreamCreate(&Strm)); + HIP_CHECK(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, Strm)); + HIP_CHECK(hipStreamSynchronize(Strm)); + HIP_CHECK(hipStreamDestroy(Strm)); + } else { + HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, numW, numH)); + } + HIP_CHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, + hipMemcpyDeviceToHost)); + + for (size_t i = 0; i < elements; i++) { + if (A_h[i] != memsetval) { + INFO("Memset2D mismatch at index:" << i << " computed:" + << A_h[i] << " memsetval:" << memsetval); + REQUIRE(false); + } + } + + HIP_CHECK(hipFree(A_d)); + free(A_h); +} + + +void PerThrdDefltStrm_Memset3D(int Async) { + constexpr int memsetval = 0x22; + constexpr size_t numH = 256; + constexpr size_t numW = 256; + constexpr size_t depth = 10; + size_t width = numW * sizeof(char); + size_t sizeElements = width * numH * depth; + size_t elements = numW * numH * depth; + char *A_h; + + hipExtent extent = make_hipExtent(width, numH, depth); + hipPitchedPtr devPitchedPtr; + + HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); + A_h = reinterpret_cast(malloc(sizeElements)); + REQUIRE(A_h != nullptr); + + for (size_t i = 0; i < elements; i++) { + A_h[i] = 1; + } + + if (Async) { + hipStream_t Strm; + HIP_CHECK(hipStreamCreate(&Strm)); + HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, Strm)); + HIP_CHECK(hipStreamSynchronize(Strm)); + HIP_CHECK(hipStreamDestroy(Strm)); + } else { + HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent)); + } + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH); + myparms.srcPtr = devPitchedPtr; + myparms.extent = extent; +#if HT_NVIDIA + myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost); +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + HIP_CHECK(hipMemcpy3D(&myparms)); + + for (size_t i = 0; i < elements; i++) { + if (A_h[i] != memsetval) { + INFO("Memset3D mismatch at index:" << i << " computed:" + << A_h[i] << " memsetval:" << memsetval); + REQUIRE(false); + } + } + HIP_CHECK(hipFree(devPitchedPtr.ptr)); + free(A_h); +} + + +void DefaultPT2_StrmQuery() { + HIP_CHECK(hipDeviceGetAttribute(&(DefltStrmPT::clockrate), + hipDeviceAttributeMemoryClockRate, 0)); + HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm))); + // StreamQuery with null stream + Wait_Kernel3<<<1, 1>>>(DefltStrmPT::clockrate, 3); + REQUIRE((hipErrorNotReady == hipStreamQuery(0))); + // StreamQuery with user created stream + Wait_Kernel3<<<1, 1, 0, DefltStrmPT::Strm>>>(DefltStrmPT::clockrate, 3); + REQUIRE((hipErrorNotReady == hipStreamQuery(DefltStrmPT::Strm))); + HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm)); +} + + +void DefaultPT2_StreamSync() { + HIP_CHECK(hipDeviceGetAttribute(&(DefltStrmPT::clockrate), + hipDeviceAttributeMemoryClockRate, 0)); + HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm))); + // Calling hipStreamSync on user created stream object + Wait_Kernel3<<<1, 1, 0, DefltStrmPT::Strm>>>(DefltStrmPT::clockrate, 1); + HIP_CHECK(hipStreamSynchronize(DefltStrmPT::Strm)); + // Calling hipStreamSync on null stream + Wait_Kernel3<<<1, 1>>>(DefltStrmPT::clockrate, 1); + HIP_CHECK(hipStreamSynchronize(0)); + HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm)); +} + + +void DefaultPT2_StrmWaitEvent() { + hipEvent_t evt; + hipStream_t Strm1; + HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm))); + HIP_CHECK(hipStreamCreate(&Strm1)); + HIP_CHECK(hipEventCreate(&evt)); + Wait_Kernel3<<<1, 1, 0, DefltStrmPT::Strm>>>(DefltStrmPT::clockrate, 3, 1); + HIP_CHECK(hipEventRecord(evt, DefltStrmPT::Strm)); + HIP_CHECK(hipStreamWaitEvent(Strm1, evt, 0)); + Wait_Kernel3<<<1, 1, 0, Strm1>>>(DefltStrmPT::clockrate, 1); + // By the time control reaches the below point SigComplte is expected + // to be still zero + if (SigComplte) { + REQUIRE(false); + } + HIP_CHECK(hipStreamSynchronize(Strm1)); + HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm)); + HIP_CHECK(hipStreamDestroy(Strm1)); + HIP_CHECK(hipEventDestroy(evt)); +} + +void DefaultPT2_EvtQuery() { + hipEvent_t evt, evt1; + hipError_t err; + HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm))); + HIP_CHECK(hipEventCreate(&evt)); + HIP_CHECK(hipEventCreate(&evt1)); + Wait_Kernel3<<<1, 1, 0, DefltStrmPT::Strm>>>(DefltStrmPT::clockrate, 3); + HIP_CHECK(hipEventRecord(evt, DefltStrmPT::Strm)); + err = hipEventQuery(evt); + if (err != hipErrorNotReady) { + REQUIRE(false); + } + // Testing for Null or default stream + HIP_CHECK(hipEventRecord(evt1, 0)); + std::chrono::time_point start = std::chrono::steady_clock::now(); + int Got_hipSuccess = 0; // 0 for no, 1 for yes + while (true) { + err = hipEventQuery(evt1); + if (err == hipSuccess) { + Got_hipSuccess = 1; + break; + } + if (std::chrono::steady_clock::now() - start > std::chrono::seconds(60)) { + break; + } + } + if (!Got_hipSuccess) { + REQUIRE(false); + } + HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm)); + HIP_CHECK(hipEventDestroy(evt)); + HIP_CHECK(hipEventDestroy(evt1)); +} + + +void Default_LaunchKernel(int NullStrm) { + DefltStrmPT::N = DefltStrmPT::N/4; + DefltStrmPT::Sz = DefltStrmPT::N * sizeof(int64_t); + DefltStrmPT::HstA = reinterpret_cast (malloc(DefltStrmPT::Sz)); + HIP_CHECK(hipMalloc(&(DefltStrmPT::DevA), DefltStrmPT::Sz)); + for (int64_t i = 0; i < DefltStrmPT::N; ++i) { + DefltStrmPT::HstA[i] = DefltStrmPT::CONST; + } + HIP_CHECK(hipMemcpy(DefltStrmPT::DevA, DefltStrmPT::HstA, DefltStrmPT::Sz, + hipMemcpyHostToDevice)); + HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm))); + unsigned ThrdsPerBlk = 32; + unsigned Blocks = ((DefltStrmPT::N + ThrdsPerBlk - 1)/ThrdsPerBlk); + void *Args[] = {&(DefltStrmPT::DevA), &(DefltStrmPT::N)}; + // launch Kernel + if (NullStrm) { + HIP_CHECK(hipLaunchKernel((const void*)DefltStrmPT_Square, + dim3(Blocks, 1, 1), dim3(ThrdsPerBlk, 1, 1), Args, 0, + 0)); + HIP_CHECK(hipStreamSynchronize(0)); + } else { + HIP_CHECK(hipLaunchKernel((const void*)DefltStrmPT_Square, + dim3(Blocks, 1, 1), dim3(ThrdsPerBlk, 1, 1), Args, 0, + DefltStrmPT::Strm)); + HIP_CHECK(hipStreamSynchronize(DefltStrmPT::Strm)); + } + HIP_CHECK(hipMemcpy(DefltStrmPT::HstA, DefltStrmPT::DevA, DefltStrmPT::Sz, + hipMemcpyDeviceToHost)); + for (int64_t i = 0; i < DefltStrmPT::N; ++i) { + if (DefltStrmPT::HstA[i] != (DefltStrmPT::CONST * DefltStrmPT::CONST)) { + REQUIRE(false); + } + } + HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm)); + HIP_CHECK(hipFree(DefltStrmPT::DevA)); + free(DefltStrmPT::HstA); +} + + +void DefaultPT2_LaunchCooperativeKernel(int NullStrm) { + bool IfTestPassed = true; + uint32_t *dA; + int64_t *dB, *dC; + uint32_t BufferSizeInDwords = 448 * 1024 * 1024; + uint32_t* init = new uint32_t[BufferSizeInDwords]; + for (uint32_t i = 0; i < BufferSizeInDwords; ++i) { + init[i] = i; + } + size_t SIZE = BufferSizeInDwords * sizeof(uint); + HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm))); + hipDeviceProp_t deviceProp; + HIP_CHECK(hipGetDeviceProperties(&deviceProp, 0)); + HIPCHECK(hipMalloc(reinterpret_cast(&dA), SIZE)); + HIPCHECK(hipHostMalloc(reinterpret_cast(&dC), sizeof(int64_t))); + HIPCHECK(hipMemcpy(dA, init, SIZE, hipMemcpyHostToDevice)); + + dim3 dimBlock = dim3(1); + dim3 dimGrid = dim3(1); + int numBlocks = 0; + dimBlock.x = 32; + // Calculate the device occupancy to know how many blocks can be run + // concurrently + HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks, + DefltStrmPT_Test_gws, + dimBlock.x * dimBlock.y * dimBlock.z, + dimBlock.x * sizeof(int64_t))); + dimGrid.x = deviceProp.multiProcessorCount * std::min(numBlocks, 32); + HIPCHECK(hipMalloc(reinterpret_cast(&dB), + dimGrid.x * sizeof(int64_t))); + + void *params[4]; + params[0] = reinterpret_cast(&dA); + params[1] = reinterpret_cast(&BufferSizeInDwords); + params[2] = reinterpret_cast(&dB); + params[3] = reinterpret_cast(&dC); + if (NullStrm) { + HIPCHECK(hipLaunchCooperativeKernel( + reinterpret_cast(DefltStrmPT_Test_gws), + dimGrid, dimBlock, params, dimBlock.x * sizeof(int64_t), 0)); + HIP_CHECK(hipStreamSynchronize(0)); + } else { + HIPCHECK(hipLaunchCooperativeKernel( + reinterpret_cast(DefltStrmPT_Test_gws), + dimGrid, dimBlock, params, dimBlock.x * sizeof(int64_t), + DefltStrmPT::Strm)); + HIP_CHECK(hipStreamSynchronize(DefltStrmPT::Strm)); + } + HIPCHECK(hipMemcpy(init, dC, sizeof(int64_t), hipMemcpyDeviceToHost)); + + if (*dC != (((int64_t)(BufferSizeInDwords) * (BufferSizeInDwords - 1)) / 2)) { + std::cout << "Data validation failed for grid size = " << dimGrid.x << + " and block size = " << dimBlock.x << "\n"; + std::cout << "Test failed! \n"; + IfTestPassed = false; + } + HIPCHECK(hipStreamDestroy(DefltStrmPT::Strm)); + HIPCHECK(hipHostFree(dC)); + HIPCHECK(hipFree(dB)); + HIPCHECK(hipFree(dA)); + delete [] init; + REQUIRE(IfTestPassed); +} + + +void DefaultPT2_StrmGetFlag() { + HIP_CHECK(hipStreamCreateWithFlags(&(DefltStrmPT::Strm), hipStreamDefault)); + unsigned int flag = 9999; + HIP_CHECK(hipStreamGetFlags(DefltStrmPT::Strm, &flag)); + if (flag != 0) { + INFO("Expected flag value: 0, Received flag value: %u\n" << flag); + REQUIRE(false); + } + HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm)); + HIP_CHECK(hipStreamCreateWithFlags(&(DefltStrmPT::Strm), + hipStreamNonBlocking)); + flag = 9999; + HIP_CHECK(hipStreamGetFlags(DefltStrmPT::Strm, &flag)); + if (flag != 1) { + INFO("Expected flag value: 1, Received flag value: %u\n" << flag); + REQUIRE(false); + } + HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm)); +} + +void DefaultPT2_StrmGetPriority() { + int low, high, ObsrvdPriority; + HIP_CHECK(hipDeviceGetStreamPriorityRange(&low, &high)); + INFO("Lowest possible priority: %d\n" << low); + INFO("Highest possible priority: %d\n" << high); + INFO("Creating streams with flag hipStreamNonBlocking\n"); + // hipStrmFlg = 0 = hipStreamDefault + // hipStrmFlg = 1 = hipStreamNonBlocking + for (int hipStrmFlg = 0; hipStrmFlg < 2; ++hipStrmFlg) { + for (int Priority = low; Priority <= high; ++Priority) { + if (hipStrmFlg == 0) { + HIP_CHECK(hipStreamCreateWithPriority(&(DefltStrmPT::Strm), + hipStreamDefault, Priority)); + } else { + HIP_CHECK(hipStreamCreateWithPriority(&(DefltStrmPT::Strm), + hipStreamNonBlocking, Priority)); + } + HIP_CHECK(hipStreamGetPriority(DefltStrmPT::Strm, &ObsrvdPriority)); + if (ObsrvdPriority != Priority) { + INFO("Expected priority: %d" << Priority << " Observed Priority: %d\n" + << ObsrvdPriority); + INFO("Test Failed!\n\n"); + REQUIRE(false); + } + HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm)); + } + } + INFO("Checking priority on null stream!!\n"); + HIP_CHECK(hipStreamGetPriority(0, &ObsrvdPriority)); + if (ObsrvdPriority != 0) { + INFO("Expected priority: 0, Observed Priority: %d\n" + << ObsrvdPriority); + INFO("Test Failed!\n\n"); + REQUIRE(false); + } +} + + +void DefaultPT2_hipMemcpyFromSymbol() { + int64_t *Hst = nullptr; + HIP_CHECK(hipHostMalloc(&(Hst), DefltStrmPT::OneMBSz)); + HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm))); + for (int i = 0; i < DefltStrmPT::OneMB; ++i) { + Hst[i] = DefltStrmPT::CONST; + } + HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(globalInDStrmPT), Hst, + DefltStrmPT::OneMBSz, 0, hipMemcpyHostToDevice)); + for (int i = 0; i < DefltStrmPT::OneMB; ++i) { + Hst[i] = 0; + } + HIP_CHECK(hipMemcpyFromSymbol(Hst, HIP_SYMBOL(globalInDStrmPT), + DefltStrmPT::OneMBSz, 0, hipMemcpyDeviceToHost)); + for (int i = 0; i < DefltStrmPT::OneMB; ++i) { + if (Hst[i] != DefltStrmPT::CONST) { + REQUIRE(false); + } + } + HIP_CHECK(hipHostFree(Hst)); + HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm)); +} + +void DefaultPT2_hipMemcpy2D(int Async) { + DefltStrmPT::numH = 1024; + DefltStrmPT::numW = 1024; + DefltStrmPT::width = DefltStrmPT::numW * sizeof(int64_t); + HIP_CHECK(hipHostMalloc(&(DefltStrmPT::HstA), + (DefltStrmPT::numH * DefltStrmPT::numW * sizeof(int64_t)))); + HIP_CHECK(hipHostMalloc(&(DefltStrmPT::HstRes), + (DefltStrmPT::numH * DefltStrmPT::numW * sizeof(int64_t)))); + DefltStrmPT::width = DefltStrmPT::numW * sizeof(int64_t); + for (size_t row = 0; row < DefltStrmPT::numH; ++row) { + for (size_t column = 0; column < DefltStrmPT::numW; ++column) { + DefltStrmPT::HstA[(row * DefltStrmPT::numW) + column] = + DefltStrmPT::CONST; + DefltStrmPT::HstRes[(row * DefltStrmPT::numW) + column] = 0; + } + } + HIP_CHECK(hipMallocPitch(reinterpret_cast(&(DefltStrmPT::DevA)), + &(DefltStrmPT::pitch_A), DefltStrmPT::width, DefltStrmPT::numH)); + if (Async) { + HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm))); + HIP_CHECK(hipMemcpy2DAsync(DefltStrmPT::DevA, DefltStrmPT::pitch_A, + DefltStrmPT::HstA, DefltStrmPT::numW*sizeof(int64_t), + DefltStrmPT::numW*sizeof(int64_t), DefltStrmPT::numH, + hipMemcpyHostToDevice, DefltStrmPT::Strm)); + HIP_CHECK(hipMemcpy2DAsync(DefltStrmPT::HstRes, + DefltStrmPT::numW*sizeof(int64_t), + DefltStrmPT::DevA, DefltStrmPT::pitch_A, + DefltStrmPT::numW*sizeof(int64_t), DefltStrmPT::numH, + hipMemcpyDeviceToHost, DefltStrmPT::Strm)); + HIP_CHECK(hipStreamSynchronize(DefltStrmPT::Strm)); + HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm)); + } else { + HIP_CHECK(hipMemcpy2D(DefltStrmPT::DevA, DefltStrmPT::pitch_A, + DefltStrmPT::HstA, DefltStrmPT::numW*sizeof(int64_t), + DefltStrmPT::numW*sizeof(int64_t), DefltStrmPT::numH, + hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy2D(DefltStrmPT::HstRes, + DefltStrmPT::numW*sizeof(int64_t), DefltStrmPT::DevA, + DefltStrmPT::pitch_A, DefltStrmPT::numW*sizeof(int64_t), + DefltStrmPT::numH, hipMemcpyDeviceToHost)); + } + for (size_t row = 0; row < DefltStrmPT::numH; ++row) { + for (size_t column = 0; column < DefltStrmPT::numW; ++column) { + if (DefltStrmPT::HstRes[(row * DefltStrmPT::numW) + column] + != DefltStrmPT::CONST) { + REQUIRE(false); + } + } + } + HIP_CHECK(hipFree(DefltStrmPT::DevA)); + HIP_CHECK(hipHostFree(DefltStrmPT::HstA)); +} + + +void DefaultPT2_hipMemcpy2DToArray() { + hipArray *Dptr = nullptr; + float *Hptr = nullptr, *HRes = nullptr; + DefltStrmPT::numH = 1024; + DefltStrmPT::numW = 1024; + DefltStrmPT::width = DefltStrmPT::numW * sizeof(float); + Hptr = new float[DefltStrmPT::width * DefltStrmPT::numH]; + HRes = new float[DefltStrmPT::width * DefltStrmPT::numH]; + for (size_t i = 0; i < DefltStrmPT::width * DefltStrmPT::numH; ++i) { + Hptr[i] = DefltStrmPT::CONST; + } + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&(Dptr), &desc, DefltStrmPT::numW, + DefltStrmPT::numH, hipArrayDefault)); + HIP_CHECK(hipMemcpy2DToArray(Dptr, 0, 0, Hptr, DefltStrmPT::width, + DefltStrmPT::width, DefltStrmPT::numH, + hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy2DFromArray(HRes, DefltStrmPT::width, Dptr, 0, 0, + DefltStrmPT::width, DefltStrmPT::numH, hipMemcpyDeviceToHost)); + // verifying the result + for (size_t i = 0; i < DefltStrmPT::numW * DefltStrmPT::numH; ++i) { + if (HRes[i] != DefltStrmPT::CONST) { + REQUIRE(false); + } + } + delete[] Hptr; + delete[] HRes; + HIP_CHECK(hipFreeArray(Dptr)); +} + + +float DefaultPT2_hipMemcpy2DFromArray() { + hipArray *Dptr = nullptr; + float *Hptr_A = nullptr, *Hptr_B = nullptr; + DefltStrmPT::numH = 1024; + DefltStrmPT::numW = 1024; + DefltStrmPT::width = DefltStrmPT::numW * sizeof(float); + HIP_CHECK(hipDeviceGetAttribute(&(DefltStrmPT::clockrate), + hipDeviceAttributeMemoryClockRate, 0)); + HIP_CHECK(hipHostMalloc(&(Hptr_A), + (DefltStrmPT::width * DefltStrmPT::numH * sizeof(float)))); + HIP_CHECK(hipHostMalloc(&(Hptr_B), + (DefltStrmPT::width * DefltStrmPT::numH * sizeof(float)))); + for (size_t i = 0; i < (DefltStrmPT::width * DefltStrmPT::numH); ++i) { + Hptr_A[i] = DefltStrmPT::CONST; + } + hipChannelFormatDesc desc = hipCreateChannelDesc(); + HIP_CHECK(hipMallocArray(&(Dptr), &desc, DefltStrmPT::numW, + DefltStrmPT::numH, hipArrayDefault)); + HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm))); + HIP_CHECK(hipMemcpy2DToArray(Dptr, 0, 0, Hptr_A, DefltStrmPT::width, + DefltStrmPT::width, DefltStrmPT::numH, + hipMemcpyHostToDevice)); + Wait_Kernel3 <<< 1, 1, 0, DefltStrmPT::Strm >>> (DefltStrmPT::clockrate, + 1); + HIP_CHECK(hipMemcpy2DFromArray(Hptr_B, DefltStrmPT::width, Dptr, 0, 0, + DefltStrmPT::width, DefltStrmPT::numH, hipMemcpyDeviceToHost)); + HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm)); + HIP_CHECK(hipFreeArray(Dptr)); + HIP_CHECK(hipHostFree(Hptr_A)); + HIP_CHECK(hipHostFree(Hptr_B)); + return true; +} + +void DefaultPT2_hipMemcpy3D() { + int width = 8, height = 8, depth = 8; + unsigned int size = width * height * depth * sizeof(float); + float* Hptr = reinterpret_cast(malloc(size)); + float* HRes = reinterpret_cast(malloc(size)); + memset(Hptr, 0, size); + memset(HRes, 0, size); + hipExtent extent = make_hipExtent(width, height, depth); + + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + Hptr[i*width*height + j*width +k] = i*width*height + j*width + k; + } + } + } + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0, + 0, 0, hipChannelFormatKindFloat); + hipArray *arr; + + HIP_CHECK(hipMalloc3DArray(&arr, &channelDesc, + make_hipExtent(width, height, depth), hipArrayDefault)); + hipMemcpy3DParms myparms{0, {0, 0, 0}, {0, 0, 0, 0}, 0, {0, 0, 0}, + {0, 0, 0, 0}, {0, 0, 0}, hipMemcpyDefault}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.srcPtr = make_hipPitchedPtr(Hptr, width * sizeof(float), width, + height); + myparms.dstArray = arr; + myparms.extent = extent; + +#ifdef __HIP_PLATFORM_NVIDIA__ + myparms.kind = cudaMemcpyHostToDevice; +#else + myparms.kind = hipMemcpyHostToDevice; +#endif + // Host to Device copy + HIP_CHECK(hipMemcpy3D(&myparms)); + + // Device to Host copy + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(HRes, width * sizeof(float), width, + height); + myparms.srcArray = arr; + myparms.extent = extent; +#ifdef __HIP_PLATFORM_NVIDIA__ + myparms.kind = cudaMemcpyDeviceToHost; +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + HIP_CHECK(hipMemcpy3D(&myparms)); + + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + if (HRes[i*width*height + j*width +k] != i*width*height + j*width + k) { + REQUIRE(false); + } + } + } + } + HIP_CHECK(hipFreeArray(arr)); + free(Hptr); + free(HRes); +} + + +TEST_CASE("Unit_hipStrmPerThrdDefault") { + SECTION("Testing hipMemset/Memcpy() and their async version") { + REQUIRE(DefaultPT2_Memcpy_MemSet(1, 0)); + REQUIRE(DefaultPT2_Memcpy_MemSet(1, 1)); + REQUIRE(DefaultPT2_Memcpy_MemSet(0, 1)); + REQUIRE(DefaultPT2_Memcpy_MemSet(0, 0)); + } + SECTION("Testing hipMemset2D() and its async version") { + DefaultPT2_Memset2D(0); + DefaultPT2_Memset2D(1); + } + + SECTION("Testing_hipMemset3D() and its async version") { + PerThrdDefltStrm_Memset3D(0); + PerThrdDefltStrm_Memset3D(1); + } + + SECTION("Testing_hipStreamQuery()") { + DefaultPT2_StrmQuery(); + } + + SECTION("Testing_hipStreamSynchronize()") { + DefaultPT2_StreamSync(); + } + + SECTION("Testing_hipLaunchKernel()") { + // launch with null stream + Default_LaunchKernel(1); + // launch with user created stream + Default_LaunchKernel(0); + } + + hipDeviceProp_t deviceProp; + HIP_CHECK(hipGetDeviceProperties(&deviceProp, 0)); + if (deviceProp.cooperativeLaunch) { + SECTION("Testing_hipLaunchCooperativeKernel()") { + // launching hipLaunchCooperativeKernel() with Null stream + DefaultPT2_LaunchCooperativeKernel(1); + // launching hipLaunchCooperativeKernel() with user created stream + DefaultPT2_LaunchCooperativeKernel(0); + } + } else { + INFO("Cooperative Launch feature is not supported, therefore skipping"); + INFO(" the test Testing_hipLaunchCooperativeKernel()"); + } + + SECTION("Testing_StrmWaitEvent()") { + DefaultPT2_StrmWaitEvent(); + } + + SECTION("Testing_hipStreamGetFlag()") { + DefaultPT2_StrmGetFlag(); + } + + SECTION("Testing_hipStreamGetPriority()") { + DefaultPT2_StrmGetPriority(); + } + + SECTION("Testing_hipMemcpyFrom/To/Symbol()") { + DefaultPT2_hipMemcpyFromSymbol(); + } + + SECTION("Testing_hipMemcpy2D() & its Async version") { + DefaultPT2_hipMemcpy2D(0); + DefaultPT2_hipMemcpy2D(1); + } + + SECTION("Testing_hipMemcpy2DToArray()") { + DefaultPT2_hipMemcpy2DToArray(); + } + + SECTION("Testing_hipMemcpy3D()") { + DefaultPT2_hipMemcpy3D(); + } + + SECTION("Testing_hipEventQuery()") { + DefaultPT2_EvtQuery(); + } +} diff --git a/catch/unit/texture/hipBindTex2DPitch.cc b/catch/unit/texture/hipBindTex2DPitch.cc index 4bd671af00..26e073ed83 100644 --- a/catch/unit/texture/hipBindTex2DPitch.cc +++ b/catch/unit/texture/hipBindTex2DPitch.cc @@ -20,6 +20,9 @@ THE SOFTWARE. #include #include + +#if CUDA_VERSION < CUDA_12000 + #define SIZE_H 8 #define SIZE_W 12 #define TYPE_t float @@ -77,3 +80,6 @@ TEST_CASE("Unit_hipBindTexture2D_Pitch") { HIP_CHECK(hipFree(devPtrA)); HIP_CHECK(hipFree(devPtrB)); } + +#endif // CUDA_VERSION < CUDA_12000 + diff --git a/catch/unit/texture/hipBindTexRef1DFetch.cc b/catch/unit/texture/hipBindTexRef1DFetch.cc index 7291b786bb..a223d27711 100644 --- a/catch/unit/texture/hipBindTexRef1DFetch.cc +++ b/catch/unit/texture/hipBindTexRef1DFetch.cc @@ -19,6 +19,8 @@ THE SOFTWARE. #include +#if CUDA_VERSION < CUDA_12000 + #define N 512 texture tex; @@ -79,3 +81,7 @@ TEST_CASE("Unit_hipBindTexture_tex1DfetchVerification") { HIP_CHECK(hipFree(texBuf)); HIP_CHECK(hipFree(devBuf)); } + + +#endif // CUDA_VERSION < CUDA_12000 + diff --git a/catch/unit/texture/hipNormalizedFloatValueTex.cc b/catch/unit/texture/hipNormalizedFloatValueTex.cc index 8e660c41cf..3bc1d8ca7d 100644 --- a/catch/unit/texture/hipNormalizedFloatValueTex.cc +++ b/catch/unit/texture/hipNormalizedFloatValueTex.cc @@ -19,6 +19,9 @@ THE SOFTWARE. #include + +#if CUDA_VERSION < CUDA_12000 + #define SIZE 10 #define EPSILON 0.00001 #define THRESH_HOLD 0.01 // For filter mode @@ -159,3 +162,5 @@ TEST_CASE("Unit_hipNormalizedFloatValueTex_CheckModes") { runTest_hipTextureFilterMode(); } } + +#endif // CUDA_VERSION < CUDA_12000 diff --git a/catch/unit/texture/hipSimpleTexture2DLayered.cc b/catch/unit/texture/hipSimpleTexture2DLayered.cc index 1bf6ea3af7..73d2167d5c 100644 --- a/catch/unit/texture/hipSimpleTexture2DLayered.cc +++ b/catch/unit/texture/hipSimpleTexture2DLayered.cc @@ -20,6 +20,8 @@ THE SOFTWARE. #include #include +#if CUDA_VERSION < CUDA_12000 + typedef float T; // Texture reference for 2D Layered texture @@ -107,3 +109,6 @@ TEST_CASE("Unit_hipSimpleTexture2DLayered_Check") { free(hData); free(hOutputData); } + +#endif // CUDA_VERSION < CUDA_12000 + diff --git a/catch/unit/texture/hipSimpleTexture3D.cc b/catch/unit/texture/hipSimpleTexture3D.cc index 0de3aec096..c17d218a53 100644 --- a/catch/unit/texture/hipSimpleTexture3D.cc +++ b/catch/unit/texture/hipSimpleTexture3D.cc @@ -20,6 +20,8 @@ THE SOFTWARE. #include #include +#if CUDA_VERSION < CUDA_12000 + // Texture reference for 3D texture texture texf; texture texi; @@ -119,3 +121,7 @@ TEST_CASE("Unit_hipSimpleTexture3D_Check_DataTypes") { runSimpleTexture3D_Check(i, i+1, i, &texc); } } + + + +#endif // CUDA_VERSION < CUDA_12000 diff --git a/catch/unit/texture/hipTextureRef2D.cc b/catch/unit/texture/hipTextureRef2D.cc index d0a31644be..9ecaab9365 100644 --- a/catch/unit/texture/hipTextureRef2D.cc +++ b/catch/unit/texture/hipTextureRef2D.cc @@ -19,6 +19,9 @@ THE SOFTWARE. #include + +#if CUDA_VERSION < CUDA_12000 + texture tex; __global__ void tex2DKernel(float* outputData, int width) { @@ -90,3 +93,7 @@ TEST_CASE("Unit_hipTextureRef2D_Check") { HIP_CHECK(hipFree(dData)); HIP_CHECK(hipFreeArray(hipArray)); } + + +#endif // CUDA_VERSION < CUDA_12000 + diff --git a/samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt b/samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt index 248dd355b0..7045a891e1 100644 --- a/samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt +++ b/samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt @@ -18,7 +18,7 @@ # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN # THE SOFTWARE. -cmake_minimum_required(VERSION 2.8.3) +cmake_minimum_required(VERSION 3.10) if (NOT DEFINED ROCM_PATH ) set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." ) endif () @@ -45,16 +45,15 @@ endif() set(MY_SOURCE_FILES MatrixTranspose.cpp) set(MY_TARGET_NAME MatrixTranspose) set(MY_HIPCC_OPTIONS) -set(MY_HCC_OPTIONS) set(MY_CLANG_OPTIONS) set(MY_NVCC_OPTIONS) set_source_files_properties(${MY_SOURCE_FILES} PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) -hip_add_executable(${MY_TARGET_NAME} ${MY_SOURCE_FILES} HIPCC_OPTIONS ${MY_HIPCC_OPTIONS} HCC_OPTIONS ${MY_HCC_OPTIONS} CLANG_OPTIONS ${MY_CLANG_OPTIONS} NVCC_OPTIONS ${MY_NVCC_OPTIONS}) +hip_add_executable(${MY_TARGET_NAME} ${MY_SOURCE_FILES} HIPCC_OPTIONS ${MY_HIPCC_OPTIONS} CLANG_OPTIONS ${MY_CLANG_OPTIONS} NVCC_OPTIONS ${MY_NVCC_OPTIONS}) # Search for rocm in common locations list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH}) -find_package(hip QUIET) +find_package(hip QUIET CONFIG) if(TARGET hip::host) message(STATUS "Found hip::host at ${hip_DIR}") target_link_libraries(${MY_TARGET_NAME} hip::host)