From e2ae92502ce3c8ea165ece4e0ad8917f581204b3 Mon Sep 17 00:00:00 2001 From: agunashe <86270081+agunashe@users.noreply.github.com> Date: Tue, 8 Mar 2022 16:08:25 -0800 Subject: [PATCH 1/9] SWDEV-313590 - appending rocmpath at the end of rpath after rebase (#2526) --- bin/hipcc.pl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/bin/hipcc.pl b/bin/hipcc.pl index 7b84067202..8431d1621e 100755 --- a/bin/hipcc.pl +++ b/bin/hipcc.pl @@ -707,7 +707,7 @@ if ($HIP_PLATFORM eq "amd") { if ($linkType eq 0) { $toolArgs = " -L$HIP_LIB_PATH -lamdhip64 -L$ROCM_PATH/lib -lhsa-runtime64 -ldl -lnuma " . ${toolArgs}; } else { - $toolArgs = " -Wl,--enable-new-dtags -Wl,-rpath=$HIP_LIB_PATH:$ROCM_PATH/lib -lamdhip64 " . ${toolArgs}; + $toolArgs = ${toolArgs} . " -Wl,--enable-new-dtags -Wl,-rpath=$HIP_LIB_PATH:$ROCM_PATH/lib -lamdhip64 "; } # To support __fp16 and _Float16, explicitly link with compiler-rt $HIP_CLANG_BUILTIN_LIB="$HIP_CLANG_PATH/../lib/clang/$HIP_CLANG_VERSION/lib/$HIP_CLANG_TARGET/libclang_rt.builtins.a"; From 69ef2e37a2128d45127b1def6dba644cd085bffd Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 10 Mar 2022 15:01:43 +0530 Subject: [PATCH 2/9] Add mising p2p checks in graph instantiate tests (#2525) Change-Id: I37ed539ccb9fc487872807170dc62eddc598f1f6 --- .../graph/hipGraphInstantiateWithFlags.cc | 28 +++++++++++++++++-- 1 file changed, 26 insertions(+), 2 deletions(-) diff --git a/tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc b/tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc index b8fabb69ee..5e53d0357c 100644 --- a/tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc +++ b/tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc @@ -248,7 +248,19 @@ by creating dependency graph and instantiate, launching and verifying the result */ TEST_CASE("Unit_hipGraphInstantiateWithFlags_DependencyGraph") { - GraphInstantiateWithFlags_DependencyGraph(); + int numDevices = 0; + int canAccessPeer = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + GraphInstantiateWithFlags_DependencyGraph(); + } else { + SUCCEED("Machine does not seem to have P2P"); + } + } else { + SUCCEED("skipped the testcase as no of devices is less than 2"); + } } /* This testcase verifies hipGraphInstantiateWithFlags API @@ -278,7 +290,19 @@ by creating capture graph and instantiate, launching and verifying the result */ TEST_CASE("Unit_hipGraphInstantiateWithFlags_StreamCapture") { - GraphInstantiateWithFlags_StreamCapture(); + int numDevices = 0; + int canAccessPeer = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + GraphInstantiateWithFlags_StreamCapture(); + } else { + SUCCEED("Machine does not seem to have P2P"); + } + } else { + SUCCEED("skipped the testcase as no of devices is less than 2"); + } } /* From 6c74d44a1f8bf66477b48c3c6db5a3597726a6bc Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 10 Mar 2022 15:02:31 +0530 Subject: [PATCH 3/9] SWDEV-323848 - Fix for p2p failed tests (#2537) Added peer check for the p2p related testcases Change-Id: I9131b284baf33bef18eceaa7ed462d733f91be92 --- tests/catch/multiproc/hipIpcMemAccessTest.cc | 10 +-- tests/catch/unit/memory/hipMemcpy3DAsync.cc | 84 +++++++++++--------- 2 files changed, 50 insertions(+), 44 deletions(-) diff --git a/tests/catch/multiproc/hipIpcMemAccessTest.cc b/tests/catch/multiproc/hipIpcMemAccessTest.cc index 4e143111c4..183da51c24 100644 --- a/tests/catch/multiproc/hipIpcMemAccessTest.cc +++ b/tests/catch/multiproc/hipIpcMemAccessTest.cc @@ -118,12 +118,12 @@ TEST_CASE("Unit_hipIpcMemAccess_Semaphores") { } for (int i = 0; i < Num_devices; ++i) { HIP_CHECK(hipSetDevice(i)); - HIP_CHECK(hipMalloc(&C_d, Nbytes)); - HIP_CHECK(hipIpcOpenMemHandle(reinterpret_cast(&B_d), - shrd_mem->memHandle, - hipIpcMemLazyEnablePeerAccess)); HIP_CHECK(hipDeviceCanAccessPeer(&CanAccessPeer, i, shrd_mem->device)); if (CanAccessPeer == 1) { + HIP_CHECK(hipMalloc(&C_d, Nbytes)); + HIP_CHECK(hipIpcOpenMemHandle(reinterpret_cast(&B_d), + shrd_mem->memHandle, + hipIpcMemLazyEnablePeerAccess)); HIP_CHECK(hipMemcpy(C_d, B_d, Nbytes, hipMemcpyDeviceToDevice)); HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); HipTest::checkTest(A_h, C_h, N); @@ -131,9 +131,9 @@ TEST_CASE("Unit_hipIpcMemAccess_Semaphores") { // Checking if the data obtained from Ipc shared memory is consistent HIP_CHECK(hipMemcpy(C_h, B_d, Nbytes, hipMemcpyDeviceToHost)); HipTest::checkTest(A_h, C_h, N); - } HIP_CHECK(hipIpcCloseMemHandle(reinterpret_cast(B_d))); HIP_CHECK(hipFree(C_d)); + } } if ((sem_post(sem_ob2)) == -1) { shrd_mem->IfTestPassed = false; diff --git a/tests/catch/unit/memory/hipMemcpy3DAsync.cc b/tests/catch/unit/memory/hipMemcpy3DAsync.cc index 6ebc480844..5918a287c9 100644 --- a/tests/catch/unit/memory/hipMemcpy3DAsync.cc +++ b/tests/catch/unit/memory/hipMemcpy3DAsync.cc @@ -481,58 +481,64 @@ template void Memcpy3DAsync::D2D_SameDeviceMem_StreamDiffDevice() { HIP_CHECK(hipSetDevice(0)); // Allocating the Memory - AllocateMemory(); - HIP_CHECK(hipSetDevice(1)); - HIP_CHECK(hipStreamCreate(&stream)); - memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); - SetDefaultData(); + int peerAccess = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 0, 1)); + if (peerAccess) { + AllocateMemory(); + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipStreamCreate(&stream)); + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); - // Host to Device - myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height); - myparms.dstArray = arr; + // Host to Device + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height); + myparms.dstArray = arr; #ifdef __HIP_PLATFORM_NVCC__ - myparms.kind = cudaMemcpyHostToDevice; + myparms.kind = cudaMemcpyHostToDevice; #else - myparms.kind = hipMemcpyHostToDevice; + myparms.kind = hipMemcpyHostToDevice; #endif - REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); - HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); - // Array to Array - memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); - SetDefaultData(); - myparms.srcArray = arr; - myparms.dstArray = arr1; + // Array to Array + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); + myparms.srcArray = arr; + myparms.dstArray = arr1; #ifdef __HIP_PLATFORM_NVCC__ - myparms.kind = cudaMemcpyDeviceToDevice; + myparms.kind = cudaMemcpyDeviceToDevice; #else - myparms.kind = hipMemcpyDeviceToDevice; + myparms.kind = hipMemcpyDeviceToDevice; #endif - REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); - HIP_CHECK(hipStreamSynchronize(stream)); - T *hOutputData = reinterpret_cast(malloc(size)); - memset(hOutputData, 0, size); + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + T *hOutputData = reinterpret_cast(malloc(size)); + memset(hOutputData, 0, size); - // Device to host - memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); - SetDefaultData(); - myparms.dstPtr = make_hipPitchedPtr(hOutputData, - width * sizeof(T), width, height); - myparms.srcArray = arr1; + // Device to host + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); + myparms.dstPtr = make_hipPitchedPtr(hOutputData, + width * sizeof(T), width, height); + myparms.srcArray = arr1; #ifdef __HIP_PLATFORM_NVCC__ - myparms.kind = cudaMemcpyDeviceToHost; + myparms.kind = cudaMemcpyDeviceToHost; #else - myparms.kind = hipMemcpyDeviceToHost; + myparms.kind = hipMemcpyDeviceToHost; #endif - REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); - HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); - // Validating the result - HipTest::checkArray(hData, hOutputData, width, height, depth); + // Validating the result + HipTest::checkArray(hData, hOutputData, width, height, depth); - // Deallocating the memory - free(hOutputData); - DeAllocateMemory(); + // Deallocating the memory + free(hOutputData); + DeAllocateMemory(); + } else { + SUCCEED("Skipped the test as there is no peer access"); + } } /* @@ -630,7 +636,7 @@ void Memcpy3DAsync::simple_Memcpy3DAsync() { memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); SetDefaultData(); myparms.dstPtr = make_hipPitchedPtr(hOutputData, - width * sizeof(T), width, height); + width * sizeof(T), width, height); myparms.srcArray = arr1; #ifdef __HIP_PLATFORM_NVCC__ myparms.kind = cudaMemcpyDeviceToHost; From 8e2aee5dbe3ac7cfb75faeec25f858ad925860ac Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 10 Mar 2022 15:03:29 +0530 Subject: [PATCH 4/9] SWDEV-324081 - Fix for dtest (#2540) Change-Id: I114cb00d0bf2f899f1f299891ccbd4c0bea29150 --- tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp b/tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp index 8950516b41..e6f5c49a9b 100644 --- a/tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp +++ b/tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp @@ -1,5 +1,5 @@ /* - Copyright (c) 2020 - 2021 Advanced Micro Devices, Inc. All rights reserved. + Copyright (c) 2020 - 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 @@ -185,7 +185,6 @@ int main(int argc, char* argv[]) { bool testStatus = true; HipTest::parseStandardArguments(argc, argv, false); if (p_tests == 1) { - testStatus &= ConcurencyCheck_GlobalVar(1); testStatus &= ConcurencyCheck_GlobalVar(0); } else if (p_tests == 2) { testStatus &= KernelTimeExecution(); From b61eb2e79c8c6abb88a579d18dcae9e338841656 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 10 Mar 2022 15:04:01 +0530 Subject: [PATCH 5/9] SWDEV-323167 - Fix for test hang (#2541) Change-Id: Ic038ceb2368839f67af65d213694820f881f4193 --- tests/src/ipc/hipMultiProcIpcMem.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/src/ipc/hipMultiProcIpcMem.cpp b/tests/src/ipc/hipMultiProcIpcMem.cpp index 268d610180..3829853774 100644 --- a/tests/src/ipc/hipMultiProcIpcMem.cpp +++ b/tests/src/ipc/hipMultiProcIpcMem.cpp @@ -19,7 +19,7 @@ THE SOFTWARE. /* HIT_START * BUILD: %t %s ../test_common.cpp - * TEST: %t + * TEST: %t --N 4 * HIT_END */ From 72af2f3299bc7d977d59995693d1b21ec7133a71 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 10 Mar 2022 15:04:48 +0530 Subject: [PATCH 6/9] SWDEV-323849 - Fix for dtests failures (#2539) Change-Id: Iffd78656d7f67b11d3eb6d1e7859bb6168b68a57 --- tests/catch/multiproc/hipMemCoherencyTstMProc.cc | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/tests/catch/multiproc/hipMemCoherencyTstMProc.cc b/tests/catch/multiproc/hipMemCoherencyTstMProc.cc index 0f842464b9..8e36c4331d 100644 --- a/tests/catch/multiproc/hipMemCoherencyTstMProc.cc +++ b/tests/catch/multiproc/hipMemCoherencyTstMProc.cc @@ -1,5 +1,5 @@ /* - Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. + Copyright (c) 2021 - 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 @@ -128,7 +128,7 @@ TEST_CASE("Unit_malloc_CoherentTst") { // Read string from child and close reading end. read(fd1[0], GpuId, 2 * sizeof(int)); close(fd1[0]); - if ((GpuId[0] == 1) || (GpuId[0] == 1)) { + if ((GpuId[0] == 1) || (GpuId[1] == 1)) { WARN("This test is not applicable on MI60 & MI100." "Skipping the test!!"); exit(0); @@ -223,7 +223,7 @@ TEST_CASE("Unit_malloc_CoherentTstWthAdvise") { // Read string from child and close reading end. read(fd1[0], GpuId, 2 * sizeof(int)); close(fd1[0]); - if ((GpuId[0] == 1) || (GpuId[0] == 1)) { + if ((GpuId[0] == 1) || (GpuId[1] == 1)) { WARN("This test is not applicable on MI60 & MI100." "Skipping the test!!"); exit(0); @@ -320,7 +320,7 @@ TEST_CASE("Unit_mmap_CoherentTst") { // Read string from child and close reading end. read(fd1[0], GpuId, 2 * sizeof(int)); close(fd1[0]); - if ((GpuId[0] == 1) || (GpuId[0] == 1)) { + if ((GpuId[0] == 1) || (GpuId[1] == 1)) { WARN("This test is not applicable on MI60 & MI100." "Skipping the test!!"); exit(0); @@ -418,7 +418,7 @@ TEST_CASE("Unit_mmap_CoherentTstWthAdvise") { // Read string from child and close reading end. read(fd1[0], GpuId, 2 * sizeof(int)); close(fd1[0]); - if ((GpuId[0] == 1) || (GpuId[0] == 1)) { + if ((GpuId[0] == 1) || (GpuId[1] == 1)) { WARN("This test is not applicable on MI60 & MI100." "Skipping the test!!"); exit(0); From 17105259fea829b07a22d4bc5223580c130f6836 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 10 Mar 2022 16:45:54 +0530 Subject: [PATCH 7/9] SWDEV-321511 - Enable device related disabled tests on Nvidia platform (#2518) Change-Id: I95a62e9562a638cb3525862b9a5d2e640c10882c --- tests/catch/unit/device/hipChooseDevice.cc | 5 +---- tests/catch/unit/device/hipGetDeviceProperties.cc | 5 +---- 2 files changed, 2 insertions(+), 8 deletions(-) diff --git a/tests/catch/unit/device/hipChooseDevice.cc b/tests/catch/unit/device/hipChooseDevice.cc index fa5e5089f7..0cbcc68a4a 100644 --- a/tests/catch/unit/device/hipChooseDevice.cc +++ b/tests/catch/unit/device/hipChooseDevice.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +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 @@ -41,8 +41,6 @@ TEST_CASE("Unit_hipChooseDevice_ValidateDevId") { * Scenario1: Validates if dev = nullptr returns error code * Scenario2: Validates if prop = nullptr returns error code */ -#if HT_AMD -// These test scenarios fail on NVIDIA. TEST_CASE("Unit_hipChooseDevice_NegTst") { hipDeviceProp_t prop; int dev = -1; @@ -57,4 +55,3 @@ TEST_CASE("Unit_hipChooseDevice_NegTst") { REQUIRE_FALSE(hipSuccess == hipChooseDevice(&dev, nullptr)); } } -#endif diff --git a/tests/catch/unit/device/hipGetDeviceProperties.cc b/tests/catch/unit/device/hipGetDeviceProperties.cc index e644f3d397..fe5ba6ffd6 100644 --- a/tests/catch/unit/device/hipGetDeviceProperties.cc +++ b/tests/catch/unit/device/hipGetDeviceProperties.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +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 @@ -182,14 +182,11 @@ TEST_CASE("Unit_hipGetDeviceProperties_ArchPropertiesTst") { TEST_CASE("Unit_hipGetDeviceProperties_NegTst") { hipDeviceProp_t prop; -#if HT_AMD SECTION("props is nullptr") { int device; HIP_CHECK(hipGetDevice(&device)); - // this test case results in segmentation fault on NVCC REQUIRE_FALSE(hipSuccess == hipGetDeviceProperties(nullptr, device)); } -#endif SECTION("device is -1") { REQUIRE_FALSE(hipSuccess == hipGetDeviceProperties(&prop, -1)); From a55d684e78b2e8e2656692c27ca6c1e69a78afd9 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 10 Mar 2022 16:46:23 +0530 Subject: [PATCH 8/9] SWDEV-317716 - Add support for hipDeviceGetUuid (#2528) Change-Id: I488edcd595b42e8119c923d3e7b3643b06c190de --- include/hip/hip_runtime_api.h | 15 +++++++++++++++ 1 file changed, 15 insertions(+) diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index c55e509256..28f035a8b9 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -73,6 +73,9 @@ typedef struct { unsigned hasDynamicParallelism : 1; ///< Dynamic parallelism. } hipDeviceArch_t; +typedef struct hipUUID_t { + char bytes[16]; +} hipUUID; //--- // Common headers for both NVCC and HCC paths: @@ -998,6 +1001,18 @@ hipError_t hipDeviceComputeCapability(int* major, int* minor, hipDevice_t device * @returns #hipSuccess, #hipErrorInavlidDevice */ hipError_t hipDeviceGetName(char* name, int len, hipDevice_t device); +/** + * @brief Returns an UUID for the device.[BETA] + * @param [out] uuid + * @param [in] device + * + * @beta This API is marked as beta, meaning, while this is feature complete, + * it is still open to changes and may have outstanding issues. + * + * @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue, #hipErrorNotInitialized, + * #hipErrorDeInitialized + */ +hipError_t hipDeviceGetUuid(hipUUID* uuid, hipDevice_t device); /** * @brief Returns a value for attr of link between two devices * @param [out] value From c52a81e96f2039d606f5788cab16bb736714c7c8 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 10 Mar 2022 16:47:03 +0530 Subject: [PATCH 9/9] SWDEV-318833 - Get and Set attribute for kernel nodes (#2533) Signed-off-by: sdashmiz Change-Id: Ieacd1fd9aa6d6e98a3a5574ba374c201f05d0e51 --- include/hip/hip_runtime_api.h | 49 +++++++++++++++++++++++++++++++++++ 1 file changed, 49 insertions(+) diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 28f035a8b9..b34b4489b3 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -165,6 +165,31 @@ typedef enum hipMemoryType { hipMemoryTypeUnified ///< Not used currently } hipMemoryType; +/** + * @brief hipKernelNodeAttrID + * @enum + * + */ +typedef enum hipKernelNodeAttrID { + hipKernelNodeAttributeAccessPolicyWindow = 1, + hipKernelNodeAttributeCooperative = 2, +} hipKernelNodeAttrID; +typedef enum hipAccessProperty { + hipAccessPropertyNormal = 0, + hipAccessPropertyStreaming = 1, + hipAccessPropertyPersisting = 2, +} hipAccessProperty; +typedef struct hipAccessPolicyWindow { + void* base_ptr; + hipAccessProperty hitProp; + float hitRatio; + hipAccessProperty missProp; + size_t num_bytes; +} hipAccessPolicyWindow; +typedef union hipKernelNodeAttrValue { + hipAccessPolicyWindow accessPolicyWindow; + int cooperative; +} hipKernelNodeAttrValue; /** * Pointer attributes @@ -4887,6 +4912,30 @@ hipError_t hipGraphMemcpyNodeGetParams(hipGraphNode_t node, hipMemcpy3DParms* pN */ hipError_t hipGraphMemcpyNodeSetParams(hipGraphNode_t node, const hipMemcpy3DParms* pNodeParams); +/** + * @brief Sets a node attribute. + * + * @param [in] hNode - instance of the node to set parameters to. + * @param [in] attr - the attribute node is set to. + * @param [in] value - const pointer to the parameters. + * @returns #hipSuccess, #hipErrorInvalidValue + * @warning : This API is marked as beta, meaning, while this is feature complete, + * it is still open to changes and may have outstanding issues. + */ +hipError_t hipGraphKernelNodeSetAttribute(hipGraphNode_t hNode, hipKernelNodeAttrID attr, + const hipKernelNodeAttrValue* value); +/** + * @brief Gets a node attribute. + * + * @param [in] hNode - instance of the node to set parameters to. + * @param [in] attr - the attribute node is set to. + * @param [in] value - const pointer to the parameters. + * @returns #hipSuccess, #hipErrorInvalidValue + * @warning : This API is marked as beta, meaning, while this is feature complete, + * it is still open to changes and may have outstanding issues. + */ +hipError_t hipGraphKernelNodeGetAttribute(hipGraphNode_t hNode, hipKernelNodeAttrID attr, + hipKernelNodeAttrValue* value); /** * @brief Sets the parameters for a memcpy node in the given graphExec. *