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"; diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index c55e509256..b34b4489b3 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: @@ -162,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 @@ -998,6 +1026,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 @@ -4872,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. * 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/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); 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)); 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"); + } } /* 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; 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 */ 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();