From fc106647e38237a66c00c822336b2d4e81b37f6f Mon Sep 17 00:00:00 2001 From: sumanthtg <90063301+sumanthtg@users.noreply.github.com> Date: Fri, 24 Sep 2021 16:27:17 +0530 Subject: [PATCH 1/9] SWDEV-303262 - Catch2 test fix for -Wunused-but-set-parameter compilation error. (#2357) Change-Id: I102d80763103fa792ab1282745bd89f8d24454d2 [ROCm/hip commit: e2dd25333a762f7a5fb6acacff5e60ab0a26b5a6] --- projects/hip/tests/catch/unit/stream/hipStreamWithCUMask.cc | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/projects/hip/tests/catch/unit/stream/hipStreamWithCUMask.cc b/projects/hip/tests/catch/unit/stream/hipStreamWithCUMask.cc index f78ef9ee2c..0dd855702d 100644 --- a/projects/hip/tests/catch/unit/stream/hipStreamWithCUMask.cc +++ b/projects/hip/tests/catch/unit/stream/hipStreamWithCUMask.cc @@ -95,10 +95,9 @@ void createDisabledCUMask(std::vector *pdisabledCUMask, } } -void Callback(hipStream_t stream, hipError_t status, +void Callback(hipStream_t, hipError_t status, void* userData) { isPassed = true; - stream = 0; HIP_CHECK(status); REQUIRE(userData == nullptr); for (size_t i = 0; i < N; i++) { From 2788cb12430657cf5d062a82c31cce55cabffb23 Mon Sep 17 00:00:00 2001 From: sumanthtg <90063301+sumanthtg@users.noreply.github.com> Date: Fri, 24 Sep 2021 16:27:39 +0530 Subject: [PATCH 2/9] SWDEV-238517 - [dtest] New Catch2 Tests for texture object apis. (#2358) Seperate Catch2 files are created for each resource type. Future test additions of respective resource type will be added to resource type specific files. Change-Id: I7ba09f5bd31605bb670ec70ed7f12904e34cf40d [ROCm/hip commit: abe851ad757bcb518951cf75b568ef61f80df4ff] --- .../tests/catch/hipTestMain/CMakeLists.txt | 21 ++ projects/hip/tests/catch/unit/CMakeLists.txt | 21 ++ .../tests/catch/unit/texture/CMakeLists.txt | 33 +++ .../hipCreateTextureObject_ArgValidation.cc | 81 +++++++ .../texture/hipCreateTextureObject_Array.cc | 67 ++++++ .../texture/hipCreateTextureObject_Linear.cc | 133 +++++++++++ .../texture/hipCreateTextureObject_Pitch2D.cc | 216 ++++++++++++++++++ 7 files changed, 572 insertions(+) create mode 100644 projects/hip/tests/catch/unit/texture/CMakeLists.txt create mode 100644 projects/hip/tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc create mode 100644 projects/hip/tests/catch/unit/texture/hipCreateTextureObject_Array.cc create mode 100644 projects/hip/tests/catch/unit/texture/hipCreateTextureObject_Linear.cc create mode 100644 projects/hip/tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc diff --git a/projects/hip/tests/catch/hipTestMain/CMakeLists.txt b/projects/hip/tests/catch/hipTestMain/CMakeLists.txt index a70169efb0..9b7457cd62 100644 --- a/projects/hip/tests/catch/hipTestMain/CMakeLists.txt +++ b/projects/hip/tests/catch/hipTestMain/CMakeLists.txt @@ -1,3 +1,23 @@ +# Copyright (c) 2021 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. + if(CMAKE_BUILD_TYPE MATCHES "^Debug$") add_definitions(-DHT_LOG_ENABLE) endif() @@ -16,6 +36,7 @@ target_link_libraries(UnitTests PRIVATE UnitDeviceTests OccupancyTest DeviceTest RTC + TextureTest stdc++fs) if(HIP_PLATFORM MATCHES "nvidia") diff --git a/projects/hip/tests/catch/unit/CMakeLists.txt b/projects/hip/tests/catch/unit/CMakeLists.txt index f0ae8b6b93..3ab31d0957 100644 --- a/projects/hip/tests/catch/unit/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/CMakeLists.txt @@ -1,3 +1,23 @@ +# Copyright (c) 2021 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. + add_subdirectory(memory) add_subdirectory(deviceLib) add_subdirectory(stream) @@ -5,3 +25,4 @@ add_subdirectory(event) add_subdirectory(occupancy) add_subdirectory(device) add_subdirectory(rtc) +add_subdirectory(texture) \ No newline at end of file diff --git a/projects/hip/tests/catch/unit/texture/CMakeLists.txt b/projects/hip/tests/catch/unit/texture/CMakeLists.txt new file mode 100644 index 0000000000..f91ae84fd3 --- /dev/null +++ b/projects/hip/tests/catch/unit/texture/CMakeLists.txt @@ -0,0 +1,33 @@ +# Copyright (c) 2021 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. + +# Common Tests - Test independent of all platforms +set(TEST_SRC + hipCreateTextureObject_ArgValidation.cc + hipCreateTextureObject_Linear.cc + hipCreateTextureObject_Pitch2D.cc + hipCreateTextureObject_Array.cc +) + +# Create shared lib of all tests +add_library(TextureTest SHARED EXCLUDE_FROM_ALL ${TEST_SRC}) + +# Add dependency on build_tests to build it on this custom target +add_dependencies(build_tests TextureTest) diff --git a/projects/hip/tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc b/projects/hip/tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc new file mode 100644 index 0000000000..91a0827127 --- /dev/null +++ b/projects/hip/tests/catch/unit/texture/hipCreateTextureObject_ArgValidation.cc @@ -0,0 +1,81 @@ +/* +Copyright (c) 2021 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 + +#define N 512 + +/* + * Validate argument list of texture object api. + */ +TEST_CASE("Unit_hipCreateTextureObject_ArgValidation") { + float *texBuf; + hipError_t ret; + constexpr int xsize = 32; + hipResourceDesc resDesc; + hipTextureDesc texDesc; + hipTextureObject_t texObj; + + /** Initialization */ + HIP_CHECK(hipMalloc(&texBuf, N * sizeof(float))); + // Populate resource descriptor + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = hipResourceTypeLinear; + resDesc.res.linear.devPtr = texBuf; + resDesc.res.linear.desc = hipCreateChannelDesc(xsize, 0, 0, 0, + hipChannelFormatKindFloat); + resDesc.res.linear.sizeInBytes = N * sizeof(float); + + // Populate texture descriptor + memset(&texDesc, 0, sizeof(texDesc)); + texDesc.readMode = hipReadModeElementType; + + + /** Sections */ + SECTION("TextureObject as nullptr") { + ret = hipCreateTextureObject(nullptr, &resDesc, &texDesc, nullptr); + REQUIRE(ret != hipSuccess); + } + + SECTION("Resouce Descriptor as nullptr") { + ret = hipCreateTextureObject(&texObj, nullptr, &texDesc, nullptr); + REQUIRE(ret != hipSuccess); + } + + SECTION("Texture Descriptor as nullptr") { + if ((TestContext::get()).isAmd()) { + ret = hipCreateTextureObject(&texObj, &resDesc, nullptr, nullptr); + REQUIRE(ret != hipSuccess); + } else { + // API expected to return failure. Test skipped + // on nvidia as api returns success and would lead + // to unexpected behavior with app. + WARN("Texture Desc(nullptr) skipped on nvidia"); + } + } + + SECTION("Destroy TextureObject with nullptr") { + ret = hipDestroyTextureObject((hipTextureObject_t)nullptr); + // api to return success and no crash seen. + REQUIRE(ret == hipSuccess); + } + + /** De-Initialization */ + HIP_CHECK(hipFree(texBuf)); +} diff --git a/projects/hip/tests/catch/unit/texture/hipCreateTextureObject_Array.cc b/projects/hip/tests/catch/unit/texture/hipCreateTextureObject_Array.cc new file mode 100644 index 0000000000..037efea1af --- /dev/null +++ b/projects/hip/tests/catch/unit/texture/hipCreateTextureObject_Array.cc @@ -0,0 +1,67 @@ +/* +Copyright (c) 2021 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 + +/* + * Validates Array Resource texture object with negative/functional tests. + */ +TEST_CASE("Unit_hipCreateTextureObject_ArrayResource") { + hipError_t ret; + hipResourceDesc resDesc; + hipTextureDesc texDesc; + hipTextureObject_t texObj; + + /* set resource type as hipResourceTypeArray and array(nullptr) */ + // Populate resource descriptor + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = hipResourceTypeArray; + resDesc.res.array.array = nullptr; + + // Populate texture descriptor + memset(&texDesc, 0, sizeof(texDesc)); + texDesc.readMode = hipReadModeElementType; + + ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr); + REQUIRE(ret != hipSuccess); +} + +/* + * Validates MipMappedArray Resource texture object + * with negative/functional tests. + */ +TEST_CASE("Unit_hipCreateTextureObject_MmArrayResource") { + hipError_t ret; + hipResourceDesc resDesc; + hipTextureDesc texDesc; + hipTextureObject_t texObj; + + /* set resource type as hipResourceTypeMipmappedArray and mipmap(nullptr) */ + // Populate resource descriptor + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = hipResourceTypeMipmappedArray; + resDesc.res.mipmap.mipmap = nullptr; + + // Populate texture descriptor + memset(&texDesc, 0, sizeof(texDesc)); + texDesc.readMode = hipReadModeElementType; + + ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr); + REQUIRE(ret != hipSuccess); +} diff --git a/projects/hip/tests/catch/unit/texture/hipCreateTextureObject_Linear.cc b/projects/hip/tests/catch/unit/texture/hipCreateTextureObject_Linear.cc new file mode 100644 index 0000000000..29bd35a821 --- /dev/null +++ b/projects/hip/tests/catch/unit/texture/hipCreateTextureObject_Linear.cc @@ -0,0 +1,133 @@ +/* +Copyright (c) 2021 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 + +#define UNALIGN_OFFSET 1 +#define N 512 + +/* + * Validates Linear Resource texture object with negative/functional tests. + */ +TEST_CASE("Unit_hipCreateTextureObject_LinearResource") { + float *texBuf; + hipError_t ret; + constexpr int xsize = 32; + hipResourceDesc resDesc; + hipTextureDesc texDesc; + hipResourceViewDesc resViewDesc; + hipTextureObject_t texObj; + hipDeviceProp_t devProp; + + /** Initialization */ + HIP_CHECK(hipMalloc(&texBuf, N * sizeof(float))); + HIP_CHECK(hipGetDeviceProperties(&devProp, 0)); + memset(&resDesc, 0, sizeof(resDesc)); + memset(&texDesc, 0, sizeof(texDesc)); + resDesc.resType = hipResourceTypeLinear; + + /** Sections */ + SECTION("hipResourceTypeLinear and devPtr(nullptr)") { + // Populate resource descriptor + resDesc.res.linear.devPtr = nullptr; + resDesc.res.linear.desc = hipCreateChannelDesc(xsize, 0, 0, 0, + hipChannelFormatKindFloat); + resDesc.res.linear.sizeInBytes = N * sizeof(float); + + // Populate texture descriptor + texDesc.readMode = hipReadModeElementType; + ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipResourceTypeLinear and sizeInBytes(0)") { + if ((TestContext::get()).isAmd()) { + // Populate resource descriptor + resDesc.res.linear.devPtr = texBuf; + resDesc.res.linear.desc = hipCreateChannelDesc(xsize, 0, 0, 0, + hipChannelFormatKindFloat); + resDesc.res.linear.sizeInBytes = 0; + + // Populate texture descriptor + texDesc.readMode = hipReadModeElementType; + ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr); + REQUIRE(ret != hipSuccess); + } else { + // API expected to return failure. Test skipped + // on nvidia as api returns success and would lead + // to unexpected behavior with app. + WARN("Resource type Linear/sizeInBytes(0) skipped on nvidia"); + } + } + + SECTION("hipResourceTypeLinear and sizeInBytes(max(size_t))") { + // Populate resource descriptor + resDesc.res.linear.devPtr = texBuf; + resDesc.res.linear.desc = hipCreateChannelDesc(xsize, 0, 0, 0, + hipChannelFormatKindFloat); + resDesc.res.linear.sizeInBytes = std::numeric_limits::max(); + + // Populate texture descriptor + texDesc.readMode = hipReadModeElementType; + ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipResourceTypeLinear and valid resource view descriptor") { +#if HT_AMD + // Populate resource descriptor + resDesc.res.linear.devPtr = texBuf; + resDesc.res.linear.desc = hipCreateChannelDesc(xsize, 0, 0, 0, + hipChannelFormatKindFloat); + resDesc.res.linear.sizeInBytes = N * sizeof(float); + + // Populate texture descriptor + texDesc.readMode = hipReadModeElementType; + + // Populate resourceview descriptor + memset(&resViewDesc, 0, sizeof(resViewDesc)); + resViewDesc.format = hipResViewFormatFloat1; + resViewDesc.width = N * sizeof(float); + ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, &resViewDesc); + REQUIRE(ret != hipSuccess); +#else + // API expected to return error according to cuda documentation. + WARN("Resource view descriptor test skipped on nvidia"); +#endif + } + + SECTION("hipResourceTypeLinear and devicePtr un-aligned") { + if (devProp.textureAlignment > UNALIGN_OFFSET) { + // Populate resource descriptor + resDesc.res.linear.devPtr = reinterpret_cast(texBuf) + + UNALIGN_OFFSET; + resDesc.res.linear.desc = hipCreateChannelDesc(xsize, 0, 0, 0, + hipChannelFormatKindFloat); + resDesc.res.linear.sizeInBytes = N * sizeof(float); + + // Populate texture descriptor + texDesc.readMode = hipReadModeElementType; + ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr); + REQUIRE(ret != hipSuccess); + } + } + + /** De-Initialization */ + HIP_CHECK(hipFree(texBuf)); +} diff --git a/projects/hip/tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc b/projects/hip/tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc new file mode 100644 index 0000000000..89eb936c3c --- /dev/null +++ b/projects/hip/tests/catch/unit/texture/hipCreateTextureObject_Pitch2D.cc @@ -0,0 +1,216 @@ +/* +Copyright (c) 2021 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 + +#define UNALIGN_OFFSET 1 +#define SIZE_H 20 +#define SIZE_W 30 +#define N 512 + + +/* + * Validates Pitch2D Resource texture object with negative and functional tests + */ +TEST_CASE("Unit_hipCreateTextureObject_Pitch2DResource") { + hipError_t ret; + hipResourceDesc resDesc; + hipTextureDesc texDesc; + hipTextureObject_t texObj; + hipDeviceProp_t devProp; + size_t devPitchA; + float *devPtrA; + + /** Initialization */ + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devPtrA), &devPitchA, + SIZE_W*sizeof(float), SIZE_H)); + HIP_CHECK(hipGetDeviceProperties(&devProp, 0)); + memset(&resDesc, 0, sizeof(resDesc)); + memset(&texDesc, 0, sizeof(texDesc)); + resDesc.resType = hipResourceTypePitch2D; + + /** Sections */ + SECTION("hipResourceTypePitch2D and devPtr(nullptr)") { + // Populate resource descriptor + resDesc.res.pitch2D.devPtr = nullptr; + resDesc.res.pitch2D.height = SIZE_H; + resDesc.res.pitch2D.width = SIZE_W; + resDesc.res.pitch2D.pitchInBytes = devPitchA; + resDesc.res.pitch2D.desc = hipCreateChannelDesc(); + + // Populate texture descriptor + texDesc.readMode = hipReadModeElementType; + + ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipResourceTypePitch2D and devicePtr un-aligned") { + if (devProp.textureAlignment > UNALIGN_OFFSET) { + // Populate resource descriptor + resDesc.res.pitch2D.devPtr = reinterpret_cast(devPtrA) + + UNALIGN_OFFSET; + resDesc.res.pitch2D.height = SIZE_H; + resDesc.res.pitch2D.width = SIZE_W; + resDesc.res.pitch2D.pitchInBytes = devPitchA; + resDesc.res.pitch2D.desc = hipCreateChannelDesc(); + + // Populate texture descriptor + memset(&texDesc, 0, sizeof(texDesc)); + texDesc.readMode = hipReadModeElementType; + ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr); + REQUIRE(ret != hipSuccess); + } + } + + SECTION("hipResourceTypePitch2D and pitch(un-aligned)") { + // Populate resource descriptor + resDesc.res.pitch2D.devPtr = devPtrA; + resDesc.res.pitch2D.height = SIZE_H; + resDesc.res.pitch2D.width = SIZE_W; + resDesc.res.pitch2D.pitchInBytes = UNALIGN_OFFSET; + resDesc.res.pitch2D.desc = hipCreateChannelDesc(); + + // Populate texture descriptor + texDesc.readMode = hipReadModeElementType; + + ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipResourceTypePitch2D and height(0)") { + if ((TestContext::get()).isAmd()) { + // Populate resource descriptor + resDesc.res.pitch2D.devPtr = devPtrA; + resDesc.res.pitch2D.height = 0; + resDesc.res.pitch2D.width = SIZE_W; + resDesc.res.pitch2D.pitchInBytes = devPitchA; + resDesc.res.pitch2D.desc = hipCreateChannelDesc(); + + // Populate texture descriptor + texDesc.readMode = hipReadModeElementType; + + ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr); + REQUIRE(ret != hipSuccess); + } else { + // Test expected to return error with height(0). + WARN("Resourcetype Pitch2D/height(0) skipped on nvidia"); + } + } + + SECTION("hipResourceTypePitch2D and height(0)/devptr(nullptr)") { + // Populate resource descriptor + resDesc.res.pitch2D.devPtr = nullptr; + resDesc.res.pitch2D.height = 0; + resDesc.res.pitch2D.width = SIZE_W; + resDesc.res.pitch2D.pitchInBytes = devPitchA; + resDesc.res.pitch2D.desc = hipCreateChannelDesc(); + + // Populate texture descriptor + texDesc.readMode = hipReadModeElementType; + + ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipResourceTypePitch2D and height(max(size_t))") { + // Populate resource descriptor + resDesc.res.pitch2D.devPtr = devPtrA; + resDesc.res.pitch2D.height = std::numeric_limits::max(); + resDesc.res.pitch2D.width = SIZE_W; + resDesc.res.pitch2D.pitchInBytes = devPitchA; + resDesc.res.pitch2D.desc = hipCreateChannelDesc(); + + // Populate texture descriptor + texDesc.readMode = hipReadModeElementType; + + ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipResourceTypePitch2D and width(0)") { + if ((TestContext::get()).isAmd()) { + // Populate resource descriptor + resDesc.resType = hipResourceTypePitch2D; + resDesc.res.pitch2D.devPtr = devPtrA; + resDesc.res.pitch2D.height = SIZE_H; + resDesc.res.pitch2D.width = 0; + resDesc.res.pitch2D.pitchInBytes = devPitchA; + resDesc.res.pitch2D.desc = hipCreateChannelDesc(); + + // Populate texture descriptor + texDesc.readMode = hipReadModeElementType; + + ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr); + REQUIRE(ret != hipSuccess); + } else { + // api expected to return failure when width(0) is passed. + WARN("ResourceType Pitch2D/width(0) skipped on nvidia"); + } + } + + SECTION("hipResourceTypePitch2D and width(0)/devPtr(nullptr)") { + // Populate resource descriptor + resDesc.res.pitch2D.devPtr = nullptr; + resDesc.res.pitch2D.height = SIZE_H; + resDesc.res.pitch2D.width = 0; + resDesc.res.pitch2D.pitchInBytes = devPitchA; + resDesc.res.pitch2D.desc = hipCreateChannelDesc(); + + // Populate texture descriptor + texDesc.readMode = hipReadModeElementType; + + ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipResourceTypePitch2D and width(max(size_t))") { + // Populate resource descriptor + resDesc.res.pitch2D.devPtr = devPtrA; + resDesc.res.pitch2D.height = SIZE_H; + resDesc.res.pitch2D.width = std::numeric_limits::max(); + resDesc.res.pitch2D.pitchInBytes = devPitchA; + resDesc.res.pitch2D.desc = hipCreateChannelDesc(); + + // Populate texture descriptor + texDesc.readMode = hipReadModeElementType; + + ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr); + REQUIRE(ret != hipSuccess); + } + + SECTION("hipResourceTypePitch2D and pitch(max(size_t))") { + // Populate resource descriptor + resDesc.res.pitch2D.devPtr = devPtrA; + resDesc.res.pitch2D.height = SIZE_H; + resDesc.res.pitch2D.width = SIZE_W; + resDesc.res.pitch2D.pitchInBytes = std::numeric_limits::max(); + resDesc.res.pitch2D.desc = hipCreateChannelDesc(); + + // Populate texture descriptor + texDesc.readMode = hipReadModeElementType; + + ret = hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr); + REQUIRE(ret != hipSuccess); + } + + /** De-Initialization */ + HIP_CHECK(hipFree(devPtrA)); +} + From e3c8349e2ea97cc51d96445afe31371210cdc305 Mon Sep 17 00:00:00 2001 From: Julia Jiang <56359287+jujiang-del@users.noreply.github.com> Date: Fri, 24 Sep 2021 06:57:51 -0400 Subject: [PATCH 3/9] SWDEV-286322 - clean up trailing space (#2361) Change-Id: I03c07e67a8d1fa1a874718ffba43eb396c2aa05c [ROCm/hip commit: 44581b4d3ced8612cb515efcd684084a6b11be65] --- .../hipDispatchLatency/hipDispatchLatency.cpp | 22 +++++----- .../2_Cookbook/13_occupancy/occupancy.cpp | 44 +++++++++---------- .../Negative/memory/hipMemcpyFromSymbol.cpp | 6 +-- .../memory/hipMemcpyFromSymbolAsync.cpp | 6 +-- .../src/Negative/memory/hipMemcpyToSymbol.cpp | 6 +-- .../memory/hipMemcpyToSymbolAsync.cpp | 10 ++--- .../tests/src/Negative/memory/hipMemory.cpp | 4 +- .../hip/tests/src/compiler/hipClassKernel.cpp | 16 +++---- .../hip/tests/src/deviceLib/hipBfloat16.cpp | 12 ++--- .../tests/src/deviceLib/hipMathFunctions.cpp | 28 ++++++------ .../hip/tests/src/deviceLib/hipStdComplex.cpp | 8 ++-- .../hip/tests/src/deviceLib/hipTestHalf.cpp | 2 +- .../hip/tests/src/deviceLib/hipTestNew.cpp | 2 +- projects/hip/tests/src/hipEnvVarDriver.cpp | 4 +- 14 files changed, 85 insertions(+), 85 deletions(-) diff --git a/projects/hip/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp b/projects/hip/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp index 1e04cb9c52..5b96f3a199 100644 --- a/projects/hip/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp +++ b/projects/hip/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp @@ -38,9 +38,9 @@ THE SOFTWARE. __global__ void EmptyKernel() { } void print_timing(std::string test, const std::array &results, int batch = 1) { - + float total_us = 0.0f, mean_us = 0.0f, stddev_us = 0.0f; - + // skip warm-up runs auto start_iter = std::next(results.begin(), WARMUP_RUN_COUNT); auto end_iter = results.end(); @@ -48,7 +48,7 @@ void print_timing(std::string test, const std::array &re // mean std::for_each(start_iter, end_iter, [&](const float &run_ms) { total_us += (run_ms * 1000) / batch; - }); + }); mean_us = total_us / TIMING_RUN_COUNT; // stddev @@ -63,18 +63,18 @@ void print_timing(std::string test, const std::array &re printf("\n %s: %.1f us, std: %.1f us\n", test.c_str(), mean_us, stddev_us); } -int main() { +int main() { hipStream_t stream0 = 0; hipDevice_t device; hipDeviceGet(&device, 0); - hipCtx_t context; - hipCtxCreate(&context, 0, device); + hipCtx_t context; + hipCtxCreate(&context, 0, device); hipModule_t module; hipFunction_t function; hipModuleLoad(&module, FILE_NAME); hipModuleGetFunction(&function, module, KERNEL_NAME); void* params = nullptr; - + std::array results; hipEvent_t start, stop; hipEventCreate(&start); @@ -83,7 +83,7 @@ int main() { /************************************************************************************/ /* HIP kernel launch enqueue rate: */ /* Measure time taken to enqueue a kernel on the GPU */ - /************************************************************************************/ + /************************************************************************************/ // Timing hipModuleLaunchKernel for (auto i = 0; i < TOTAL_RUN_COUNT; ++i) { @@ -104,8 +104,8 @@ int main() { print_timing("hipLaunchKernelGGL enqueue rate", results); /***********************************************************************************/ - /* Single dispatch execution latency using HIP events: */ - /* Measures latency to start & finish executing a kernel with GPU-scope visibility */ + /* Single dispatch execution latency using HIP events: */ + /* Measures latency to start & finish executing a kernel with GPU-scope visibility */ /***********************************************************************************/ //Timing around the dispatch @@ -120,7 +120,7 @@ int main() { /*********************************************************************************/ /* Batch dispatch execution latency using HIP events: */ - /* Measures latency to start & finish executing each dispatch in a batch */ + /* Measures latency to start & finish executing each dispatch in a batch */ /*********************************************************************************/ for (auto i = 0; i < TOTAL_RUN_COUNT; ++i) { diff --git a/projects/hip/samples/2_Cookbook/13_occupancy/occupancy.cpp b/projects/hip/samples/2_Cookbook/13_occupancy/occupancy.cpp index 98728f8124..4f51b61c6e 100644 --- a/projects/hip/samples/2_Cookbook/13_occupancy/occupancy.cpp +++ b/projects/hip/samples/2_Cookbook/13_occupancy/occupancy.cpp @@ -29,23 +29,23 @@ THE SOFTWARE. // Device (Kernel) function __global__ void multiply(float* C, float* A, float* B, int N){ - + int tx = blockDim.x*blockIdx.x+threadIdx.x; - + if (tx < N){ C[tx] = A[tx] * B[tx]; } } // CPU implementation void multiplyCPU(float* C, float* A, float* B, int N){ - - for(unsigned int i=0; i eps) { errors++; } } - + if (errors != 0){ printf("\nManual Test FAILED: %d errors\n", errors); errors=0; } else { printf("\nManual Test PASSED!\n"); } - + for (i = 0; i < NUM; i++) { if (std::abs(C1[i] - cpuC[i]) > eps) { errors++; } } - + if (errors != 0){ printf("\n Automatic Test FAILED: %d errors\n", errors); } else { diff --git a/projects/hip/tests/src/Negative/memory/hipMemcpyFromSymbol.cpp b/projects/hip/tests/src/Negative/memory/hipMemcpyFromSymbol.cpp index 524a7b7c7a..6a96db06e6 100644 --- a/projects/hip/tests/src/Negative/memory/hipMemcpyFromSymbol.cpp +++ b/projects/hip/tests/src/Negative/memory/hipMemcpyFromSymbol.cpp @@ -36,11 +36,11 @@ int main(){ e = hipMemcpyFromSymbol(S, HIP_SYMBOL(Sd), SIZE, 0, hipMemcpyDeviceToHost); HIPASSERT(e==hipErrorInvalidSymbol); - + e = hipMemcpyFromSymbol(S, NULL, SIZE, 0, hipMemcpyDeviceToHost); HIPASSERT(e==hipErrorInvalidSymbol); - HIPCHECK(hipFree(Sd)); - + HIPCHECK(hipFree(Sd)); + passed(); } diff --git a/projects/hip/tests/src/Negative/memory/hipMemcpyFromSymbolAsync.cpp b/projects/hip/tests/src/Negative/memory/hipMemcpyFromSymbolAsync.cpp index 81f738dc96..e9b64a16a8 100644 --- a/projects/hip/tests/src/Negative/memory/hipMemcpyFromSymbolAsync.cpp +++ b/projects/hip/tests/src/Negative/memory/hipMemcpyFromSymbolAsync.cpp @@ -39,11 +39,11 @@ int main(){ e = hipMemcpyFromSymbolAsync(S, HIP_SYMBOL(Sd), SIZE, 0, hipMemcpyDeviceToHost, stream); HIPASSERT(e==hipErrorInvalidSymbol); - + e = hipMemcpyFromSymbolAsync(S, NULL, SIZE, 0, hipMemcpyDeviceToHost, stream); HIPASSERT(e==hipErrorInvalidSymbol); - HIPCHECK(hipFree(Sd)); - + HIPCHECK(hipFree(Sd)); + passed(); } diff --git a/projects/hip/tests/src/Negative/memory/hipMemcpyToSymbol.cpp b/projects/hip/tests/src/Negative/memory/hipMemcpyToSymbol.cpp index 0a4797bef7..3a9e5dfdb9 100644 --- a/projects/hip/tests/src/Negative/memory/hipMemcpyToSymbol.cpp +++ b/projects/hip/tests/src/Negative/memory/hipMemcpyToSymbol.cpp @@ -36,11 +36,11 @@ int main(){ e = hipMemcpyToSymbol(HIP_SYMBOL(Sd), S, SIZE, 0, hipMemcpyHostToDevice); HIPASSERT(e==hipErrorInvalidSymbol); - + e = hipMemcpyToSymbol(NULL, S, SIZE, 0, hipMemcpyHostToDevice); HIPASSERT(e==hipErrorInvalidSymbol); - HIPCHECK(hipFree(Sd)); - + HIPCHECK(hipFree(Sd)); + passed(); } diff --git a/projects/hip/tests/src/Negative/memory/hipMemcpyToSymbolAsync.cpp b/projects/hip/tests/src/Negative/memory/hipMemcpyToSymbolAsync.cpp index 3ad207082f..8b7d479dc8 100644 --- a/projects/hip/tests/src/Negative/memory/hipMemcpyToSymbolAsync.cpp +++ b/projects/hip/tests/src/Negative/memory/hipMemcpyToSymbolAsync.cpp @@ -31,7 +31,7 @@ int main(){ void *Sd; hipError_t e; char S[SIZE]="This is not a device symbol"; - + HIPCHECK(hipMalloc(&Sd,SIZE)); hipStream_t stream; @@ -39,11 +39,11 @@ int main(){ e = hipMemcpyToSymbolAsync(HIP_SYMBOL(Sd), S, SIZE, 0, hipMemcpyHostToDevice, stream); HIPASSERT(e==hipErrorInvalidSymbol); - + e = hipMemcpyToSymbolAsync(NULL, S, SIZE, 0, hipMemcpyHostToDevice, stream); HIPASSERT(e==hipErrorInvalidSymbol); - - HIPCHECK(hipFree(Sd)); - + + HIPCHECK(hipFree(Sd)); + passed(); } diff --git a/projects/hip/tests/src/Negative/memory/hipMemory.cpp b/projects/hip/tests/src/Negative/memory/hipMemory.cpp index 5274b94d94..34825d1e5b 100644 --- a/projects/hip/tests/src/Negative/memory/hipMemory.cpp +++ b/projects/hip/tests/src/Negative/memory/hipMemory.cpp @@ -1,4 +1,4 @@ -/* +/* Copyright (c) 2015 - 2021 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 @@ -32,7 +32,7 @@ int main(){ e = hipMemcpy(0, str, SIZE, hipMemcpyHostToDevice); HIPASSERT(e==hipErrorInvalidValue); - + e = hipMemcpy(NULL, str, SIZE, hipMemcpyHostToDevice); HIPASSERT(e==hipErrorInvalidValue); diff --git a/projects/hip/tests/src/compiler/hipClassKernel.cpp b/projects/hip/tests/src/compiler/hipClassKernel.cpp index e91fb97cd7..5e368dd3b2 100644 --- a/projects/hip/tests/src/compiler/hipClassKernel.cpp +++ b/projects/hip/tests/src/compiler/hipClassKernel.cpp @@ -45,7 +45,7 @@ void HipClassTests::TestForOverride(void){ 0, 0, result_ecd); - + HipClassTests::VerifyResult(result_ech,result_ecd); HipClassTests::FreeMem(result_ech,result_ecd); } @@ -70,13 +70,13 @@ void HipClassTests::TestForOverload(void){ 0, 0, result_ecd); - + HipClassTests::VerifyResult(result_ech,result_ecd); HipClassTests::FreeMem(result_ech,result_ecd); } #endif -#ifdef ENABLE_FRIEND_TEST +#ifdef ENABLE_FRIEND_TEST // check for friend __global__ void friendClassKernel(bool* result_ecd){ @@ -106,7 +106,7 @@ void HipClassTests::TestForEmptyClass(void){ 0, 0, result_ecd); - + HipClassTests::VerifyResult(result_ech,result_ecd); HipClassTests::FreeMem(result_ech,result_ecd); } @@ -157,7 +157,7 @@ void HipClassTests::TestForClassSize(void){ 0, 0, result_ecd); - + HipClassTests::VerifyResult(result_ech,result_ecd); HipClassTests::FreeMem(result_ech,result_ecd); } @@ -217,7 +217,7 @@ void HipClassTests::TestForPassByValue(void){ HipClassTests::VerifyResult(result_ech,result_ecd); HipClassTests::FreeMem(result_ech,result_ecd); } - + // check obj created with hipMalloc __global__ void mallocObjKernel(testPassByValue *obj, bool* result_ecd) { @@ -292,7 +292,7 @@ bool* HipClassTests::AllocateHostMemory(void){ } bool* HipClassTests::AllocateDeviceMemory(void){ - bool* result_ecd; + bool* result_ecd; HIPCHECK(hipMalloc(&result_ecd, NBOOL)); HIPCHECK(hipMemset(result_ecd, @@ -351,5 +351,5 @@ int main(){ #ifdef ENABLE_DESTRUCTOR_TEST classTests.TestForConsrtDesrt(); test_passed(TestForConsrtDesrt); -#endif +#endif } diff --git a/projects/hip/tests/src/deviceLib/hipBfloat16.cpp b/projects/hip/tests/src/deviceLib/hipBfloat16.cpp index 42627ea593..55de397b40 100644 --- a/projects/hip/tests/src/deviceLib/hipBfloat16.cpp +++ b/projects/hip/tests/src/deviceLib/hipBfloat16.cpp @@ -55,17 +55,17 @@ __host__ __device__ void testOperations(float &fa, float &fb) { hip_bfloat16 bf_a(fa); hip_bfloat16 bf_b(fb); float fc = float(bf_a); - float fd = float(bf_b); + float fd = float(bf_b); assert(testRelativeAccuracy(fa, bf_a)); assert(testRelativeAccuracy(fb, bf_b)); assert(testRelativeAccuracy(fc + fd, bf_a + bf_b)); - //when checked as above for add, operation sub fails on GPU + //when checked as above for add, operation sub fails on GPU assert(hip_bfloat16(fc - fd) == (bf_a - bf_b)); assert(testRelativeAccuracy(fc * fd, bf_a * bf_b)); assert(testRelativeAccuracy(fc / fd, bf_a / bf_b)); - + hip_bfloat16 bf_opNegate = -bf_a; assert(bf_opNegate == -bf_a); @@ -75,7 +75,7 @@ __host__ __device__ void testOperations(float &fa, float &fb) { bf_x--; ++bf_x; --bf_x; - //hip_bfloat16 is converted to float and then inc/decremented, hence check with reduced precision + //hip_bfloat16 is converted to float and then inc/decremented, hence check with reduced precision assert(testRelativeAccuracy(bf_x,bf_a)); bf_x = bf_a; @@ -95,7 +95,7 @@ __host__ __device__ void testOperations(float &fa, float &fb) { if (isnan(bf_rounded)) { assert(isnan(bf_rounded) || isinf(bf_rounded)); } -} +} __global__ void testOperationsGPU(float* d_a, float* d_b) { @@ -126,7 +126,7 @@ int main(){ hipLaunchKernelGGL(testOperationsGPU, 1, SIZE, 0, 0, d_fa, d_fb); hipDeviceSynchronize(); - cout<<"Device bfloat16 Operations Successful!!"< __global__ void kernel_simple(F f, T *out) { @@ -191,7 +191,7 @@ int main(int argc, char* argv[]) { check_abs_int64(); // check_lgamma_double(); - + test_fp16(); test_pown(); diff --git a/projects/hip/tests/src/deviceLib/hipStdComplex.cpp b/projects/hip/tests/src/deviceLib/hipStdComplex.cpp index 55a3ef264e..4870a77a5b 100644 --- a/projects/hip/tests/src/deviceLib/hipStdComplex.cpp +++ b/projects/hip/tests/src/deviceLib/hipStdComplex.cpp @@ -82,7 +82,7 @@ __device__ __host__ complex calc(complex A, return A * B; case CK_div: return A / B; - + ONE_ARG(abs) ONE_ARG(arg) ONE_ARG(sin) @@ -111,7 +111,7 @@ void test() { hipMalloc((void**)&Ad, sizeof(ComplexT)*LEN); hipMalloc((void**)&Bd, sizeof(ComplexT)*LEN); hipMalloc((void**)&Cd, sizeof(ComplexT)*LEN); - + for (uint32_t i = 0; i < LEN; i++) { A[i] = ComplexT((i + 1) * 1.0f, (i + 2) * 1.0f); B[i] = A[i]; @@ -119,7 +119,7 @@ void test() { } hipMemcpy(Ad, A, sizeof(ComplexT)*LEN, hipMemcpyHostToDevice); hipMemcpy(Bd, B, sizeof(ComplexT)*LEN, hipMemcpyHostToDevice); - + // Run kernel for a calculation kind and verify by comparing with host // calculation result. Returns false if fails. auto test_fun = [&](enum CalcKind CK) { @@ -145,7 +145,7 @@ void test() { } return true; }; - + #define OP(x) assert(test_fun(CK_##x)); ALL_FUN #undef OP diff --git a/projects/hip/tests/src/deviceLib/hipTestHalf.cpp b/projects/hip/tests/src/deviceLib/hipTestHalf.cpp index 2e027bf491..a418080550 100644 --- a/projects/hip/tests/src/deviceLib/hipTestHalf.cpp +++ b/projects/hip/tests/src/deviceLib/hipTestHalf.cpp @@ -84,7 +84,7 @@ void kernel_hisnan(__half* input, int* output) { } __global__ -void kernel_hisinf(__half* input, int* output) { +void kernel_hisinf(__half* input, int* output) { int tx = threadIdx.x; output[tx] = __hisinf(input[tx]); } diff --git a/projects/hip/tests/src/deviceLib/hipTestNew.cpp b/projects/hip/tests/src/deviceLib/hipTestNew.cpp index a2e7732348..6d5e67e5fe 100644 --- a/projects/hip/tests/src/deviceLib/hipTestNew.cpp +++ b/projects/hip/tests/src/deviceLib/hipTestNew.cpp @@ -41,7 +41,7 @@ THE SOFTWARE. private: int a; }; - + static __global__ void kernel(int* Ad) { int tid = threadIdx.x + blockIdx.x * blockDim.x; new(Ad+tid) A(); diff --git a/projects/hip/tests/src/hipEnvVarDriver.cpp b/projects/hip/tests/src/hipEnvVarDriver.cpp index 5ec6ead939..2d303e08fe 100644 --- a/projects/hip/tests/src/hipEnvVarDriver.cpp +++ b/projects/hip/tests/src/hipEnvVarDriver.cpp @@ -41,7 +41,7 @@ int readHipEnvVar(string flags, char* buff){ std::cout << "\nFinding hipEnvVar in " << directed_dir << "...\n"; FILE* directed_in = popen((directed_dir + flags).c_str(), "r"); - + if(fgets(buff, 512, directed_in) == NULL){ std::cout << "Finding hipEnvVar in " << dir << "...\n"; FILE* in = popen((dir + flags).c_str(), "r"); @@ -74,7 +74,7 @@ int getDeviceNumber(bool print_err=true) { } // Query the current device ID remotely to hipEnvVar -void getDevicePCIBusNumRemote(int deviceID, char* pciBusID) { +void getDevicePCIBusNumRemote(int deviceID, char* pciBusID) { std::this_thread::sleep_for(std::chrono::milliseconds(10)); if (readHipEnvVar((" -d " + std::to_string(deviceID)), pciBusID)){ std::cerr << "The system cannot find hipEnvVar\n"; From 3d53651d310dfdb171d1fa61bb4113a279173311 Mon Sep 17 00:00:00 2001 From: Julia Jiang <56359287+jujiang-del@users.noreply.github.com> Date: Fri, 24 Sep 2021 06:58:03 -0400 Subject: [PATCH 4/9] SWDEV-286322 - remove unnecessary executable permission (#2362) Change-Id: Iebcac3f717ed916554d57a4b839d27b0427efd84 [ROCm/hip commit: c61bbcce358934ac68ce29d1ccebe5aa51e7c096] --- projects/hip/tests/src/ipc/hipMultiProcIpcEvent.cpp | 0 projects/hip/tests/src/ipc/hipMultiProcIpcMem.cpp | 0 projects/hip/tests/src/kernel/hipLaunchParmFunctor.cpp | 0 projects/hip/tests/src/printf/hipPrintfManyDevices.cpp | 0 .../runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaBaseType.cpp | 0 .../runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaPublicApi.cpp | 0 .../src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp | 0 .../tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockType.cpp | 0 .../cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cpp | 0 .../cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cpp | 0 projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoD.cpp | 0 projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp | 0 projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp | 0 projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp | 0 projects/hip/tests/src/runtimeApi/memory/hipMemcpyWithStream.cpp | 0 .../hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp | 0 projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp | 0 17 files changed, 0 insertions(+), 0 deletions(-) mode change 100755 => 100644 projects/hip/tests/src/ipc/hipMultiProcIpcEvent.cpp mode change 100755 => 100644 projects/hip/tests/src/ipc/hipMultiProcIpcMem.cpp mode change 100755 => 100644 projects/hip/tests/src/kernel/hipLaunchParmFunctor.cpp mode change 100755 => 100644 projects/hip/tests/src/printf/hipPrintfManyDevices.cpp mode change 100755 => 100644 projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaBaseType.cpp mode change 100755 => 100644 projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaPublicApi.cpp mode change 100755 => 100644 projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp mode change 100755 => 100644 projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockType.cpp mode change 100755 => 100644 projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cpp mode change 100755 => 100644 projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cpp mode change 100755 => 100644 projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoD.cpp mode change 100755 => 100644 projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp mode change 100755 => 100644 projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp mode change 100755 => 100644 projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp mode change 100755 => 100644 projects/hip/tests/src/runtimeApi/memory/hipMemcpyWithStream.cpp mode change 100755 => 100644 projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp mode change 100755 => 100644 projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp diff --git a/projects/hip/tests/src/ipc/hipMultiProcIpcEvent.cpp b/projects/hip/tests/src/ipc/hipMultiProcIpcEvent.cpp old mode 100755 new mode 100644 diff --git a/projects/hip/tests/src/ipc/hipMultiProcIpcMem.cpp b/projects/hip/tests/src/ipc/hipMultiProcIpcMem.cpp old mode 100755 new mode 100644 diff --git a/projects/hip/tests/src/kernel/hipLaunchParmFunctor.cpp b/projects/hip/tests/src/kernel/hipLaunchParmFunctor.cpp old mode 100755 new mode 100644 diff --git a/projects/hip/tests/src/printf/hipPrintfManyDevices.cpp b/projects/hip/tests/src/printf/hipPrintfManyDevices.cpp old mode 100755 new mode 100644 diff --git a/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaBaseType.cpp b/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaBaseType.cpp old mode 100755 new mode 100644 diff --git a/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaPublicApi.cpp b/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaPublicApi.cpp old mode 100755 new mode 100644 diff --git a/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp b/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp old mode 100755 new mode 100644 diff --git a/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockType.cpp b/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockType.cpp old mode 100755 new mode 100644 diff --git a/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cpp b/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cpp old mode 100755 new mode 100644 diff --git a/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cpp b/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cpp old mode 100755 new mode 100644 diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoD.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoD.cpp old mode 100755 new mode 100644 diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp old mode 100755 new mode 100644 diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp old mode 100755 new mode 100644 diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp old mode 100755 new mode 100644 diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyWithStream.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyWithStream.cpp old mode 100755 new mode 100644 diff --git a/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp b/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp old mode 100755 new mode 100644 diff --git a/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp b/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp old mode 100755 new mode 100644 From 86a0af8a0437d8d64e669e67145b51c0c187286d Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary <51944368+cjatin@users.noreply.github.com> Date: Fri, 24 Sep 2021 16:28:25 +0530 Subject: [PATCH 5/9] SWDEV-294226 - Add -x c++ if openmp targets are mentioned (#2363) Change-Id: Ie15d63eebe2b0b86487ac59845b4fc7f5e5a9e86 [ROCm/hip commit: 63dfe6e8a1271f555b6bca1a1b211a74580b5f7f] --- projects/hip/bin/hipcc | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index c72badcf69..840aaafd41 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -229,6 +229,7 @@ my $needCXXFLAGS = 0; # need to add CXX flags to compile step my $needCFLAGS = 0; # need to add C flags to compile step my $needLDFLAGS = 1; # need to add LDFLAGS to compile step. my $fileTypeFlag = 0; # to see if -x flag is mentioned +my $hasOMPTargets = 0; # If OMP targets is mentioned my $hasC = 0; # options contain a c-style file my $hasCXX = 0; # options contain a cpp-style file (NVCC must force recognition as GPU file) my $hasCU = 0; # options contain a cu-style file (HCC must force recognition as GPU file) @@ -517,6 +518,8 @@ foreach $arg (@ARGV) $hasC = 0; $hasCXX = 0; $hasHIP = 1; + } elsif ($arg =~ '-fopenmp-targets=') { + $hasOMPTargets = 1; } elsif ($arg =~ m/^-/) { # options start with - if ($arg eq '-fgpu-rdc') { @@ -555,7 +558,7 @@ foreach $arg (@ARGV) $toolArgs .= " -x c"; } elsif (($arg =~ /\.cpp$/) or ($arg =~ /\.cxx$/) or ($arg =~ /\.cc$/) or ($arg =~ /\.C$/)) { $needCXXFLAGS = 1; - if ($HIP_COMPILE_CXX_AS_HIP eq '0' or $HIP_PLATFORM ne "amd") { + if ($HIP_COMPILE_CXX_AS_HIP eq '0' or $HIP_PLATFORM ne "amd" or $hasOMPTargets eq 1) { $hasCXX = 1; } elsif ($HIP_PLATFORM eq "amd") { $hasHIP = 1; From 6b654c0bf345a052b41803ac9cacd1de8731f8f5 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary <51944368+cjatin@users.noreply.github.com> Date: Fri, 24 Sep 2021 16:28:38 +0530 Subject: [PATCH 6/9] SWDEV-297706 - Set HIP_ARCHITECTURE Property to off. Since cmake 3.21 cmake has added this property(https://gitlab.kitware.com/cmake/cmake/-/blob/master/Help/prop_tgt/HIP_ARCHITECTURES.rst) (#2364) Change-Id: I33580ac0171ac3744341fcbf25cc3421a1512166 [ROCm/hip commit: 6e27c48da741ca536b30f2cc611263bb3fa2c798] --- projects/hip/tests/catch/CMakeLists.txt | 33 ++++++++++++++++++++ projects/hip/tests/hit/HIT.cmake | 40 ++++++++++++++++++++++++- 2 files changed, 72 insertions(+), 1 deletion(-) diff --git a/projects/hip/tests/catch/CMakeLists.txt b/projects/hip/tests/catch/CMakeLists.txt index 434dc92b2b..4d17fc037b 100644 --- a/projects/hip/tests/catch/CMakeLists.txt +++ b/projects/hip/tests/catch/CMakeLists.txt @@ -54,6 +54,39 @@ if(POLICY CMP0037) cmake_policy(SET CMP0037 OLD) endif() +# Turn off CMAKE_HIP_ARCHITECTURES Feature if cmake version is 3.21+ +if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.21.0) + set(CMAKE_HIP_ARCHITECTURES OFF) +endif() +message(STATUS "CMAKE HIP ARCHITECTURES: ${CMAKE_HIP_ARCHITECTURES}") + +# Identify the GPU Targets. +# This is done due to limitation of rocm_agent_enumerator +# While building test parallelly, rocm_agent_enumerator can fail and give out an empty target +# That results in hipcc building the test for gfx803 (the default target) +if(NOT DEFINED OFFLOAD_ARCH_STR AND EXISTS "${ROCM_PATH}/bin/rocm_agent_enumerator" + AND HIP_PLATFORM STREQUAL "amd" AND UNIX) + execute_process(COMMAND ${ROCM_PATH}/bin/rocm_agent_enumerator OUTPUT_VARIABLE HIP_GPU_ARCH + RESULT_VARIABLE ROCM_AGENT_ENUM_RESULT) + # Trim out gfx000 + string(REPLACE "gfx000\n" "" HIP_GPU_ARCH ${HIP_GPU_ARCH}) + string(LENGTH ${HIP_GPU_ARCH} HIP_GPU_ARCH_LEN) + + # If string has more gfx target except gfx000 + if(${HIP_GPU_ARCH_LEN} GREATER_EQUAL 1) + string(REGEX REPLACE "\n" ";" HIP_GPU_ARCH_LIST "${HIP_GPU_ARCH}") + set(OFFLOAD_ARCH_STR "") + foreach(_hip_gpu_arch ${HIP_GPU_ARCH_LIST}) + set(OFFLOAD_ARCH_STR " ${OFFLOAD_ARCH_STR} --offload-arch=${_hip_gpu_arch} ") + endforeach() + message(STATUS "Using offload arch string: ${OFFLOAD_ARCH_STR}") + endif() +endif() + +if(DEFINED OFFLOAD_ARCH_STR) + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OFFLOAD_ARCH_STR} ") +endif() + # Use clang as host compiler with nvcc if(HIP_COMPILER MATCHES "nvcc") set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -ccbin clang") diff --git a/projects/hip/tests/hit/HIT.cmake b/projects/hip/tests/hit/HIT.cmake index f5192963ea..f831d5965f 100755 --- a/projects/hip/tests/hit/HIT.cmake +++ b/projects/hip/tests/hit/HIT.cmake @@ -29,12 +29,50 @@ if (NOT ${BUILD_SHARED_LIBS}) endif() message(STATUS "HIP runtime lib type - ${HIP_LIB_TYPE}") message(STATUS "CMAKE_TESTING_TOOL: ${CMAKE_TESTING_TOOL}") + +# Turn off CMAKE_HIP_ARCHITECTURES Feature if cmake version is 3.21+ +if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.21.0) + set(CMAKE_HIP_ARCHITECTURES OFF) +endif() +message(STATUS "CMAKE HIP ARCHITECTURES: ${CMAKE_HIP_ARCHITECTURES}") + +# Identify the GPU Targets. +# This is done due to limitation of rocm_agent_enumerator +# While building test parallelly, rocm_agent_enumerator can fail and give out an empty target +# That results in hipcc building the test for gfx803 (the default target) +if(NOT DEFINED OFFLOAD_ARCH_STR AND EXISTS "${ROCM_PATH}/bin/rocm_agent_enumerator" + AND HIP_PLATFORM STREQUAL "amd" AND UNIX) + execute_process(COMMAND ${ROCM_PATH}/bin/rocm_agent_enumerator OUTPUT_VARIABLE HIP_GPU_ARCH + RESULT_VARIABLE ROCM_AGENT_ENUM_RESULT) + message(STATUS "ROCm Agent Enumurator Result: ${ROCM_AGENT_ENUM_RESULT}") + + # Trim out gfx000 + string(REPLACE "gfx000\n" "" HIP_GPU_ARCH ${HIP_GPU_ARCH}) + string(LENGTH ${HIP_GPU_ARCH} HIP_GPU_ARCH_LEN) + + # If string has more gfx target except gfx000 + if(${HIP_GPU_ARCH_LEN} GREATER_EQUAL 1) + string(REGEX REPLACE "\n" ";" HIP_GPU_ARCH_LIST "${HIP_GPU_ARCH}") + set(OFFLOAD_ARCH_STR "") + foreach(_hip_gpu_arch ${HIP_GPU_ARCH_LIST}) + set(OFFLOAD_ARCH_STR " ${OFFLOAD_ARCH_STR} --offload-arch=${_hip_gpu_arch} ") + endforeach() + message(STATUS "Using offload arch string: ${OFFLOAD_ARCH_STR}") + endif() +else() + message(STATUS "ROCm Agent Enumurator Not Found") +endif() + #------------------------------------------------------------------------------- # Helper macro to parse BUILD instructions macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _clang_options _nvcc_options _link_options _exclude_platforms _exclude_runtime _exclude_compiler _exclude_lib_type _depends _dir) set(${_target}) set(${_sources}) - set(${_hipcc_options}) + if(DEFINED OFFLOAD_ARCH_STR) + set(${_hipcc_options} "${OFFLOAD_ARCH_STR}") + else() + set(${_hipcc_options}) + endif() set(${_clang_options}) set(${_nvcc_options}) set(${_link_options}) From 349a189e7af78f11487cd90e04bf1b591f0f7d7d Mon Sep 17 00:00:00 2001 From: lthakur Date: Fri, 24 Sep 2021 16:29:39 +0530 Subject: [PATCH 7/9] SWDEV-298757 - hipMallocManagedFlagsTst.cc Added tests to test flags of hipMallocManaged() api. (#2366) Change-Id: I4294a4e5c3176c9ece8ed6b35cb83e4d1a3e4773 [ROCm/hip commit: 9ae69c59359b3aa49716cdbac82da9f6971e1ddc] --- .../tests/catch/unit/memory/CMakeLists.txt | 2 + .../unit/memory/hipMallocManagedFlagsTst.cc | 340 ++++++++++++++++++ 2 files changed, 342 insertions(+) create mode 100644 projects/hip/tests/catch/unit/memory/hipMallocManagedFlagsTst.cc diff --git a/projects/hip/tests/catch/unit/memory/CMakeLists.txt b/projects/hip/tests/catch/unit/memory/CMakeLists.txt index a4619410b6..dea1415c8f 100644 --- a/projects/hip/tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/memory/CMakeLists.txt @@ -44,6 +44,7 @@ set(TEST_SRC hipMemset3DFunctional.cc hipMemset3DNegative.cc hipMemset3DRegressMultiThread.cc + hipMallocManagedFlagsTst.cc ) else() set(TEST_SRC @@ -88,6 +89,7 @@ set(TEST_SRC hipMemset3DFunctional.cc hipMemset3DNegative.cc hipMemset3DRegressMultiThread.cc + hipMallocManagedFlagsTst.cc ) endif() # Create shared lib of all tests diff --git a/projects/hip/tests/catch/unit/memory/hipMallocManagedFlagsTst.cc b/projects/hip/tests/catch/unit/memory/hipMallocManagedFlagsTst.cc new file mode 100644 index 0000000000..e0d13b764e --- /dev/null +++ b/projects/hip/tests/catch/unit/memory/hipMallocManagedFlagsTst.cc @@ -0,0 +1,340 @@ +/* +Copyright (c) 2021 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 + +// Kernel function +__global__ void MallcMangdFlgTst(int n, float *x, float *y) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + int stride = blockDim.x * gridDim.x; + for (int i = index; i < n; i += stride) + y[i] = x[i] * x[i]; +} + + +// The following function prints info on attributes related to HMM +static int HmmAttrPrint() { + int managed = 0; + INFO("The following are the attribute values related to HMM for" + " device 0:\n"); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributeDirectManagedMemAccessFromHost, 0)); + INFO("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributeConcurrentManagedAccess, 0)); + INFO("hipDeviceAttributeConcurrentManagedAccess: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributePageableMemoryAccess, 0)); + INFO("hipDeviceAttributePageableMemoryAccess: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0)); + INFO("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:" + << managed); + + HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, + 0)); + INFO("hipDeviceAttributeManagedMemory: " << managed); + return managed; +} + +// The following section tests working of hipMallocManaged with flag parameters +TEST_CASE("Unit_hipMallocManaged_FlgParam") { + int managed = HmmAttrPrint(); + if (managed == 1) { + std::atomic DataMismatch{0}; + bool IfTestPassed = true; + float *HmmAG = NULL, *HmmAH1 = NULL, *HmmAH2 = NULL, INIT_VAL = 2.5; + int NumDevs = 0, NUM_ELMS = 4096; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + float *Ad = NULL, *Ah = NULL; + Ah = new float[NUM_ELMS]; + // Testing hipMemAttachGlobal Flag + HIP_CHECK(hipMallocManaged(&HmmAG, NUM_ELMS * sizeof(float), + hipMemAttachGlobal)); + + // Initializing HmmAG memory + for (int i = 0; i < NUM_ELMS; i++) { + HmmAG[i] = INIT_VAL; + Ah[i] = 0; + } + + int blockSize = 256; + int numBlocks = (NUM_ELMS + blockSize - 1) / blockSize; + dim3 dimGrid(numBlocks, 1, 1); + dim3 dimBlock(blockSize, 1, 1); + hipStream_t strm; + for (int i = 0; i < NumDevs; i++) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMalloc(&Ad, NUM_ELMS * sizeof(float))); + HIP_CHECK(hipMemset(Ad, 0, NUM_ELMS * sizeof(float))); + MallcMangdFlgTst<<>>(NUM_ELMS, HmmAG, Ad); + HIP_CHECK(hipStreamSynchronize(strm)); + HIP_CHECK(hipMemcpy(Ah, Ad, NUM_ELMS * sizeof(float), + hipMemcpyDeviceToHost)); + for (int j = 0; j < NUM_ELMS; ++j) { + if (Ah[j] != (INIT_VAL * INIT_VAL)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + WARN("Data Mismatch observed when kernel launched on"); + WARN(" device: " << i); + IfTestPassed = false; + } + DataMismatch = 0; + + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipStreamDestroy(strm)); + } + delete[] Ah; + HIP_CHECK(hipFree(HmmAG)); + + DataMismatch = 0; + HIP_CHECK(hipMallocManaged(&HmmAH1, NUM_ELMS * sizeof(float), + hipMemAttachHost)); + HIP_CHECK(hipMallocManaged(&HmmAH2, NUM_ELMS * sizeof(float), + hipMemAttachHost)); + + // Initializing HmmAH memory + for (int i = 0; i < NUM_ELMS; i++) { + HmmAH1[i] = INIT_VAL; + HmmAH2[i] = 0; + } + for (int i = 0; i < NumDevs; i++) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMemset(HmmAH2, 0, NUM_ELMS * sizeof(float))); + MallcMangdFlgTst<<>>(NUM_ELMS, + HmmAH1, HmmAH2); + HIP_CHECK(hipStreamSynchronize(strm)); + for (int j = 0; j < NUM_ELMS; ++j) { + if (HmmAH2[j] != (INIT_VAL * INIT_VAL)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + WARN("Data Mismatch observed when kernel launched on"); + WARN(" device: " << i); + IfTestPassed = false; + } + HIP_CHECK(hipStreamDestroy(strm)); + } + HIP_CHECK(hipFree(HmmAH1)); + HIP_CHECK(hipFree(HmmAH2)); + REQUIRE(IfTestPassed); + } else { + SUCCEED("Gpu doesnt support HMM! Hence skipping the test with PASS result"); + } +} + +// The following function tests Memory access allocated using hipMallocManaged +// in multiple streams +TEST_CASE("Unit_hipMallocManaged_AccessMultiStream") { + int managed = HmmAttrPrint(); + if (managed == 1) { + std::atomic DataMismatch{0}; + bool IfTestPassed = true; + float *HmmAG = NULL, *HmmAH1 = NULL, *HmmAH2 = NULL, INIT_VAL = 2.5; + int NumStrms = 0, MultiDevice = 0, NUM_ELMS = 4096; + HIP_CHECK(hipGetDeviceCount(&MultiDevice)); + if (MultiDevice >= 2) { + HIP_CHECK(hipGetDeviceCount(&NumStrms)); + } else { + NumStrms = 4; + } + hipStream_t **Stream = new hipStream_t*[NumStrms]; + for (int i = 0; i < NumStrms; ++i) { + Stream[i] = reinterpret_cast(malloc(sizeof(hipStream_t))); + } + float *Ad = NULL, *Ah = NULL; + Ah = new float[NUM_ELMS]; + for (int i = 0; i < NumStrms; ++i) { + if (MultiDevice >= 2) { + HIP_CHECK(hipSetDevice(i)); + } + HIP_CHECK(hipStreamCreate(Stream[i])); + } + HIP_CHECK(hipSetDevice(0)); + // Testing hipMemAttachGlobal Flag + HIP_CHECK(hipMallocManaged(&HmmAG, NUM_ELMS * sizeof(float), + hipMemAttachGlobal)); + + // Initializing HmmAG memory + for (int i = 0; i < NUM_ELMS; i++) { + HmmAG[i] = INIT_VAL; + Ah[i] = 0; + } + + int blockSize = 256; + int numBlocks = (NUM_ELMS + blockSize - 1) / blockSize; + dim3 dimGrid(numBlocks, 1, 1); + dim3 dimBlock(blockSize, 1, 1); + for (int i = 0; i < NumStrms; i++) { + if (MultiDevice >= 2) { + HIP_CHECK(hipSetDevice(i)); + } + HIP_CHECK(hipMalloc(&Ad, NUM_ELMS * sizeof(float))); + HIP_CHECK(hipMemset(Ad, 0, NUM_ELMS * sizeof(float))); + MallcMangdFlgTst<<>>(NUM_ELMS, + HmmAG, Ad); + HIP_CHECK(hipStreamSynchronize(*(Stream[i]))); + // Validating the results + HIP_CHECK(hipMemcpy(Ah, Ad, NUM_ELMS * sizeof(float), + hipMemcpyDeviceToHost)); + for (int j = 0; j < NUM_ELMS; ++j) { + if (Ah[j] != (INIT_VAL * INIT_VAL)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + WARN("Data Mismatch observed when kernel launched on"); + WARN(" device: " << i); + IfTestPassed = false; + } + DataMismatch = 0; + + HIP_CHECK(hipFree(Ad)); + } + delete[] Ah; + HIP_CHECK(hipFree(HmmAG)); + + DataMismatch = 0; + HIP_CHECK(hipMallocManaged(&HmmAH1, NUM_ELMS * sizeof(float), + hipMemAttachHost)); + HIP_CHECK(hipMallocManaged(&HmmAH2, NUM_ELMS * sizeof(float), + hipMemAttachHost)); + + // Initializing HmmAH memory + for (int i = 0; i < NUM_ELMS; i++) { + HmmAH1[i] = INIT_VAL; + HmmAH2[i] = 0; + } + for (int i = 0; i < NumStrms; i++) { + if (MultiDevice >= 2) { + HIP_CHECK(hipSetDevice(i)); + } + HIP_CHECK(hipMemset(HmmAH2, 0, NUM_ELMS * sizeof(float))); + MallcMangdFlgTst<<>>(NUM_ELMS, + HmmAH1, HmmAH2); + HIP_CHECK(hipStreamSynchronize(*(Stream[i]))); + for (int j = 0; j < NUM_ELMS; ++j) { + if (HmmAH2[j] != (INIT_VAL * INIT_VAL)) { + DataMismatch++; + break; + } + } + if (DataMismatch != 0) { + WARN("Data Mismatch observed when kernel launched on"); + WARN(" device: " << i); + IfTestPassed = false; + } + } + + HIP_CHECK(hipFree(HmmAH1)); + HIP_CHECK(hipFree(HmmAH2)); + for (int i = 0; i < NumStrms; ++i) { + HIP_CHECK(hipStreamDestroy(*(Stream[i]))); + } + REQUIRE(IfTestPassed); + } else { + SUCCEED("Gpu doesnt support HMM! Hence skipping the test with PASS result"); + } +} + +TEST_CASE("Unit_hipMallocManaged_ExtremeSizes") { + int managed = HmmAttrPrint(); + if (managed == 1) { + bool IfTestPassed = true; + hipError_t err; + void *Hmm = NULL; + size_t totalDevMem = 0, freeDevMem = 0; + int NumDevs = 0; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + // Testing allocation of extreme and unusual mem values + for (int i = 0; i < NumDevs; i++) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipMemGetInfo(&freeDevMem, &totalDevMem)); + err = hipMallocManaged(&Hmm, 1, hipMemAttachGlobal); + if (hipSuccess == err) { + HIP_CHECK(hipFree(Hmm)); + } else { + WARN("Observed error while allocating memory on GPU: " << i); + WARN(" size 1 with"); + WARN(" hipMallocManaged() api with flag 'hipMemAttachGlobal'\n"); + WARN("Error: " << hipGetErrorString(err)); + IfTestPassed = false; + } + err = hipMallocManaged(&Hmm, freeDevMem, hipMemAttachGlobal); + if (hipSuccess == err) { + HIP_CHECK(hipFree(Hmm)); + } else { + WARN("Observed error while allocating max free memory on GPU: " << i); + WARN(" with hipMallocManaged() api with flag 'hipMemAttachGlobal'\n"); + WARN("Error: " << hipGetErrorString(err)); + IfTestPassed = false; + } + err = hipMallocManaged(&Hmm, (freeDevMem - 1), hipMemAttachGlobal); + if (hipSuccess == err) { + HIP_CHECK(hipFree(Hmm)); + } else { + WARN("Observed error while allocating max (free - 1) memory on "); + WARN("GPU: " << i); + WARN(" using hipMallocManaged() api with flag 'hipMemAttachGlobal'\n"); + WARN("Error: " << hipGetErrorString(err)); + IfTestPassed = false; + } + err = hipMallocManaged(&Hmm, 1, hipMemAttachHost); + if (hipSuccess == err) { + HIP_CHECK(hipFree(Hmm)); + } else { + WARN("Observed error while allocating memory size 1 on GPU: " << i); + WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n"); + WARN("Error: " << hipGetErrorString(err)); + IfTestPassed = false; + } + err = hipMallocManaged(&Hmm, freeDevMem, hipMemAttachHost); + if (hipSuccess == err) { + HIP_CHECK(hipFree(Hmm)); + } else { + WARN("Observed error while allocating max free memory on GPU: " << i); + WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n"); + WARN("Error: " << hipGetErrorString(err)); + IfTestPassed = false; + } + err = hipMallocManaged(&Hmm, (freeDevMem - 1), hipMemAttachHost); + if (hipSuccess == err) { + HIP_CHECK(hipFree(Hmm)); + } else { + WARN("Observed error while allocating max (freeDevMem - 1) memory" + " on GPU: " << i); + WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n"); + WARN("Error: " << hipGetErrorString(err)); + IfTestPassed = false; + } + } + REQUIRE(IfTestPassed); + } else { + SUCCEED("Gpu doesnt support HMM! Hence skipping the test with PASS result"); + } +} From ec17bf5601c4c264c49e82a2f30f4a4763dd6b4f Mon Sep 17 00:00:00 2001 From: ansurya <50609411+ansurya@users.noreply.github.com> Date: Fri, 24 Sep 2021 16:30:15 +0530 Subject: [PATCH 8/9] SWDEV-240806 - Added few more Graph APIs (#2367) Change-Id: Ibae3a640edf45a90517bee9ebc364ac01731418a [ROCm/hip commit: b4e84a333d7c36d0275804a67426b7e6d08b39f4] --- projects/hip/include/hip/hip_runtime_api.h | 1083 ++++++++++++++++---- 1 file changed, 879 insertions(+), 204 deletions(-) diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index d1ba129917..7718a8be4f 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -254,6 +254,7 @@ typedef enum __HIP_NODISCARD hipError_t { hipErrorInvalidHandle = 400, // Deprecated hipErrorInvalidResourceHandle = 400, ///< Resource handle (hipEvent_t or hipStream_t) invalid. + hipErrorIllegalState = 401, ///< Resource required is not in a valid state to perform operation. hipErrorNotFound = 500, hipErrorNotReady = 600, ///< Indicates that asynchronous operations enqueued earlier are not ///< ready. This is not actually an error, but is used to distinguish @@ -302,6 +303,10 @@ typedef enum __HIP_NODISCARD hipError_t { ///< the hipStreamCaptureModeRelaxed argument to ///< hipStreamBeginCapture was passed to ///< hipStreamEndCapture in a different thread. + hipErrorGraphExecUpdateFailure = 910, ///< This error indicates that the graph update + ///< not performed because it included changes which + ///< violated constraintsspecific to instantiated graph + ///< update. hipErrorUnknown = 999, //< Unknown error. // HSA Runtime Error Codes start here. hipErrorRuntimeMemory = 1052, ///< HSA runtime memory call returned error. Typically not seen @@ -4142,6 +4147,14 @@ const char* hipKernelNameRef(const hipFunction_t f); const char* hipKernelNameRefByPtr(const void* hostFunction, hipStream_t stream); int hipGetStreamDeviceId(hipStream_t stream); +/** + *------------------------------------------------------------------------------------------------- + *------------------------------------------------------------------------------------------------- + * @defgroup Graph Management + * @{ + * This section describes the graph management types & functions of HIP runtime API. + */ + /** * An opaque value that represents a hip graph */ @@ -4228,23 +4241,119 @@ typedef enum hipStreamCaptureStatus { hipStreamCaptureStatusInvalidated ///< Stream is part of a capture sequence that has been ///< invalidated, but not terminated } hipStreamCaptureStatus; + +typedef enum hipStreamUpdateCaptureDependenciesFlags { + hipStreamAddCaptureDependencies = 0, ///< Add new nodes to the dependency set + hipStreamSetCaptureDependencies, ///< Replace the dependency set with the new nodes +} hipStreamUpdateCaptureDependenciesFlags; + +/** + * @brief Begins graph capture on a stream. + * + * @param [in] stream - Stream to initiate capture. + * @param [in] mode - Controls the interaction of this capture sequence with other API calls that + * are not safe. + * + * @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 hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode); + +/** + * @brief Ends capture on a stream, returning the captured graph. + * + * @param [in] stream - Stream to end capture. + * @param [out] pGraph - returns the graph captured. + * + * @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 hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph); /** - *------------------------------------------------------------------------------------------------- - *------------------------------------------------------------------------------------------------- - * @defgroup Graph Graph Management - * @{ - * This section describes the graph management functions of HIP runtime API. + * @brief Get capture status of a stream. + * + * @param [in] stream - Stream under capture. + * @param [out] pCaptureStatus - returns current status of the capture. + * @param [out] pId - unique ID of the capture. + * + * @returns #hipSuccess, #hipErrorStreamCaptureImplicit + * + * @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 hipStreamGetCaptureInfo(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus, + unsigned long long* pId); + +/** + * @brief Get stream's capture state + * + * @param [in] stream - Stream under capture. + * @param [out] captureStatus_out - returns current status of the capture. + * @param [out] id_out - unique ID of the capture. + * @param [in] graph_out - returns the graph being captured into. + * @param [out] dependencies_out - returns pointer to an array of nodes. + * @param [out] numDependencies_out - returns size of the array returned in dependencies_out. + * + * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorStreamCaptureImplicit + * + * @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 hipStreamGetCaptureInfo_v2(hipStream_t stream, hipStreamCaptureStatus* captureStatus_out, + unsigned long long* id_out __dparm(0), + hipGraph_t* graph_out __dparm(0), + const hipGraphNode_t** dependencies_out __dparm(0), + size_t* numDependencies_out __dparm(0)); + +/** + * @brief Get stream's capture state + * + * @param [in] stream - Stream under capture. + * @param [out] pCaptureStatus - returns current status of the capture. + * + * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorStreamCaptureImplicit + * + * @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 hipStreamIsCapturing(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus); + +/** + * @brief Update the set of dependencies in a capturing stream + * + * @param [in] stream - Stream under capture. + * @param [in] dependencies - pointer to an array of nodes to Add/Replace. + * @param [in] numDependencies - size of the array in dependencies. + * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorIllegalState + * + * @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 hipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t* dependencies, + size_t numDependencies, + unsigned int flags __dparm(0)); + /** * @brief Creates a graph * * @param [out] pGraph - pointer to graph to create. * @param [in] flags - flags for graph creation, must be 0. * - * @returns #hipSuccess. + * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryAllocation + * + * @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 hipGraphCreate(hipGraph_t* pGraph, unsigned int flags); @@ -4254,206 +4363,14 @@ hipError_t hipGraphCreate(hipGraph_t* pGraph, unsigned int flags); * * @param [in] graph - instance of graph to destroy. * - * @returns #hipSuccess. + * @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 hipGraphDestroy(hipGraph_t graph); -/** - * @brief Destroys an executable graph - * - * @param [in] pGraphExec - instance of executable graph to destry. - * - * @returns #hipSuccess. - * - */ -hipError_t hipGraphExecDestroy(hipGraphExec_t pGraphExec); - -/** - * @brief Creates an executable graph from a graph - * - * @param [out] pGraphExec - pointer to instantiated executable graph to create. - * @param [in] graph - instance of graph to instantiate. - * @param [out] pErrorNode - pointer to error node in case error occured in graph instantiation, - * it could modify the correponding node. - * @param [out] pLogBuffer - pointer to log buffer. - * @param [in] bufferSize - the size of log buffer. - * - * @returns #hipSuccess, #hipErrorOutOfMemory. - * - */ -hipError_t hipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph, - hipGraphNode_t* pErrorNode, char* pLogBuffer, size_t bufferSize); - -// Launches an executable graph in a stream. -/** - * @brief launches an executable graph in a stream - * - * @param [in] graphExec - instance of executable graph to launch. - * @param [in] stream - instance of stream in which to launch executable graph. - * @returns #hipSuccess, #hipErrorOutOfMemory, #hipErrorInvalidHandle, #hipErrorInvalidValue - * - */ -hipError_t hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream); - -/** - * @brief Creates a kernel execution node and adds it to a graph. - * - * @param [out] pGraphNode - pointer to graph node to create. - * @param [in,out] graph - instance of graph to add the created node. - * @param [in] pDependencies - pointer to the dependencies on the kernel execution node. - * @param [in] numDependencies - the number of the dependencies. - * @param [in] pNodeParams - pointer to the parameters to the kernel execution node on the GPU. - * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDeviceFunction - * - */ -hipError_t hipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, - const hipGraphNode_t* pDependencies, size_t numDependencies, - const hipKernelNodeParams* pNodeParams); - - -/** - * @brief Creates a memcpy node and adds it to a graph. - * - * @param [out] pGraphNode - pointer to graph node to create. - * @param [in,out] graph - instance of graph to add the created node. - * @param [in] pDependencies - const pointer to the dependencies on the kernel execution node. - * @param [in] numDependencies - the number of the dependencies. - * @param [in] pCopyParams - const pointer to the parameters for the memory copy. - * @returns #hipSuccess, #hipErrorInvalidValue - * - */ -hipError_t hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, - const hipGraphNode_t* pDependencies, size_t numDependencies, - const hipMemcpy3DParms* pCopyParams); - -/** - * @brief Creates a 1D memcpy node and adds it to a graph. - * - * @param [out] pGraphNode - pointer to graph node to create. - * @param [in,out] graph - instance of the graph to add the created node. - * @param [in] pDependencies - const pointer to the dependencies on the kernel execution node. - * @param [in] numDependencies - the number of the dependencies. - * @param [in] dst - pointer to memory address to the destination. - * @param [in] src - pointer to memory address to the source. - * @param [in] count - the size of the memory to copy. - * @param [in] kind - the type of memory copy. - * @returns #hipSuccess, #hipErrorInvalidValue - * - */ -hipError_t hipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph, - const hipGraphNode_t* pDependencies, size_t numDependencies, - void* dst, const void* src, size_t count, hipMemcpyKind kind); - -/** - * @brief Creates a memset node and adds it to a graph. - * - * @param [out] pGraphNode - pointer to the graph node to create. - * @param [in,out] graph - instance of the graph to add the created node. - * @param [in] pDependencies - const pointer to the dependencies on the kernel execution node. - * @param [in] numDependencies - the number of the dependencies. - * @param [in] pMemsetParams - const pointer to the parameters for the memory set. - * @returns #hipSuccess, #hipErrorInvalidValue - * - */ -hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, - const hipGraphNode_t* pDependencies, size_t numDependencies, - const hipMemsetParams* pMemsetParams); - -/** - * @brief Returns graph nodes. - * - * @param [in] graph - instance of graph to get the nodes. - * @param [out] nodes - pointer to the graph nodes. - * @param [out] numNodes - the number of graph nodes. - * @returns #hipSuccess, #hipErrorInvalidValue - * - */ -hipError_t hipGraphGetNodes(hipGraph_t graph, hipGraphNode_t* nodes, size_t* numNodes); - -/** - * @brief Returns graph's root nodes. - * - * @param [in] graph - instance of the graph to get the nodes. - * @param [out] pRootNodes - pointer to the graph's root nodes. - * @param [out] pNumRootNodes - the number of graph's root nodes. - * @returns #hipSuccess, #hipErrorInvalidValue - * - */ -hipError_t hipGraphGetRootNodes(hipGraph_t graph, hipGraphNode_t* pRootNodes, - size_t* pNumRootNodes); - -/** - * @brief Gets kernel node's parameters. - * - * @param [in] node - instance of the node to get parameters from. - * @param [out] pNodeParams - pointer to the parameters - * @returns #hipSuccess, #hipErrorInvalidValue - * - */ -hipError_t hipGraphKernelNodeGetParams(hipGraphNode_t node, hipKernelNodeParams* pNodeParams); - -/** - * @brief Sets a kernel node's parameters. - * - * @param [in] node - instance of the node to set parameters to. - * @param [in] pNodeParams - const pointer to the parameters. - * @returns #hipSuccess, #hipErrorInvalidValue - * - */ -hipError_t hipGraphKernelNodeSetParams(hipGraphNode_t node, const hipKernelNodeParams* pNodeParams); - -/** - * @brief Gets a memcpy node's parameters. - * - * @param [in] node - instance of the node to get parameters from. - * @param [out] pNodeParams - pointer to the parameters. - * @returns #hipSuccess, #hipErrorInvalidValue - * - */ -hipError_t hipGraphMemcpyNodeGetParams(hipGraphNode_t node, hipMemcpy3DParms* pNodeParams); - -/** - * @brief Sets a memcpy node's parameters. - * - * @param [in] node - instance of the node to set parameters to. - * @param [in] pNodeParams - const pointer to the parameters. - * @returns #hipSuccess, #hipErrorInvalidValue - * - */ -hipError_t hipGraphMemcpyNodeSetParams(hipGraphNode_t node, const hipMemcpy3DParms* pNodeParams); - -/** - * @brief Gets a memset node's parameters. - * - * @param [in] node - instane of the node to get parameters from. - * @param [out] pNodeParams - pointer to the parameters. - * @returns #hipSuccess, #hipErrorInvalidValue - * - */ -hipError_t hipGraphMemsetNodeGetParams(hipGraphNode_t node, hipMemsetParams* pNodeParams); - -/** - * @brief Sets a memset node's parameters. - * - * @param [in] node - instance of the node to set parameters to. - * @param [out] pNodeParams - pointer to the parameters. - * @returns #hipSuccess, #hipErrorInvalidValue - * - */ -hipError_t hipGraphMemsetNodeSetParams(hipGraphNode_t node, const hipMemsetParams* pNodeParams); - -/** - * @brief Sets the parameters for a kernel node in the given graphExec. - * - * @param [in] hGraphExec - instance of the executable graph with the node. - * @param [in] node - instance of the node to set parameters to. - * @param [in] pNodeParams - const pointer to the kernel node parameters. - * @returns #hipSuccess, #hipErrorInvalidValue - * - */ -hipError_t hipGraphExecKernelNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node, - const hipKernelNodeParams* pNodeParams); /** * @brief Adds dependency edges to a graph. * @@ -4463,23 +4380,781 @@ hipError_t hipGraphExecKernelNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNo * @param [in] numDependencies - the number of dependencies to add. * @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 hipGraphAddDependencies(hipGraph_t graph, const hipGraphNode_t* from, const hipGraphNode_t* to, size_t numDependencies); +/** + * @brief Removes dependency edges from a graph. + * + * @param [in] graph - instance of the graph to remove dependencies. + * @param [in] from - Array of nodes that provide the dependencies. + * @param [in] to - Array of dependent nodes. + * @param [in] numDependencies - the number of dependencies to remove. + * @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 hipGraphRemoveDependencies(hipGraph_t graph, const hipGraphNode_t* from, + const hipGraphNode_t* to, size_t numDependencies); + +/** + * @brief Returns a graph's dependency edges. + * + * @param [in] graph - instance of the graph to get the edges from. + * @param [out] from - pointer to the graph nodes to return edge endpoints. + * @param [out] to - pointer to the graph nodes to return edge endpoints. + * @param [out] numEdges - returns number of edges. + * @returns #hipSuccess, #hipErrorInvalidValue + * + * from and to may both be NULL, in which case this function only returns the number of edges in + * numEdges. Otherwise, numEdges entries will be filled in. If numEdges is higher than the actual + * number of edges, the remaining entries in from and to will be set to NULL, and the number of + * edges actually returned will be written to numEdges + * @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 hipGraphGetEdges(hipGraph_t graph, hipGraphNode_t* from, hipGraphNode_t* to, + size_t* numEdges); + +/** + * @brief Returns graph nodes. + * + * @param [in] graph - instance of graph to get the nodes. + * @param [out] nodes - pointer to return the graph nodes. + * @param [out] numNodes - returns number of graph nodes. + * @returns #hipSuccess, #hipErrorInvalidValue + * + * nodes may be NULL, in which case this function will return the number of nodes in numNodes. + * Otherwise, numNodes entries will be filled in. If numNodes is higher than the actual number of + * nodes, the remaining entries in nodes will be set to NULL, and the number of nodes actually + * obtained will be returned in numNodes. + * @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 hipGraphGetNodes(hipGraph_t graph, hipGraphNode_t* nodes, size_t* numNodes); + +/** + * @brief Returns graph's root nodes. + * + * @param [in] graph - instance of the graph to get the nodes. + * @param [out] pRootNodes - pointer to return the graph's root nodes. + * @param [out] pNumRootNodes - returns the number of graph's root nodes. + * @returns #hipSuccess, #hipErrorInvalidValue + * + * pRootNodes may be NULL, in which case this function will return the number of root nodes in + * pNumRootNodes. Otherwise, pNumRootNodes entries will be filled in. If pNumRootNodes is higher + * than the actual number of root nodes, the remaining entries in pRootNodes will be set to NULL, + * and the number of nodes actually obtained will be returned in pNumRootNodes. + * @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 hipGraphGetRootNodes(hipGraph_t graph, hipGraphNode_t* pRootNodes, + size_t* pNumRootNodes); + +/** + * @brief Returns a node's dependencies. + * + * @param [in] node - graph node to get the dependencies from. + * @param [out] pDependencies - pointer to to return the dependencies. + * @param [out] pNumDependencies - returns the number of graph node dependencies. + * @returns #hipSuccess, #hipErrorInvalidValue + * + * pDependencies may be NULL, in which case this function will return the number of dependencies in + * pNumDependencies. Otherwise, pNumDependencies entries will be filled in. If pNumDependencies is + * higher than the actual number of dependencies, the remaining entries in pDependencies will be set + * to NULL, and the number of nodes actually obtained will be returned in pNumDependencies. + * @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 hipGraphNodeGetDependencies(hipGraphNode_t node, hipGraphNode_t* pDependencies, + size_t* pNumDependencies); + +/** + * @brief Returns a node's dependent nodes. + * + * @param [in] node - graph node to get the Dependent nodes from. + * @param [out] pDependentNodes - pointer to return the graph dependent nodes. + * @param [out] pNumDependentNodes - returns the number of graph node dependent nodes. + * @returns #hipSuccess, #hipErrorInvalidValue + * + * DependentNodes may be NULL, in which case this function will return the number of dependent nodes + * in pNumDependentNodes. Otherwise, pNumDependentNodes entries will be filled in. If + * pNumDependentNodes is higher than the actual number of dependent nodes, the remaining entries in + * pDependentNodes will be set to NULL, and the number of nodes actually obtained will be returned + * in pNumDependentNodes. + * @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 hipGraphNodeGetDependentNodes(hipGraphNode_t node, hipGraphNode_t* pDependentNodes, + size_t* pNumDependentNodes); + +/** + * @brief Returns a node's type. + * + * @param [in] node - instance of the graph to add dependencies. + * @param [out] pType - pointer to the return the type + * @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 hipGraphNodeGetType(hipGraphNode_t node, hipGraphNodeType* pType); + +/** + * @brief Remove a node from the graph. + * + * @param [in] node - graph node to remove + * @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 hipGraphDestroyNode(hipGraphNode_t node); + +/** + * @brief Clones a graph. + * + * @param [out] pGraphClone - Returns newly created cloned graph. + * @param [in] originalGraph - original graph to clone from. + * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryAllocation + * + * @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 hipGraphClone(hipGraph_t* pGraphClone, hipGraph_t originalGraph); + +/** + * @brief Finds a cloned version of a node. + * + * @param [out] pNode - Returns the cloned node. + * @param [in] originalNode - original node handle. + * @param [in] clonedGraph - Cloned graph to query. + * @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 hipGraphNodeFindInClone(hipGraphNode_t* pNode, hipGraphNode_t originalNode, + hipGraph_t clonedGraph); + +/** + * @brief Creates an executable graph from a graph + * + * @param [out] pGraphExec - pointer to instantiated executable graph that is created. + * @param [in] graph - instance of graph to instantiate. + * @param [out] pErrorNode - pointer to error node in case error occured in graph instantiation, + * it could modify the correponding node. + * @param [out] pLogBuffer - pointer to log buffer. + * @param [out] bufferSize - the size of log buffer. + * + * @returns #hipSuccess, #hipErrorOutOfMemory + * + * @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 hipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph, + hipGraphNode_t* pErrorNode, char* pLogBuffer, size_t bufferSize); + +/** + * @brief Creates an executable graph from a graph. + * + * @param [out] pGraphExec - pointer to instantiated executable graph that is created. + * @param [in] graph - instance of graph to instantiate. + * @param [in] flags - Flags to control instantiation. + * @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 hipGraphInstantiateWithFlags(hipGraphExec_t* pGraphExec, hipGraph_t graph, + unsigned long long flags); + +/** + * @brief launches an executable graph in a stream + * + * @param [in] graphExec - instance of executable graph to launch. + * @param [in] stream - instance of stream in which to launch executable graph. + * @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 hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream); + +/** + * @brief Destroys an executable graph + * + * @param [in] pGraphExec - instance of executable graph to destry. + * + * @returns #hipSuccess. + * + * @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 hipGraphExecDestroy(hipGraphExec_t graphExec); + +// Check whether an executable graph can be updated with a graph and perform the update if possible. +/** + * @brief Check whether an executable graph can be updated with a graph and perform the update if * + * possible. + * + * @param [in] hGraphExec - instance of executable graph to update. + * @param [in] hGraph - graph that contains the updated parameters. + * @param [in] hErrorNode_out - node which caused the permissibility check to forbid the update. + * @param [in] updateResult_out - Whether the graph update was permitted. + * @returns #hipSuccess, #hipErrorGraphExecUpdateFailure + * + * @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 hipGraphExecUpdate(hipGraphExec_t hGraphExec, hipGraph_t hGraph, + hipGraphNode_t* hErrorNode_out, + hipGraphExecUpdateResult* updateResult_out); + +/** + * @brief Creates a kernel execution node and adds it to a graph. + * + * @param [out] pGraphNode - pointer to graph node to create. + * @param [in] graph - instance of graph to add the created node. + * @param [in] pDependencies - pointer to the dependencies on the kernel execution node. + * @param [in] numDependencies - the number of the dependencies. + * @param [in] pNodeParams - pointer to the parameters to the kernel execution node on the GPU. + * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDeviceFunction + * @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 hipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + const hipKernelNodeParams* pNodeParams); + +/** + * @brief Gets kernel node's parameters. + * + * @param [in] node - instance of the node to get parameters from. + * @param [out] pNodeParams - 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 hipGraphKernelNodeGetParams(hipGraphNode_t node, hipKernelNodeParams* pNodeParams); + +/** + * @brief Sets a kernel node's parameters. + * + * @param [in] node - instance of the node to set parameters to. + * @param [in] pNodeParams - 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 hipGraphKernelNodeSetParams(hipGraphNode_t node, const hipKernelNodeParams* pNodeParams); + +/** + * @brief Sets the parameters for a kernel node in the given graphExec. + * + * @param [in] hGraphExec - instance of the executable graph with the node. + * @param [in] node - instance of the node to set parameters to. + * @param [in] pNodeParams - const pointer to the kernel node 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 hipGraphExecKernelNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node, + const hipKernelNodeParams* pNodeParams); + +/** + * @brief Creates a memcpy node and adds it to a graph. + * + * @param [out] pGraphNode - pointer to graph node to create. + * @param [in] graph - instance of graph to add the created node. + * @param [in] pDependencies - const pointer to the dependencies on the memcpy execution node. + * @param [in] numDependencies - the number of the dependencies. + * @param [in] pCopyParams - const pointer to the parameters for the memory copy. + * @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 hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + const hipMemcpy3DParms* pCopyParams); +/** + * @brief Gets a memcpy node's parameters. + * + * @param [in] node - instance of the node to get parameters from. + * @param [out] pNodeParams - 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 hipGraphMemcpyNodeGetParams(hipGraphNode_t node, hipMemcpy3DParms* pNodeParams); + +/** + * @brief Sets a memcpy node's parameters. + * + * @param [in] node - instance of the node to set parameters to. + * @param [in] pNodeParams - 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 hipGraphMemcpyNodeSetParams(hipGraphNode_t node, const hipMemcpy3DParms* pNodeParams); + +/** + * @brief Sets the parameters for a memcpy node in the given graphExec. + * + * @param [in] hGraphExec - instance of the executable graph with the node. + * @param [in] node - instance of the node to set parameters to. + * @param [in] pNodeParams - const pointer to the kernel node 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 hipGraphExecMemcpyNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node, + hipMemcpy3DParms* pNodeParams); + +/** + * @brief Creates a 1D memcpy node and adds it to a graph. + * + * @param [out] pGraphNode - pointer to graph node to create. + * @param [in] graph - instance of graph to add the created node. + * @param [in] pDependencies - const pointer to the dependencies on the memcpy execution node. + * @param [in] numDependencies - the number of the dependencies. + * @param [in] dst - pointer to memory address to the destination. + * @param [in] src - pointer to memory address to the source. + * @param [in] count - the size of the memory to copy. + * @param [in] kind - the type of memory copy. + * @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 hipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + void* dst, const void* src, size_t count, hipMemcpyKind kind); + +/** + * @brief Sets a memcpy node's parameters to perform a 1-dimensional copy. + * + * @param [in] node - instance of the node to set parameters to. + * @param [in] dst - pointer to memory address to the destination. + * @param [in] src - pointer to memory address to the source. + * @param [in] count - the size of the memory to copy. + * @param [in] kind - the type of memory copy. + * @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 hipGraphMemcpyNodeSetParams1D(hipGraphNode_t node, void* dst, const void* src, + size_t count, hipMemcpyKind kind); + +/** + * @brief Sets the parameters for a memcpy node in the given graphExec to perform a 1-dimensional + * copy. + * + * @param [in] hGraphExec - instance of the executable graph with the node. + * @param [in] node - instance of the node to set parameters to. + * @param [in] dst - pointer to memory address to the destination. + * @param [in] src - pointer to memory address to the source. + * @param [in] count - the size of the memory to copy. + * @param [in] kind - the type of memory copy. + * @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 hipGraphExecMemcpyNodeSetParams1D(hipGraphExec_t hGraphExec, hipGraphNode_t node, + void* dst, const void* src, size_t count, + hipMemcpyKind kind); + +/** + * @brief Creates a memcpy node to copy from a symbol on the device and adds it to a graph. + * + * @param [out] pGraphNode - pointer to graph node to create. + * @param [in] graph - instance of graph to add the created node. + * @param [in] pDependencies - const pointer to the dependencies on the memcpy execution node. + * @param [in] numDependencies - the number of the dependencies. + * @param [in] dst - pointer to memory address to the destination. + * @param [in] symbol - Device symbol address. + * @param [in] count - the size of the memory to copy. + * @param [in] offset - Offset from start of symbol in bytes. + * @param [in] kind - the type of memory copy. + * @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 hipGraphAddMemcpyNodeFromSymbol(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, + size_t numDependencies, void* dst, const void* symbol, + size_t count, size_t offset, hipMemcpyKind kind); + +/** + * @brief Sets a memcpy node's parameters to copy from a symbol on the device. + * + * @param [in] node - instance of the node to set parameters to. + * @param [in] dst - pointer to memory address to the destination. + * @param [in] symbol - Device symbol address. + * @param [in] count - the size of the memory to copy. + * @param [in] offset - Offset from start of symbol in bytes. + * @param [in] kind - the type of memory copy. + * @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 hipGraphMemcpyNodeSetParamsFromSymbol(hipGraphNode_t node, void* dst, const void* symbol, + size_t count, size_t offset, hipMemcpyKind kind); + +/** + * @brief Sets the parameters for a memcpy node in the given graphExec to copy from a symbol on the + * * device. + * + * @param [in] hGraphExec - instance of the executable graph with the node. + * @param [in] node - instance of the node to set parameters to. + * @param [in] dst - pointer to memory address to the destination. + * @param [in] symbol - Device symbol address. + * @param [in] count - the size of the memory to copy. + * @param [in] offset - Offset from start of symbol in bytes. + * @param [in] kind - the type of memory copy. + * @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 hipGraphExecMemcpyNodeSetParamsFromSymbol(hipGraphExec_t hGraphExec, hipGraphNode_t node, + void* dst, const void* symbol, size_t count, + size_t offset, hipMemcpyKind kind); + +/** + * @brief Creates a memcpy node to copy to a symbol on the device and adds it to a graph. + * + * @param [out] pGraphNode - pointer to graph node to create. + * @param [in] graph - instance of graph to add the created node. + * @param [in] pDependencies - const pointer to the dependencies on the memcpy execution node. + * @param [in] numDependencies - the number of the dependencies. + * @param [in] symbol - Device symbol address. + * @param [in] src - pointer to memory address of the src. + * @param [in] count - the size of the memory to copy. + * @param [in] offset - Offset from start of symbol in bytes. + * @param [in] kind - the type of memory copy. + * @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 hipGraphAddMemcpyNodeToSymbol(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, + size_t numDependencies, const void* symbol, + const void* src, size_t count, size_t offset, + hipMemcpyKind kind); + +/** + * @brief Sets a memcpy node's parameters to copy to a symbol on the device. + * + * @param [in] node - instance of the node to set parameters to. + * @param [in] symbol - Device symbol address. + * @param [in] src - pointer to memory address of the src. + * @param [in] count - the size of the memory to copy. + * @param [in] offset - Offset from start of symbol in bytes. + * @param [in] kind - the type of memory copy. + * @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 hipGraphMemcpyNodeSetParamsToSymbol(hipGraphNode_t node, const void* symbol, + const void* src, size_t count, size_t offset, + hipMemcpyKind kind); + + +/** + * @brief Sets the parameters for a memcpy node in the given graphExec to copy to a symbol on the + * device. + * @param [in] hGraphExec - instance of the executable graph with the node. + * @param [in] node - instance of the node to set parameters to. + * @param [in] symbol - Device symbol address. + * @param [in] src - pointer to memory address of the src. + * @param [in] count - the size of the memory to copy. + * @param [in] offset - Offset from start of symbol in bytes. + * @param [in] kind - the type of memory copy. + * @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 hipGraphExecMemcpyNodeSetParamsToSymbol(hipGraphExec_t hGraphExec, hipGraphNode_t node, + const void* symbol, const void* src, + size_t count, size_t offset, hipMemcpyKind kind); + +/** + * @brief Creates a memset node and adds it to a graph. + * + * @param [out] pGraphNode - pointer to the graph node to create. + * @param [in] graph - instance of the graph to add the created node. + * @param [in] pDependencies - const pointer to the dependencies on the memset execution node. + * @param [in] numDependencies - the number of the dependencies. + * @param [in] pMemsetParams - const pointer to the parameters for the memory set. + * @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 hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + const hipMemsetParams* pMemsetParams); + +/** + * @brief Gets a memset node's parameters. + * + * @param [in] node - instane of the node to get parameters from. + * @param [out] pNodeParams - 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 hipGraphMemsetNodeGetParams(hipGraphNode_t node, hipMemsetParams* pNodeParams); + +/** + * @brief Sets a memset node's parameters. + * + * @param [in] node - instance of the node to set parameters to. + * @param [in] pNodeParams - 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 hipGraphMemsetNodeSetParams(hipGraphNode_t node, const hipMemsetParams* pNodeParams); + +/** + * @brief Sets the parameters for a memset node in the given graphExec. + * + * @param [in] hGraphExec - instance of the executable graph with the node. + * @param [in] node - instance of the node to set parameters to. + * @param [in] pNodeParams - 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 hipGraphExecMemsetNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node, + const hipMemsetParams* pNodeParams); + +/** + * @brief Creates a host execution node and adds it to a graph. + * + * @param [out] pGraphNode - pointer to the graph node to create. + * @param [in] graph - instance of the graph to add the created node. + * @param [in] pDependencies - const pointer to the dependencies on the memset execution node. + * @param [in] numDependencies - the number of the dependencies. + * @param [in] pNodeParams -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 hipGraphAddHostNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + const hipHostNodeParams* pNodeParams); + +/** + * @brief Returns a host node's parameters. + * + * @param [in] node - instane of the node to get parameters from. + * @param [out] pNodeParams - 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 hipGraphHostNodeGetParams(hipGraphNode_t node, hipHostNodeParams* pNodeParams); + +/** + * @brief Sets a host node's parameters. + * + * @param [in] node - instance of the node to set parameters to. + * @param [in] pNodeParams - 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 hipGraphHostNodeSetParams(hipGraphNode_t node, const hipHostNodeParams* pNodeParams); + +/** + * @brief Sets the parameters for a host node in the given graphExec. + * + * @param [in] hGraphExec - instance of the executable graph with the node. + * @param [in] node - instance of the node to set parameters to. + * @param [in] pNodeParams - 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 hipGraphExecHostNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node, + const hipHostNodeParams* pNodeParams); + +/** + * @brief Creates a child graph node and adds it to a graph. + * + * @param [out] pGraphNode - pointer to the graph node to create. + * @param [in] graph - instance of the graph to add the created node. + * @param [in] pDependencies - const pointer to the dependencies on the memset execution node. + * @param [in] numDependencies - the number of the dependencies. + * @param [in] childGraph - the graph to clone into this node + * @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 hipGraphAddChildGraphNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + hipGraph_t childGraph); + +/** + * @brief Gets a handle to the embedded graph of a child graph node. + * + * @param [in] node - instane of the node to get child graph. + * @param [out] pGraph - pointer to get the graph. + * @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 hipGraphChildGraphNodeGetGraph(hipGraphNode_t node, hipGraph_t* pGraph); + +/** + * @brief Updates node parameters in the child graph node in the given graphExec. + * + * @param [in] hGraphExec - instance of the executable graph with the node. + * @param [in] node - node from the graph which was used to instantiate graphExec. + * @param [in] childGraph - child graph with updated 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 hipGraphExecChildGraphNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node, + hipGraph_t childGraph); + /** * @brief Creates an empty node and adds it to a graph. * * @param [out] pGraphNode - pointer to the graph node to create and add to the graph. - * @param [in,out] graph - instane of the graph the node is add to. + * @param [in] graph - instane of the graph the node is add to. * @param [in] pDependencies - const pointer to the node dependenties. * @param [in] numDependencies - the number of dependencies. * @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 hipGraphAddEmptyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, const hipGraphNode_t* pDependencies, size_t numDependencies); + +/** + * @brief Creates an event record node and adds it to a graph. + * + * @param [out] pGraphNode - pointer to the graph node to create and add to the graph. + * @param [in] graph - instane of the graph the node to be added. + * @param [in] pDependencies - const pointer to the node dependenties. + * @param [in] numDependencies - the number of dependencies. + * @param [in] event - Event for the node. + * @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 hipGraphAddEventRecordNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + hipEvent_t event); + +/** + * @brief Returns the event associated with an event record node. + * + * @param [in] node - instane of the node to get event from. + * @param [out] event_out - Pointer to return the event. + * @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 hipGraphEventRecordNodeGetEvent(hipGraphNode_t node, hipEvent_t* event_out); + +/** + * @brief Sets an event record node's event. + * + * @param [in] node - instane of the node to set event to. + * @param [in] event - pointer to the event. + * @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 hipGraphEventRecordNodeSetEvent(hipGraphNode_t node, hipEvent_t event); + +/** + * @brief Sets the event for an event record node in the given graphExec. + * + * @param [in] hGraphExec - instance of the executable graph with the node. + * @param [in] hNode - node from the graph which was used to instantiate graphExec. + * @param [in] event - pointer to the event. + * @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 hipGraphExecEventRecordNodeSetEvent(hipGraphExec_t hGraphExec, hipGraphNode_t hNode, + hipEvent_t event); + +/** + * @brief Creates an event wait node and adds it to a graph. + * + * @param [out] pGraphNode - pointer to the graph node to create and add to the graph. + * @param [in] graph - instane of the graph the node to be added. + * @param [in] pDependencies - const pointer to the node dependenties. + * @param [in] numDependencies - the number of dependencies. + * @param [in] event - Event for the node. + * @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 hipGraphAddEventWaitNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + hipEvent_t event); + + +/** + * @brief Returns the event associated with an event wait node. + * + * @param [in] node - instane of the node to get event from. + * @param [out] event_out - Pointer to return the event. + * @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 hipGraphEventWaitNodeGetEvent(hipGraphNode_t node, hipEvent_t* event_out); + +/** + * @brief Sets an event wait node's event. + * + * @param [in] node - instane of the node to set event to. + * @param [in] event - pointer to the event. + * @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 hipGraphEventWaitNodeSetEvent(hipGraphNode_t node, hipEvent_t event); + +/** + * @brief Sets the event for an event record node in the given graphExec. + * + * @param [in] hGraphExec - instance of the executable graph with the node. + * @param [in] hNode - node from the graph which was used to instantiate graphExec. + * @param [in] event - pointer to the event. + * @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 hipGraphExecEventWaitNodeSetEvent(hipGraphExec_t hGraphExec, hipGraphNode_t hNode, + hipEvent_t event); + // doxygen end graph API /** * @} From 266355ffb3b28811854cceabfd26af080689bb5c Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 30 Sep 2021 13:17:15 +0530 Subject: [PATCH 9/9] [dtest] Remove unused variables in hipMemset2DAsyncMultiThreadAndKernel.cc [ROCm/hip commit: 865b40d8bd8b26a7dfe4d9719e8dcf26f4b3afc6] --- .../unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc | 6 +----- 1 file changed, 1 insertion(+), 5 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc b/projects/hip/tests/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc index 04240a4104..479e163c26 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc @@ -112,9 +112,6 @@ TEST_CASE("Unit_hipMemset2DAsync_WithKernel") { * hipMemSet2DAsync execution in multiple threads. */ TEST_CASE("Unit_hipMemset2DAsync_MultiThread") { - constexpr auto N = 4 * 1024 * 1024; - constexpr auto blocksPerCU = 6; // to hide latency - constexpr auto threadsPerBlock = 256; constexpr auto memPerThread = 200; constexpr int memsetval = 0x22; char *A_d, *A_h, *B_d, *B_h, *C_d; @@ -122,11 +119,10 @@ TEST_CASE("Unit_hipMemset2DAsync_MultiThread") { size_t width = NUM_W * sizeof(char); size_t sizeElements = width * NUM_H; size_t elements = NUM_W * NUM_H; - unsigned blocks{}; + int validateCount{}; hipStream_t stream; - blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); auto thread_count = HipTest::getHostThreadCount(memPerThread, NUM_THREADS); if (thread_count == 0) { WARN("Resources not available for thread creation");