From cd90931c0d878456c595e3279b6c890da8c05a80 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Fri, 1 Jul 2022 07:00:02 +0100 Subject: [PATCH 01/30] EXSWCPHIPT-102 - Adding hipEventRecord Tests (#2722) --- tests/catch/unit/event/Unit_hipEventRecord.cc | 120 +++++++++++------- 1 file changed, 75 insertions(+), 45 deletions(-) diff --git a/tests/catch/unit/event/Unit_hipEventRecord.cc b/tests/catch/unit/event/Unit_hipEventRecord.cc index 408f802d41..41c793d23e 100644 --- a/tests/catch/unit/event/Unit_hipEventRecord.cc +++ b/tests/catch/unit/event/Unit_hipEventRecord.cc @@ -19,72 +19,102 @@ 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. */ + // Test hipEventRecord serialization behavior. -// Through manual inspection of the reported timestamps, can determine if recording a NULL event -// forces synchronization : set -#include -#include -#include + #include +#include +#include +#include + TEST_CASE("Unit_hipEventRecord") { - size_t N = 4 * 1024 * 1024; - unsigned threadsPerBlock = 256; - int iterations = 1; + constexpr size_t N = 1024; + constexpr int iterations = 1; - unsigned blocks = (N + threadsPerBlock - 1) / threadsPerBlock; - if (blocks > 1024) blocks = 1024; - if (blocks == 0) blocks = 1; + constexpr int blocks = 1024; - printf("N=%zu (A+B+C= %6.1f MB total) blocks=%u threadsPerBlock=%u iterations=%d\n", N, - ((double)3 * N * sizeof(float)) / 1024 / 1024, blocks, threadsPerBlock, iterations); - printf("iterations=%d\n", iterations); + constexpr size_t Nbytes = N * sizeof(float); - size_t Nbytes = N * sizeof(float); + float *A_h, *B_h, *C_h; + float *A_d, *B_d, *C_d; + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N); - float *A_h, *B_h, *C_h; - float *A_d, *B_d, *C_d; - HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N); + enum TestType { + WithFlags_Default = hipEventDefault, + WithFlags_Blocking = hipEventBlockingSync, + WithFlags_DisableTiming = hipEventDisableTiming, +#if HT_AMD + WithFlags_ReleaseToDevice = hipEventReleaseToDevice, + WithFlags_ReleaseToSystem = hipEventReleaseToSystem, +#endif + WithoutFlags + }; - hipEvent_t start, stop; +#if HT_AMD + auto flags = GENERATE(WithFlags_Default, WithFlags_Blocking, WithFlags_DisableTiming, + WithFlags_ReleaseToDevice, WithFlags_ReleaseToSystem, WithoutFlags); +#endif - // NULL stream check: +#if HT_NVIDIA + auto flags = + GENERATE(WithFlags_Default, WithFlags_Blocking, WithFlags_DisableTiming, WithoutFlags); +#endif + + + hipEvent_t start{}, stop{}; + + if (flags == WithoutFlags) { HIP_CHECK(hipEventCreate(&start)); HIP_CHECK(hipEventCreate(&stop)); + } else { + HIP_CHECK(hipEventCreateWithFlags(&start, flags)); + HIP_CHECK(hipEventCreateWithFlags(&stop, flags)); + } - HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); - for (int i = 0; i < iterations; i++) { - //--- START TIMED REGION - long long hostStart = HipTest::get_time(); - // Record the start event - HIP_CHECK(hipEventRecord(start, NULL)); + for (int i = 0; i < iterations; i++) { + //--- START TIMED REGION + long long hostStart = HipTest::get_time(); + // Record the start event + HIP_CHECK(hipEventRecord(start, NULL)); - HipTest::launchKernel(HipTest::vectorADD, blocks, threadsPerBlock, 0, 0, -static_cast(A_d), static_cast(B_d), C_d, N); + HipTest::launchKernel(HipTest::vectorADD, blocks, 1, 0, 0, + static_cast(A_d), static_cast(B_d), + C_d, N); - HIP_CHECK(hipEventRecord(stop, NULL)); - HIP_CHECK(hipEventSynchronize(stop)); - long long hostStop = HipTest::get_time(); - //--- STOP TIMED REGION + HIP_CHECK(hipEventRecord(stop, NULL)); + HIP_CHECK(hipEventSynchronize(stop)); + long long hostStop = HipTest::get_time(); + //--- STOP TIMED REGION - float eventMs = 1.0f; - HIP_CHECK(hipEventElapsedTime(&eventMs, start, stop)); - float hostMs = HipTest::elapsed_time(hostStart, hostStop); + float hostMs = HipTest::elapsed_time(hostStart, hostStop); - printf("host_time (chrono) =%6.3fms\n", hostMs); - printf("kernel_time (hipEventElapsedTime) =%6.3fms\n", eventMs); - printf("\n"); + INFO("host_time (chrono) = " << hostMs); - // Make sure timer is timing something... - REQUIRE(eventMs > 0.0f); + // Make sure timer is timing something... + if (flags != WithFlags_DisableTiming) { + float eventMs = 1.0f; + HIP_CHECK(hipEventElapsedTime(&eventMs, start, stop)); + INFO("kernel_time (hipEventElapsedTime) = " << eventMs); + REQUIRE(eventMs > 0.0f); } + } - HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); - HIP_CHECK(hipEventDestroy(start)); - HIP_CHECK(hipEventDestroy(stop)); + HIP_CHECK(hipEventDestroy(start)); + HIP_CHECK(hipEventDestroy(stop)); - HipTest::checkVectorADD(A_h, B_h, C_h, N, true); + HipTest::checkVectorADD(A_h, B_h, C_h, N, true); + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + TestContext::get().cleanContext(); } + +TEST_CASE("Unit_hipEventRecord_Negative") { + SECTION("Nullptr event") { + HIP_CHECK_ERROR(hipEventRecord(nullptr, nullptr), hipErrorInvalidResourceHandle); + } +} \ No newline at end of file From a598eeebb0a6144ea3f197536fef7d783d800462 Mon Sep 17 00:00:00 2001 From: Finlay Date: Fri, 1 Jul 2022 09:17:47 +0100 Subject: [PATCH 02/30] EXSWCPHIPT-95 - More comprehensive tests for hipArrayCreate (#2702) --- tests/catch/unit/memory/CMakeLists.txt | 7 +- tests/catch/unit/memory/DriverContext.cc | 40 +++ tests/catch/unit/memory/DriverContext.hh | 41 +++ tests/catch/unit/memory/hipArrayCommon.hh | 124 +++++++ tests/catch/unit/memory/hipArrayCreate.cc | 387 ++++++++++++++++++---- tests/catch/unit/memory/hipMallocArray.cc | 79 +---- 6 files changed, 538 insertions(+), 140 deletions(-) create mode 100644 tests/catch/unit/memory/DriverContext.cc create mode 100644 tests/catch/unit/memory/DriverContext.hh create mode 100644 tests/catch/unit/memory/hipArrayCommon.hh diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index af72155549..1028010aa5 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -1,4 +1,4 @@ -# Copyright (c) 2021 Advanced Micro Devices, Inc. All Rights Reserved. +# Copyright (c) 2022 Advanced Micro Devices, Inc. All Rights Reserved. # # Permission is hereby granted, free of charge, to any person obtaining a copy # of this software and associated documentation files (the "Software"), to deal @@ -21,6 +21,7 @@ # Common Tests - Test independent of all platforms if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC + DriverContext.cc memset.cc malloc.cc hipMemcpy2DToArray.cc @@ -88,6 +89,7 @@ set(TEST_SRC ) else() set(TEST_SRC + DriverContext.cc memset.cc malloc.cc hipMemcpy2DToArray.cc @@ -159,4 +161,5 @@ endif() hip_add_exe_to_target(NAME MemoryTest TEST_SRC ${TEST_SRC} - TEST_TARGET_NAME build_tests) + TEST_TARGET_NAME build_tests + COMPILE_OPTIONS -std=c++14) diff --git a/tests/catch/unit/memory/DriverContext.cc b/tests/catch/unit/memory/DriverContext.cc new file mode 100644 index 0000000000..6791650567 --- /dev/null +++ b/tests/catch/unit/memory/DriverContext.cc @@ -0,0 +1,40 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "DriverContext.hh" +#include + +DriverContext::DriverContext() { +#if HT_NVIDIA + HIP_CHECK(hipInit(0)); + HIP_CHECK(hipDeviceGet(&device, 0)); + HIP_CHECK(hipDevicePrimaryCtxRetain(&ctx, device)); + HIP_CHECK(hipCtxPushCurrent(ctx)); +#endif +} + +DriverContext::~DriverContext() { +#if HT_NVIDIA + HIP_CHECK(hipCtxPopCurrent(&ctx)); + HIP_CHECK(hipDevicePrimaryCtxRelease(device)); +#endif +} diff --git a/tests/catch/unit/memory/DriverContext.hh b/tests/catch/unit/memory/DriverContext.hh new file mode 100644 index 0000000000..76afe28f44 --- /dev/null +++ b/tests/catch/unit/memory/DriverContext.hh @@ -0,0 +1,41 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +#include + +class DriverContext { + private: +#if HT_NVIDIA + hipCtx_t ctx; + hipDevice_t device; +#endif + + public: + DriverContext(); + ~DriverContext(); + + // Rule of three + DriverContext(const DriverContext& other) = delete; + DriverContext(DriverContext&& other) noexcept = delete; +}; diff --git a/tests/catch/unit/memory/hipArrayCommon.hh b/tests/catch/unit/memory/hipArrayCommon.hh new file mode 100644 index 0000000000..1c6100a6e9 --- /dev/null +++ b/tests/catch/unit/memory/hipArrayCommon.hh @@ -0,0 +1,124 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +#include + +constexpr size_t BlockSize = 16; + +template struct type_and_size_and_format { + using type = T; + static constexpr size_t size = N; + static constexpr hipArray_Format format = Format; +}; + +// Create a map of type to scalar type, vector size and scalar type format enum. +// This is useful for creating simpler function that depend on the vector size. +template struct vector_info; +template <> +struct vector_info : type_and_size_and_format {}; +template <> struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; + +template <> +struct vector_info : type_and_size_and_format {}; +template <> struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; + +template <> +struct vector_info : type_and_size_and_format {}; +template <> struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; +template <> +struct vector_info + : type_and_size_and_format {}; + +// read from a texture using normalized coordinates +constexpr size_t ChannelToRead = 1; +template +__global__ void readFromTexture(T* output, hipTextureObject_t texObj, size_t width, size_t height, + bool textureGather) { + // Calculate normalized texture coordinates + const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; + const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; + const float u = x / (float)width; + + // Read from texture and write to global memory + if (height == 0) { + output[x] = tex1D(texObj, u); + } else { + const float v = y / (float)height; + output[y * width + x] = + textureGather ? tex2Dgather(texObj, u, v, ChannelToRead) : tex2D(texObj, u, v); + } +} + +template void checkDataIsAscending(const std::vector& hostData) { + bool allMatch = true; + size_t i = 0; + for (; i < hostData.size(); ++i) { + allMatch = allMatch && hostData[i] == static_cast(i); + if (!allMatch) break; + } + INFO("hostData[" << i << "] == " << static_cast(hostData[i])); + REQUIRE(allMatch); +} + +inline size_t getFreeMem() { + size_t free = 0, total = 0; + HIP_CHECK(hipMemGetInfo(&free, &total)); + return free; +} diff --git a/tests/catch/unit/memory/hipArrayCreate.cc b/tests/catch/unit/memory/hipArrayCreate.cc index 3bbf6cbd54..9ef637821e 100644 --- a/tests/catch/unit/memory/hipArrayCreate.cc +++ b/tests/catch/unit/memory/hipArrayCreate.cc @@ -24,7 +24,11 @@ hipArrayCreate API test scenarios 3. Multithreaded scenario */ +#include +#include #include +#include "hipArrayCommon.hh" +#include "DriverContext.hh" static constexpr auto NUM_W{4}; static constexpr auto BIGNUM_W{100}; @@ -48,76 +52,31 @@ static constexpr auto ARRAY_LOOP{100}; static void ArrayCreate_DiffSizes(int gpu) { HIP_CHECK(hipSetDevice(gpu)); - std::vector array_size; - array_size.push_back(NUM_W); - array_size.push_back(BIGNUM_W); - for (auto &size : array_size) { - HIP_ARRAY array[ARRAY_LOOP]; - size_t tot, avail, ptot, pavail; - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + std::vector> array_size{{NUM_W, NUM_H}, {BIGNUM_W, BIGNUM_H}}; + for (auto& size : array_size) { + std::array array; + const size_t pavail = getFreeMem(); + HIP_ARRAY_DESCRIPTOR desc; + desc.NumChannels = 1; + desc.Width = std::get<0>(size); + desc.Height = std::get<1>(size); + desc.Format = HIP_AD_FORMAT_FLOAT; for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_ARRAY_DESCRIPTOR desc; - desc.NumChannels = 1; - if (size == NUM_W) { - desc.Width = NUM_W; - desc.Height = NUM_H; - } else { - desc.Width = BIGNUM_W; - desc.Height = BIGNUM_H; - } - desc.Format = HIP_AD_FORMAT_FLOAT; HIP_CHECK(hipArrayCreate(&array[i], &desc)); } for (int i = 0; i < ARRAY_LOOP; i++) { - ARRAY_DESTROY(array[i]); + HIP_CHECK(hipArrayDestroy(array[i])); } - HIP_CHECK(hipMemGetInfo(&avail, &tot)); - if ((pavail != avail)) { + const size_t avail = getFreeMem(); + if (pavail != avail) { HIPASSERT(false); } } } -/*Thread function*/ -static void ArrayCreateThreadFunc(int gpu) { - ArrayCreate_DiffSizes(gpu); -} - /* This testcase verifies hipArrayCreate API for small and big chunks data*/ -TEST_CASE("Unit_hipArrayCreate_DiffSizes") { - ArrayCreate_DiffSizes(0); -} +TEST_CASE("Unit_hipArrayCreate_DiffSizes") { ArrayCreate_DiffSizes(0); } - -/* This testcase verifies the negative scenarios of - * hipArrayCreate API - */ -TEST_CASE("Unit_hipArrayCreate_Negative") { - HIP_ARRAY_DESCRIPTOR desc; - HIP_ARRAY array; - desc.Format = HIP_AD_FORMAT_FLOAT; - desc.NumChannels = 1; - desc.Width = NUM_W; - desc.Height = NUM_H; -#if HT_NVIDIA - SECTION("NullPointer to Array") { - REQUIRE(hipArrayCreate(nullptr, &desc) != hipSuccess); - } - - SECTION("NullPointer to Channel Descriptor") { - REQUIRE(hipArrayCreate(&array, nullptr) != hipSuccess); - } -#endif - SECTION("Width 0 for Array Descriptor") { - desc.Width = 0; - REQUIRE(hipArrayCreate(&array, &desc) != hipSuccess); - } - - SECTION("Invalid NumChannels") { - desc.NumChannels = 3; - REQUIRE(hipArrayCreate(&array, &desc) != hipSuccess); - } -} /* This testcase verifies the hipArrayCreate API in multithreaded scenario by launching threads in parallel on multiple GPUs @@ -129,16 +88,16 @@ TEST_CASE("Unit_hipArrayCreate_MultiThread") { devCnt = HipTest::getDeviceCount(); - size_t tot, avail, ptot, pavail; - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + const size_t pavail = getFreeMem(); for (int i = 0; i < devCnt; i++) { - threadlist.push_back(std::thread(ArrayCreateThreadFunc, i)); + // FIXME: the HIP_CHECK and HIPASSERT are not threadsafe so this test is broken. + threadlist.push_back(std::thread(ArrayCreate_DiffSizes, i)); } - for (auto &t : threadlist) { + for (auto& t : threadlist) { t.join(); } - HIP_CHECK(hipMemGetInfo(&avail, &tot)); + const size_t avail = getFreeMem(); if (pavail != avail) { WARN("Memory leak of hipMalloc3D API in multithreaded scenario"); @@ -146,3 +105,305 @@ TEST_CASE("Unit_hipArrayCreate_MultiThread") { } } + +// All the possible formats for channel data in an array. +static const std::vector formats{ + HIP_AD_FORMAT_UNSIGNED_INT8, HIP_AD_FORMAT_UNSIGNED_INT16, HIP_AD_FORMAT_UNSIGNED_INT32, + HIP_AD_FORMAT_SIGNED_INT8, HIP_AD_FORMAT_SIGNED_INT16, HIP_AD_FORMAT_SIGNED_INT32, + HIP_AD_FORMAT_HALF, HIP_AD_FORMAT_FLOAT}; + +// Helpful for printing errors +const char* formatToString(hipArray_Format f) { + switch (f) { + case HIP_AD_FORMAT_UNSIGNED_INT8: + return "Unsigned Int 8"; + case HIP_AD_FORMAT_UNSIGNED_INT16: + return "Unsigned Int 16"; + case HIP_AD_FORMAT_UNSIGNED_INT32: + return "Unsigned Int 32"; + case HIP_AD_FORMAT_SIGNED_INT8: + return "Signed Int 8"; + case HIP_AD_FORMAT_SIGNED_INT16: + return "Signed Int 16"; + case HIP_AD_FORMAT_SIGNED_INT32: + return "Signed Int 32"; + case HIP_AD_FORMAT_HALF: + return "Float 16"; + case HIP_AD_FORMAT_FLOAT: + return "Float 32"; + default: + return "not found"; + } +} + +// Tests ///////////////////////////////////////// + +#if HT_AMD +constexpr auto MemoryTypeHost = hipMemoryTypeHost; +constexpr auto MemoryTypeArray = hipMemoryTypeArray; +constexpr auto NORMALIZED_COORDINATES = HIP_TRSF_NORMALIZED_COORDINATES; +constexpr auto READ_AS_INTEGER = HIP_TRSF_READ_AS_INTEGER; +#else +constexpr auto MemoryTypeHost = CU_MEMORYTYPE_HOST; +constexpr auto MemoryTypeArray = CU_MEMORYTYPE_ARRAY; +// (EXSWCPHIPT-92) HIP equivalents not defined for CUDA backend. +constexpr auto NORMALIZED_COORDINATES = CU_TRSF_NORMALIZED_COORDINATES; +constexpr auto READ_AS_INTEGER = CU_TRSF_READ_AS_INTEGER; +#endif + +// Copy data from host to the hiparray, accounting 1D or 2D arrays +template +void copyToArray(hiparray dst, const std::vector& src, const size_t height) { + const auto sizeInBytes = src.size() * sizeof(T); + if (height == 0) { + // FIXME(EXSWCPHIPT-64) remove cast when API is fixed (will require major version change) + HIP_CHECK(hipMemcpyHtoA(reinterpret_cast(dst), 0, src.data(), sizeInBytes)); + } else { + const auto pitch = sizeInBytes / height; + hip_Memcpy2D copyParams{}; + copyParams.srcMemoryType = MemoryTypeHost; + copyParams.srcXInBytes = 0; // x offset + copyParams.srcY = 0; // y offset + copyParams.srcHost = src.data(); + copyParams.srcPitch = pitch; + + + copyParams.dstMemoryType = MemoryTypeArray; + copyParams.dstXInBytes = 0; // x offset + copyParams.dstY = 0; // y offset + copyParams.dstArray = dst; + + copyParams.WidthInBytes = pitch; + copyParams.Height = height; + + HIP_CHECK(hipMemcpyParam2D(©Params)); + } +} + +// Test the allocated array by generating a texture from it then reading from that texture. +// Textures are read-only, so write to the array then copy that into normal device memory. +template +void testArrayAsTexture(hiparray array, const size_t width, const size_t height) { + using vec_info = vector_info; + using scalar_type = typename vec_info::type; + const auto h = height ? height : 1; + const auto size = sizeof(T) * width * h; + + // set hip array + std::vector hostData(width * h * vec_info::size); + // assigned ascending values to the data array to show indexing is working + std::iota(std::begin(hostData), std::end(hostData), 0); + + copyToArray(array, hostData, height); + + // create texture + hipTextureObject_t textObj{}; + + HIP_RESOURCE_DESC resDesc{}; + memset(&resDesc, 0, sizeof(HIP_RESOURCE_DESC)); + resDesc.resType = HIP_RESOURCE_TYPE_ARRAY; + resDesc.res.array.hArray = array; + resDesc.flags = 0; + + HIP_TEXTURE_DESC texDesc{}; + memset(&texDesc, 0, sizeof(HIP_TEXTURE_DESC)); + // use the actual values in the texture, not normalized data + texDesc.filterMode = HIP_TR_FILTER_MODE_POINT; + // Use normalized coordinates and also read the data in the original data type + texDesc.flags |= NORMALIZED_COORDINATES | READ_AS_INTEGER; + + HIP_CHECK(hipTexObjectCreate(&textObj, &resDesc, &texDesc, nullptr)); + + // run kernel + T* device_data{}; + HIP_CHECK(hipMalloc(&device_data, size)); + readFromTexture<<>>(device_data, textObj, width, + height, false); + HIP_CHECK(hipGetLastError()); // check for errors when running the kernel + + // copy data back and then test it + std::fill(std::begin(hostData), std::end(hostData), 0); + HIP_CHECK(hipMemcpy(hostData.data(), device_data, size, hipMemcpyDeviceToHost)); + + checkDataIsAscending(hostData); + + // clean up + HIP_CHECK(hipTexObjectDestroy(textObj)); + HIP_CHECK(hipFree(device_data)); +} + +// Selection of types chosen since trying all types would be slow to compile +// Test the happy path of the hipArrayCreate +TEMPLATE_TEST_CASE("Unit_hipArrayCreate_happy", "", uint, int, int4, ushort, short2, char, uchar2, + char4, float, float2, float4) { +#if HT_AMD + if (std::is_same::value || std::is_same::value || + std::is_same::value) { + HipTest::HIP_SKIP_TEST("Probably EXSWCPHIPT-62"); + return; + } +#endif + using vec_info = vector_info; + DriverContext ctx; + + HIP_ARRAY_DESCRIPTOR desc; + desc.Format = vec_info::format; + desc.NumChannels = vec_info::size; + desc.Width = 1024; + desc.Height = GENERATE(0, 1024); + + size_t initFree = getFreeMem(); + + // pointer to the array in device memory + hiparray array{}; + + HIP_CHECK(hipArrayCreate(&array, &desc)); + + testArrayAsTexture(array, desc.Width, desc.Height); + + size_t finalFree = getFreeMem(); + + const size_t allocSize = sizeof(TestType) * desc.Width * (desc.Height ? desc.Height : 1); + // will be aligned to some size, so this is not exact + REQUIRE(initFree - finalFree >= allocSize); + + HIP_CHECK(hipArrayDestroy(array)); +} + + +// Only widths and Heights up to the maxTexture size is supported +TEMPLATE_TEST_CASE("Unit_hipArrayCreate_maxTexture", "", uint, int, int4, ushort, short2, char, + uchar2, char4, float, float2, float4) { + using vec_info = vector_info; + DriverContext ctx; + + HIP_ARRAY_DESCRIPTOR desc; + desc.Format = vec_info::format; + desc.NumChannels = vec_info::size; + + int device; + HIP_CHECK(hipGetDevice(&device)); + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + + hiparray array{}; + SECTION("Happy") { + SECTION("1D - Max") { + desc.Width = prop.maxTexture1D; + desc.Height = 0; + } + SECTION("2D - Max Width") { + desc.Width = prop.maxTexture2D[0]; + desc.Height = 64; + } + SECTION("2D - Max Height") { + desc.Width = 64; + desc.Height = prop.maxTexture2D[1]; + } + SECTION("2D - Max Width and Height") { + desc.Width = prop.maxTexture2D[0]; + desc.Height = prop.maxTexture2D[1]; + } + auto maxArrayCreateError = hipArrayCreate(&array, &desc); + // this can try to alloc many GB of memory, so out of memory is acceptable + // return to avoid destroy + if (maxArrayCreateError == hipErrorOutOfMemory) return; + HIP_CHECK(maxArrayCreateError); + HIP_CHECK(hipArrayDestroy(array)); + } + SECTION("Negative") { + SECTION("1D - More Than Max") { + desc.Width = prop.maxTexture1D + 1; + desc.Height = 0; + } + SECTION("2D - More Than Max Width") { + desc.Width = prop.maxTexture2D[0] + 1; + desc.Height = 64; + } + SECTION("2D - More Than Max Height") { + desc.Width = 64; + desc.Height = prop.maxTexture2D[1] + 1; + } + SECTION("2D - More Than Max Width and Height") { + desc.Width = prop.maxTexture2D[0] + 1; + desc.Height = prop.maxTexture2D[1] + 1; + } + HIP_CHECK_ERROR(hipArrayCreate(&array, &desc), hipErrorInvalidValue); + } +} + +// zero-width array is not supported +TEST_CASE("Unit_hipArrayCreate_ZeroWidth") { + DriverContext ctx; + HIP_ARRAY_DESCRIPTOR desc; + desc.Format = formats[0]; + desc.NumChannels = 4; + desc.Width = 0; + desc.Height = GENERATE(0, 1024); + + // pointer to the array in device memory + hiparray array; + HIP_CHECK_ERROR(hipArrayCreate(&array, &desc), hipErrorInvalidValue); +} + +// HipArrayCreate will return an error when nullptr is used as the array argument +TEST_CASE("Unit_hipArrayCreate_Nullptr") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("Probably EXSWCPHIPT-45"); + return; +#endif + DriverContext ctx; + SECTION("Null array") { + HIP_ARRAY_DESCRIPTOR desc; + desc.Format = formats[0]; + desc.NumChannels = 4; + desc.Width = 1024; + desc.Height = 1024; + + HIP_CHECK_ERROR(hipArrayCreate(nullptr, &desc), hipErrorInvalidValue); + } + SECTION("Null Description") { + hiparray array; + HIP_CHECK_ERROR(hipArrayCreate(&array, nullptr), hipErrorInvalidValue); + } +} + +// Only elements with 1,2, or 4 channels is supported +TEST_CASE("Unit_hipArrayCreate_BadNumberChannelElement") { + DriverContext ctx; + HIP_ARRAY_DESCRIPTOR desc; + desc.Format = GENERATE(from_range(std::begin(formats), std::end(formats))); + desc.NumChannels = GENERATE(-1, 0, 3, 5, 8); + desc.Width = 1024; + desc.Height = GENERATE(0, 1024); + + hiparray array; + + INFO("Format: " << formatToString(desc.Format) << " NumChannels: " << desc.NumChannels + << " Height: " << desc.Height) + HIP_CHECK_ERROR(hipArrayCreate(&array, &desc), hipErrorInvalidValue); +} + +// Only certain channel formats are acceptable. +TEST_CASE("Unit_hipArrayCreate_BadChannelFormat") { + DriverContext ctx; + HIP_ARRAY_DESCRIPTOR desc; + + // create a bad format + desc.Format = + std::accumulate(std::begin(formats), std::end(formats), formats[0], + [](auto i, auto f) { return static_cast(i + f); }); + for (auto&& format : formats) { + REQUIRE(desc.Format != format); + } + + desc.NumChannels = 4; + desc.Width = 1024; + desc.Height = GENERATE(0, 1024); + + hiparray array; + + INFO("Format: " << formatToString(desc.Format) << " Height: " << desc.Height) + HIP_CHECK_ERROR(hipArrayCreate(&array, &desc), hipErrorInvalidValue); +} diff --git a/tests/catch/unit/memory/hipMallocArray.cc b/tests/catch/unit/memory/hipMallocArray.cc index 22e1ce9864..8859aadb45 100644 --- a/tests/catch/unit/memory/hipMallocArray.cc +++ b/tests/catch/unit/memory/hipMallocArray.cc @@ -27,9 +27,8 @@ hipMallocArray API test scenarios #include #include -#if defined(_WIN32) || defined(_WIN64) #include -#endif +#include "hipArrayCommon.hh" static constexpr auto NUM_W{4}; static constexpr auto BIGNUM_W{100}; @@ -86,7 +85,7 @@ TEST_CASE("Unit_hipMallocArray_MultiThread") { size_t tot, avail, ptot, pavail; HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); for (int i = 0; i < devCnt; i++) { - // TODO the HIP_CHECK and HIPASSERT are not threadsafe so this test is broken. + // FIXME: the HIP_CHECK and HIPASSERT are not threadsafe so this test is broken. threadlist.push_back(std::thread(MallocArray_DiffSizes, i)); } @@ -101,63 +100,8 @@ TEST_CASE("Unit_hipMallocArray_MultiThread") { } } - -constexpr size_t BlockSize = 16; - -template struct type_and_size { - using type = T; - static constexpr size_t size = N; -}; - -// scalars are interpreted as a vector of 1 length. -// template using int_constant = std::integral_constant; -template struct vector_info; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; - -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; - -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; -template <> struct vector_info : type_and_size {}; - // Kernels /////////////////////////////////////// -// read from a texture using normalized coordinates -constexpr size_t ChannelToRead = 1; -template -__global__ void readFromTexture(T* output, hipTextureObject_t texObj, size_t width, size_t height, - bool textureGather) { - // Calculate normalized texture coordinates - const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; - const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; - const float u = x / (float)width; - - // Read from texture and write to global memory - if (height == 0) { - output[x] = tex1D(texObj, u); - } else { - const float v = y / (float)height; - output[y * width + x] = - textureGather ? tex2Dgather(texObj, u, v, ChannelToRead) : tex2D(texObj, u, v); - } -} - template __device__ void addOne(T* a) { using scalar_type = typename vector_info::type; auto as = reinterpret_cast(a); @@ -190,16 +134,6 @@ template size_t getAllocSize(const size_t width, const size_t heigh return sizeof(T) * width * (height ? height : 1); } -template void checkDataIsAscending(const std::vector& hostData) { - bool allMatch = true; - size_t i = 0; - for (; i < hostData.size(); ++i) { - allMatch = allMatch && hostData[i] == static_cast(i); - if (!allMatch) break; - } - INFO("hostData[" << i << "] == " << static_cast(hostData[i])); - REQUIRE(allMatch); -} const char* channelFormatString(hipChannelFormatKind formatKind) noexcept { switch (formatKind) { @@ -458,12 +392,6 @@ void testArrayAsSurface(hipArray_t arrayPtr, const size_t width, const size_t he HIP_CHECK(hipFree(device_data)); } -size_t getFreeMem() { - size_t free = 0, total = 0; - HIP_CHECK(hipMemGetInfo(&free, &total)); - return free; -} - // The happy path of a default array and a SurfaceLoadStore array should work // Selection of types chosen to reduce compile times TEMPLATE_TEST_CASE("Unit_hipMallocArray_happy", "", uint, int, int4, ushort, short2, char, uchar2, @@ -526,6 +454,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_MaxTexture_Default", "", uint, int4, ush HIP_CHECK(hipGetDevice(&device)); hipDeviceProp_t prop; HIP_CHECK(hipGetDeviceProperties(&prop, device)); + size_t width, height; hipArray_t array{}; hipChannelFormatDesc desc = hipCreateChannelDesc(); @@ -549,7 +478,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_MaxTexture_Default", "", uint, int4, ush height = prop.maxTexture2D[1]; } auto maxArrayCreateError = hipMallocArray(&array, &desc, width, height, flag); - // this can try to alloc many GB of memory, so out of memory is fair + // this can try to alloc many GB of memory, so out of memory is acceptable if (maxArrayCreateError == hipErrorOutOfMemory) return; HIP_CHECK(maxArrayCreateError); HIP_CHECK(hipFreeArray(array)); From 5f7a820b1ee0be14988949e644a0993cc0645199 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio?= Date: Fri, 1 Jul 2022 09:54:17 +0100 Subject: [PATCH 03/30] EXSWCPHIPT-50 - Testing HipMalloc3D / HipMallocPitch / HipMemAllocPitch (#2613) --- tests/catch/unit/memory/CMakeLists.txt | 2 + tests/catch/unit/memory/hipMalloc3D.cc | 20 -- tests/catch/unit/memory/hipMallocPitch.cc | 282 +++++++++++++++++++++- 3 files changed, 273 insertions(+), 31 deletions(-) diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index 1028010aa5..9f31f65ea5 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -39,6 +39,7 @@ set(TEST_SRC hipMemcpyAllApiNegative.cc hipMemcpy_MultiThread.cc hipHostRegister.cc + hipMallocPitch.cc hipMemPtrGetInfo.cc hipPointerGetAttributes.cc hipHostGetFlags.cc @@ -107,6 +108,7 @@ set(TEST_SRC hipMemcpyAllApiNegative.cc hipMemcpy_MultiThread.cc hipHostRegister.cc + hipMallocPitch.cc hipHostGetFlags.cc hipHostGetDevicePointer.cc hipMallocManaged_MultiScenario.cc diff --git a/tests/catch/unit/memory/hipMalloc3D.cc b/tests/catch/unit/memory/hipMalloc3D.cc index 2efc9ff7db..386eff0a51 100644 --- a/tests/catch/unit/memory/hipMalloc3D.cc +++ b/tests/catch/unit/memory/hipMalloc3D.cc @@ -64,26 +64,6 @@ static void Malloc3DThreadFunc(int gpu) { MemoryAlloc3DDiffSizes(gpu); } -/* - * This verifies the negative scenarios of hipMalloc3D API - */ -TEST_CASE("Unit_hipMalloc3D_Negative") { - size_t width = SMALL_SIZE * sizeof(char); - size_t height{SMALL_SIZE}, depth{SMALL_SIZE}; - hipPitchedPtr devPitchedPtr; - - SECTION("Passing nullptr to device pitched pointer") { - hipExtent extent = make_hipExtent(width, height, depth); - REQUIRE(hipMalloc3D(nullptr, extent) != hipSuccess); - } - - SECTION("Passing Max values to extent") { - hipExtent extent = make_hipExtent(std::numeric_limits::max(), - std::numeric_limits::max(), - std::numeric_limits::max()); - REQUIRE(hipMalloc3D(&devPitchedPtr, extent) != hipSuccess); - } -} /* * This verifies the hipMalloc3D API by * assigning width,height and depth as 10 diff --git a/tests/catch/unit/memory/hipMallocPitch.cc b/tests/catch/unit/memory/hipMallocPitch.cc index 08691b2045..0a63db52eb 100644 --- a/tests/catch/unit/memory/hipMallocPitch.cc +++ b/tests/catch/unit/memory/hipMallocPitch.cc @@ -1,13 +1,16 @@ /* Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions: + The above copyright notice and this permission notice shall be included in all copies or substantial portions of the Software. + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE @@ -17,20 +20,278 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/* -Test Scenarios of hipMallocPitch API -1. Negative Scenarios -2. Basic Functionality Scenario -3. Allocate memory using hipMallocPitch API, Launch Kernel validate result. -4. Allocate Memory in small chunks and large chunks and check for possible memory leaks -5. Allocate Memory using hipMallocPitch API, Memcpy2D on the allocated variables. -6. Multithreaded scenario -*/ - #include +#include +#include +#include "hip/driver_types.h" +#include +#include +#include #include #include +/** + * @brief Test hipMalloc3D, hipMallocPitch and hipMemAllocPitch with multiple input values. + * Checks that the memory has been allocated with the specified pitch and extent sizes. + */ + +struct MemoryInfo { + size_t freeMem; + size_t totalMem; +}; + +inline static MemoryInfo createMemoryInfo() { + MemoryInfo memoryInfo{}; + HIP_CHECK(hipMemGetInfo(&memoryInfo.freeMem, &memoryInfo.totalMem)); + return memoryInfo; +} + +static void validateMemory(void* devPtr, hipExtent extent, size_t pitch, + MemoryInfo memBeforeAllocation) { + INFO("Width: " << extent.width << " Height: " << extent.height << " Depth: " << extent.depth); + + MemoryInfo memAfterAllocation{createMemoryInfo()}; + const size_t theoreticalAllocatedMemory{pitch * extent.height * extent.depth}; + const size_t allocatedMemory = memBeforeAllocation.freeMem - memAfterAllocation.freeMem; + + if (theoreticalAllocatedMemory == 0) { + REQUIRE(theoreticalAllocatedMemory == allocatedMemory); + return; /* If there was no memory allocated then we don't need to do further checks. */ + } else { + REQUIRE(theoreticalAllocatedMemory <= allocatedMemory); + } + + std::unique_ptr hostPtr{new char[theoreticalAllocatedMemory]}; + std::memset(hostPtr.get(), 2, theoreticalAllocatedMemory); + + hipPitchedPtr devPitchedPtr{devPtr, pitch, extent.width, extent.height}; + hipPitchedPtr hostPitchedPtr{hostPtr.get(), pitch, extent.width, extent.height}; + + HIP_CHECK(hipMemset3D(devPitchedPtr, 1, extent)); + + hipMemcpy3DParms params{}; + params.srcPtr = devPitchedPtr; + params.kind = hipMemcpyKind::hipMemcpyDeviceToHost; + params.dstPtr = hostPitchedPtr; + params.extent = extent; + HIP_CHECK(hipMemcpy3D(¶ms)) + + bool mismatch = false; + for (size_t width = 0; width < extent.width; ++width) { + for (size_t height = 0; height < extent.height; ++height) { + for (size_t depth = 0; depth < extent.depth; ++depth) { + char* reinterpretedPtr = reinterpret_cast(hostPtr.get()); + size_t index = (pitch * extent.height * depth) + (pitch * height) + width; + if (*(reinterpretedPtr + index) != 1) { + mismatch = true; + } + } + } + } + REQUIRE(!mismatch); +} + +class ExtentGenerator { + public: + static constexpr size_t totalRandomValues{20}; + static constexpr size_t seed{1337}; + + std::uniform_int_distribution width_distribution{1, 1024}; + std::uniform_int_distribution height_distribution{1, 100}; + std::uniform_int_distribution depth_distribution{1, 100}; + + std::vector extents2D{}; + std::vector extents3D{}; + + static ExtentGenerator& getInstance() { + static ExtentGenerator instance; + return instance; + } + + private: + ExtentGenerator() { + std::mt19937 randomGenerator{seed}; + extents3D = std::vector{hipExtent{0, 0, 0}, hipExtent{1, 0, 0}, hipExtent{0, 1, 0}, + hipExtent{0, 0, 1}}; + + for (size_t i = 0; i < totalRandomValues; ++i) { + extents3D.push_back(hipExtent{width_distribution(randomGenerator), + height_distribution(randomGenerator), + depth_distribution(randomGenerator)}); + } + + extents2D = std::vector{hipExtent{0, 0, 1}, hipExtent{1, 0, 1}, hipExtent{0, 1, 1}}; + + for (size_t i = 0; i < totalRandomValues; ++i) { + extents2D.push_back( + hipExtent{width_distribution(randomGenerator), height_distribution(randomGenerator), 1}); + } + } +}; + +enum class AllocationApi { hipMalloc3D, hipMallocPitch, hipMemAllocPitch }; + +hipExtent generateExtent(AllocationApi api) { + hipExtent extent; + if (api == AllocationApi::hipMalloc3D) { + auto& extents3D = ExtentGenerator::getInstance().extents3D; + extent = GENERATE_REF(from_range(extents3D.begin(), extents3D.end())); + } else { + auto& extents2D = ExtentGenerator::getInstance().extents2D; + extent = GENERATE_REF(from_range(extents2D.begin(), extents2D.end())); + } + + return extent; +} + + +TEST_CASE("Unit_hipMalloc3D_ValidatePitch") { + hipPitchedPtr hipPitchedPtr; + hipExtent validExtent{generateExtent(AllocationApi::hipMalloc3D)}; + + MemoryInfo memBeforeAllocation{createMemoryInfo()}; + HIP_CHECK(hipMalloc3D(&hipPitchedPtr, validExtent)); + validateMemory(hipPitchedPtr.ptr, validExtent, hipPitchedPtr.pitch, memBeforeAllocation); + HIP_CHECK(hipFree(hipPitchedPtr.ptr)); +} + +TEST_CASE("Unit_hipMemAllocPitch_ValidatePitch") { + size_t pitch; + hipDeviceptr_t ptr; + hipExtent validExtent{generateExtent(AllocationApi::hipMemAllocPitch)}; + MemoryInfo memBeforeAllocation{createMemoryInfo()}; + unsigned int elementSizeBytes = GENERATE(4, 8, 16); + + if (validExtent.width == 0 || validExtent.height == 0) { + return; + } + + HIP_CHECK( + hipMemAllocPitch(&ptr, &pitch, validExtent.width, validExtent.height, elementSizeBytes)); + validateMemory(reinterpret_cast(ptr), validExtent, pitch, memBeforeAllocation); + HIP_CHECK(hipFree(reinterpret_cast(ptr))); +} + +TEST_CASE("Unit_hipMallocPitch_ValidatePitch") { + size_t pitch; + void* ptr; + hipExtent validExtent{generateExtent(AllocationApi::hipMemAllocPitch)}; + MemoryInfo memBeforeAllocation{createMemoryInfo()}; + HIP_CHECK(hipMallocPitch(&ptr, &pitch, validExtent.width, validExtent.height)); + validateMemory(ptr, validExtent, pitch, memBeforeAllocation); + HIP_CHECK(hipFree(ptr)); +} + +TEST_CASE("Unit_hipMalloc3D_Negative") { + SECTION("Invalid ptr") { + hipExtent validExtent{1, 1, 1}; + HIP_CHECK_ERROR(hipMalloc3D(nullptr, validExtent), hipErrorInvalidValue); + } + + hipPitchedPtr ptr; + constexpr size_t maxSizeT = std::numeric_limits::max(); + + SECTION("Max size_t width") { + hipExtent validExtent{maxSizeT, 1, 1}; + HIP_CHECK_ERROR(hipMalloc3D(&ptr, validExtent), hipErrorInvalidValue); + } + + SECTION("Max size_t height") { + hipExtent validExtent{1, maxSizeT, 1}; + HIP_CHECK_ERROR(hipMalloc3D(&ptr, validExtent), hipErrorOutOfMemory); + } + + SECTION("Max size_t depth") { + hipExtent validExtent{1, 1, maxSizeT}; + HIP_CHECK_ERROR(hipMalloc3D(&ptr, validExtent), hipErrorOutOfMemory); + } + + SECTION("Max size_t all dimensions") { + hipExtent validExtent{maxSizeT, maxSizeT, maxSizeT}; + HIP_CHECK_ERROR(hipMalloc3D(&ptr, validExtent), hipErrorInvalidValue); + } +} + +TEST_CASE("Unit_hipMallocPitch_Negative") { + size_t pitch; + void* ptr; + constexpr size_t maxSizeT = std::numeric_limits::max(); + + SECTION("Invalid ptr") { + HIP_CHECK_ERROR(hipMallocPitch(nullptr, &pitch, 1, 1), hipErrorInvalidValue); + } + + SECTION("Invalid pitch") { + HIP_CHECK_ERROR(hipMallocPitch(&ptr, nullptr, 1, 1), hipErrorInvalidValue); + } + + SECTION("Max size_t width") { + HIP_CHECK_ERROR(hipMallocPitch(&ptr, &pitch, maxSizeT, 1), hipErrorInvalidValue); + } + + SECTION("Max size_t height") { + HIP_CHECK_ERROR(hipMallocPitch(&ptr, &pitch, 1, maxSizeT), hipErrorOutOfMemory); + } +} + +TEST_CASE("Unit_hipMemAllocPitch_Negative") { + size_t pitch; + hipDeviceptr_t ptr{}; + unsigned int validElementSizeBytes{4}; + constexpr size_t maxSizeT = std::numeric_limits::max(); + +#if HT_NVIDIA + /* Device synchronize is used here to initialize the device. + * Nvidia does not implicitly do it for this Api. And hipInit(0) does not work either. + */ + HIP_CHECK(hipDeviceSynchronize()); + + SECTION("Invalid elementSizeBytes") { + unsigned int invalidElementSizeBytes = GENERATE(0, 7, 12, 17); + HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, &pitch, 1, 1, invalidElementSizeBytes), + hipErrorInvalidValue); + } + + SECTION("Zero width") { + HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, &pitch, 0, 1, validElementSizeBytes), + hipErrorInvalidValue); + } + SECTION("Zero height") { + HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, &pitch, 1, 0, validElementSizeBytes), + hipErrorInvalidValue); + } +#endif + + SECTION("Invalid dptr") { + HIP_CHECK_ERROR(hipMemAllocPitch(nullptr, &pitch, 1, 1, validElementSizeBytes), + hipErrorInvalidValue); + } + + SECTION("Invalid pitch") { + HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, nullptr, 1, 1, validElementSizeBytes), + hipErrorInvalidValue); + } + + SECTION("Max size_t width") { + HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, &pitch, maxSizeT, 1, validElementSizeBytes), + hipErrorInvalidValue); + } + + SECTION("Max size_t height") { + HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, &pitch, 1, maxSizeT, validElementSizeBytes), + hipErrorOutOfMemory); + } +} + +/* +Test Scenarios of hipMallocPitch API +1. Basic Functionality Scenario +2. Allocate memory using hipMallocPitch API, Launch Kernel validate result. +3. Allocate Memory in small chunks and large chunks and check for possible memory leaks +4. Allocate Memory using hipMallocPitch API, Memcpy2D on the allocated variables. +5. Multithreaded scenario +*/ + static constexpr auto SMALLCHUNK_NUMW{4}; static constexpr auto SMALLCHUNK_NUMH{4}; static constexpr auto LARGECHUNK_NUMW{1025}; @@ -295,4 +556,3 @@ TEMPLATE_TEST_CASE("Unit_hipMallocPitch_KernelLaunch", "" A_h, B_h, C_h, false); } - From f5685158ad11c45c4c568135b865b69201a11904 Mon Sep 17 00:00:00 2001 From: Finlay Date: Fri, 1 Jul 2022 10:41:35 +0100 Subject: [PATCH 04/30] EXSWCPHIPT-91 - Added hipHostUnregister tests (#2579) --- tests/catch/unit/memory/CMakeLists.txt | 2 + tests/catch/unit/memory/hipHostUnregister.cc | 95 ++++++++++++++++++++ 2 files changed, 97 insertions(+) create mode 100644 tests/catch/unit/memory/hipHostUnregister.cc diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index 9f31f65ea5..20531a5e8a 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -39,6 +39,7 @@ set(TEST_SRC hipMemcpyAllApiNegative.cc hipMemcpy_MultiThread.cc hipHostRegister.cc + hipHostUnregister.cc hipMallocPitch.cc hipMemPtrGetInfo.cc hipPointerGetAttributes.cc @@ -108,6 +109,7 @@ set(TEST_SRC hipMemcpyAllApiNegative.cc hipMemcpy_MultiThread.cc hipHostRegister.cc + hipHostUnregister.cc hipMallocPitch.cc hipHostGetFlags.cc hipHostGetDevicePointer.cc diff --git a/tests/catch/unit/memory/hipHostUnregister.cc b/tests/catch/unit/memory/hipHostUnregister.cc new file mode 100644 index 0000000000..2721c29ad1 --- /dev/null +++ b/tests/catch/unit/memory/hipHostUnregister.cc @@ -0,0 +1,95 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +namespace hipHostUnregisterTests { +constexpr unsigned int allFlags = hipHostRegisterDefault & // 0 + hipHostRegisterPortable & // 1 + hipHostRegisterMapped & // 2 + hipHostRegisterIoMemory // 4 +#if HT_NVIDIA + & cudaHostRegisterReadOnly; // 8 +#else + ; +#endif + +inline bool hipHostRegisterSupported() { +#if HT_NVIDIA + // unable to query for cudaDevAttrHostRegisterSupported equivalent + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-40"); + HipTest::HIP_SKIP_TEST("hipHostRegister is not supported on this device"); + return false; +#else + return true; +#endif +} + + +TEST_CASE("Unit_hipHostUnregister_MemoryNotAccessableAfterUnregister") { + if (!hipHostRegisterSupported()) { + return; + } + // try all combinations of flags + for (unsigned int flag = 0; flag <= allFlags; ++flag) { + DYNAMIC_SECTION("Using flag: " << flag) { + auto x = std::unique_ptr(new int); + HIP_CHECK(hipHostRegister(x.get(), sizeof(int), flag)); + + void* device_memory; + HIP_CHECK(hipHostGetDevicePointer(&device_memory, x.get(), 0)); + + HIP_CHECK(hipHostUnregister(x.get())); + HIP_CHECK_ERROR(hipHostGetDevicePointer(&device_memory, x.get(), 0), hipErrorInvalidValue); + } + } +} + +TEST_CASE("Unit_hipHostUnregister_NullPtr") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-90"); + return; +#endif + HIP_CHECK_ERROR(hipHostUnregister(nullptr), hipErrorInvalidValue); +} + +TEST_CASE("Unit_hipHostUnregister_NotRegisteredPointer") { + auto x = std::unique_ptr(new int); + HIP_CHECK_ERROR(hipHostUnregister(x.get()), hipErrorHostMemoryNotRegistered); +} + +TEST_CASE("Unit_hipHostUnregister_AlreadyUnregisteredPointer") { + if (!hipHostRegisterSupported()) { + return; + } + // try all combinations of flags + for (unsigned int flag = 0; flag <= allFlags; ++flag) { + DYNAMIC_SECTION("Using flag: " << flag) { + auto x = std::unique_ptr(new int); + HIP_CHECK(hipHostRegister(x.get(), sizeof(int), flag)); + HIP_CHECK(hipHostUnregister(x.get())); + HIP_CHECK_ERROR(hipHostUnregister(x.get()), hipErrorHostMemoryNotRegistered); + } + } +} + +} // namespace hipHostUnregisterTests From 4d2787749c63af1cab8470d9b0c35b8c6e59311d Mon Sep 17 00:00:00 2001 From: Aiden Nibali Date: Thu, 7 Jul 2022 16:33:19 +1000 Subject: [PATCH 05/30] Fix use of blockDim in 'square' sample (#2789) --- samples/0_Intro/square/square.hipref.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/samples/0_Intro/square/square.hipref.cpp b/samples/0_Intro/square/square.hipref.cpp index b8213e0074..80ff2a00bf 100644 --- a/samples/0_Intro/square/square.hipref.cpp +++ b/samples/0_Intro/square/square.hipref.cpp @@ -38,7 +38,7 @@ THE SOFTWARE. */ template __global__ void vector_square(T* C_d, const T* A_d, size_t N) { - size_t offset = (blockIdx.x * blockDim_x + threadIdx.x); + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); size_t stride = blockDim.x * gridDim.x; for (size_t i = offset; i < N; i += stride) { From 76ce7131d989f446cb7048fd1d841741b0f4ef0a Mon Sep 17 00:00:00 2001 From: Sarbojit2019 <52527887+SarbojitAMD@users.noreply.github.com> Date: Thu, 7 Jul 2022 12:07:52 +0530 Subject: [PATCH 06/30] SWDEV-343159 - corrected macro value (#2785) --- include/hip/hip_runtime_api.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index f8a244d046..868f0de44a 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -657,7 +657,7 @@ enum hipLimit_t { #define hipDeviceScheduleBlockingSync 0x4 #define hipDeviceScheduleMask 0x7 #define hipDeviceMapHost 0x8 -#define hipDeviceLmemResizeToMax 0x16 +#define hipDeviceLmemResizeToMax 0x10 /** Default HIP array allocation flag.*/ #define hipArrayDefault 0x00 #define hipArrayLayered 0x01 From 4514f350849b1090954295f8f87a5f8d78bd781b Mon Sep 17 00:00:00 2001 From: Sarbojit2019 <52527887+SarbojitAMD@users.noreply.github.com> Date: Thu, 7 Jul 2022 12:08:01 +0530 Subject: [PATCH 07/30] Updated correct API for reading error string (#2786) --- tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp b/tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp index 1ce5c52bff..3468a1fa53 100644 --- a/tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp +++ b/tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp @@ -83,12 +83,12 @@ int main(int argc, char* argv[]) { defaultCUMask.push_back(temp); } - str_out = hipGetErrorString(hipExtStreamGetCUMask(0, cuMask.size(), 0)); + str_out = hipGetErrorName(hipExtStreamGetCUMask(0, cuMask.size(), 0)); if ((str_err.compare(str_out)) != 0) { failed("hipExtStreamGetCUMask returned wrong error code!"); } - str_out = hipGetErrorString(hipExtStreamGetCUMask(0, 0, &cuMask[0])); + str_out = hipGetErrorName(hipExtStreamGetCUMask(0, 0, &cuMask[0])); if ((str_err.compare(str_out)) != 0) { failed("hipExtStreamGetCUMask returned wrong error code!"); } From 982a3cf3879b2725fba9a3b65b7492ffb401b3cf Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 8 Jul 2022 13:35:50 +0530 Subject: [PATCH 08/30] SWDEV-306122 - [catch2][dtest] hipGraphAddMemsetNode API (#2754) Change-Id: I19e01a0f801af411b3e9ac4b57f53c41b79014a5 --- tests/catch/unit/graph/CMakeLists.txt | 1 + .../catch/unit/graph/hipGraphAddMemsetNode.cc | 123 ++++++++++++++++++ 2 files changed, 124 insertions(+) create mode 100644 tests/catch/unit/graph/hipGraphAddMemsetNode.cc diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index a8a37733d4..ba18e7fcc6 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -64,6 +64,7 @@ set(TEST_SRC hipStreamEndCapture.cc hipGraphMemcpyNodeSetParamsFromSymbol.cc hipGraphExecEventWaitNodeSetEvent.cc + hipGraphAddMemsetNode.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphAddMemsetNode.cc b/tests/catch/unit/graph/hipGraphAddMemsetNode.cc new file mode 100644 index 0000000000..0984e194c8 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphAddMemsetNode.cc @@ -0,0 +1,123 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Negative Testcase Scenarios for api hipGraphAddMemsetNode : +1) Pass pGraphNode as nullptr and check if api returns error. +2) Pass pGraphNode as un-initialize object and check. +3) Pass Graph as nullptr and check if api returns error. +4) Pass Graph as empty object(skipping graph creation), api should return error code. +5) Pass pDependencies as nullptr, api should return success. +6) Pass numDependencies is max(size_t) and pDependencies is not valid ptr, api expected to return error code. +7) Pass pDependencies is nullptr, but numDependencies is non-zero, api expected to return error. +8) Pass pMemsetParams as nullptr and check if api returns error code. +9) Pass pMemsetParams as un-initialize object and check if api returns error code. +10) Pass hipMemsetParams::dst as nullptr should return error code. +11) Pass hipMemsetParams::element size other than 1, 2, or 4 and check api should return error code. +12) Pass hipMemsetParams::height as zero and check api should return error code. +*/ + +#include + +/** + * Negative Test for API hipGraphAddMemsetNode + */ + +TEST_CASE("Unit_hipGraphAddMemsetNode_Negative") { + hipError_t ret; + hipGraph_t graph; + hipGraphNode_t memsetNode; + char *devData; + + HIP_CHECK(hipMalloc(&devData, 1024)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(devData); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = 1024; + memsetParams.height = 1; + + SECTION("Pass pGraphNode as nullptr") { + ret = hipGraphAddMemsetNode(nullptr, graph, nullptr, 0, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass pGraphNode as un-initialize object") { + hipGraphNode_t memsetNode_1; + ret = hipGraphAddMemsetNode(&memsetNode_1, graph, + nullptr, 0, &memsetParams); + REQUIRE(hipSuccess == ret); + } + SECTION("Pass graph as nullptr") { + ret = hipGraphAddMemsetNode(&memsetNode, nullptr, + nullptr, 0, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass Graph as empty object") { + hipGraph_t graph_1{}; + ret = hipGraphAddMemsetNode(&memsetNode, graph_1, + nullptr, 0, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass pDependencies as nullptr") { + ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams); + REQUIRE(hipSuccess == ret); + } + SECTION("Pass numDependencies is max and pDependencies is not valid ptr") { + ret = hipGraphAddMemsetNode(&memsetNode, graph, + nullptr, INT_MAX, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass pDependencies as nullptr, but numDependencies is non-zero") { + ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 9, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass pMemsetParams as nullptr") { + ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass pMemsetParams as un-initialize object") { + hipMemsetParams memsetParams1; + ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams1); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass hipMemsetParams::dst as nullptr") { + memsetParams.dst = nullptr; + ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass hipMemsetParams::element size other than 1, 2, or 4") { + memsetParams.dst = reinterpret_cast(devData); + memsetParams.elementSize = 9; + ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass hipMemsetParams::height as zero") { + memsetParams.elementSize = sizeof(char); + memsetParams.height = 0; + ret = hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams); + REQUIRE(hipErrorInvalidValue == ret); + } + HIP_CHECK(hipFree(devData)); + HIP_CHECK(hipGraphDestroy(graph)); +} From bcd2b7800c01b75bf1a890b6cad0a9070f7d77d4 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 8 Jul 2022 13:36:06 +0530 Subject: [PATCH 09/30] SWDEV-342404 - Update coordinates in perf test kernels (#2756) Change-Id: I33aefc54f1cf2453e77fe6a7a6b9140b343d9895 --- .../performance/compute/hipPerfDotProduct.cpp | 118 +++++++++--------- 1 file changed, 59 insertions(+), 59 deletions(-) diff --git a/tests/performance/compute/hipPerfDotProduct.cpp b/tests/performance/compute/hipPerfDotProduct.cpp index 3b2e0c72dd..e30d5ab039 100644 --- a/tests/performance/compute/hipPerfDotProduct.cpp +++ b/tests/performance/compute/hipPerfDotProduct.cpp @@ -38,7 +38,7 @@ __global__ void vectors_not_equal(int n, const double* __restrict__ x, const double* __restrict__ y, double* __restrict__ workspace) { - int gid = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; + int gid = blockIdx.x * blockDim.x + threadIdx.x; double sum = 0.0; for(int idx = gid; idx < n; idx += hipGridDim_x * hipBlockDim_x) { @@ -46,51 +46,51 @@ __global__ void vectors_not_equal(int n, } __shared__ double sdata[BLOCKSIZE]; - sdata[hipThreadIdx_x] = sum; + sdata[threadIdx.x] = sum; __syncthreads(); - if(hipThreadIdx_x < 128) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 128]; + if(threadIdx.x < 128) { + sdata[threadIdx.x] += sdata[threadIdx.x + 128]; } __syncthreads(); - if(hipThreadIdx_x < 64){ - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 64]; + if(threadIdx.x < 64){ + sdata[threadIdx.x] += sdata[threadIdx.x + 64]; } __syncthreads(); - if(hipThreadIdx_x < 32){ - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 32]; + if(threadIdx.x < 32){ + sdata[threadIdx.x] += sdata[threadIdx.x + 32]; } __syncthreads(); - if(hipThreadIdx_x < 16) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 16]; + if(threadIdx.x < 16) { + sdata[threadIdx.x] += sdata[threadIdx.x + 16]; } __syncthreads(); - if(hipThreadIdx_x < 8) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 8]; + if(threadIdx.x < 8) { + sdata[threadIdx.x] += sdata[threadIdx.x + 8]; } __syncthreads(); - if(hipThreadIdx_x < 4) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 4]; + if(threadIdx.x < 4) { + sdata[threadIdx.x] += sdata[threadIdx.x + 4]; } __syncthreads(); - if(hipThreadIdx_x < 2) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 2]; + if(threadIdx.x < 2) { + sdata[threadIdx.x] += sdata[threadIdx.x + 2]; } __syncthreads(); - if(hipThreadIdx_x < 1) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 1]; + if(threadIdx.x < 1) { + sdata[threadIdx.x] += sdata[threadIdx.x + 1]; } - if(hipThreadIdx_x == 0) { - workspace[hipBlockIdx_x] = sdata[0]; + if(threadIdx.x == 0) { + workspace[blockIdx.x] = sdata[0]; } } @@ -99,59 +99,59 @@ template __launch_bounds__(BLOCKSIZE) __global__ void vectors_equal(int n, const double* __restrict__ x, double* __restrict__ workspace) { - int gid = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; + int gid = blockIdx.x * blockDim.x + threadIdx.x; double sum = 0.0; - for(int idx = gid; idx < n; idx += hipGridDim_x * hipBlockDim_x) { + for(int idx = gid; idx < n; idx += hipGridDim_x * blockDim.x) { sum = fma(x[idx], x[idx], sum); } __shared__ double sdata[BLOCKSIZE]; - sdata[hipThreadIdx_x] = sum; + sdata[threadIdx.x] = sum; __syncthreads(); - if(hipThreadIdx_x < 128) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 128]; + if(threadIdx.x < 128) { + sdata[threadIdx.x] += sdata[threadIdx.x + 128]; } __syncthreads(); - if(hipThreadIdx_x < 64) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 64]; + if(threadIdx.x < 64) { + sdata[threadIdx.x] += sdata[threadIdx.x + 64]; } __syncthreads(); - if(hipThreadIdx_x < 32) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 32]; + if(threadIdx.x < 32) { + sdata[threadIdx.x] += sdata[threadIdx.x + 32]; } __syncthreads(); - if(hipThreadIdx_x < 16) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 16]; + if(threadIdx.x < 16) { + sdata[threadIdx.x] += sdata[threadIdx.x + 16]; } __syncthreads(); - if(hipThreadIdx_x < 8) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 8]; + if(threadIdx.x < 8) { + sdata[threadIdx.x] += sdata[threadIdx.x + 8]; } __syncthreads(); - if(hipThreadIdx_x < 4) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 4]; + if(threadIdx.x < 4) { + sdata[threadIdx.x] += sdata[threadIdx.x + 4]; } __syncthreads(); - if(hipThreadIdx_x < 2) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 2]; + if(threadIdx.x < 2) { + sdata[threadIdx.x] += sdata[threadIdx.x + 2]; } __syncthreads(); - if(hipThreadIdx_x < 1) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 1]; + if(threadIdx.x < 1) { + sdata[threadIdx.x] += sdata[threadIdx.x + 1]; } - if(hipThreadIdx_x == 0) { - workspace[hipBlockIdx_x] = sdata[0]; + if(threadIdx.x == 0) { + workspace[blockIdx.x] = sdata[0]; } } @@ -161,49 +161,49 @@ __global__ void dot_reduction(double* __restrict__ workspace) { __shared__ double sdata[BLOCKSIZE]; - sdata[hipThreadIdx_x] = workspace[hipThreadIdx_x]; + sdata[threadIdx.x] = workspace[threadIdx.x]; __syncthreads(); - if(hipThreadIdx_x < 128) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 128]; + if(threadIdx.x < 128) { + sdata[threadIdx.x] += sdata[threadIdx.x + 128]; } __syncthreads(); - if(hipThreadIdx_x < 64) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 64]; + if(threadIdx.x < 64) { + sdata[threadIdx.x] += sdata[threadIdx.x + 64]; } __syncthreads(); - if(hipThreadIdx_x < 32) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 32]; + if(threadIdx.x < 32) { + sdata[threadIdx.x] += sdata[threadIdx.x + 32]; } __syncthreads(); - if(hipThreadIdx_x < 16) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 16]; + if(threadIdx.x < 16) { + sdata[threadIdx.x] += sdata[threadIdx.x + 16]; } __syncthreads(); - if(hipThreadIdx_x < 8) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 8]; + if(threadIdx.x < 8) { + sdata[threadIdx.x] += sdata[threadIdx.x + 8]; } __syncthreads(); - if(hipThreadIdx_x < 4) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 4]; + if(threadIdx.x < 4) { + sdata[threadIdx.x] += sdata[threadIdx.x + 4]; } __syncthreads(); - if(hipThreadIdx_x < 2) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 2]; + if(threadIdx.x < 2) { + sdata[threadIdx.x] += sdata[threadIdx.x + 2]; } __syncthreads(); - if(hipThreadIdx_x < 1) { - sdata[hipThreadIdx_x] += sdata[hipThreadIdx_x + 1]; + if(threadIdx.x < 1) { + sdata[threadIdx.x] += sdata[threadIdx.x + 1]; } - if(hipThreadIdx_x == 0) { + if(threadIdx.x == 0) { workspace[0] = sdata[0]; } From 0d274931ee1e681036691e3a757e306d2dd84fcb Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 8 Jul 2022 13:36:27 +0530 Subject: [PATCH 10/30] SWDEV-322620 - Virtual Memory Management (#2772) Adding simple VM test. Change-Id: I68bb55c9439a8d0e77035833d19d72d2a21f0753 --- tests/catch/unit/memory/CMakeLists.txt | 1 + tests/catch/unit/memory/hipMemVmm.cc | 85 ++++++++++++++++++++++++++ 2 files changed, 86 insertions(+) create mode 100644 tests/catch/unit/memory/hipMemVmm.cc diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index 20531a5e8a..695dae8792 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -88,6 +88,7 @@ set(TEST_SRC hipMallocMngdMultiThread.cc hipMemPrefetchAsync.cc hipArray.cc + hipMemVmm.cc ) else() set(TEST_SRC diff --git a/tests/catch/unit/memory/hipMemVmm.cc b/tests/catch/unit/memory/hipMemVmm.cc new file mode 100644 index 0000000000..1d3041eae5 --- /dev/null +++ b/tests/catch/unit/memory/hipMemVmm.cc @@ -0,0 +1,85 @@ +/* + Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR + IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. + */ + +/* Test Case Description: + 1) This testcase verifies the basic scenario - supported on + all devices +*/ + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +/* + This testcase verifies HIP Mem VMM API basic scenario - supported on all devices + */ + +TEST_CASE("Unit_hipMemVmm_Basic") { + int vmm = 0; + HIP_CHECK(hipDeviceGetAttribute(&vmm, hipDeviceAttributeVirtualMemoryManagementSupported, 0)); + INFO("hipDeviceAttributeVirtualMemoryManagementSupported: " << vmm); + + if (vmm == 0) { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeVirtualMemoryManagement " + "attribute. Hence skipping the testing with Pass result.\n"); + return; + } + + size_t granularity = 0; + + hipMemAllocationProp memAllocationProp; + memAllocationProp.type = hipMemAllocationTypePinned; + memAllocationProp.location.id = 0; + memAllocationProp.location.type = hipMemLocationTypeDevice; + + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &memAllocationProp, hipMemAllocationGranularityRecommended)); + + size_t size = 4 * 1024; + void* reservedAddress{nullptr}; + HIP_CHECK(hipMemAddressReserve(&reservedAddress, size, granularity, nullptr, 0)); + + hipMemGenericAllocationHandle_t gaHandle{nullptr}; + HIP_CHECK(hipMemCreate(&gaHandle, size, &memAllocationProp, 0)); + + HIP_CHECK(hipMemMap(reservedAddress, size, 0, gaHandle, 0)); + + hipMemAccessDesc desc; + std::vector values(size); + const char value = 1; + + HIP_CHECK(hipMemSetAccess(reservedAddress, size, &desc, 1)); + HIP_CHECK(hipMemset(reservedAddress, value, size)); + HIP_CHECK(hipMemcpy(&values[0], reservedAddress, size, hipMemcpyDeviceToHost)); + + for (size_t i=0; i < size; ++i) { + REQUIRE(values[i] == value); + } + + HIP_CHECK(hipMemUnmap(reservedAddress, size)); + + HIP_CHECK(hipMemRelease(gaHandle)); + HIP_CHECK(hipMemAddressFree(reservedAddress, size)); +} + From a2313ade4c81071a38cb51817ad0606ff3590634 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 8 Jul 2022 13:36:41 +0530 Subject: [PATCH 11/30] SWDEV-332251 - packaging in windows and multiple arch in catchInfo.txt (#2779) Change-Id: I627662b59de002e9042f8218478404de0aa73ea5 --- tests/catch/CMakeLists.txt | 3 +++ tests/catch/packaging/hip-tests.txt | 2 +- 2 files changed, 4 insertions(+), 1 deletion(-) diff --git a/tests/catch/CMakeLists.txt b/tests/catch/CMakeLists.txt index ff12be86e3..a8a488002d 100644 --- a/tests/catch/CMakeLists.txt +++ b/tests/catch/CMakeLists.txt @@ -41,6 +41,7 @@ if(NOT DEFINED ROCM_PATH) endif() endif() endif() +file(TO_CMAKE_PATH "${ROCM_PATH}" ROCM_PATH) message(STATUS "ROCM Path: ${ROCM_PATH}") @@ -179,6 +180,8 @@ if(NOT DEFINED OFFLOAD_ARCH_STR else() message(STATUS "ROCm Agent Enumurator found no valid architectures") endif() +elseif(DEFINED OFFLOAD_ARCH_STR) + string(REPLACE "--offload-arch=" "" HIP_GPU_ARCH_LIST ${OFFLOAD_ARCH_STR}) endif() if(DEFINED OFFLOAD_ARCH_STR) diff --git a/tests/catch/packaging/hip-tests.txt b/tests/catch/packaging/hip-tests.txt index bbf3ec0ade..77b3a0e46c 100644 --- a/tests/catch/packaging/hip-tests.txt +++ b/tests/catch/packaging/hip-tests.txt @@ -30,7 +30,7 @@ install(FILES @PROJECT_BINARY_DIR@/CTestTestfile.cmake # Packaging steps ############################# set(CPACK_SET_DESTDIR TRUE) -set(CPACK_INSTALL_PREFIX @ROCM_PATH@/test/hip/) +set(CPACK_INSTALL_PREFIX "@ROCM_PATH@/test/hip/") set(PKG_NAME catch-@HIP_PLATFORM@) set(CPACK_PACKAGE_NAME ${PKG_NAME}) set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [CATCH TESTS]") From 8b90789c3a9676d8fa26d1b7b29768cf347ff821 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 8 Jul 2022 13:36:55 +0530 Subject: [PATCH 12/30] SWDEV-342570, SWDEV-340007 - update HIP documents (#2780) Change-Id: If6c2de4dd8468984cd718aea4a8613d98269a736 --- docs/markdown/hip_build.md | 2 +- docs/markdown/hip_deprecated_api_list.md | 52 +++++++++++++++++++++++- docs/markdown/hip_programming_guide.md | 3 ++ 3 files changed, 54 insertions(+), 3 deletions(-) diff --git a/docs/markdown/hip_build.md b/docs/markdown/hip_build.md index 17bc1af955..54d85fa540 100755 --- a/docs/markdown/hip_build.md +++ b/docs/markdown/hip_build.md @@ -141,7 +141,7 @@ After build and install HIP commands, catch tests can be built via the following ``` cd "$HIP_DIR" mkdir -p build; cd build -export HIP_PATH=$HIPAMD_DIR/build +export HIP_PATH=$HIPAMD_DIR/build/install cmake ../tests/catch/ -DHIP_PLATFORM=amd make -j$(nproc) build_tests ctest # run tests diff --git a/docs/markdown/hip_deprecated_api_list.md b/docs/markdown/hip_deprecated_api_list.md index 783ea4a390..b26aceecee 100644 --- a/docs/markdown/hip_deprecated_api_list.md +++ b/docs/markdown/hip_deprecated_api_list.md @@ -1,9 +1,11 @@ # HIP Deprecated APIs + ## HIP Context Management APIs CUDA supports cuCtx API, the Driver API that defines "Context" and "Devices" as separate entities. Contexts contain a single device, and a device can theoretically have multiple contexts. HIP initially added limited support for these API to facilitate easy porting from existing driver codes. These API are marked as deprecated now since there are better alternate interface (such as hipSetDevice or the stream API) to achieve the required functions. - +### hipCtxCreate +### hipCtxDestroy ### hipCtxPopCurrent ### hipCtxPushCurrent ### hipCtxSetCurrent @@ -19,6 +21,7 @@ CUDA supports cuCtx API, the Driver API that defines "Context" and "Devices" as ### hipCtxEnablePeerAccess ### hipCtxDisablePeerAccess + ## HIP Memory Management APIs ### hipMallocHost @@ -31,4 +34,49 @@ Should use "hipHostMalloc" instead. Should use "hipHostMalloc" instead. ### hipFreeHost -Should use "hipHostFree" instead. \ No newline at end of file +Should use "hipHostFree" instead. + +### hipMemcpyToArray +### hipMemcpyFromArray + + +## HIP Profiler Control APIs + +### hipProfilerStart +Should use roctracer/rocTX instead + +### hipProfilerStop +Should use roctracer/rocTX instead + + +## HIP Texture Management APIs + +###hipGetTextureReference +###hipTexRefSetAddressMode +###hipTexRefSetArray +###hipTexRefSetFilterMode +###hipTexRefSetFlags +###hipTexRefSetFormat +###hipBindTexture +###hipBindTexture2D +###hipBindTextureToArray +###hipGetTextureAlignmentOffset +###hipUnbindTexture +###hipTexRefGetAddress +###hipTexRefGetAddressMode +###hipTexRefGetFilterMode +###hipTexRefGetFlags +###hipTexRefGetFormat +###hipTexRefGetMaxAnisotropy +###hipTexRefGetMipmapFilterMode +###hipTexRefGetMipmapLevelBias +###hipTexRefGetMipmapLevelClamp +###hipTexRefGetMipMappedArray +###hipTexRefSetAddress +###hipTexRefSetAddress2D +###hipTexRefSetMaxAnisotropy +###hipTexRefSetBorderColor +###hipTexRefSetMipmapFilterMode +###hipTexRefSetMipmapLevelBias +###hipTexRefSetMipmapLevelClamp +###hipTexRefSetMipmappedArray \ No newline at end of file diff --git a/docs/markdown/hip_programming_guide.md b/docs/markdown/hip_programming_guide.md index e9d25b06bc..9f48870300 100644 --- a/docs/markdown/hip_programming_guide.md +++ b/docs/markdown/hip_programming_guide.md @@ -129,6 +129,9 @@ For more details on hipRTC APIs, refer to HIP-API.pdf in GitHub (https://github. The link here(https://github.com/ROCm-Developer-Tools/HIP/blob/main/tests/src/hiprtc/saxpy.cpp) shows an example how to program HIP application using runtime compilation mechanism, and detail hipRTC programming guide is also available in Github (https://github.com/ROCm-Developer-Tools/HIP/blob/main/docs/markdown/hip_rtc.md). +## HIP Graph +HIP graph is supported. For more details, refer to the HIP API Guide. + ## Device-Side Malloc HIP-Clang now supports device-side malloc and free. From ab89350ab68717330d04ee4ec8f40a7f120c4b9a Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 8 Jul 2022 15:18:09 +0530 Subject: [PATCH 13/30] SWDEV-306122 - [catch2][dtest] hipGraphAddKernelNode Added negative test (#2734) Change-Id: I0d6353f77c3c2abda7557f8268a024b1b4d8c1d7 --- tests/catch/unit/graph/CMakeLists.txt | 1 + .../catch/unit/graph/hipGraphAddKernelNode.cc | 101 ++++++++++++++++++ 2 files changed, 102 insertions(+) create mode 100644 tests/catch/unit/graph/hipGraphAddKernelNode.cc diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index ba18e7fcc6..dc2344638b 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -65,6 +65,7 @@ set(TEST_SRC hipGraphMemcpyNodeSetParamsFromSymbol.cc hipGraphExecEventWaitNodeSetEvent.cc hipGraphAddMemsetNode.cc + hipGraphAddKernelNode.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphAddKernelNode.cc b/tests/catch/unit/graph/hipGraphAddKernelNode.cc new file mode 100644 index 0000000000..79ed3c76ea --- /dev/null +++ b/tests/catch/unit/graph/hipGraphAddKernelNode.cc @@ -0,0 +1,101 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +/* Test verifies hipGraphAddKernelNode API Negative scenarios. + */ + +TEST_CASE("Unit_hipGraphAddKernelNode_Negative") { + constexpr int N = 1024; + size_t NElem{N}; + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + int *A_d, *B_d, *C_d; + hipGraph_t graph; + hipError_t ret; + hipGraphNode_t kNode; + hipKernelNodeParams kNodeParams{}; + std::vector dependencies; + + HIP_CHECK(hipMalloc(&A_d, sizeof(int) * N)); + HIP_CHECK(hipMalloc(&B_d, sizeof(int) * N)); + HIP_CHECK(hipMalloc(&C_d, sizeof(int) * N)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.gridDim = dim3(blocks); + kNodeParams.blockDim = dim3(threadsPerBlock); + kNodeParams.sharedMemBytes = 0; + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kNodeParams.extra = nullptr; + + SECTION("Pass pGraphNode as nullptr") { + ret = hipGraphAddKernelNode(nullptr, graph, nullptr, 0, &kNodeParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass Graph as nullptr") { + ret = hipGraphAddKernelNode(&kNode, nullptr, nullptr, 0, &kNodeParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass invalid numDependencies") { + ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 11, &kNodeParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass invalid numDependencies and valid list for dependencies") { + HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams)); + dependencies.push_back(kNode); + ret = hipGraphAddKernelNode(&kNode, graph, + dependencies.data(), dependencies.size()+1, &kNodeParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass NodeParams as nullptr") { + ret = hipGraphAddKernelNode(&kNode, graph, + dependencies.data(), dependencies.size(), nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass NodeParams func datamember as nullptr") { + kNodeParams.func = nullptr; + ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams); + REQUIRE(hipSuccess != ret); + } + SECTION("Pass kernelParams datamember as nullptr") { + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.kernelParams = nullptr; + ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams); + REQUIRE(hipErrorInvalidValue == ret); + } +#if HT_AMD +// On Cuda setup this test case getting failed + SECTION("Try adding kernel node after destroy the already created graph") { + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); + HIP_CHECK(hipGraphDestroy(graph)); + ret = hipGraphAddKernelNode(&kNode, graph, nullptr, 0, &kNodeParams); + REQUIRE(hipErrorInvalidValue == ret); + } +#endif + + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(B_d)); + HIP_CHECK(hipFree(C_d)); +} + From 778bf8631d360481a96f7911c38dd5760dff4d74 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 8 Jul 2022 15:19:02 +0530 Subject: [PATCH 14/30] SWDEV-306122 - [catch2][dtest] Added tests for hipGraphMemcpyNodeGetParams/hipGraphMemcpyNodeSetParams API. (#2775) Change-Id: Ie0e323a545c3bac17a35814a90b93ecd44cec537 --- .../config/config_amd_windows.json | 3 +- tests/catch/unit/graph/CMakeLists.txt | 2 + .../unit/graph/hipGraphMemcpyNodeGetParams.cc | 232 ++++++++++++++++++ .../unit/graph/hipGraphMemcpyNodeSetParams.cc | 215 ++++++++++++++++ 4 files changed, 451 insertions(+), 1 deletion(-) create mode 100644 tests/catch/unit/graph/hipGraphMemcpyNodeGetParams.cc create mode 100644 tests/catch/unit/graph/hipGraphMemcpyNodeSetParams.cc diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 2d92fd5752..58f6439f5e 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -29,6 +29,7 @@ "Unit_hipDeviceGetUuid", "Unit_hipGraphMemcpyNodeSetParamsFromSymbol_Functional", "Unit_hipGraphExecEventWaitNodeSetEvent_Negative", - "Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory" + "Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory", + "Unit_hipGraphMemcpyNodeSetParams_Functional" ] } diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index dc2344638b..33a6ce3cd5 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -66,6 +66,8 @@ set(TEST_SRC hipGraphExecEventWaitNodeSetEvent.cc hipGraphAddMemsetNode.cc hipGraphAddKernelNode.cc + hipGraphMemcpyNodeGetParams.cc + hipGraphMemcpyNodeSetParams.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphMemcpyNodeGetParams.cc b/tests/catch/unit/graph/hipGraphMemcpyNodeGetParams.cc new file mode 100644 index 0000000000..6393358bb2 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphMemcpyNodeGetParams.cc @@ -0,0 +1,232 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : +Negative - +1) Pass node as nullptr and verify api returns error code. +2) Pass un-initialize node and verify api returns error code. +3) Pass pNodeParams as nullptr and verify api returns error code. +Functional - +1) Create a graph, add Memcpy node to graph with desired node params. + Verify api fetches the node params mentioned while adding Memcpy node. +2) Set Memcpy node params with hipGraphMemcpyNodeSetParams, + now get the params and verify both are same. +*/ + +#include +#include + +#define SIZE 10 +#define UPDATESIZE 8 + +/* Test verifies hipGraphMemcpyNodeGetParams API Negative scenarios. + */ +TEST_CASE("Unit_hipGraphMemcpyNodeGetParams_Negative") { + constexpr int width{SIZE}, height{SIZE}, depth{SIZE}; + hipArray *devArray; + hipChannelFormatKind formatKind = hipChannelFormatKindSigned; + hipMemcpy3DParms myparms; + int* hData; + uint32_t size = width * height * depth * sizeof(int); + hData = reinterpret_cast(malloc(size)); + REQUIRE(hData != nullptr); + memset(hData, 0, size); + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + hData[i*width*height + j*width + k] = i*width*height + j*width + k; + } + } + } + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8, + 0, 0, 0, formatKind); + HIP_CHECK(hipMalloc3DArray(&devArray, &channelDesc, make_hipExtent(width, + height, depth), hipArrayDefault)); + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = make_hipExtent(width , height, depth); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int), + width, height); + myparms.dstArray = devArray; + myparms.kind = hipMemcpyHostToDevice; + + hipGraph_t graph; + hipError_t ret; + hipGraphNode_t memcpyNode; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, NULL, 0, &myparms)); + + SECTION("Pass node as nullptr") { + ret = hipGraphMemcpyNodeGetParams(nullptr, &myparms); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass un-initilize node") { + hipGraphNode_t memcpyNode_uninit{}; + ret = hipGraphMemcpyNodeGetParams(memcpyNode_uninit, &myparms); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass GetNodeParams as nullptr") { + ret = hipGraphMemcpyNodeGetParams(memcpyNode, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } + HIP_CHECK(hipFreeArray(devArray)); + free(hData); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/* Test verifies hipGraphMemcpyNodeGetParams API Functional scenarios. + */ + +static bool compareHipPos(hipPos hPos1, hipPos hPos2) { + if ((hPos1.x == hPos2.x) && (hPos1.y == hPos2.y) && (hPos1.z == hPos2.z)) + return true; + else + return false; +} +static bool compareHipExtent(hipExtent hExt1, hipExtent hExt2) { + if ((hExt1.width == hExt2.width) && (hExt1.height == hExt2.height) && + (hExt1.depth == hExt2.depth)) + return true; + else + return false; +} +static bool compareHipPitchedPtr(hipPitchedPtr hpPtr1, hipPitchedPtr hpPtr2) { + if ((reinterpret_cast(hpPtr1.ptr) == + reinterpret_cast(hpPtr2.ptr)) + && (hpPtr1.pitch == hpPtr2.pitch) + #if HT_AMD + && (hpPtr1.xsize == hpPtr2.xsize) + /* xsize check below is disabled on nvidia as xsize value + * is not being updated properly due to issue with CUDA api */ + #endif + && (hpPtr1.ysize == hpPtr2.ysize)) + return true; + else + return false; +} + +static bool memcpyNodeCompare(hipMemcpy3DParms *mNode1, + hipMemcpy3DParms *mNode2) { + if (mNode1->srcArray != mNode2->srcArray) + return false; + if (!compareHipPos(mNode1->srcPos, mNode2->srcPos)) + return false; + if (!compareHipPitchedPtr(mNode1->srcPtr, mNode2->srcPtr)) + return false; + if (mNode1->dstArray != mNode2->dstArray) + return false; + if (!compareHipPos(mNode1->dstPos, mNode2->dstPos)) + return false; + if (!compareHipPitchedPtr(mNode1->dstPtr, mNode2->dstPtr)) + return false; + if (!compareHipExtent(mNode1->extent, mNode2->extent)) + return false; + if (mNode1->kind != mNode2->kind) + return false; + return true; +} + +TEST_CASE("Unit_hipGraphMemcpyNodeGetParams_Functional") { + constexpr int width{SIZE}, height{SIZE}, depth{SIZE}; + hipArray *devArray; + hipChannelFormatKind formatKind = hipChannelFormatKindSigned; + hipMemcpy3DParms myparms; + int* hData; + uint32_t size = width * height * depth * sizeof(int); + hData = reinterpret_cast(malloc(size)); + REQUIRE(hData != nullptr); + memset(hData, 0, size); + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + hData[i*width*height + j*width + k] = i*width*height + j*width + k; + } + } + } + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8, + 0, 0, 0, formatKind); + HIP_CHECK(hipMalloc3DArray(&devArray, &channelDesc, make_hipExtent(width, + height, depth), hipArrayDefault)); + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = make_hipExtent(width , height, depth); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int), + width, height); + myparms.dstArray = devArray; + myparms.kind = hipMemcpyHostToDevice; + + hipGraph_t graph; + hipGraphNode_t memcpyNode; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, NULL, 0, &myparms)); + + SECTION("Get Memcpy Param and verify.") { + hipMemcpy3DParms m3DGetParams; + REQUIRE(hipSuccess == hipGraphMemcpyNodeGetParams(memcpyNode, + &m3DGetParams)); + // Validating the result + REQUIRE(true == memcpyNodeCompare(&myparms, &m3DGetParams)); + } + + SECTION("Set memcpy params and Get param and verify.") { + hipMemcpy3DParms myparms1, m3DGetParams1; + constexpr int width1{UPDATESIZE}, height1{UPDATESIZE}, depth1{UPDATESIZE}; + hipArray *devArray1; + hipChannelFormatKind formatKind1 = hipChannelFormatKindSigned; + int* hData1; + uint32_t size1 = width1 * height1 * depth1 * sizeof(int); + hData1 = reinterpret_cast(malloc(size1)); + REQUIRE(hData1 != nullptr); + memset(hData1, 0, size1); + for (int i = 0; i < depth1; i++) { + for (int j = 0; j < height1; j++) { + for (int k = 0; k < width1; k++) { + hData1[i*width1*height1 + j*width1 + k] = i*width1*height1 + + j*width1 + k; + } + } + } + hipChannelFormatDesc channelDesc1 = hipCreateChannelDesc(sizeof(int)*8, + 0, 0, 0, formatKind1); + HIP_CHECK(hipMalloc3DArray(&devArray1, &channelDesc1, + make_hipExtent(width1, height1, depth1), hipArrayDefault)); + memset(&myparms1, 0x0, sizeof(hipMemcpy3DParms)); + myparms1.srcPos = make_hipPos(0, 0, 0); + myparms1.dstPos = make_hipPos(0, 0, 0); + myparms1.extent = make_hipExtent(width1 , height1, depth1); + myparms1.srcPtr = make_hipPitchedPtr(hData1, width1 * sizeof(int), + width1, height1); + myparms1.dstArray = devArray1; + myparms1.kind = hipMemcpyHostToDevice; + + REQUIRE(hipSuccess == hipGraphMemcpyNodeSetParams(memcpyNode, &myparms1)); + REQUIRE(hipSuccess == hipGraphMemcpyNodeGetParams(memcpyNode, + &m3DGetParams1)); + REQUIRE(true == memcpyNodeCompare(&myparms1, &m3DGetParams1)); + + HIP_CHECK(hipFreeArray(devArray1)); + free(hData1); + } + HIP_CHECK(hipFreeArray(devArray)); + free(hData); + HIP_CHECK(hipGraphDestroy(graph)); +} diff --git a/tests/catch/unit/graph/hipGraphMemcpyNodeSetParams.cc b/tests/catch/unit/graph/hipGraphMemcpyNodeSetParams.cc new file mode 100644 index 0000000000..12cf7e22eb --- /dev/null +++ b/tests/catch/unit/graph/hipGraphMemcpyNodeSetParams.cc @@ -0,0 +1,215 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : +Negative - +1) Pass node as nullptr and verify api returns error code. +2) Pass un-initialize node and verify api returns error code. +3) Pass pNodeParams as nullptr and verify api returns error code. +Functional - +1) Add Memcpy node to graph, update the Memcpy node params with set and + launch the graph and check updated params are taking effect. +2) Add Memcpy node to graph, launch graph, then update the Memcpy node params + with set and launch the graph and check updated params are taking effect. +*/ + +#include +#include + +#define SIZE 10 + +/* Test verifies hipGraphMemcpyNodeSetParams API Negative scenarios. + */ +TEST_CASE("Unit_hipGraphMemcpyNodeSetParams_Negative") { + constexpr int width{SIZE}, height{SIZE}, depth{SIZE}; + hipArray *devArray; + hipChannelFormatKind formatKind = hipChannelFormatKindSigned; + hipMemcpy3DParms myparms; + int* hData; + uint32_t size = width * height * depth * sizeof(int); + hData = reinterpret_cast(malloc(size)); + REQUIRE(hData != nullptr); + memset(hData, 0, size); + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + hData[i*width*height + j*width + k] = i*width*height + j*width + k; + } + } + } + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8, + 0, 0, 0, formatKind); + HIP_CHECK(hipMalloc3DArray(&devArray, &channelDesc, make_hipExtent(width, + height, depth), hipArrayDefault)); + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = make_hipExtent(width , height, depth); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int), + width, height); + myparms.dstArray = devArray; + myparms.kind = hipMemcpyHostToDevice; + + hipGraph_t graph; + hipError_t ret; + hipGraphNode_t memcpyNode; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, NULL, 0, &myparms)); + + SECTION("Pass node as nullptr") { + ret = hipGraphMemcpyNodeSetParams(nullptr, &myparms); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass un-initialize node") { + hipGraphNode_t memcpyNode_uninit{}; + ret = hipGraphMemcpyNodeSetParams(memcpyNode_uninit, &myparms); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass SetNodeParams as nullptr") { + ret = hipGraphMemcpyNodeSetParams(memcpyNode, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } + HIP_CHECK(hipFreeArray(devArray)); + free(hData); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/* Test verifies hipGraphMemcpyNodeSetParams API Functional scenarios. + */ +TEST_CASE("Unit_hipGraphMemcpyNodeSetParams_Functional") { + constexpr int width{SIZE}, height{SIZE}, depth{SIZE}; + hipArray *devArray; + hipChannelFormatKind formatKind = hipChannelFormatKindSigned; + hipMemcpy3DParms myparms, myparms1; + uint32_t size = width * height * depth * sizeof(int); + + int *hData = reinterpret_cast(malloc(size)); + REQUIRE(hData != nullptr); + memset(hData, 0, size); + int *hDataTemp = reinterpret_cast(malloc(size)); + REQUIRE(hDataTemp != nullptr); + memset(hDataTemp, 0, size); + int *hOutputData = reinterpret_cast(malloc(size)); + REQUIRE(hOutputData != nullptr); + memset(hOutputData, 0, size); + int *hOutputData1 = reinterpret_cast(malloc(size)); + REQUIRE(hOutputData1 != nullptr); + memset(hOutputData1, 0, size); + + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + hData[i*width*height + j*width + k] = i*width*height + j*width + k; + } + } + } + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8, + 0, 0, 0, formatKind); + HIP_CHECK(hipMalloc3DArray(&devArray, &channelDesc, make_hipExtent(width, + height, depth), hipArrayDefault)); + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + + // Host to Device + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = make_hipExtent(width , height, depth); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int), + width, height); + myparms.dstArray = devArray; + myparms.kind = hipMemcpyHostToDevice; + + hipGraph_t graph; + hipGraphNode_t memcpyNode; + std::vector dependencies; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, NULL, 0, &myparms)); + dependencies.push_back(memcpyNode); + + // Device to host + memset(&myparms1, 0x0, sizeof(hipMemcpy3DParms)); + myparms1.srcPos = make_hipPos(0, 0, 0); + myparms1.dstPos = make_hipPos(0, 0, 0); + myparms1.dstPtr = make_hipPitchedPtr(hDataTemp, width * sizeof(int), + width, height); + myparms1.srcArray = devArray; + myparms1.extent = make_hipExtent(width, height, depth); + myparms1.kind = hipMemcpyDeviceToHost; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(), + dependencies.size(), &myparms1)); + + SECTION("Update the memcpyNode and check") { + // Device to host with updated host ptr hDataTemp -> hOutputData + memset(&myparms1, 0x0, sizeof(hipMemcpy3DParms)); + myparms1.srcPos = make_hipPos(0, 0, 0); + myparms1.dstPos = make_hipPos(0, 0, 0); + myparms1.dstPtr = make_hipPitchedPtr(hOutputData, width * sizeof(int), + width, height); + myparms1.srcArray = devArray; + myparms1.extent = make_hipExtent(width, height, depth); + myparms1.kind = hipMemcpyDeviceToHost; + + HIP_CHECK(hipGraphMemcpyNodeSetParams(memcpyNode, &myparms1)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Check result + HipTest::checkArray(hData, hOutputData, width, height, depth); + } + + SECTION("Update the memcpyNode again and check") { + // Device to host with updated host ptr hOutputData -> hOutputData1 + memset(&myparms1, 0x0, sizeof(hipMemcpy3DParms)); + myparms1.srcPos = make_hipPos(0, 0, 0); + myparms1.dstPos = make_hipPos(0, 0, 0); + myparms1.dstPtr = make_hipPitchedPtr(hOutputData1, width * sizeof(int), + width, height); + myparms1.srcArray = devArray; + myparms1.extent = make_hipExtent(width, height, depth); + myparms1.kind = hipMemcpyDeviceToHost; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(), + dependencies.size(), &myparms1)); + HIP_CHECK(hipGraphMemcpyNodeSetParams(memcpyNode, &myparms1)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Check result + HipTest::checkArray(hData, hOutputData1, width, height, depth); + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipFreeArray(devArray)); + free(hData); + free(hDataTemp); + free(hOutputData); + free(hOutputData1); +} From 903035a0affb5bf1cb635a4b69b924aa2776a4ef Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 8 Jul 2022 15:19:24 +0530 Subject: [PATCH 15/30] SWDEV-314080 - Adding test cases based on hipStreamPerThread stream object. (#2781) Change-Id: I9b60f4ecdc3196d5dfdb3d7f484f0b5429aae8b3 --- .../catch/unit/streamperthread/CMakeLists.txt | 1 + .../streamperthread/hipStreamPerThrdTsts.cc | 596 ++++++++++++++++++ 2 files changed, 597 insertions(+) create mode 100644 tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc diff --git a/tests/catch/unit/streamperthread/CMakeLists.txt b/tests/catch/unit/streamperthread/CMakeLists.txt index a47e1e8c69..dd5b54dabc 100644 --- a/tests/catch/unit/streamperthread/CMakeLists.txt +++ b/tests/catch/unit/streamperthread/CMakeLists.txt @@ -4,6 +4,7 @@ set(TEST_SRC hipStreamPerThread_Event.cc hipStreamPerThread_MultiThread.cc hipStreamPerThread_DeviceReset.cc + hipStreamPerThrdTsts.cc ) hip_add_exe_to_target(NAME StreamPerThreadTest diff --git a/tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc b/tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc new file mode 100644 index 0000000000..730dcbef6a --- /dev/null +++ b/tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc @@ -0,0 +1,596 @@ +/* +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. +*/ + + +/* Test Description: + Scenario-1: Launch a kernel in hipStreamPerThread, while it is in flight + check for hipStreamQuery(hipStreamPerThread) it should return + hipErrorNotReady. + Scenario-2: Testing hipStreamPerThread stream object with hipMallocManaged() + memory + Scenario-3: To check the working of hipStreamPerThread in forked process + Scenario-4: The following test case tests the working of hipEventSynchronize + in multiple threads which are launched in quick succession + Scenario-5: The following test case checks the working of + hipStreamWaitEvent() with hipStreamWaitEvent() + Scenario-6: Testing hipLaunchCooperativeKernel() api with hipStreamPerThread + Scenario-7: Testing hipLaunchCooperativeKernelMultiDevice() with + hipStreamPerThread +*/ +#include +#include +#include +#ifdef _WIN32 + #include +#endif +#ifdef __linux__ + #include + #include + #include +#endif + +#include +#ifdef HT_AMD + #include "hip/hip_cooperative_groups.h" +#endif +using namespace std::chrono; +using namespace cooperative_groups; +#if HT_AMD +#define HIPRT_CB +#endif + + +static bool IfTestPassed = false; +// kernel +__global__ void StreamPerThrd(int *Ad, int *Ad1, size_t n, int Pk_Clk, + int Wait, int WaitEvnt = 0) { + size_t index = blockIdx.x * blockDim.x + threadIdx.x; + if (index < n) { + Ad[index] = Ad[index] + 10; + } + if (Wait) { + int64_t GpuFrq = (Pk_Clk * 1000); + int64_t StrtTck = clock64(); + if (index == 0) { + // The following while loop checks the value in ptr for around 4 seconds + while ((clock64() - StrtTck) <= (6 * GpuFrq)) { + } + if (WaitEvnt == 1) { + *Ad1 = 1; + } + } + } +} + + +__global__ void StreamPerThrd1(int *A, int Pk_Clk) { + int64_t GpuFrq = (Pk_Clk * 1000); + int64_t StrtTck = clock64(); + // The following while loop checks the value in ptr for around 3-4 seconds + while ((clock64() - StrtTck) <= (3 * GpuFrq)) { + } + *A = 1; +} + +__global__ void MiniKernel(int *A) { + if (*A == 0) { + *A = 2; // Fail condition + } else if (*A == 1) { + *A = 3; // Pass condition + } else { + *A = 4; // Garbage value found in A + } +} + +__global__ void StreamPerThrdCoopKrnl(int *Ad, int *n) { + int NumElms = (*n); + int index = blockIdx.x * blockDim.x + threadIdx.x; + if (index < NumElms) { + Ad[index] = Ad[index] + 10; + } +} + +#if HT_AMD +__global__ void test_gwsPerThrd(uint* buf, uint bufSize, int64_t* tmpBuf, + int64_t* result) { + extern __shared__ int64_t tmp[]; + uint groups = gridDim.x; + uint group_id = blockIdx.x; + uint local_id = threadIdx.x; + uint chunk = gridDim.x * blockDim.x; + + uint i = group_id * blockDim.x + local_id; + int64_t sum = 0; + while (i < bufSize) { + sum += buf[i]; + i += chunk; + } + tmp[local_id] = sum; + __syncthreads(); + i = 0; + if (local_id == 0) { + sum = 0; + while (i < blockDim.x) { + sum += tmp[i]; + i++; + } + tmpBuf[group_id] = sum; + } + + // wait + cooperative_groups::this_grid().sync(); + + if (((blockIdx.x * blockDim.x) + threadIdx.x) == 0) { + for (uint i = 1; i < groups; ++i) { + sum += tmpBuf[i]; + } + // *result = sum; + result[1 + cooperative_groups::this_multi_grid().grid_rank()] = sum; + } + cooperative_groups::this_multi_grid().sync(); + if (cooperative_groups::this_multi_grid().grid_rank() == 0) { + sum = 0; + for (uint i = 1; i <= cooperative_groups::this_multi_grid().num_grids(); + ++i) { + sum += result[i]; + } + *result = sum; + } +} +#endif +static const uint BufferSizeInDwords = 256 * 1024 * 1024; +static constexpr uint NumKernelArgs = 4; +static constexpr uint MaxGPUs = 8; +// callback function +static void HIPRT_CB CallBackFunctn(hipStream_t strm, hipError_t err, + void *ChkVal) { + // The following HIPASSERT() is just to satisfy catch2 framework. + // As it ensures the use of all the variables. + HIPASSERT(strm); + HIPCHECK(err); + if (*(reinterpret_cast(ChkVal)) == 1) { + IfTestPassed = true; + } else { + IfTestPassed = false; + } +} + +static void EventSync() { + int *Ad = nullptr, *Ah = nullptr, NumElms = 4096, CONST = 123; + int blockSize = 32, peak_clk; + HIP_CHECK(hipMalloc(&Ad, NumElms * sizeof(int))); + Ah = new int[NumElms]; + for (int i = 0; i < NumElms; ++i) { + Ah[i] = CONST; + } + // creating event objects + hipEvent_t start, end; + HIP_CHECK(hipEventCreate(&start)); + HIP_CHECK(hipEventCreate(&end)); + HIP_CHECK(hipMemcpy(Ad, Ah, NumElms * sizeof(int), hipMemcpyHostToDevice)); + HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0)); + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + HIP_CHECK(hipEventRecord(start, hipStreamPerThread)); + StreamPerThrd<<>>(Ad, NULL, NumElms, + peak_clk, 0); + HIP_CHECK(hipEventRecord(end, hipStreamPerThread)); + HIP_CHECK(hipEventSynchronize(end)); + HIP_CHECK(hipMemcpy(Ah, Ad, NumElms * sizeof(int), hipMemcpyDeviceToHost)); + int MisMatch = 0; + for (int i = 0; i < NumElms; ++i) { + if (Ah[i] != (CONST + 10)) { + MisMatch++; + } + } + delete[] Ah; + HIP_CHECK(hipFree(Ad)); + if (MisMatch) { + WARN("Data Mismatch observed!!\n"); + IfTestPassed = false; + } else { + IfTestPassed = true; + } +} + +/* Launch a kernel in hipStreamPerThread, while it is in flight check for + hipStreamQuery(hipStreamPerThread) it should return hipErrorNotReady.*/ +TEST_CASE("Unit_hipStreamPerThreadTst_StrmQuery") { + int *Ad = nullptr, *Ah = nullptr, NumElms = 4096, CONST = 123; + int blockSize = 32, peak_clk; + hipError_t err; + HIP_CHECK(hipMalloc(&Ad, NumElms * sizeof(int))); + Ah = new int[NumElms]; + for (int i = 0; i < NumElms; ++i) { + Ah[i] = CONST; + } + HIP_CHECK(hipMemcpy(Ad, Ah, NumElms * sizeof(int), hipMemcpyHostToDevice)); + HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0)); + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + SECTION("Test working of hipStreamQuery") { + StreamPerThrd<<>>(Ad, NULL, + NumElms, peak_clk, 1); + err = hipStreamQuery(hipStreamPerThread); + if (err != hipErrorNotReady) { + WARN("hipStreamQuery on hipStreamPerThread didnt return expected error!"); + IfTestPassed = false; + } else { + IfTestPassed = true; + } + } + SECTION("check working of hipStreamAddCallback() with hipStreamPerThread") { + int *Hptr = nullptr, *A_d = nullptr; + HIP_CHECK(hipHostMalloc(&Hptr, sizeof(int))); + *Hptr = 0; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), Hptr, 0)); + StreamPerThrd1<<<1, 1, 0, hipStreamPerThread>>>(A_d, peak_clk); + HIP_CHECK(hipStreamAddCallback(hipStreamPerThread, CallBackFunctn, A_d, 0)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + HIP_CHECK(hipHostFree(Hptr)); + } + HIP_CHECK(hipFree(Ad)); + delete[] Ah; + REQUIRE(IfTestPassed); +} + +/* Testing hipStreamPerThread stream object with hipMallocManaged() memory*/ +TEST_CASE("Unit_hipStreamPerThread_MangdMem") { + int managed = 0; + HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, + 0)); + if (managed == 1) { + int *Hmm = nullptr, NumElms = 4096, CONST = 123, blockSize = 32; + SECTION("Using Managed memory") { + HIP_CHECK(hipMallocManaged(&Hmm, NumElms * sizeof(int))); + for (int i = 0; i < NumElms; ++i) { + Hmm[i] = CONST; + } + } + SECTION("Prefetching Managed memory to device") { + HIP_CHECK(hipMallocManaged(&Hmm, NumElms * sizeof(int))); + for (int i = 0; i < NumElms; ++i) { + Hmm[i] = CONST; + } + HIP_CHECK(hipMemPrefetchAsync(Hmm, NumElms * sizeof(int), 0, + hipStreamPerThread)); + } + int peak_clk; + HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0)); + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + StreamPerThrd<<>>(Hmm, NULL, + NumElms, peak_clk, 0); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + // Validating the result + int MisMatch = 0; + for (int i = 0; i < NumElms; ++i) { + if (Hmm[i] != (CONST + 10)) { + MisMatch++; + } + } + HIP_CHECK(hipFree(Hmm)); + if (MisMatch) { + WARN("Data mismatch observed!!\n"); + REQUIRE(false); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +/* To check the working of hipStreamPerThread in forked process*/ +#ifdef __linux__ +TEST_CASE("Unit_hipStreamPerThread_ChildProc") { + if (fork() == 0) { // child process + int *Ad = nullptr, *Ah = nullptr, NumElms = 4096, CONST = 123; + int blockSize = 32, peak_clk; + HIP_CHECK(hipMalloc(&Ad, NumElms * sizeof(int))); + Ah = new int[NumElms]; + for (int i = 0; i < NumElms; ++i) { + Ah[i] = CONST; + } + HIP_CHECK(hipMemcpy(Ad, Ah, NumElms * sizeof(int), hipMemcpyHostToDevice)); + HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0)); + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + StreamPerThrd<<>>(Ad, NULL, + NumElms, peak_clk, 0); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + HIP_CHECK(hipMemcpy(Ah, Ad, NumElms * sizeof(int), hipMemcpyDeviceToHost)); + int MisMatch = 0; + for (int i = 0; i < NumElms; ++i) { + if (Ah[i] != (CONST + 10)) { + MisMatch++; + } + } + delete[] Ah; + HIP_CHECK(hipFree(Ad)); + if (MisMatch) { + WARN("Data Mismatch observed!!\n"); + exit(9); + } else { + exit(10); + } + } else { // Parent process + int stat; + wait(&stat); + int Result = WEXITSTATUS(stat); + if (Result != 10) { + REQUIRE(false); + } + } +} +#endif + +/* The following test case tests the working of hipEventSynchronize in + multiple threads which are launched in quick succession*/ +TEST_CASE("Unit_hipStreamPerThread_EvtRcrdMThrd") { + IfTestPassed = true; + int MAX_THREAD_CNT = 20; + std::vector threads(MAX_THREAD_CNT); + for (auto &th : threads) { + th = std::thread(EventSync); + } + for (auto& th : threads) { + th.join(); + } + REQUIRE(IfTestPassed); +} + +/* The following test case checks the working of hipStreamWaitEvent() with + hipStreamWaitEvent()*/ +TEST_CASE("Unit_hipStreamPerThread_StrmWaitEvt") { + IfTestPassed = true; + int *Ad = nullptr, NumElms = 4096, CONST = 123, blockSize = 32, *Ah = nullptr; + int *Ad1 = nullptr, *Ah1 = nullptr; + Ah = new int[NumElms]; + Ah1 = new int; + hipStream_t Strm; + HIP_CHECK(hipStreamCreate(&Strm)); + for (int i = 0; i < NumElms; ++i) { + Ah[i] = CONST; + } + Ah1[0] = 0; + HIP_CHECK(hipMalloc(&Ad, NumElms * sizeof(int))); + HIP_CHECK(hipMemcpy(Ad, Ah, NumElms * sizeof(int), hipMemcpyHostToDevice)); + memset(Ah, 0, NumElms * sizeof(int)); + HIP_CHECK(hipMalloc(&Ad1, sizeof(int))); + HIP_CHECK(hipMemset(Ad1, 0, sizeof(int))); + int peak_clk; + HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0)); + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + hipEvent_t e1; + HIPCHECK(hipEventCreate(&e1)); + StreamPerThrd<<>>(Ad, Ad1, NumElms, + peak_clk, 1, 1); + HIP_CHECK(hipEventRecord(e1, Strm)); + HIP_CHECK(hipStreamWaitEvent(hipStreamPerThread, e1, 0 /*flags*/)); + MiniKernel<<<1, 1, 0, hipStreamPerThread>>>(Ad1); + sleep(1); + HIP_CHECK(hipMemcpy(Ah1, Ad1, sizeof(int), hipMemcpyDeviceToHost)); + if (*Ah1 != 3) { + IfTestPassed = false; + if (*Ah1 == 2) { + WARN("hipStreamPerThread didn't honour hipStreamWaitEvent()"); + } else if (*Ah1 == 4) { + WARN("Unexpected behavior observed with hipStreamPerThread"); + } + } + // Validating the result + HIP_CHECK(hipMemcpy(Ah, Ad, NumElms * sizeof(int), hipMemcpyDeviceToHost)); + int MisMatch = 0; + for (int i = 0; i < NumElms; ++i) { + if (Ah[i] != (CONST + 10)) { + MisMatch++; + } + } + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(Ad1)); + HIP_CHECK(hipStreamDestroy(Strm)); + delete[] Ah; + delete Ah1; + if (MisMatch) { + WARN("Data mismatch observed!!\n"); + IfTestPassed = false; + } + REQUIRE(IfTestPassed); +} + + +/* Testing hipLaunchCooperativeKernel() api with hipStreamPerThread*/ +TEST_CASE("Unit_hipStreamPerThread_CoopLaunch") { + hipDeviceProp_t device_properties; + HIPCHECK(hipGetDeviceProperties(&device_properties, 0)); + /* Test whether target device supports cooperative groups ****************/ + if (device_properties.cooperativeLaunch == 0) { + SUCCEED("Cooperative group support not available..."); + } else { + /* We will launch enough waves to fill up all of the GPU *****************/ + int warp_size = device_properties.warpSize; + int num_sms = device_properties.multiProcessorCount; + // long long totalTicks = device_properties.clockRate ; + int max_blocks_per_sm = 0; + // Calculate the device occupancy to know how many blocks can be run. + HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_sm, + StreamPerThrdCoopKrnl, + warp_size, 0)); + int max_active_blocks = max_blocks_per_sm * num_sms; + int *Ad = nullptr, *Ah = nullptr, *DNumElms = nullptr, NumElms = 4096; + int Const = 123; + Ah = new int[NumElms]; + for (int i = 0; i < NumElms; ++i) { + Ah[i] = Const; + } + HIP_CHECK(hipMalloc(&Ad, sizeof(int) * NumElms)); + HIP_CHECK(hipMalloc(&DNumElms, sizeof(int))); + HIP_CHECK(hipMemcpyAsync(Ad, Ah, sizeof(int) * NumElms, + hipMemcpyHostToDevice, hipStreamPerThread)); + HIP_CHECK(hipMemcpyAsync(DNumElms, &NumElms, sizeof(int), + hipMemcpyHostToDevice, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + + void *coop_params[2]; + coop_params[0] = reinterpret_cast(&Ad); + coop_params[1] = reinterpret_cast(&DNumElms); + HIP_CHECK(hipLaunchCooperativeKernel( + reinterpret_cast(StreamPerThrdCoopKrnl), + max_active_blocks, warp_size, + coop_params, 0, hipStreamPerThread)); + HIP_CHECK(hipMemcpy(Ah, Ad, sizeof(int) * NumElms, hipMemcpyDeviceToHost)); + // Verifying the result + int DataMismatch = 0; + for (int i = 0; i < NumElms; ++i) { + if (Ah[i] != (Const + 10)) { + DataMismatch++; + } + } + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(DNumElms)); + delete[] Ah; + if (DataMismatch > 0) { + REQUIRE(false); + } + } +} + +/* Testing hipLaunchCooperativeKernelMultiDevice() with hipStreamPerThread*/ +#if HT_AMD +TEST_CASE("Unit_hipStreamPerThread_CoopLaunchMDev") { + uint* dA[MaxGPUs]; + int64_t* dB[MaxGPUs]; + int64_t* dC; + + uint32_t* init = new uint32_t[BufferSizeInDwords]; + for (uint32_t i = 0; i < BufferSizeInDwords; ++i) { + init[i] = i; + } + + int nGpu = 0; + HIPCHECK(hipGetDeviceCount(&nGpu)); + size_t copySizeInDwords = BufferSizeInDwords / nGpu; + hipDeviceProp_t deviceProp[MaxGPUs]; + + for (int i = 0; i < nGpu; i++) { + HIPCHECK(hipSetDevice(i)); + + // Calculate the device occupancy to know how many blocks can be + // run concurrently + hipGetDeviceProperties(&deviceProp[i], 0); + if (!deviceProp[i].cooperativeMultiDeviceLaunch) { + WARN("Device doesn't support cooperative launch!"); + SUCCEED(""); + } + size_t SIZE = copySizeInDwords * sizeof(uint); + + HIPCHECK(hipMalloc(reinterpret_cast(&dA[i]), SIZE)); + HIPCHECK(hipMalloc(reinterpret_cast(&dB[i]), + 64 * deviceProp[i].multiProcessorCount * sizeof(int64_t))); + if (i == 0) { + HIPCHECK(hipHostMalloc(reinterpret_cast(&dC), + (nGpu + 1) * sizeof(int64_t))); + } + HIPCHECK(hipMemcpy(dA[i], &init[i * copySizeInDwords] , SIZE, + hipMemcpyHostToDevice)); + hipDeviceSynchronize(); + } + + dim3 dimBlock; + dim3 dimGrid; + dimGrid.x = 1; + dimGrid.y = 1; + dimGrid.z = 1; + dimBlock.x = 64; + dimBlock.y = 1; + dimBlock.z = 1; + + int numBlocks = 0; + uint workgroups[3] = {64, 128, 256}; + + hipLaunchParams* launchParamsList = new hipLaunchParams[nGpu]; + std::time_t end_time; + double time = 0; + for (uint set = 0; set < 3; ++set) { + void* args[MaxGPUs * NumKernelArgs]; + WARN("---------- Test#" << set << ", size: "<< BufferSizeInDwords << + " dwords ---------------\n"); + for (int i = 0; i < nGpu; i++) { + HIPCHECK(hipSetDevice(i)); + dimBlock.x = workgroups[set]; + HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks, + test_gwsPerThrd, dimBlock.x * dimBlock.y * dimBlock.z, + dimBlock.x * sizeof(int64_t))); + + WARN("GPU(" << i << ") Block size: " << dimBlock.x << + " Num blocks per CU: " << numBlocks << "\n"); + + dimGrid.x = deviceProp[i].multiProcessorCount * std::min(numBlocks, 32); + + args[i * NumKernelArgs] = reinterpret_cast(&dA[i]); + args[i * NumKernelArgs + 1] = reinterpret_cast(©SizeInDwords); + args[i * NumKernelArgs + 2] = reinterpret_cast(&dB[i]); + args[i * NumKernelArgs + 3] = reinterpret_cast(&dC); + + launchParamsList[i].func = reinterpret_cast(test_gwsPerThrd); + launchParamsList[i].gridDim = dimGrid; + launchParamsList[i].blockDim = dimBlock; + launchParamsList[i].sharedMem = dimBlock.x * sizeof(int64_t); + + launchParamsList[i].stream = hipStreamPerThread; + launchParamsList[i].args = &args[i * NumKernelArgs]; + } + + system_clock::time_point start = system_clock::now(); + hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0); + for (int i = 0; i < nGpu; i++) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipDeviceSynchronize()); + } + system_clock::time_point end = system_clock::now(); + std::chrono::duration elapsed_seconds = end - start; + end_time = std::chrono::system_clock::to_time_t(end); + + time += elapsed_seconds.count(); + + size_t processedDwords = copySizeInDwords * nGpu; + if (*dC != (((int64_t)(processedDwords) * (processedDwords - 1)) / 2)) { + WARN("Data validation failed ("<< *dC << " != " << + (((int64_t)(BufferSizeInDwords) * (BufferSizeInDwords - 1)) / 2) << + ") for grid size = " << dimGrid.x << " and block size = " << + dimBlock.x << "\n"); + WARN("Test failed!"); + } + } + + delete [] launchParamsList; + + WARN("finished computation at " << std::ctime(&end_time)); + WARN("elapsed time: " << time << "s\n"); + + hipSetDevice(0); + hipFree(dC); + for (int i = 0; i < nGpu; i++) { + hipFree(dA[i]); + hipFree(dB[i]); + } + delete [] init; +} +#endif From e7034719be25dad8dfb7316fc3224072b983da3e Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 8 Jul 2022 15:20:49 +0530 Subject: [PATCH 16/30] Update Jenkinsfile --- .jenkins/Jenkinsfile | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/.jenkins/Jenkinsfile b/.jenkins/Jenkinsfile index d64775c026..a4c767be38 100644 --- a/.jenkins/Jenkinsfile +++ b/.jenkins/Jenkinsfile @@ -58,8 +58,8 @@ def hipBuildTest(String backendLabel) { set -x # Check if backend label contains string "amd" or backend host is a server with amd gpu if [[ $backendLabel =~ amd ]]; then - LLVM_PATH=/opt/rocm/llvm ctest -E 'cooperative_streams_least_capacity.tst|cooperative_streams_half_capacity.tst|cooperative_streams_full_capacity.tst|grid_group_data_sharing.tst|hipIpcMemAccessTest.tst|p2p_copy_coherency.tst' sleep 120 + LLVM_PATH=/opt/rocm/llvm ctest -E 'cooperative_streams_least_capacity.tst|cooperative_streams_half_capacity.tst|cooperative_streams_full_capacity.tst|grid_group_data_sharing.tst|hipIpcMemAccessTest.tst|p2p_copy_coherency.tst' else make test fi @@ -100,9 +100,9 @@ def hipBuildTest(String backendLabel) { set -x # Check if backend label contains string "amd" or backend host is a server with amd gpu if [[ $backendLabel =~ amd ]]; then + sleep 120 export HT_CONFIG_FILE="$HIP_DIR/tests/catch/hipTestMain/config/config_amd_linux.json" LLVM_PATH=/opt/rocm/llvm ctest -E 'Unit_hipGraphChildGraphNodeGetGraph_Functional|Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Negative|Unit_hipPtrGetAttribute_Simple|Unit_hipStreamPerThread_DeviceReset_2' - sleep 120 else make test fi From 1b576c9fa66750068ef60af420e56aaf5a7d789a Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 8 Jul 2022 17:49:53 +0530 Subject: [PATCH 17/30] Disable newly added failing tests (#2790) * Update hipHostUnregister.cc - Disable Unit_hipHostUnregister_NotRegisteredPointer - Disable Unit_hipHostUnregister_AlreadyUnregisteredPointer * Update hipMallocPitch.cc - Disable Unit_hipMallocPitch_ValidatePitch - Disable failing sections of Unit_hipMalloc3D_Negative - Disable failing sections of Unit_hipMallocPitch_Negative - Disable failing sections of Unit_hipMemAllocPitch_Negative * Update config_amd_windows.json - Disabled Unit_hipMalloc3D_ValidatePitch - Disabled Unit_hipArrayCreate_happy --- .../hipTestMain/config/config_amd_windows.json | 4 +++- tests/catch/unit/memory/hipHostUnregister.cc | 8 ++++++++ tests/catch/unit/memory/hipMallocPitch.cc | 16 ++++++++++++++++ 3 files changed, 27 insertions(+), 1 deletion(-) diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 58f6439f5e..3c12aee9f9 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -30,6 +30,8 @@ "Unit_hipGraphMemcpyNodeSetParamsFromSymbol_Functional", "Unit_hipGraphExecEventWaitNodeSetEvent_Negative", "Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory", - "Unit_hipGraphMemcpyNodeSetParams_Functional" + "Unit_hipGraphMemcpyNodeSetParams_Functional", + "Unit_hipMalloc3D_ValidatePitch", + "Unit_hipArrayCreate_happy" ] } diff --git a/tests/catch/unit/memory/hipHostUnregister.cc b/tests/catch/unit/memory/hipHostUnregister.cc index 2721c29ad1..68a34122b9 100644 --- a/tests/catch/unit/memory/hipHostUnregister.cc +++ b/tests/catch/unit/memory/hipHostUnregister.cc @@ -73,11 +73,19 @@ TEST_CASE("Unit_hipHostUnregister_NullPtr") { } TEST_CASE("Unit_hipHostUnregister_NotRegisteredPointer") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("TODO-MATCH-ERRORCODE"); + return; +#endif auto x = std::unique_ptr(new int); HIP_CHECK_ERROR(hipHostUnregister(x.get()), hipErrorHostMemoryNotRegistered); } TEST_CASE("Unit_hipHostUnregister_AlreadyUnregisteredPointer") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("TODO-MATCH-ERRORCODE"); + return; +#endif if (!hipHostRegisterSupported()) { return; } diff --git a/tests/catch/unit/memory/hipMallocPitch.cc b/tests/catch/unit/memory/hipMallocPitch.cc index 0a63db52eb..33133f4a52 100644 --- a/tests/catch/unit/memory/hipMallocPitch.cc +++ b/tests/catch/unit/memory/hipMallocPitch.cc @@ -173,6 +173,10 @@ TEST_CASE("Unit_hipMemAllocPitch_ValidatePitch") { } TEST_CASE("Unit_hipMallocPitch_ValidatePitch") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("TODO-FIX-EXTENT-GENERATOR"); + return; +#endif size_t pitch; void* ptr; hipExtent validExtent{generateExtent(AllocationApi::hipMemAllocPitch)}; @@ -191,10 +195,13 @@ TEST_CASE("Unit_hipMalloc3D_Negative") { hipPitchedPtr ptr; constexpr size_t maxSizeT = std::numeric_limits::max(); +#if HT_NVIDIA + //TODO-MATCH-ERRORCODE SECTION("Max size_t width") { hipExtent validExtent{maxSizeT, 1, 1}; HIP_CHECK_ERROR(hipMalloc3D(&ptr, validExtent), hipErrorInvalidValue); } +#endif SECTION("Max size_t height") { hipExtent validExtent{1, maxSizeT, 1}; @@ -206,10 +213,13 @@ TEST_CASE("Unit_hipMalloc3D_Negative") { HIP_CHECK_ERROR(hipMalloc3D(&ptr, validExtent), hipErrorOutOfMemory); } +#if HT_NVIDIA + //TODO-MATCH-ERRORCODE SECTION("Max size_t all dimensions") { hipExtent validExtent{maxSizeT, maxSizeT, maxSizeT}; HIP_CHECK_ERROR(hipMalloc3D(&ptr, validExtent), hipErrorInvalidValue); } +#endif } TEST_CASE("Unit_hipMallocPitch_Negative") { @@ -225,9 +235,12 @@ TEST_CASE("Unit_hipMallocPitch_Negative") { HIP_CHECK_ERROR(hipMallocPitch(&ptr, nullptr, 1, 1), hipErrorInvalidValue); } +#if HT_NVIDIA + //TODO-MATCH-ERRORCODE SECTION("Max size_t width") { HIP_CHECK_ERROR(hipMallocPitch(&ptr, &pitch, maxSizeT, 1), hipErrorInvalidValue); } +#endif SECTION("Max size_t height") { HIP_CHECK_ERROR(hipMallocPitch(&ptr, &pitch, 1, maxSizeT), hipErrorOutOfMemory); @@ -272,10 +285,13 @@ TEST_CASE("Unit_hipMemAllocPitch_Negative") { hipErrorInvalidValue); } +#if HT_NVIDIA + //TODO-MATCH-ERRORCODE SECTION("Max size_t width") { HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, &pitch, maxSizeT, 1, validElementSizeBytes), hipErrorInvalidValue); } +#endif SECTION("Max size_t height") { HIP_CHECK_ERROR(hipMemAllocPitch(&ptr, &pitch, 1, maxSizeT, validElementSizeBytes), From c6092028d3569f13e53d2ab761416bc9fa3ee3c9 Mon Sep 17 00:00:00 2001 From: agunashe <86270081+agunashe@users.noreply.github.com> Date: Sun, 10 Jul 2022 22:22:04 -0700 Subject: [PATCH 18/30] SWDEV-327563 - Fix hipStreamPerThrdTsts test failures on windows (#2792) --- .../streamperthread/hipStreamPerThrdTsts.cc | 33 ++++++++++--------- 1 file changed, 17 insertions(+), 16 deletions(-) diff --git a/tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc b/tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc index 730dcbef6a..8aba0fa0f5 100644 --- a/tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc +++ b/tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc @@ -38,6 +38,7 @@ THE SOFTWARE. #include #ifdef _WIN32 #include + #define sleep(x) _sleep(x) #endif #ifdef __linux__ #include @@ -172,12 +173,12 @@ static void HIPRT_CB CallBackFunctn(hipStream_t strm, hipError_t err, } static void EventSync() { - int *Ad = nullptr, *Ah = nullptr, NumElms = 4096, CONST = 123; + int *Ad = nullptr, *Ah = nullptr, NumElms = 4096, CONST_NUM = 123; int blockSize = 32, peak_clk; HIP_CHECK(hipMalloc(&Ad, NumElms * sizeof(int))); Ah = new int[NumElms]; for (int i = 0; i < NumElms; ++i) { - Ah[i] = CONST; + Ah[i] = CONST_NUM; } // creating event objects hipEvent_t start, end; @@ -195,7 +196,7 @@ static void EventSync() { HIP_CHECK(hipMemcpy(Ah, Ad, NumElms * sizeof(int), hipMemcpyDeviceToHost)); int MisMatch = 0; for (int i = 0; i < NumElms; ++i) { - if (Ah[i] != (CONST + 10)) { + if (Ah[i] != (CONST_NUM + 10)) { MisMatch++; } } @@ -212,13 +213,13 @@ static void EventSync() { /* Launch a kernel in hipStreamPerThread, while it is in flight check for hipStreamQuery(hipStreamPerThread) it should return hipErrorNotReady.*/ TEST_CASE("Unit_hipStreamPerThreadTst_StrmQuery") { - int *Ad = nullptr, *Ah = nullptr, NumElms = 4096, CONST = 123; + int *Ad = nullptr, *Ah = nullptr, NumElms = 4096, CONST_NUM = 123; int blockSize = 32, peak_clk; hipError_t err; HIP_CHECK(hipMalloc(&Ad, NumElms * sizeof(int))); Ah = new int[NumElms]; for (int i = 0; i < NumElms; ++i) { - Ah[i] = CONST; + Ah[i] = CONST_NUM; } HIP_CHECK(hipMemcpy(Ad, Ah, NumElms * sizeof(int), hipMemcpyHostToDevice)); HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0)); @@ -256,17 +257,17 @@ TEST_CASE("Unit_hipStreamPerThread_MangdMem") { HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, 0)); if (managed == 1) { - int *Hmm = nullptr, NumElms = 4096, CONST = 123, blockSize = 32; + int *Hmm = nullptr, NumElms = 4096, CONST_NUM = 123, blockSize = 32; SECTION("Using Managed memory") { HIP_CHECK(hipMallocManaged(&Hmm, NumElms * sizeof(int))); for (int i = 0; i < NumElms; ++i) { - Hmm[i] = CONST; + Hmm[i] = CONST_NUM; } } SECTION("Prefetching Managed memory to device") { HIP_CHECK(hipMallocManaged(&Hmm, NumElms * sizeof(int))); for (int i = 0; i < NumElms; ++i) { - Hmm[i] = CONST; + Hmm[i] = CONST_NUM; } HIP_CHECK(hipMemPrefetchAsync(Hmm, NumElms * sizeof(int), 0, hipStreamPerThread)); @@ -281,7 +282,7 @@ TEST_CASE("Unit_hipStreamPerThread_MangdMem") { // Validating the result int MisMatch = 0; for (int i = 0; i < NumElms; ++i) { - if (Hmm[i] != (CONST + 10)) { + if (Hmm[i] != (CONST_NUM + 10)) { MisMatch++; } } @@ -300,12 +301,12 @@ TEST_CASE("Unit_hipStreamPerThread_MangdMem") { #ifdef __linux__ TEST_CASE("Unit_hipStreamPerThread_ChildProc") { if (fork() == 0) { // child process - int *Ad = nullptr, *Ah = nullptr, NumElms = 4096, CONST = 123; + int *Ad = nullptr, *Ah = nullptr, NumElms = 4096, CONST_NUM = 123; int blockSize = 32, peak_clk; HIP_CHECK(hipMalloc(&Ad, NumElms * sizeof(int))); Ah = new int[NumElms]; for (int i = 0; i < NumElms; ++i) { - Ah[i] = CONST; + Ah[i] = CONST_NUM; } HIP_CHECK(hipMemcpy(Ad, Ah, NumElms * sizeof(int), hipMemcpyHostToDevice)); HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0)); @@ -317,7 +318,7 @@ TEST_CASE("Unit_hipStreamPerThread_ChildProc") { HIP_CHECK(hipMemcpy(Ah, Ad, NumElms * sizeof(int), hipMemcpyDeviceToHost)); int MisMatch = 0; for (int i = 0; i < NumElms; ++i) { - if (Ah[i] != (CONST + 10)) { + if (Ah[i] != (CONST_NUM + 10)) { MisMatch++; } } @@ -359,14 +360,14 @@ TEST_CASE("Unit_hipStreamPerThread_EvtRcrdMThrd") { hipStreamWaitEvent()*/ TEST_CASE("Unit_hipStreamPerThread_StrmWaitEvt") { IfTestPassed = true; - int *Ad = nullptr, NumElms = 4096, CONST = 123, blockSize = 32, *Ah = nullptr; + int *Ad = nullptr, NumElms = 4096, CONST_NUM = 123, blockSize = 32, *Ah = nullptr; int *Ad1 = nullptr, *Ah1 = nullptr; Ah = new int[NumElms]; Ah1 = new int; hipStream_t Strm; HIP_CHECK(hipStreamCreate(&Strm)); for (int i = 0; i < NumElms; ++i) { - Ah[i] = CONST; + Ah[i] = CONST_NUM; } Ah1[0] = 0; HIP_CHECK(hipMalloc(&Ad, NumElms * sizeof(int))); @@ -399,7 +400,7 @@ TEST_CASE("Unit_hipStreamPerThread_StrmWaitEvt") { HIP_CHECK(hipMemcpy(Ah, Ad, NumElms * sizeof(int), hipMemcpyDeviceToHost)); int MisMatch = 0; for (int i = 0; i < NumElms; ++i) { - if (Ah[i] != (CONST + 10)) { + if (Ah[i] != (CONST_NUM + 10)) { MisMatch++; } } @@ -542,7 +543,7 @@ TEST_CASE("Unit_hipStreamPerThread_CoopLaunchMDev") { WARN("GPU(" << i << ") Block size: " << dimBlock.x << " Num blocks per CU: " << numBlocks << "\n"); - dimGrid.x = deviceProp[i].multiProcessorCount * std::min(numBlocks, 32); + dimGrid.x = deviceProp[i].multiProcessorCount * (std::min)(numBlocks, 32); args[i * NumKernelArgs] = reinterpret_cast(&dA[i]); args[i * NumKernelArgs + 1] = reinterpret_cast(©SizeInDwords); From 02747bcdb1c6944c1e845823323692f9ca9aa80f Mon Sep 17 00:00:00 2001 From: Dylan Angus <61192377+dylan-angus-codeplay@users.noreply.github.com> Date: Mon, 11 Jul 2022 07:48:19 +0100 Subject: [PATCH 19/30] EXSWCPHIPT-77 - Extending tests for hipHostRegister (#2609) --- tests/catch/include/hip_test_helper.hh | 2 +- tests/catch/unit/memory/hipHostRegister.cc | 118 ++++++++++++++------- 2 files changed, 80 insertions(+), 40 deletions(-) diff --git a/tests/catch/include/hip_test_helper.hh b/tests/catch/include/hip_test_helper.hh index fbbcb6cb06..9d4cbcd73f 100644 --- a/tests/catch/include/hip_test_helper.hh +++ b/tests/catch/include/hip_test_helper.hh @@ -38,7 +38,7 @@ static inline int getGeviceCount() { } // Get Free Memory from the system -static size_t getMemoryAmount() { +static inline size_t getMemoryAmount() { #ifdef __linux__ struct sysinfo info{}; sysinfo(&info); diff --git a/tests/catch/unit/memory/hipHostRegister.cc b/tests/catch/unit/memory/hipHostRegister.cc index 8312cf7f50..c2fa6ed009 100644 --- a/tests/catch/unit/memory/hipHostRegister.cc +++ b/tests/catch/unit/memory/hipHostRegister.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -28,51 +28,49 @@ This testfile verifies the following scenarios of hipHostRegister API */ #include -#include -#include +#include #define OFFSET 128 -static constexpr auto LEN{1024*1024}; +static constexpr auto LEN{1024 * 1024}; -template -__global__ void Inc(T* Ad) { - int tx = threadIdx.x + blockIdx.x * blockDim.x; - Ad[tx] = Ad[tx] + static_cast(1); +template __global__ void Inc(T* Ad) { + int tx = threadIdx.x + blockIdx.x * blockDim.x; + Ad[tx] = Ad[tx] + static_cast(1); } template -void doMemCopy(size_t numElements, int offset, T* A, T* Bh, T* Bd, - bool internalRegister) { - constexpr auto memsetval = 13.0f; - A = A + offset; - numElements -= offset; +void doMemCopy(size_t numElements, int offset, T* A, T* Bh, T* Bd, bool internalRegister) { + constexpr auto memsetval = 13.0f; + A = A + offset; + numElements -= offset; - size_t sizeBytes = numElements * sizeof(T); + size_t sizeBytes = numElements * sizeof(T); - if (internalRegister) { - HIP_CHECK(hipHostRegister(A, sizeBytes, 0)); - } + if (internalRegister) { + HIP_CHECK(hipHostRegister(A, sizeBytes, 0)); + } - // Reset - for (size_t i = 0; i < numElements; i++) { - A[i] = static_cast(i); - Bh[i] = 0.0f; - } + // Reset + for (size_t i = 0; i < numElements; i++) { + A[i] = static_cast(i); + Bh[i] = 0.0f; + } - HIP_CHECK(hipMemset(Bd, memsetval, sizeBytes)); + HIP_CHECK(hipMemset(Bd, memsetval, sizeBytes)); - HIP_CHECK(hipMemcpy(Bd, A, sizeBytes, hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy(Bh, Bd, sizeBytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(Bd, A, sizeBytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(Bh, Bd, sizeBytes, hipMemcpyDeviceToHost)); - // Make sure the copy worked - for (size_t i = 0; i < numElements; i++) { - REQUIRE(Bh[i] == A[i]); - } + // Make sure the copy worked + for (size_t i = 0; i < numElements; i++) { + REQUIRE(Bh[i] == A[i]); + } - if (internalRegister) { - HIP_CHECK(hipHostUnregister(A)); - } + if (internalRegister) { + HIP_CHECK(hipHostUnregister(A)); + } } + /* This testcase verifies the hipHostRegister API by 1. Allocating the memory using malloc @@ -81,9 +79,7 @@ This testcase verifies the hipHostRegister API by 4. Launching kernel and access the device pointer variable 5. performing hipMemset on the device pointer variable */ -TEMPLATE_TEST_CASE("Unit_hipHostRegister_ReferenceFromKernelandhipMemset", - "", int, - float, double) { +TEMPLATE_TEST_CASE("Unit_hipHostRegister_ReferenceFromKernelandhipMemset", "", int, float, double) { size_t sizeBytes{LEN * sizeof(TestType)}; TestType *A, **Ad; int num_devices; @@ -118,14 +114,14 @@ TEMPLATE_TEST_CASE("Unit_hipHostRegister_ReferenceFromKernelandhipMemset", HIP_CHECK(hipHostUnregister(A)); free(A); - delete [] Ad; + delete[] Ad; } + /* This testcase verifies hipHostRegister API by performing memcpy on the hipHostRegistered variable. */ -TEMPLATE_TEST_CASE("Unit_hipHostRegister_Memcpy", "", - int, float, double) { +TEMPLATE_TEST_CASE("Unit_hipHostRegister_Memcpy", "", int, float, double) { // 1 refers to hipHostRegister // 0 refers to malloc auto mem_type = GENERATE(0, 1); @@ -156,5 +152,49 @@ TEMPLATE_TEST_CASE("Unit_hipHostRegister_Memcpy", "", free(A); free(Bh); - hipFree(Bd); + HIP_CHECK(hipFree(Bd)); +} + +template __global__ void fill_kernel(T* dataPtr, T value) { + size_t tid{blockIdx.x * blockDim.x + threadIdx.x}; + dataPtr[tid] = value; +} + +TEMPLATE_TEST_CASE("Unit_hipHostRegister_Negative", "", int, float, double) { + TestType* hostPtr = nullptr; + + size_t sizeBytes = 1 * sizeof(TestType); + SECTION("hipHostRegister Negative Test - nullptr") { + HIP_CHECK_ERROR(hipHostRegister(hostPtr, 1, 0), hipErrorInvalidValue); + } + + hostPtr = reinterpret_cast(malloc(sizeBytes)); + SECTION("hipHostRegister Negative Test - zero size") { + HIP_CHECK_ERROR(hipHostRegister(hostPtr, 0, 0), hipErrorInvalidValue); + } + +#if HT_NVIDIA + // Flags aren't used for AMD devices currently + SECTION("hipHostRegister Negative Test - invalid flag") { + HIP_CHECK_ERROR(hipHostRegister(hostPtr, sizeBytes, 0b11111111), hipErrorInvalidValue); + } +#endif + + size_t devMemAvail{0}, devMemFree{0}; + HIP_CHECK(hipMemGetInfo(&devMemFree, &devMemAvail)); + auto hostMemFree = HipTest::getMemoryAmount() /* In MB */ * 1024 * 1024; // In bytes + REQUIRE(devMemFree > 0); + REQUIRE(devMemAvail > 0); + REQUIRE(hostMemFree > 0); + + size_t memFree = std::min(devMemFree, hostMemFree); // which is the limiter cpu or gpu + + SECTION("hipHostRegister Negative Test - invalid memory size") { + HIP_CHECK_ERROR(hipHostRegister(hostPtr, memFree, 0), hipErrorInvalidValue); + } + + free(hostPtr); + SECTION("hipHostRegister Negative Test - freed memory") { + HIP_CHECK_ERROR(hipHostRegister(hostPtr, 0, 0), hipErrorInvalidValue); + } } From 25f2b260beb51f6c1040ef19da8e3d89935248a7 Mon Sep 17 00:00:00 2001 From: Dylan Angus <61192377+dylan-angus-codeplay@users.noreply.github.com> Date: Mon, 11 Jul 2022 07:48:43 +0100 Subject: [PATCH 20/30] EXSWCPHIPT-76 - Adding hipMemset functional tests (#2669) --- tests/catch/unit/memory/CMakeLists.txt | 2 + .../catch/unit/memory/hipMemsetFunctional.cc | 560 ++++++++++++++++++ 2 files changed, 562 insertions(+) create mode 100644 tests/catch/unit/memory/hipMemsetFunctional.cc diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index 695dae8792..c262620f53 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -76,6 +76,7 @@ set(TEST_SRC hipHostMalloc.cc hipMemcpy.cc hipMemcpyAsync.cc + hipMemsetFunctional.cc hipMallocPitch.cc hipMallocArray.cc hipMalloc3D.cc @@ -143,6 +144,7 @@ set(TEST_SRC hipHostMalloc.cc hipMemcpy.cc hipMemcpyAsync.cc + hipMemsetFunctional.cc hipMallocPitch.cc hipMallocArray.cc hipMalloc3D.cc diff --git a/tests/catch/unit/memory/hipMemsetFunctional.cc b/tests/catch/unit/memory/hipMemsetFunctional.cc new file mode 100644 index 0000000000..274074266d --- /dev/null +++ b/tests/catch/unit/memory/hipMemsetFunctional.cc @@ -0,0 +1,560 @@ +/* + * Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +/** + * Testcase Scenarios: + * For hipMemset, hipMemsetD8, hipMemsetD16, hipMemsetD32, hipMemset2D, hipMemset3D and all async + * counterparts + * 1) (ZeroValue) - Test setting a specified range to zero. + * 2) (SmallSize) - Test setting a unique memset value for small sizes. + * 3) (ZeroSize) - Test that trying to set memory with a zero dimension does not fail and doesn't + * affect the memory. + * 4) (PartialSet) - Test setting a partial range of total allocated memory and + * ensure the full range isn't affected. + */ +#include + +constexpr size_t FULL_DIM = 10; + +// Enum used to determine which 1D memset function to use. +enum MemsetType { + hipMemsetTypeDefault = 1, + hipMemsetTypeD8 = 2, + hipMemsetTypeD16 = 3, + hipMemsetTypeD32 = 4 +}; + +// Macro used to assert all elements in a flat vector range is equal to a specified value +#define HIP_ASSERT_VEC_EQ(ptr, value, N) \ + for (size_t i = 0; i < N; i++) { \ + CAPTURE(N, i, ptr[i], value); \ + HIP_ASSERT(ptr[i] == value); \ + } + +// Copies device data to host and checks that each element is equal to the +// specified value +template void check_device_data(T* devPtr, T value, size_t numElems) { + std::unique_ptr hostPtr(new T[numElems]); + HIP_CHECK(hipMemcpy(hostPtr.get(), devPtr, numElems * sizeof(T), hipMemcpyDeviceToHost)); + HIP_ASSERT_VEC_EQ(hostPtr.get(), value, numElems); +} + +// Macro to assist calling and then checking the result of the 1D memset API with the necessary +// manipulation to the arguments. +#define HIP_MEMSET_CHECK(hipMemsetFunc, devPtr, value, count, async) \ + using scalar_t = decltype(value); \ + size_t sizeBytes = count * sizeof(scalar_t); \ + HIP_CHECK(hipMemsetFunc(devPtr, value, sizeBytes)); \ + if (async) { \ + HIP_CHECK(hipStreamSynchronize(stream)); \ + } \ + check_device_data(devPtr, value, count); + +#define HIP_MEMSET_CHECK_DTYPE(hipMemsetFunc, devPtr, value, count, async) \ + HIP_CHECK(hipMemsetFunc(reinterpret_cast(devPtr), value, count)); \ + if (async) { \ + HIP_CHECK(hipStreamSynchronize(stream)); \ + } \ + check_device_data(devPtr, value, count); + +// Enum for specifying wether to allocate the data using hipMalloc, hipHostMalloc or not at all. +enum MemsetMallocType { hipDeviceMalloc_t = 1, hipHostMalloc_t = 2, hipNoMalloc_t }; + +// Helper function for allocating memory, setting data with the specified 1D memset API and then +// checking result of operation. +template +void checkMemset(T value, size_t count, MemsetType memsetType, bool async = false, + MemsetMallocType mallocType = hipDeviceMalloc_t, T* devPtr = nullptr) { + hipStream_t stream{nullptr}; + if (async) { + hipStreamCreate(&stream); + } + + // Allocate Memory + if (mallocType == hipDeviceMalloc_t) { + HIP_CHECK(hipMalloc(&devPtr, count * sizeof(T))); + } else if (mallocType == hipHostMalloc_t) { + HIP_CHECK(hipHostMalloc(&devPtr, count * sizeof(T))); + } + + // memset API calls + switch (memsetType) { + case hipMemsetTypeDefault: + if (!async) { + INFO("Testing hipMemset call") + HIP_MEMSET_CHECK(hipMemset, devPtr, value, count, false); + } else { + INFO("Testing hipMemsetAsync call") + HIP_MEMSET_CHECK(hipMemsetAsync, devPtr, value, count, true); + } + break; + case hipMemsetTypeD8: + if (!async) { + INFO("Testing hipMemsetD8 call") + HIP_MEMSET_CHECK_DTYPE(hipMemsetD8, devPtr, value, count, false); + } else { + INFO("Testing hipMemsetD8Async call") + HIP_MEMSET_CHECK_DTYPE(hipMemsetD8Async, devPtr, value, count, true); + } + break; + case hipMemsetTypeD16: + if (!async) { + INFO("Testing hipMemsetD16 call") + HIP_MEMSET_CHECK_DTYPE(hipMemsetD16, devPtr, value, count, false); + } else { + INFO("Testing hipMemsetD16Async call") + HIP_MEMSET_CHECK_DTYPE(hipMemsetD16Async, devPtr, value, count, true); + } + break; + case hipMemsetTypeD32: + if (!async) { + INFO("Testing hipMemsetD32 call") + HIP_MEMSET_CHECK_DTYPE(hipMemsetD32, devPtr, value, count, false); + } else { + INFO("Testing hipMemsetD32Async call") + HIP_MEMSET_CHECK_DTYPE(hipMemsetD32Async, devPtr, value, count, true); + } + break; + } + + // Cleanup + if (async) { + HIP_CHECK(hipStreamDestroy(stream)); + } + + // Free memory + if (mallocType == hipDeviceMalloc_t) { + HIP_CHECK(hipFree(devPtr)); + } else if (mallocType == hipHostMalloc_t) { + HIP_CHECK(hipHostFree(devPtr)); + } +} + +// Macro which defines a TEST_CASE which calls and then checks the result of the 1D memset macros +// for all combinations of sync/async and hipMalloc/hipHostMalloc, given the value and memory range. +#define DEFINE_1D_BASIC_TEST_CASE(suffix, memsetType, T, value, count) \ + TEST_CASE("Unit_hipMemsetFunctional_" + std::string(suffix)) { \ + const std::string memsetStr = std::string(suffix); \ + SECTION(memsetStr + " - Device Malloc") { \ + checkMemset(static_cast(value), count, memsetType, false, hipDeviceMalloc_t); \ + } \ + SECTION(memsetStr + " - Host Malloc") { \ + checkMemset(static_cast(value), count, memsetType, false, hipHostMalloc_t); \ + } \ + SECTION(memsetStr + "Async - Device Malloc") { \ + checkMemset(static_cast(value), count, memsetType, true, hipDeviceMalloc_t); \ + } \ + SECTION(memsetStr + "Async - Host Malloc") { \ + checkMemset(static_cast(value), count, memsetType, true, hipHostMalloc_t); \ + } \ + } + +DEFINE_1D_BASIC_TEST_CASE("ZeroValue_hipMemset", hipMemsetTypeDefault, float, 0, 1024) +DEFINE_1D_BASIC_TEST_CASE("ZeroValue_hipMemsetD32", hipMemsetTypeD32, uint32_t, 0, 1024) +DEFINE_1D_BASIC_TEST_CASE("ZeroValue_hipMemsetD16", hipMemsetTypeD16, int16_t, 0, 1024) +DEFINE_1D_BASIC_TEST_CASE("ZeroValue_hipMemsetD8", hipMemsetTypeD8, int8_t, 0, 1024) + +DEFINE_1D_BASIC_TEST_CASE("SmallSize_hipMemset", hipMemsetTypeDefault, char, 0x42, 1) +DEFINE_1D_BASIC_TEST_CASE("SmallSize_hipMemsetD32", hipMemsetTypeD32, uint32_t, 0x101, 1) +DEFINE_1D_BASIC_TEST_CASE("SmallSize_hipMemsetD16", hipMemsetTypeD16, int16_t, 0x10, 1) +DEFINE_1D_BASIC_TEST_CASE("SmallSize_hipMemsetD8", hipMemsetTypeD8, int8_t, 0x1, 1) + +DEFINE_1D_BASIC_TEST_CASE("ZeroSize_hipMemset", hipMemsetTypeDefault, char, 0x42, 0) +DEFINE_1D_BASIC_TEST_CASE("ZeroSize_hipMemsetD32", hipMemsetTypeD32, uint32_t, 0x101, 0) +DEFINE_1D_BASIC_TEST_CASE("ZeroSize_hipMemsetD16", hipMemsetTypeD16, int16_t, 0x10, 0) +DEFINE_1D_BASIC_TEST_CASE("ZeroSize_hipMemsetD8", hipMemsetTypeD8, int8_t, 0x1, 0) + +// Helper function that sets a full region of memory with an initial value, sets a smaller subregion +// with another value and check that the memset API do not write outside of the subregion of data. +template +void partialMemsetTest(T valA, T valB, size_t count, size_t offset, MemsetType memsetType, + bool async) { + T* devPtr; + size_t subSize{count - offset}; + HIP_CHECK(hipMalloc(&devPtr, count * sizeof(T))); + + // Set entire region to be first value. + INFO("Setting full region"); + checkMemset(valA, count, memsetType, async, hipNoMalloc_t, devPtr); + + // Set partial region to be second value. + INFO("Setting partial region"); + checkMemset(valB, subSize, memsetType, async, hipNoMalloc_t, devPtr + offset); + + // Ensure the first section remains unchanged + check_device_data(devPtr, valA, offset); + HIP_CHECK(hipFree(devPtr)); +} + +TEST_CASE("Unit_hipMemsetFunctional_PartialSet_1D") { + for (auto widthOffset = 8; widthOffset <= 8; widthOffset *= 2) { + SECTION("hipMemset - Partial Set") { + partialMemsetTest(0x1, 0x42, 1024, widthOffset, hipMemsetTypeDefault, false); + } + SECTION("hipMemsetAsync - Partial Set") { + partialMemsetTest(0x1, 0x42, 1024, widthOffset, hipMemsetTypeDefault, true); + } + SECTION("hipMemsetD8 - Partial Set") { + partialMemsetTest(0x1, 0xDE, 1024, widthOffset, hipMemsetTypeD8, false); + } + SECTION("hipMemsetD8Async - Partial Set") { + partialMemsetTest(0x1, 0xDE, 1024, widthOffset, hipMemsetTypeD8, true); + } + SECTION("hipMemsetD16 - Partial Set") { + partialMemsetTest(0x1, 0xDEAD, 1024, widthOffset, hipMemsetTypeD16, false); + } + SECTION("hipMemsetD16Async - Partial Set") { + partialMemsetTest(0x1, 0xDEAD, 1024, widthOffset, hipMemsetTypeD16, true); + } + SECTION("hipMemsetD32 - Partial Set") { + partialMemsetTest(0x1, 0xDEADBEEF, 1024, widthOffset, hipMemsetTypeD32, false); + } + SECTION("hipMemsetD32Async - Partial Set") { + partialMemsetTest(0x1, 0xDEADBEEF, 1024, widthOffset, hipMemsetTypeD32, true); + } + } +} + +// Helper function that copies the device data to the host and returns a unique_ptr to that data. +template +std::unique_ptr get_device_data_2D(T* devPtr, size_t pitch, size_t width, size_t height) { + std::unique_ptr hostPtr(new T[width * height]); + constexpr size_t elementSize = sizeof(T); + HIP_CHECK(hipMemcpy2D(hostPtr.get(), width * elementSize, devPtr, pitch, width, height, + hipMemcpyDeviceToHost)); + return hostPtr; +} + +// Copies device data to host and checks that each element is equal to the +// specified value +template +void check_device_data_2D(T* devPtr, T value, size_t pitch, size_t width, size_t height) { + auto hostPtr = get_device_data_2D(devPtr, pitch, width, height); + HIP_ASSERT_VEC_EQ(hostPtr.get(), value, width * height); +} + +// Helper function for allocating memory, setting data with the specified 2D memset API and then +// checking result of operation. +template +void checkMemset2D(T value, size_t width, size_t height, bool async = false, size_t pitch = 0, + T* devPtr = nullptr) { + hipStream_t stream{nullptr}; + hipStreamCreate(&stream); + constexpr size_t elementSize = sizeof(T); + bool freeDevPtr = false; + if (devPtr == nullptr) { + freeDevPtr = true; + HIP_CHECK( + hipMallocPitch(reinterpret_cast(&devPtr), &pitch, width * elementSize, height)); + } + + if (!async) { + INFO("Testing hipMemset2D call") + HIP_CHECK(hipMemset2D(devPtr, pitch, value, width * elementSize, height)); + } else { + INFO("Testing hipMemset2DAsync call") + HIP_CHECK(hipMemset2DAsync(devPtr, pitch, value, width * elementSize, height, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + if (width * height > 0) { + check_device_data_2D(devPtr, value, pitch, width, height); + } + if (freeDevPtr) { + HIP_CHECK(hipFree(devPtr)); + } + hipStreamDestroy(stream); +} + +TEST_CASE("Unit_hipMemsetFunctional_ZeroValue_2D") { + constexpr size_t width{128}; + constexpr size_t height{128}; + constexpr char memsetVal = 0; + SECTION("hipMemset2D - Zero Value") { checkMemset2D(memsetVal, width, height, false); } + SECTION("hipMemset2DAsync - Zero Value") { checkMemset2D(memsetVal, width, height, true); } +} + +TEST_CASE("Unit_hipMemsetFunctional_SmallSize_2D") { + constexpr char memsetVal = 0x42; + SECTION("hipMemset2D - Small Size") { checkMemset2D(memsetVal, 1, 1, false); } + SECTION("hipMemset2DAsync - Small Size") { checkMemset2D(memsetVal, 1, 1, true); } +} + +TEST_CASE("Unit_hipMemsetFunctional_ZeroSize_2D") { + size_t pitch{0}; + size_t width{10}; + size_t height{10}; + char* devPtr{nullptr}; + HIP_CHECK( + hipMallocPitch(reinterpret_cast(&devPtr), &pitch, width * sizeof(char), height)); + + const char initValue = 0x1; + const char testValue = 0x11; + // Set full region to initial value + checkMemset2D(initValue, width, height, false, pitch, devPtr); + + SECTION("hipMemset2D - Zero Width") { + checkMemset2D(testValue, 0, height, false, pitch, devPtr); + check_device_data_2D(devPtr, initValue, pitch, width, height); + } + SECTION("hipMemset2DAsync - Zero Width") { + checkMemset2D(testValue, 0, height, true, pitch, devPtr); + check_device_data_2D(devPtr, initValue, pitch, width, height); + } + SECTION("hipMemset2D - Zero Height") { + checkMemset2D(testValue, width, 0, false, pitch, devPtr); + check_device_data_2D(devPtr, initValue, pitch, width, height); + } + SECTION("hipMemset2DAsync - Zero Height") { + checkMemset2D(testValue, width, 0, true, pitch, devPtr); + check_device_data_2D(devPtr, initValue, pitch, width, height); + } + SECTION("hipMemset2D - Zero Width and Height") { + checkMemset2D(testValue, 0, 0, false, pitch, devPtr); + check_device_data_2D(devPtr, initValue, pitch, width, height); + } + SECTION("hipMemset2DAsync - Zero Width and Height") { + checkMemset2D(testValue, 0, 0, true, pitch, devPtr); + check_device_data_2D(devPtr, initValue, pitch, width, height); + } + HIP_CHECK(hipFree(devPtr)); +} + +// Helper function that sets a full region of memory with an initial value, sets a smaller subregion +// with another value and check that the memset API do not write outside of the subregion of data. +template +void partialMemsetTest2D(T valA, T valB, size_t width, size_t height, size_t widthOffset, + size_t heightOffset, bool async) { + T* devPtr{nullptr}; + size_t pitch{0}; + size_t subWidth{width - widthOffset}; + size_t subHeight{height - heightOffset}; + constexpr size_t elementSize = sizeof(T); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devPtr), &pitch, width * elementSize, height)); + + // Set entire region to be first value. + INFO("Setting full square region"); + checkMemset2D(valA, width, height, async, pitch, devPtr); + + // Set partial region to be second value. + INFO("Setting partial square region") + checkMemset2D(valB, subWidth, subHeight, async, pitch, devPtr); + + auto hostPtr = get_device_data_2D(devPtr, pitch, width, height); + T comparVal{0}; + size_t idx{0}; + for (size_t i = 0; i < width; i++) { + for (size_t j = 0; j < height; j++) { + if (i < subWidth && j < subHeight) { + // Compare subregion value + comparVal = valB; + } else { + // Compare full region value + comparVal = valA; + } + idx = i * height + j; + CAPTURE(width, height, subWidth, subHeight, i, j, idx, hostPtr[idx], comparVal); + HIP_ASSERT(hostPtr[idx] == comparVal); + } + } + HIP_CHECK(hipFree(devPtr)); +} + +TEST_CASE("Unit_hipMemsetFunctional_PartialSet_2D") { + for (auto widthOffset = 8; widthOffset <= 128; widthOffset *= 2) { + for (auto heightOffset = 8; heightOffset <= 128; heightOffset *= 2) { + SECTION("hipMemset2D - Partial Set") { + partialMemsetTest2D('a', 'b', 200, 200, widthOffset, heightOffset, false); + } + SECTION("hipMemset2DAsync - Partial Set") { + partialMemsetTest2D('a', 'b', 200, 200, widthOffset, heightOffset, true); + } + } + } +} + +// Helper function that copies the device data to the host and returns a unique_ptr to that data. +template +std::unique_ptr get_device_data_3D(hipPitchedPtr& devPitchedPtr, hipExtent extent) { + constexpr size_t elementSize = sizeof(T); + std::unique_ptr hostPtr( + new T[devPitchedPtr.pitch * extent.width * extent.height / elementSize]); + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.dstPtr = make_hipPitchedPtr(hostPtr.get(), devPitchedPtr.pitch, + extent.width / elementSize, extent.height); + myparms.srcPtr = devPitchedPtr; + myparms.extent = extent; + myparms.kind = hipMemcpyDeviceToHost; + HIP_CHECK(hipMemcpy3D(&myparms)); + return hostPtr; +} + +// Copies device data to host and checks that each element is equal to the +// specified value +template +void check_device_data_3D(hipPitchedPtr& devPitchedPtr, T value, hipExtent extent) { + auto hostPtr = get_device_data_3D(devPitchedPtr, extent); + size_t width = extent.width / sizeof(T); + size_t height = extent.height; + size_t depth = extent.depth; + size_t idx; + for (size_t k = 0; k < depth; k++) { + for (size_t j = 0; j < height; j++) { + for (size_t i = 0; i < width; i++) { + idx = devPitchedPtr.pitch * height * k + devPitchedPtr.pitch * j + i; + INFO("idx=" << idx << " hostPtr[idx]=" << hostPtr[idx] << " value=" << value) + HIP_ASSERT(hostPtr[idx] == value); + } + } + } +} + +// Helper function for allocating memory, setting data with the specified 3D memset API and then +// checking result of operation. +template +void checkMemset3D(hipPitchedPtr& devPitchedPtr, T value, hipExtent extent, bool async = false) { + hipStream_t stream{nullptr}; + hipStreamCreate(&stream); + if (devPitchedPtr.ptr == nullptr) { + HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); + } + if (!async) { + INFO("Testing hipMemset3D call") + HIP_CHECK(hipMemset3D(devPitchedPtr, value, extent)); + } else { + INFO("Testing hipMemset3DAsync call") + HIP_CHECK(hipMemset3DAsync(devPitchedPtr, value, extent, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + if (extent.width * extent.height * extent.depth > 0) { + check_device_data_3D(devPitchedPtr, value, extent); + } + hipStreamDestroy(stream); +} + +void check_memset_3D(std::string sectionStr, size_t width, size_t height, size_t depth, + char value) { + hipPitchedPtr devPitchedPtr; + hipExtent fullExtent; + constexpr char fullVal = 0x21; + hipExtent extent = make_hipExtent(width, height, depth); + // Check if any of the dimensions are zero + bool anyZero = width * height * depth == 0; + if (anyZero) { + // If they are zero then set a full region with memset value to later check if it's changed. + devPitchedPtr.ptr = nullptr; + fullExtent = make_hipExtent(FULL_DIM, FULL_DIM, FULL_DIM); + checkMemset3D(devPitchedPtr, fullVal, fullExtent, false); + } + SECTION("hipMemset3D - " + sectionStr) { + if (!anyZero) { + devPitchedPtr.ptr = nullptr; + } + checkMemset3D(devPitchedPtr, value, extent, false); + if (anyZero) { + // Check to make sure memsets with a zero dimension did not affect above set region. + check_device_data_3D(devPitchedPtr, fullVal, fullExtent); + } + HIP_CHECK(hipFree(devPitchedPtr.ptr)); + } + SECTION("hipMemset3DAsync - " + sectionStr) { + if (!anyZero) { + devPitchedPtr.ptr = nullptr; + } + checkMemset3D(devPitchedPtr, value, extent, true); + if (anyZero) { + // Check to make sure memsets with a zero dimension did not affect above set region. + check_device_data_3D(devPitchedPtr, fullVal, fullExtent); + } + HIP_CHECK(hipFree(devPitchedPtr.ptr)); + } +} + +TEST_CASE("Unit_hipMemsetFunctional_ZeroValue_3D") { + check_memset_3D("Zero Value", 128, 128, 10, 0); +} + +TEST_CASE("Unit_hipMemsetFunctional_SmallSize_3D") { check_memset_3D("Small Size", 1, 1, 1, 0x42); } + +TEST_CASE("Unit_hipMemsetFunctional_ZeroSize_3D") { + constexpr size_t elementSize = sizeof(char); + check_memset_3D("Zero Width", 0, FULL_DIM, FULL_DIM, 0x23); + check_memset_3D("Zero Height", FULL_DIM * elementSize, 0, FULL_DIM, 0x23); + check_memset_3D("Zero Depth", FULL_DIM * elementSize, FULL_DIM, 0, 0x23); + check_memset_3D("Zero Width and Height", 0 * elementSize, 0, FULL_DIM, 0x23); + check_memset_3D("Zero Width and Depth", 0 * elementSize, FULL_DIM, 0, 0x23); + check_memset_3D("Zero Height and Depth", FULL_DIM * elementSize, 0, 0, 0x23); + check_memset_3D("Zero Width, Height and Depth", 0 * elementSize, 0, 0, 0x23); +} + +// Helper function that sets a full region of memory with an initial value, sets a smaller subregion +// with another value and check that the memset API do not write outside of the subregion of data. +template +void partialMemsetTest3D(T valA, T valB, size_t width, size_t height, size_t depth, + size_t widthOffset, size_t heightOffset, size_t depthOffset, bool async) { + size_t subWidth{width - widthOffset}; + size_t subHeight{height - heightOffset}; + size_t subDepth{depth - depthOffset}; + hipPitchedPtr devPitchedPtr; + devPitchedPtr.ptr = nullptr; + hipExtent extent = make_hipExtent(width * sizeof(T), height, depth); + hipExtent subExtent = make_hipExtent(subWidth * sizeof(T), subHeight, subDepth); + + // Set entire region to be first value. + INFO("Setting full cuboid region") { checkMemset3D(devPitchedPtr, valA, extent, async); } + // Set partial region to be second value. + INFO("Setting partial cuboid region") { checkMemset3D(devPitchedPtr, valB, subExtent, async); } + auto pitch = devPitchedPtr.pitch; + auto hostPtr = get_device_data_3D(devPitchedPtr, extent); + T comparVal{0}; + size_t idx{0}; + for (size_t k = 0; k < depth; k++) { + for (size_t j = 0; j < height; j++) { + for (size_t i = 0; i < width; i++) { + if (i < subWidth && j < subHeight && k < subDepth) { + comparVal = valB; + } else { + comparVal = valA; + } + idx = devPitchedPtr.pitch * height * k + devPitchedPtr.pitch * j + i; + CAPTURE(width, height, depth, pitch, subWidth, subHeight, subDepth, i, j, k, idx, + hostPtr[idx], comparVal); + HIP_ASSERT(hostPtr[idx] == comparVal); + } + } + } + HIP_CHECK(hipFree(devPitchedPtr.ptr)); +} + +TEST_CASE("Unit_hipMemsetFunctional_PartialSet_3D") { + for (auto widthOffset = 8; widthOffset <= 128; widthOffset *= 2) { + for (auto heightOffset = 8; heightOffset <= 128; heightOffset *= 2) { + for (auto depthOffset = 2; depthOffset <= 5; depthOffset++) { + SECTION("hipMemset3D - Partial Set") { + partialMemsetTest3D('a', 'b', 200, 200, 10, widthOffset, heightOffset, depthOffset, + false); + } + SECTION("hipMemset3DAsync - Partial Set") { + partialMemsetTest3D('a', 'b', 200, 200, 10, widthOffset, heightOffset, depthOffset, true); + } + } + } + } +} From c0deb17bbcd5bfdbbba0c7d456625f28e6e45015 Mon Sep 17 00:00:00 2001 From: Finlay Date: Mon, 11 Jul 2022 07:49:05 +0100 Subject: [PATCH 21/30] EXSWCPHIPT-112 - Unit test for hipMalloc3DArray for default and surface arrays (#2713) --- tests/catch/unit/memory/hipArrayCommon.hh | 36 +++ tests/catch/unit/memory/hipArrayCreate.cc | 89 +++---- tests/catch/unit/memory/hipMalloc3DArray.cc | 266 +++++++++++++------- tests/catch/unit/memory/hipMallocArray.cc | 95 +++---- 4 files changed, 302 insertions(+), 184 deletions(-) diff --git a/tests/catch/unit/memory/hipArrayCommon.hh b/tests/catch/unit/memory/hipArrayCommon.hh index 1c6100a6e9..c9823806c5 100644 --- a/tests/catch/unit/memory/hipArrayCommon.hh +++ b/tests/catch/unit/memory/hipArrayCommon.hh @@ -122,3 +122,39 @@ inline size_t getFreeMem() { HIP_CHECK(hipMemGetInfo(&free, &total)); return free; } + +struct Sizes { + int max1D; + std::array max2D; + std::array max3D; + + Sizes(unsigned int flag) { + int device; + HIP_CHECK(hipGetDevice(&device)); + switch (flag) { + case hipArrayDefault: { + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + max1D = prop.maxTexture1D; + max2D = {prop.maxTexture2D[0], prop.maxTexture2D[1]}; + max3D = {prop.maxTexture3D[0], prop.maxTexture3D[1], prop.maxTexture3D[2]}; + return; + } + case hipArraySurfaceLoadStore: { + int value; + HIP_CHECK(hipDeviceGetAttribute(&value, hipDeviceAttributeMaxSurface1D, device)); + max1D = value; + HIP_CHECK(hipDeviceGetAttribute(&value, hipDeviceAttributeMaxSurface2D, device)); + max2D = {value, value}; + HIP_CHECK(hipDeviceGetAttribute(&value, hipDeviceAttributeMaxSurface3D, device)); + max3D = {value, value, value}; + return; + } + default: { + INFO("Array flag not supported"); + REQUIRE(false); + return; + } + } + } +}; diff --git a/tests/catch/unit/memory/hipArrayCreate.cc b/tests/catch/unit/memory/hipArrayCreate.cc index 9ef637821e..d55bdae255 100644 --- a/tests/catch/unit/memory/hipArrayCreate.cc +++ b/tests/catch/unit/memory/hipArrayCreate.cc @@ -30,10 +30,10 @@ hipArrayCreate API test scenarios #include "hipArrayCommon.hh" #include "DriverContext.hh" -static constexpr auto NUM_W{4}; -static constexpr auto BIGNUM_W{100}; -static constexpr auto NUM_H{4}; -static constexpr auto BIGNUM_H{100}; +static constexpr size_t NUM_W{4}; +static constexpr size_t BIGNUM_W{100}; +static constexpr size_t NUM_H{4}; +static constexpr size_t BIGNUM_H{100}; static constexpr auto ARRAY_LOOP{100}; /* @@ -51,31 +51,34 @@ static constexpr auto ARRAY_LOOP{100}; */ static void ArrayCreate_DiffSizes(int gpu) { - HIP_CHECK(hipSetDevice(gpu)); - std::vector> array_size{{NUM_W, NUM_H}, {BIGNUM_W, BIGNUM_H}}; - for (auto& size : array_size) { - std::array array; - const size_t pavail = getFreeMem(); - HIP_ARRAY_DESCRIPTOR desc; - desc.NumChannels = 1; - desc.Width = std::get<0>(size); - desc.Height = std::get<1>(size); - desc.Format = HIP_AD_FORMAT_FLOAT; - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK(hipArrayCreate(&array[i], &desc)); - } - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK(hipArrayDestroy(array[i])); - } - const size_t avail = getFreeMem(); - if (pavail != avail) { - HIPASSERT(false); - } + HIP_CHECK_THREAD(hipSetDevice(gpu)); + std::pair size = + GENERATE(std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H)); + std::array array; + size_t pavail, avail; + HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr)); + HIP_ARRAY_DESCRIPTOR desc; + desc.NumChannels = 1; + desc.Width = std::get<0>(size); + desc.Height = std::get<1>(size); + desc.Format = HIP_AD_FORMAT_FLOAT; + + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipArrayCreate(&array[i], &desc)); } + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipArrayDestroy(array[i])); + } + + HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr)); + REQUIRE_THREAD(pavail == avail); } /* This testcase verifies hipArrayCreate API for small and big chunks data*/ -TEST_CASE("Unit_hipArrayCreate_DiffSizes") { ArrayCreate_DiffSizes(0); } +TEST_CASE("Unit_hipArrayCreate_DiffSizes") { + ArrayCreate_DiffSizes(0); + HIP_CHECK_THREAD_FINALIZE(); +} /* This testcase verifies the hipArrayCreate API in multithreaded @@ -90,13 +93,13 @@ TEST_CASE("Unit_hipArrayCreate_MultiThread") { const size_t pavail = getFreeMem(); for (int i = 0; i < devCnt; i++) { - // FIXME: the HIP_CHECK and HIPASSERT are not threadsafe so this test is broken. threadlist.push_back(std::thread(ArrayCreate_DiffSizes, i)); } for (auto& t : threadlist) { t.join(); } + HIP_CHECK_THREAD_FINALIZE(); const size_t avail = getFreeMem(); if (pavail != avail) { @@ -282,28 +285,26 @@ TEMPLATE_TEST_CASE("Unit_hipArrayCreate_maxTexture", "", uint, int, int4, ushort desc.Format = vec_info::format; desc.NumChannels = vec_info::size; - int device; - HIP_CHECK(hipGetDevice(&device)); - hipDeviceProp_t prop; - HIP_CHECK(hipGetDeviceProperties(&prop, device)); + const Sizes sizes(hipArrayDefault); + const size_t s = 64; hiparray array{}; SECTION("Happy") { SECTION("1D - Max") { - desc.Width = prop.maxTexture1D; + desc.Width = sizes.max1D; desc.Height = 0; } SECTION("2D - Max Width") { - desc.Width = prop.maxTexture2D[0]; - desc.Height = 64; + desc.Width = sizes.max2D[0]; + desc.Height = s; } SECTION("2D - Max Height") { - desc.Width = 64; - desc.Height = prop.maxTexture2D[1]; + desc.Width = s; + desc.Height = sizes.max2D[1]; } SECTION("2D - Max Width and Height") { - desc.Width = prop.maxTexture2D[0]; - desc.Height = prop.maxTexture2D[1]; + desc.Width = sizes.max2D[0]; + desc.Height = sizes.max2D[1]; } auto maxArrayCreateError = hipArrayCreate(&array, &desc); // this can try to alloc many GB of memory, so out of memory is acceptable @@ -314,20 +315,20 @@ TEMPLATE_TEST_CASE("Unit_hipArrayCreate_maxTexture", "", uint, int, int4, ushort } SECTION("Negative") { SECTION("1D - More Than Max") { - desc.Width = prop.maxTexture1D + 1; + desc.Width = sizes.max1D + 1; desc.Height = 0; } SECTION("2D - More Than Max Width") { - desc.Width = prop.maxTexture2D[0] + 1; - desc.Height = 64; + desc.Width = sizes.max2D[0] + 1; + desc.Height = s; } SECTION("2D - More Than Max Height") { - desc.Width = 64; - desc.Height = prop.maxTexture2D[1] + 1; + desc.Width = s; + desc.Height = sizes.max2D[1] + 1; } SECTION("2D - More Than Max Width and Height") { - desc.Width = prop.maxTexture2D[0] + 1; - desc.Height = prop.maxTexture2D[1] + 1; + desc.Width = sizes.max2D[0] + 1; + desc.Height = sizes.max2D[1] + 1; } HIP_CHECK_ERROR(hipArrayCreate(&array, &desc), hipErrorInvalidValue); } diff --git a/tests/catch/unit/memory/hipMalloc3DArray.cc b/tests/catch/unit/memory/hipMalloc3DArray.cc index fe8ba3f8b0..253913345e 100644 --- a/tests/catch/unit/memory/hipMalloc3DArray.cc +++ b/tests/catch/unit/memory/hipMalloc3DArray.cc @@ -25,9 +25,9 @@ hipMalloc3DArray API test scenarios 4. Multithreaded scenario */ - - +#include #include +#include "hipArrayCommon.hh" static constexpr auto ARRAY_SIZE{4}; static constexpr auto BIG_ARRAY_SIZE{100}; @@ -47,34 +47,24 @@ static constexpr auto ARRAY_LOOP{100}; * */ static void Malloc3DArray_DiffSizes(int gpu) { - HIP_CHECK(hipSetDevice(gpu)); - std::vector array_size; - array_size.push_back(ARRAY_SIZE); - array_size.push_back(BIG_ARRAY_SIZE); - for (auto &size : array_size) { - int width{size}, height{size}, depth{size}; - hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0, - 0, 0, hipChannelFormatKindFloat); - hipArray *arr[ARRAY_LOOP]; - size_t tot, avail, ptot, pavail; - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK(hipMalloc3DArray(&arr[i], &channelDesc, make_hipExtent(width, - height, depth), hipArrayDefault)); - } - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK(hipFreeArray(arr[i])); - } - HIP_CHECK(hipMemGetInfo(&avail, &tot)); - if ((pavail != avail)) { - HIPASSERT(false); - } - } -} + HIP_CHECK_THREAD(hipSetDevice(gpu)); + const int size = GENERATE(ARRAY_SIZE, BIG_ARRAY_SIZE); + int width{size}, height{size}, depth{size}; + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + std::array arr; + size_t pavail, avail; + HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr)); -/* Thread Function */ -static void Malloc3DArrayThreadFunc(int gpu) { - Malloc3DArray_DiffSizes(gpu); + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipMalloc3DArray(&arr[i], &channelDesc, make_hipExtent(width, height, depth), + hipArrayDefault)); + } + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipFreeArray(arr[i])); + } + + HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr)); + REQUIRE_THREAD(pavail == avail); } /* @@ -82,91 +72,53 @@ static void Malloc3DArrayThreadFunc(int gpu) { */ TEST_CASE("Unit_hipMalloc3DArray_Negative") { constexpr int width{ARRAY_SIZE}, height{ARRAY_SIZE}, depth{ARRAY_SIZE}; - hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0, - 0, 0, hipChannelFormatKindFloat); - hipArray *arr; + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + hipArray* arr; #if HT_NVIDIA SECTION("NullPointer to Array") { - REQUIRE(hipMalloc3DArray(nullptr, &channelDesc, make_hipExtent(width, - height, depth), hipArrayDefault) != hipSuccess); + REQUIRE(hipMalloc3DArray(nullptr, &channelDesc, make_hipExtent(width, height, depth), + hipArrayDefault) != hipSuccess); } SECTION("NullPointer to Channel Descriptor") { - REQUIRE(hipMalloc3DArray(&arr, nullptr, make_hipExtent(width, - height, depth), hipArrayDefault) != hipSuccess); + REQUIRE(hipMalloc3DArray(&arr, nullptr, make_hipExtent(width, height, depth), + hipArrayDefault) != hipSuccess); } #endif SECTION("Width 0 in hipExtent") { - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(0, - height, width), hipArrayDefault) != hipSuccess); + REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(0, height, width), + hipArrayDefault) != hipSuccess); } SECTION("Height 0 in hipExtent") { - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, - 0, width), hipArrayDefault) != hipSuccess); + REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, 0, width), + hipArrayDefault) != hipSuccess); } SECTION("Invalid Flag") { - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, - height, depth), 100) != hipSuccess); + REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), 100) != + hipSuccess); } SECTION("Width,Height & Depth 0 in hipExtent") { - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(0, - 0, 0), hipArrayDefault) != hipSuccess); + REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(0, 0, 0), hipArrayDefault) != + hipSuccess); } SECTION("Max int values to extent") { - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, - make_hipExtent(std::numeric_limits::max(), - std::numeric_limits::max(), - std::numeric_limits::max()), - hipArrayDefault) != hipSuccess); + REQUIRE(hipMalloc3DArray( + &arr, &channelDesc, + make_hipExtent(std::numeric_limits::max(), std::numeric_limits::max(), + std::numeric_limits::max()), + hipArrayDefault) != hipSuccess); } } -/* - * Verifies the extent validation scenarios - * 1. Passing depth as 0 would create 2D array - * 2. Passing height and depth as 0 would create 1D array - * from hipMalloc3DArray API - */ -TEST_CASE("Unit_hipMalloc3DArray_ExtentValidation") { - constexpr int width{ARRAY_SIZE}, height{ARRAY_SIZE}; - hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0, - 0, 0, hipChannelFormatKindFloat); - hipArray *arr; - - SECTION("Depth 0 in hipExtent") { - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, - height, 0), hipArrayDefault) == hipSuccess); - HIP_CHECK(hipFreeArray(arr)); - } - - SECTION("Height & Depth 0 in hipExtent") { - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, - 0, 0), hipArrayDefault) == hipSuccess); - HIP_CHECK(hipFreeArray(arr)); - } -} - -/* - * Verifies hipMalloc3DArray API by passing width,height - * and depth as 10 - */ -TEST_CASE("Unit_hipMalloc3DArray_Basic") { - constexpr int width{ARRAY_SIZE}, height{ARRAY_SIZE}, depth{ARRAY_SIZE}; - hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0, - 0, 0, hipChannelFormatKindFloat); - hipArray *arr; - - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, - height, depth), hipArrayDefault) == hipSuccess); - HIP_CHECK(hipFreeArray(arr)); -} TEST_CASE("Unit_hipMalloc3DArray_DiffSizes") { Malloc3DArray_DiffSizes(0); + HIP_CHECK_THREAD_FINALIZE(); } + /* This testcase verifies the hipMalloc3DArray API in multithreaded scenario by launching threads in parallel on multiple GPUs @@ -176,16 +128,16 @@ TEST_CASE("Unit_hipMalloc3DArray_MultiThread") { std::vector threadlist; int devCnt = 0; devCnt = HipTest::getDeviceCount(); - size_t tot, avail, ptot, pavail; - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + const auto pavail = getFreeMem(); for (int i = 0; i < devCnt; i++) { - threadlist.push_back(std::thread(Malloc3DArrayThreadFunc, i)); + threadlist.push_back(std::thread(Malloc3DArray_DiffSizes, i)); } - for (auto &t : threadlist) { + for (auto& t : threadlist) { t.join(); } - HIP_CHECK(hipMemGetInfo(&avail, &tot)); + HIP_CHECK_THREAD_FINALIZE(); + const auto avail = getFreeMem(); if (pavail != avail) { WARN("Memory leak of hipMalloc3D API in multithreaded scenario"); @@ -193,3 +145,131 @@ TEST_CASE("Unit_hipMalloc3DArray_MultiThread") { } } +void checkArrayIsExpected(hipArray_t array, const hipChannelFormatDesc& expected_desc, + const hipExtent& expected_extent, const unsigned int expected_flags) { +// hipArrayGetInfo doesn't currently exist (EXSWCPHIPT-87) +#if HT_AMD + std::ignore = array; + std::ignore = expected_desc; + std::ignore = expected_extent; + std::ignore = expected_flags; +#else + cudaChannelFormatDesc queried_desc; + cudaExtent queried_extent; + unsigned int queried_flags; + + cudaArrayGetInfo(&queried_desc, &queried_extent, &queried_flags, array); + + REQUIRE(expected_desc.x == queried_desc.x); + REQUIRE(expected_desc.y == queried_desc.y); + REQUIRE(expected_desc.z == queried_desc.z); + REQUIRE(expected_desc.f == queried_desc.f); + + REQUIRE(expected_extent.width == queried_extent.width); + REQUIRE(expected_extent.height == queried_extent.height); + REQUIRE(expected_extent.depth == queried_extent.depth); + + REQUIRE(expected_flags == queried_flags); +#endif +} + +TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_happy", "", char, uchar2, uint2, int4, short4, float, + float2, float4) { + hipArray_t array; + const auto desc = hipCreateChannelDesc(); +#if HT_AMD + const unsigned int flags = hipArrayDefault; +#else + const unsigned int flags = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore); +#endif + constexpr size_t size = 64; + hipExtent extent; + + SECTION("1D Array") { + extent = make_hipExtent(size, 0, 0); + HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); + } + SECTION("2D Array") { + extent = make_hipExtent(size, size, 0); + HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); + } + SECTION("3D Array") { + extent = make_hipExtent(size, size, size); + HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); + } + + checkArrayIsExpected(array, desc, extent, flags); + + HIP_CHECK(hipFreeArray(array)); +} + +TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_MaxTexture", "", int, uint4, short, ushort2, + unsigned char, float, float4) { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-97"); + return; +#endif + + hipArray_t array; + const hipChannelFormatDesc desc = hipCreateChannelDesc(); +#if HT_AMD + const unsigned int flag = hipArrayDefault; +#else + const unsigned int flag = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore); +#endif + if (flag == hipArraySurfaceLoadStore) { + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-58"); + return; + } + CAPTURE(flag); + const Sizes sizes(flag); + CAPTURE(sizes.max1D, sizes.max2D, sizes.max3D); + + const size_t s = 64; + SECTION("Happy") { + // stored in a vector so some values can be ifdef'd out + std::vector extentsToTest{ + make_hipExtent(sizes.max1D, 0, 0), // 1D max + make_hipExtent(sizes.max2D[0], s, 0), // 2D max width + make_hipExtent(s, sizes.max2D[1], 0), // 2D max height + make_hipExtent(sizes.max2D[0], sizes.max2D[1], 0), // 2D max + make_hipExtent(sizes.max3D[0], s, s), // 3D max width + make_hipExtent(s, sizes.max3D[1], s), // 3D max height + make_hipExtent(s, s, sizes.max3D[2]), // 3D max depth + make_hipExtent(s, sizes.max3D[1], sizes.max3D[2]), // 3D max height and depth + make_hipExtent(sizes.max3D[0], s, sizes.max3D[2]), // 3D max width and depth + make_hipExtent(sizes.max3D[0], sizes.max3D[1], s), // 3D max width and height + make_hipExtent(sizes.max3D[0], sizes.max3D[1], sizes.max3D[2]) // 3D max + }; + const auto extent = + GENERATE_COPY(from_range(std::begin(extentsToTest), std::end(extentsToTest))); + CAPTURE(extent.width, extent.height, extent.depth); + auto maxArrayCreateError = hipMalloc3DArray(&array, &desc, extent, flag); + // this can try to alloc many GB of memory, so out of memory is acceptable + if (maxArrayCreateError == hipErrorOutOfMemory) return; + HIP_CHECK(maxArrayCreateError); + checkArrayIsExpected(array, desc, extent, flag); + HIP_CHECK(hipFreeArray(array)); + } + SECTION("Negative") { + std::vector extentsToTest { + make_hipExtent(sizes.max1D + 1, 0, 0), // 1D max + make_hipExtent(sizes.max2D[0] + 1, s, 0), // 2D max width + make_hipExtent(s, sizes.max2D[1] + 1, 0), // 2D max height + make_hipExtent(sizes.max2D[0] + 1, sizes.max2D[1] + 1, 0), // 2D max + make_hipExtent(sizes.max3D[0] + 1, s, s), // 3D max width + make_hipExtent(s, sizes.max3D[1] + 1, s), // 3D max height +#if !HT_NVIDIA // leads to hipSuccess on NVIDIA + make_hipExtent(s, s, sizes.max3D[2] + 1), // 3D max depth +#endif + make_hipExtent(s, sizes.max3D[1] + 1, sizes.max3D[2] + 1), // 3D max height and depth + make_hipExtent(sizes.max3D[0] + 1, s, sizes.max3D[2] + 1), // 3D max width and depth + make_hipExtent(sizes.max3D[0] + 1, sizes.max3D[1] + 1, s), // 3D max width and height + make_hipExtent(sizes.max3D[0] + 1, sizes.max3D[1] + 1, sizes.max3D[2] + 1) // 3D max + }; + const auto extent = + GENERATE_COPY(from_range(std::begin(extentsToTest), std::end(extentsToTest))); + CAPTURE(extent.width, extent.height, extent.depth); + HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, extent, flag), hipErrorInvalidValue); + } +} diff --git a/tests/catch/unit/memory/hipMallocArray.cc b/tests/catch/unit/memory/hipMallocArray.cc index 8859aadb45..820e917313 100644 --- a/tests/catch/unit/memory/hipMallocArray.cc +++ b/tests/catch/unit/memory/hipMallocArray.cc @@ -30,11 +30,11 @@ hipMallocArray API test scenarios #include #include "hipArrayCommon.hh" -static constexpr auto NUM_W{4}; -static constexpr auto BIGNUM_W{100}; -static constexpr auto BIGNUM_H{100}; -static constexpr auto NUM_H{4}; -static constexpr auto ARRAY_LOOP{100}; +static constexpr size_t NUM_W{4}; +static constexpr size_t NUM_H{4}; +static constexpr size_t BIGNUM_W{100}; +static constexpr size_t BIGNUM_H{100}; +static constexpr int ARRAY_LOOP{100}; /* * This API verifies memory allocations for small and @@ -50,28 +50,30 @@ static constexpr auto ARRAY_LOOP{100}; * */ static void MallocArray_DiffSizes(int gpu) { - HIP_CHECK(hipSetDevice(gpu)); - std::vector> array_size{{NUM_W, NUM_H}, {BIGNUM_W, BIGNUM_H}}; - for (auto& size : array_size) { - std::array A_d; - size_t tot, avail, ptot, pavail; - hipChannelFormatDesc desc = hipCreateChannelDesc(); - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK( - hipMallocArray(&A_d[i], &desc, std::get<0>(size), std::get<1>(size), hipArrayDefault)); - } - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK(hipFreeArray(A_d[i])); - } - HIP_CHECK(hipMemGetInfo(&avail, &tot)); - if ((pavail != avail)) { - HIPASSERT(false); - } + HIP_CHECK_THREAD(hipSetDevice(gpu)); + std::pair size = + GENERATE(std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H)); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + std::array A_d; + size_t pavail, avail; + HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr)); + + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD( + hipMallocArray(&A_d[i], &desc, std::get<0>(size), std::get<1>(size), hipArrayDefault)); } + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipFreeArray(A_d[i])); + } + + HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr)); + REQUIRE_THREAD(pavail == avail); } -TEST_CASE("Unit_hipMallocArray_DiffSizes") { MallocArray_DiffSizes(0); } +TEST_CASE("Unit_hipMallocArray_DiffSizes") { + MallocArray_DiffSizes(0); + HIP_CHECK_THREAD_FINALIZE(); +} /* This testcase verifies the hipMallocArray API in multithreaded @@ -82,18 +84,17 @@ TEST_CASE("Unit_hipMallocArray_MultiThread") { std::vector threadlist; int devCnt = 0; devCnt = HipTest::getDeviceCount(); - size_t tot, avail, ptot, pavail; - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); + const auto pavail = getFreeMem(); for (int i = 0; i < devCnt; i++) { - // FIXME: the HIP_CHECK and HIPASSERT are not threadsafe so this test is broken. threadlist.push_back(std::thread(MallocArray_DiffSizes, i)); } for (auto& t : threadlist) { t.join(); } - HIP_CHECK(hipMemGetInfo(&avail, &tot)); + HIP_CHECK_THREAD_FINALIZE(); + const auto avail = getFreeMem(); if (pavail != avail) { WARN("Memory leak of hipMalloc3D API in multithreaded scenario"); REQUIRE(false); @@ -450,32 +451,32 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_happy", "", uint, int, int4, ushort, sho // EXSWCPHIPT-71 - no equivalent value for maxSurface and maxTexture2DGather. TEMPLATE_TEST_CASE("Unit_hipMallocArray_MaxTexture_Default", "", uint, int4, ushort, short2, char, char4, float2, float4) { - int device; - HIP_CHECK(hipGetDevice(&device)); - hipDeviceProp_t prop; - HIP_CHECK(hipGetDeviceProperties(&prop, device)); - size_t width, height; hipArray_t array{}; hipChannelFormatDesc desc = hipCreateChannelDesc(); const unsigned int flag = hipArrayDefault; + const Sizes sizes(flag); + CAPTURE(sizes.max1D, sizes.max2D, sizes.max3D); + + const size_t s = 64; + SECTION("Happy") { SECTION("1D - Max") { - width = prop.maxTexture1D; + width = sizes.max1D; height = 0; } SECTION("2D - Max Width") { - width = prop.maxTexture2D[0]; - height = 64; + width = sizes.max2D[0]; + height = s; } SECTION("2D - Max Height") { - width = 64; - height = prop.maxTexture2D[1]; + width = s; + height = sizes.max2D[1]; } SECTION("2D - Max Width and Height") { - width = prop.maxTexture2D[0]; - height = prop.maxTexture2D[1]; + width = sizes.max2D[0]; + height = sizes.max2D[1]; } auto maxArrayCreateError = hipMallocArray(&array, &desc, width, height, flag); // this can try to alloc many GB of memory, so out of memory is acceptable @@ -485,20 +486,20 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_MaxTexture_Default", "", uint, int4, ush } SECTION("Negative") { SECTION("1D - More Than Max") { - width = prop.maxTexture1D + 1; + width = sizes.max1D + 1; height = 0; } SECTION("2D - More Than Max Width") { - width = prop.maxTexture2D[0] + 1; - height = 64; + width = sizes.max2D[0] + 1; + height = s; } SECTION("2D - More Than Max Height") { - width = 64; - height = prop.maxTexture2D[1] + 1; + width = s; + height = sizes.max2D[1] + 1; } SECTION("2D - More Than Max Width and Height") { - width = prop.maxTexture2D[0] + 1; - height = prop.maxTexture2D[1] + 1; + width = sizes.max2D[0] + 1; + height = sizes.max2D[1] + 1; } HIP_CHECK_ERROR(hipMallocArray(&array, &desc, width, height, flag), hipErrorInvalidValue); } From 2b1794e5d24e24076a8d88abd4d1cfa4b32e850f Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Mon, 11 Jul 2022 10:07:22 +0100 Subject: [PATCH 22/30] EXSWCPHIPT-37 - Add test for hipStreamAttachMemAsync (#2558) --- tests/catch/unit/stream/CMakeLists.txt | 4 + .../unit/stream/hipStreamAttachMemAsync.cc | 143 ++++++++++++++++++ 2 files changed, 147 insertions(+) create mode 100644 tests/catch/unit/stream/hipStreamAttachMemAsync.cc diff --git a/tests/catch/unit/stream/CMakeLists.txt b/tests/catch/unit/stream/CMakeLists.txt index 932e76f718..229df157c1 100644 --- a/tests/catch/unit/stream/CMakeLists.txt +++ b/tests/catch/unit/stream/CMakeLists.txt @@ -25,9 +25,13 @@ set(TEST_SRC hipStreamCreateWithFlags.cc hipStreamCreateWithPriority.cc hipAPIStreamDisable.cc + # hipStreamAttachMemAsync.cc # Disabling it on nvidia due to issue in function definition of hipStreamAttachMemAsync + # Fixing would break ABI, to be re-enabled when the fix is made. streamCommon.cc hipStreamValue.cc ) + + set_source_files_properties(hipStreamAttachMemAsync.cc PROPERTIES COMPILE_FLAGS -std=c++17) endif() hip_add_exe_to_target(NAME StreamTest diff --git a/tests/catch/unit/stream/hipStreamAttachMemAsync.cc b/tests/catch/unit/stream/hipStreamAttachMemAsync.cc new file mode 100644 index 0000000000..e02754dfbb --- /dev/null +++ b/tests/catch/unit/stream/hipStreamAttachMemAsync.cc @@ -0,0 +1,143 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +// TODO Enable it after hipStreamAttachMemAsync is feature complete on HIP + +#include +#include + +__device__ __managed__ int var = 0; + +enum class StreamAttachTestType { NullStream = 0, StreamPerThread, CreatedStream }; + +TEST_CASE("Unit_hipStreamAttachMemAsync_Negative") { + hipStream_t stream{nullptr}; + + auto streamType = + GENERATE(StreamAttachTestType::NullStream, StreamAttachTestType::StreamPerThread, + StreamAttachTestType::CreatedStream); + + if (streamType == StreamAttachTestType::StreamPerThread) { + stream = hipStreamPerThread; + } else if (streamType == StreamAttachTestType::CreatedStream) { + HIP_CHECK(hipStreamCreate(&stream)); + REQUIRE(stream != nullptr); + } + + SECTION("Invalid Resource Handle") { + int definitelyNotAManagedVariable = 0; + HIP_CHECK_ERROR( + hipStreamAttachMemAsync(stream, reinterpret_cast(&definitelyNotAManagedVariable), + sizeof(int), hipMemAttachSingle), + hipErrorInvalidValue); + } + + SECTION("Invalid devptr") { + HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream, nullptr, sizeof(int), hipMemAttachSingle), + hipErrorInvalidValue); + } + + SECTION("Invalid Resource Size") { + HIP_CHECK_ERROR(hipStreamAttachMemAsync(stream, reinterpret_cast(&var), sizeof(int) - 1, + hipMemAttachSingle), + hipErrorInvalidValue); + } + + SECTION("Invalid Flags") { + HIP_CHECK_ERROR( + hipStreamAttachMemAsync(stream, reinterpret_cast(&var), sizeof(int) - 1, + hipMemAttachSingle | hipMemAttachHost | hipMemAttachGlobal), + hipErrorInvalidValue); + } + + if (streamType == StreamAttachTestType::CreatedStream) { + HIP_CHECK(hipStreamDestroy(stream)); + } +} + +__global__ void kernel(int* ptr, size_t size) { + auto i = threadIdx.x; + if (i < size) { + ptr[i] = 1024; + } +} + +constexpr size_t size = 1024; +__device__ __managed__ int m_memory[size]; + +TEST_CASE("Unit_hipStreamAttachMemAsync_UseCase") { + hipStream_t stream{nullptr}; + + auto streamType = + GENERATE(StreamAttachTestType::NullStream, StreamAttachTestType::StreamPerThread, + StreamAttachTestType::CreatedStream); + + if (streamType == StreamAttachTestType::CreatedStream) { + HIP_CHECK(hipStreamCreate(&stream)); + REQUIRE(stream != nullptr); + } + + SECTION("Size zero is valid") { + int* d_memory{nullptr}; + HIP_CHECK(hipMallocManaged(&d_memory, sizeof(int) * size, hipMemAttachHost)); + HIP_CHECK( + hipStreamAttachMemAsync(stream, reinterpret_cast(d_memory), 0, hipMemAttachHost)); + HIP_CHECK(hipStreamSynchronize(stream)); // Wait for command to complete + HIP_CHECK(hipFree(d_memory)); + } + + SECTION("Access from device and host") { + int* d_memory{nullptr}; + + HIP_CHECK(hipMallocManaged(&d_memory, sizeof(int) * size, hipMemAttachHost)); + HIP_CHECK(hipMemset(d_memory, 0, sizeof(int) * size)); + HIP_CHECK( + hipStreamAttachMemAsync(stream, reinterpret_cast(d_memory), 0, hipMemAttachHost)); + HIP_CHECK(hipStreamSynchronize(stream)); // Wait for the command to complete + + kernel<<<1, size, 0, stream>>>(d_memory, size); + HIP_CHECK(hipStreamSynchronize(stream)); // Wait for the kernel to complete + + auto ptr = std::make_unique(size); + std::copy(d_memory, d_memory + size, ptr.get()); + + HIP_CHECK(hipFree(d_memory)); + + REQUIRE(std::all_of(ptr.get(), ptr.get() + size, [](int n) { return n == size; })); + } + + SECTION("Access ManagedMemory") { + HIP_CHECK(hipMemset(m_memory, 0, sizeof(int) * size)); + HIP_CHECK( + hipStreamAttachMemAsync(stream, reinterpret_cast(m_memory), 0, hipMemAttachHost)); + HIP_CHECK(hipStreamSynchronize(stream)); // Wait for the command to complete + + kernel<<<1, size, 0, stream>>>(m_memory, size); + HIP_CHECK(hipStreamSynchronize(stream)); // Wait for the kernel to complete + + auto ptr = std::make_unique(size); + std::copy(m_memory, m_memory + size, ptr.get()); + + REQUIRE(std::all_of(ptr.get(), ptr.get() + size, [](int n) { return n == size; })); + } + + if (streamType == StreamAttachTestType::CreatedStream) { + HIP_CHECK(hipStreamDestroy(stream)); + } +} \ No newline at end of file From 5b10765c6b0324089ef77dde39880dd2068f9ff4 Mon Sep 17 00:00:00 2001 From: Anton Mitkov Date: Mon, 11 Jul 2022 10:12:18 +0100 Subject: [PATCH 23/30] EXSWCPHIPT-88 - [catch2] Testing for hipHostGetFlags API extention (#2691) --- tests/catch/unit/memory/hipHostGetFlags.cc | 215 ++++++++++++++++----- 1 file changed, 172 insertions(+), 43 deletions(-) diff --git a/tests/catch/unit/memory/hipHostGetFlags.cc b/tests/catch/unit/memory/hipHostGetFlags.cc index f150aaa5a8..25fb5b2734 100644 --- a/tests/catch/unit/memory/hipHostGetFlags.cc +++ b/tests/catch/unit/memory/hipHostGetFlags.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -27,8 +27,34 @@ This testcase verifies the basic scenario of hipHostGetFlags API #include #include #include +#include +#include -static constexpr auto LEN{1024*1024}; +std::vector FlagPart1Vec{hipHostMallocDefault, + hipHostMallocDefault | hipHostMallocPortable, + hipHostMallocDefault | hipHostMallocMapped, + hipHostMallocDefault | hipHostMallocWriteCombined, + hipHostMallocPortable, + hipHostMallocPortable | hipHostMallocMapped, + hipHostMallocPortable | hipHostMallocWriteCombined, + hipHostMallocMapped, + hipHostMallocMapped | hipHostMallocWriteCombined, + hipHostMallocWriteCombined}; +#if HT_AMD +// For cases where flags from FlagPart1Vec are not used, +// hipHostMallocDefault is the default on AMD +// and hipHostMallocMapped on Nvidia +std::vector FlagPart2Vec{0x0, + hipHostMallocNumaUser, + hipHostMallocNumaUser | hipHostMallocCoherent, + hipHostMallocNumaUser | hipHostMallocNonCoherent, + hipHostMallocCoherent, + hipHostMallocNonCoherent}; +#else +std::vector FlagPart2Vec{0x0}; +#endif + +static constexpr auto LEN{1024 * 1024}; /* This testcase verifies hipHostGetFlags API basic scenario @@ -38,57 +64,160 @@ This testcase verifies hipHostGetFlags API basic scenario 3. Validates it with the initial flags used while allocating memory */ -TEMPLATE_TEST_CASE("Unit_hipHostGetFlags_Basic", "", int, - float, double) { - constexpr auto SIZE{LEN * sizeof(TestType)}; - TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; - TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; - unsigned int FlagA, FlagB, FlagC; - FlagA = hipHostMallocWriteCombined | hipHostMallocMapped; - FlagB = hipHostMallocWriteCombined | hipHostMallocMapped; - FlagC = hipHostMallocMapped; +/* Possible host flags + * hipHostMallocDefault 0x0 + * hipHostMallocPortable 0x1 + * hipHostMallocMapped 0x2 + * hipHostMallocWriteCombined 0x4 + * NOT on Nvidia + * hipHostMallocNumaUser 0x20000000 + * hipHostMallocCoherent 0x40000000 + * hipHostMallocNonCoherent 0x80000000 + */ + +inline void checkFlags(unsigned int expected, unsigned int obtained) { + // Account for cases where flags from FlagPart1Vec do not include hipHostMallocMapped, + // on Nvidia devices it is added by default +#if HT_NVIDIA + expected = expected | hipHostMallocMapped; +#endif + REQUIRE(expected == obtained); +} + +TEST_CASE("Unit_hipHostGetFlags_flagCombos") { + + constexpr auto SIZE{LEN * sizeof(int)}; + int* A_h{nullptr}; + + const unsigned int FlagPart1 = GENERATE(from_range(FlagPart1Vec.begin(), FlagPart1Vec.end())); + const unsigned int FlagPart2 = GENERATE(from_range(FlagPart2Vec.begin(), FlagPart2Vec.end())); + + unsigned int FlagComp = FlagPart1 | FlagPart2; hipDeviceProp_t prop; - int device; + int device{}; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + + // Skip test if device does not support the property canMapHostMemory + if (prop.canMapHostMemory != 1) { + HipTest::HIP_SKIP_TEST("Device Property canMapHostMemory is not set"); + return; + } else { + // Allocate using the generated flags combos + INFO("Flag passed when allocating: 0x" << std::hex << FlagComp << "\n"); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), SIZE, FlagComp)); + unsigned int flagA{}; + + // get the flags from allocations and check if they are the same as the one set + HIP_CHECK(hipHostGetFlags(&flagA, A_h)); + + checkFlags(FlagComp, flagA); + HIP_CHECK(hipHostFree(A_h)); + } +} + +// Test Allocation with flags and getting flags in another thread +TEST_CASE("Unit_hipHostGetFlags_DifferentThreads") { + constexpr auto SIZE{LEN * sizeof(int)}; + int* A_h{nullptr}; + + const unsigned int FlagPart1 = GENERATE(from_range(FlagPart1Vec.begin(), FlagPart1Vec.end())); + const unsigned int FlagPart2 = GENERATE(from_range(FlagPart2Vec.begin(), FlagPart2Vec.end())); + + + unsigned int FlagComp = FlagPart1 | FlagPart2; + + hipDeviceProp_t prop; + int device{}; HIP_CHECK(hipGetDevice(&device)); HIP_CHECK(hipGetDeviceProperties(&prop, device)); if (prop.canMapHostMemory != 1) { - SUCCEED("Device Property canMapHostMemory is not set"); + HipTest::HIP_SKIP_TEST("Device Property canMapHostMemory is not set"); + return; } else { - HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), SIZE, - hipHostMallocWriteCombined | hipHostMallocMapped)); - HIP_CHECK(hipHostMalloc(reinterpret_cast(&B_h), SIZE, - hipHostMallocWriteCombined | hipHostMallocMapped)); - HIP_CHECK(hipHostMalloc(reinterpret_cast(&C_h), SIZE, - hipHostMallocMapped)); - - unsigned int flagA, flagB, flagC; - HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), A_h, 0)); - HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&B_d), B_h, 0)); - HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&C_d), C_h, 0)); + // Make sure we allocate before trying to get the flags + std::thread malloc_thread( + [&]() { HIP_CHECK_THREAD(hipHostMalloc(reinterpret_cast(&A_h), SIZE, FlagComp)); }); + malloc_thread.join(); + HIP_CHECK_THREAD_FINALIZE(); + unsigned int flagA{}; HIP_CHECK(hipHostGetFlags(&flagA, A_h)); - HIP_CHECK(hipHostGetFlags(&flagB, B_h)); - HIP_CHECK(hipHostGetFlags(&flagC, C_h)); - HipTest::setDefaultData(LEN, A_h, B_h, C_h); + checkFlags(FlagComp, flagA); - dim3 dimGrid(LEN / 512, 1, 1); - dim3 dimBlock(512, 1, 1); - hipLaunchKernelGGL(HipTest::vectorADD, dimGrid, dimBlock, - 0, 0, static_cast(A_d), - static_cast(B_d), C_d, LEN); - - HIP_CHECK(hipMemcpy(C_h, C_d, SIZE, hipMemcpyDeviceToHost)); - // Note this really HostToHost not - // DeviceToHost, since memory is mapped... - HipTest::checkVectorADD(A_h, B_h, C_h, LEN); - - REQUIRE(flagA == FlagA); - REQUIRE(flagB == FlagB); - REQUIRE(flagC == FlagC); HIP_CHECK(hipHostFree(A_h)); - HIP_CHECK(hipHostFree(B_h)); - HIP_CHECK(hipHostFree(C_h)); + } +} + +// Test behaviour of hipHostGetFlags with invalid args +TEST_CASE("Unit_hipHostGetFlags_InvalidArgs") { + constexpr auto SIZE{LEN * sizeof(int)}; + int* A_h{nullptr}; + + hipDeviceProp_t prop; + int device{}; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + + // Skip test if device does not support the property canMapHostMemory + if (prop.canMapHostMemory != 1) { + HipTest::HIP_SKIP_TEST("Device Property canMapHostMemory is not set"); + return; + } else { + SECTION("Invalid flag ptr being passed to hipHostGetFlags") { + // Use default flag + unsigned int FlagComp = 0x0; + + // Allocate using the generated flags combos + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), SIZE, FlagComp)); + + // use a nullptr to return flags to + unsigned int* flagA = nullptr; + + // get the flags from allocations and check if they are the same as the one set + HIP_CHECK_ERROR(hipHostGetFlags(flagA, A_h), hipErrorInvalidValue); + + HIP_CHECK(hipHostFree(A_h)); + } + + SECTION("Device ptr allocated with hipMalloc passed to hipHostGetFlags") { + unsigned int FlagComp = 0x4; + + // Allocate memory on device + HIP_CHECK(hipMalloc(reinterpret_cast(&A_h), SIZE)); + + unsigned int flagA{}; + + // get the flags from allocations and check if they are the same as the one set + HIP_CHECK_ERROR(hipHostGetFlags(&flagA, A_h), hipErrorInvalidValue); + INFO("Flag passed when allocating: " << std::hex << FlagComp << " Returned flag: " << std::hex + << flagA << "\n"); + + HIP_CHECK(hipFree(A_h)); + } + + SECTION("Ptr from hipHostGetDevicePointer passed to hipHostGetFlags") { + unsigned int FlagComp = 0x4; + + int* A_d{nullptr}; + // Allocate memory on device + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), SIZE, FlagComp)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), A_h, 0)); + + unsigned int flagA; + + // get the flags from allocations and check if they are the same as the one set + HIP_CHECK(hipHostGetFlags(&flagA, A_d)); + INFO("Flag passed when allocating: " << std::hex << FlagComp << " Returned flag: " << std::hex + << flagA << "\n"); +#if HT_NVIDIA + // on Nvidia adjust for cudaHostAllocMapped being set by default + FlagComp = FlagComp | hipHostMallocMapped; +#endif + REQUIRE(flagA == FlagComp); + HIP_CHECK(hipHostFree(A_h)); + } } } From 9595654ddf6fa3ba636c4cc9e45dfb202b4979e2 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Mon, 11 Jul 2022 10:32:51 +0100 Subject: [PATCH 24/30] EXSWCPHIPT-125 - Add tests for hipGetDeviceCount and restructure some SpawnProc tests (#2765) --- tests/catch/README.md | 6 +- tests/catch/include/hip_test_process.hh | 35 ++++++++-- tests/catch/unit/CMakeLists.txt | 1 - tests/catch/unit/device/CMakeLists.txt | 6 ++ tests/catch/unit/device/getDeviceCount_exe.cc | 68 +++++++++++++++++++ tests/catch/unit/device/hipGetDeviceCount.cc | 23 +++++++ tests/catch/unit/printf/CMakeLists.txt | 7 ++ tests/catch/unit/printf/printfFlags.cc | 18 +---- .../printfFlags_exe.cc} | 2 +- tests/catch/unit/printf/printfSpecifiers.cc | 4 +- .../printfSpecifiers_exe.cc} | 2 +- tests/catch/unit/printfExe/CMakeLists.txt | 5 -- 12 files changed, 144 insertions(+), 33 deletions(-) create mode 100644 tests/catch/unit/device/getDeviceCount_exe.cc rename tests/catch/unit/{printfExe/printfFlags.cc => printf/printfFlags_exe.cc} (95%) rename tests/catch/unit/{printfExe/printfSepcifiers.cc => printf/printfSpecifiers_exe.cc} (97%) delete mode 100644 tests/catch/unit/printfExe/CMakeLists.txt diff --git a/tests/catch/README.md b/tests/catch/README.md index 7d1ff6265b..a0c3c3adc2 100644 --- a/tests/catch/README.md +++ b/tests/catch/README.md @@ -144,7 +144,7 @@ If there arises a condition where certain flag is disabled and due to which a te ```cpp TEST_CASE("TestOnlyOnXnack") { if(!XNACKEnabled) { - HIP_SKIP_TEST("Test only runs on system with XNACK enabled"); + HipTest::HIP_SKIP_TEST("Test only runs on system with XNACK enabled"); return; } // Rest of test functionality @@ -162,11 +162,11 @@ These macros are to be called in multi process tests, inside a process which get There is a special interface available for process isolation. ```hip::SpawnProc``` in ```hip_test_process.hh```. Using this interface test can spawn a process and place passing conditions on its return value or its output to stdout. This can be useful for testing printf output. Sample Usage: ```cpp -hip::SpawnProc proc(, ); +hip::SpawnProc proc(, ); REQUIRE(0 == proc.run()); // Test of return value of the proc REQUIRE(exepctedOutput == proc.getOutput()); // Test on expected output of the process ``` -The process can be a standalone exe (see tests/catch/unit/printfExe for more information). +The process must be a standalone exe inside the same folder as other tests. ## Enabling New Tests Initially, the new tests can be enabled via using ```-DHIP_CATCH_TEST=1```. After porting existing tests, this will be turned on by default. diff --git a/tests/catch/include/hip_test_process.hh b/tests/catch/include/hip_test_process.hh index 665f5d6b97..335e6c7bba 100644 --- a/tests/catch/include/hip_test_process.hh +++ b/tests/catch/include/hip_test_process.hh @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -21,7 +21,9 @@ THE SOFTWARE. */ #pragma once + #include "hip_test_common.hh" +#include "hip_test_filesystem.hh" #include #include @@ -29,9 +31,17 @@ THE SOFTWARE. #include #include #include -#include "hip_test_filesystem.hh" namespace hip { +/* +Class to spawn a process in isolation and test its standard output and return status +Good for printf tests and environment variable tests + +How to use: +Have the stand alone exe in the same folder +Init a class using hip::SpawnProc proc("ExeName", yes_or_no_to_capture_output); +proc.run("Optional command line args"); +*/ class SpawnProc { std::string exeName; std::string resultStr; @@ -53,18 +63,31 @@ class SpawnProc { public: SpawnProc(std::string exeName_, bool captureOutput_ = false) : exeName(exeName_), captureOutput(captureOutput_) { - auto dir = fs::path(TestContext::get().currentPath()).parent_path(); + auto dir = fs::path(TestContext::get().currentPath()); dir /= exeName; exeName = dir.string(); + + INFO("Testing that exe exists: " << exeName); + REQUIRE(fs::exists(exeName)); + if (captureOutput) { auto path = fs::temp_directory_path(); path /= getRandomString(); tmpFileName = path.string(); + INFO("Testing that capture file does not exist already: " << tmpFileName); + REQUIRE(!fs::exists(tmpFileName)); } } - int run() { + int run(std::string commandLineArgs = "") { std::string execCmd = exeName; + + // Append command line args + if (commandLineArgs.size() > 0) { + execCmd += " "; // Add space for command line args + execCmd += commandLineArgs; + } + if (captureOutput) { execCmd += " > "; execCmd += tmpFileName; @@ -77,7 +100,11 @@ class SpawnProc { resultStr = std::string((std::istreambuf_iterator(t)), std::istreambuf_iterator()); } +#if HT_LINUX + return WEXITSTATUS(res); +#else return res; +#endif } std::string getOutput() { return resultStr; } diff --git a/tests/catch/unit/CMakeLists.txt b/tests/catch/unit/CMakeLists.txt index 9ebff9b9de..28c6fcb6aa 100644 --- a/tests/catch/unit/CMakeLists.txt +++ b/tests/catch/unit/CMakeLists.txt @@ -27,7 +27,6 @@ add_subdirectory(event) add_subdirectory(occupancy) add_subdirectory(device) add_subdirectory(printf) -add_subdirectory(printfExe) add_subdirectory(texture) add_subdirectory(streamperthread) add_subdirectory(kernel) diff --git a/tests/catch/unit/device/CMakeLists.txt b/tests/catch/unit/device/CMakeLists.txt index bebee8af71..247867bf5c 100644 --- a/tests/catch/unit/device/CMakeLists.txt +++ b/tests/catch/unit/device/CMakeLists.txt @@ -18,7 +18,13 @@ set(TEST_SRC hipDeviceGetUuid.cc ) +set_source_files_properties(hipGetDeviceCount.cc PROPERTIES COMPILE_FLAGS -std=c++17) + +add_executable(getDeviceCount EXCLUDE_FROM_ALL getDeviceCount_exe.cc) + hip_add_exe_to_target(NAME DeviceTest TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests COMPILE_OPTIONS -std=c++14) + +add_dependencies(DeviceTest getDeviceCount) diff --git a/tests/catch/unit/device/getDeviceCount_exe.cc b/tests/catch/unit/device/getDeviceCount_exe.cc new file mode 100644 index 0000000000..1da5f7c3a9 --- /dev/null +++ b/tests/catch/unit/device/getDeviceCount_exe.cc @@ -0,0 +1,68 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +// Expects 1 command line arg, which is the Device Visible String +int main(int argc, char** argv) { + if (argc != 2) { + std::cerr << "Invalid number of args passed.\n" + << "argc : " << argc << std::endl; + for (int i = 0; i < argc; i++) { + std::cerr << " argv[" << i << "] : " << argv[0] << std::endl; + } + std::cerr << "The program expects device visibility string i.e. 0,1,2" << std::endl; + return -1; + } + + // disable visible_devices env from shell +#ifdef __HIP_PLATFORM_NVCC__ + unsetenv("CUDA_VISIBLE_DEVICES"); + setenv("CUDA_VISIBLE_DEVICES", argv[1], 1); + auto init_res = hipInit(0); + if (hipSuccess != init_res) { + std::cerr << "CUDA INIT API returned : " << hipGetErrorString(init_res) << std::endl; + return -1; + } +#else + unsetenv("ROCR_VISIBLE_DEVICES"); + unsetenv("HIP_VISIBLE_DEVICES"); + setenv("ROCR_VISIBLE_DEVICES", argv[1], 1); + setenv("HIP_VISIBLE_DEVICES", argv[1], 1); +#endif + + int count = 0; + auto res = hipGetDeviceCount(&count); + if (hipSuccess != res) { + std::cerr << "HIP API returned : " << hipGetErrorString(res) << std::endl; + return -1; + } + +#ifdef __HIP_PLATFORM_NVCC__ + unsetenv("CUDA_VISIBLE_DEVICES"); +#else + unsetenv("ROCR_VISIBLE_DEVICES"); + unsetenv("HIP_VISIBLE_DEVICES"); +#endif + return count; +} \ No newline at end of file diff --git a/tests/catch/unit/device/hipGetDeviceCount.cc b/tests/catch/unit/device/hipGetDeviceCount.cc index 2179bdad05..2facbf9515 100644 --- a/tests/catch/unit/device/hipGetDeviceCount.cc +++ b/tests/catch/unit/device/hipGetDeviceCount.cc @@ -23,6 +23,7 @@ THE SOFTWARE. */ #include +#include /** * hipGetDeviceCount tests @@ -32,3 +33,25 @@ TEST_CASE("Unit_hipGetDeviceCount_NegTst") { // Scenario1 REQUIRE_FALSE(hipGetDeviceCount(nullptr) == hipSuccess); } + +TEST_CASE("Unit_hipGetDeviceCount_HideDevices") { + int deviceCount = HipTest::getDeviceCount(); + if (deviceCount < 2) { + HipTest::HIP_SKIP_TEST("This test requires more than 2 GPUs. Skipping."); + return; + } + + for (int i = deviceCount; i >= 1; i--) { + std::string visibleStr; + for (int j = 0; j < i; j++) { // Generate a string which has first i devices + visibleStr += std::to_string(j); + if (j != (i - 1)) { + visibleStr += ","; + } + } + + hip::SpawnProc proc("getDeviceCount", true); + INFO("Output from process : " << proc.getOutput()); + REQUIRE(proc.run(visibleStr) == i); + } +} diff --git a/tests/catch/unit/printf/CMakeLists.txt b/tests/catch/unit/printf/CMakeLists.txt index b911796d66..d2a9d6eced 100644 --- a/tests/catch/unit/printf/CMakeLists.txt +++ b/tests/catch/unit/printf/CMakeLists.txt @@ -16,3 +16,10 @@ elseif (HIP_PLATFORM MATCHES "nvidia") TEST_TARGET_NAME build_tests COMPILE_OPTIONS -std=c++17) endif() + +# Standalone exes +add_executable(printfFlags EXCLUDE_FROM_ALL printfFlags_exe.cc) +add_executable(printfSpecifiers EXCLUDE_FROM_ALL printfSpecifiers_exe.cc) + +add_dependencies(printfTests printfFlags) +add_dependencies(printfTests printfSpecifiers) diff --git a/tests/catch/unit/printf/printfFlags.cc b/tests/catch/unit/printf/printfFlags.cc index 8a9aa75a3f..b9955b0b60 100644 --- a/tests/catch/unit/printf/printfFlags.cc +++ b/tests/catch/unit/printf/printfFlags.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -23,20 +23,6 @@ THE SOFTWARE. #include #include -__global__ void test_kernel() { - printf("%08d\n", 42); - printf("%08i\n", -42); - printf("%08u\n", 42); - printf("%08g\n", 123.456); - printf("%0+8d\n", 42); - printf("%+d\n", -42); - printf("%+08d\n", 42); - printf("%-8s\n", "xyzzy"); - printf("% i\n", -42); - printf("%-16.8d\n", 42); - printf("%16.8d\n", 42); -} - TEST_CASE("Unit_printf_flags") { std::string reference(R"here(00000042 -0000042 @@ -51,7 +37,7 @@ xyzzy 00000042 )here"); - hip::SpawnProc proc("printfExe/printfFlags", true); + hip::SpawnProc proc("printfFlags", true); REQUIRE(proc.run() == 0); REQUIRE(proc.getOutput() == reference); } diff --git a/tests/catch/unit/printfExe/printfFlags.cc b/tests/catch/unit/printf/printfFlags_exe.cc similarity index 95% rename from tests/catch/unit/printfExe/printfFlags.cc rename to tests/catch/unit/printf/printfFlags_exe.cc index a96c099340..6ce9691b17 100644 --- a/tests/catch/unit/printfExe/printfFlags.cc +++ b/tests/catch/unit/printf/printfFlags_exe.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal diff --git a/tests/catch/unit/printf/printfSpecifiers.cc b/tests/catch/unit/printf/printfSpecifiers.cc index 37e8584943..223a3c338b 100644 --- a/tests/catch/unit/printf/printfSpecifiers.cc +++ b/tests/catch/unit/printf/printfSpecifiers.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -89,7 +89,7 @@ x )here"); #endif - hip::SpawnProc proc("printfExe/printfSepcifiers", true); + hip::SpawnProc proc("printfSpecifiers", true); REQUIRE(0 == proc.run()); REQUIRE(proc.getOutput() == reference); } diff --git a/tests/catch/unit/printfExe/printfSepcifiers.cc b/tests/catch/unit/printf/printfSpecifiers_exe.cc similarity index 97% rename from tests/catch/unit/printfExe/printfSepcifiers.cc rename to tests/catch/unit/printf/printfSpecifiers_exe.cc index 364376fb0a..5dcabdc4d4 100644 --- a/tests/catch/unit/printfExe/printfSepcifiers.cc +++ b/tests/catch/unit/printf/printfSpecifiers_exe.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal diff --git a/tests/catch/unit/printfExe/CMakeLists.txt b/tests/catch/unit/printfExe/CMakeLists.txt deleted file mode 100644 index 3b73810b35..0000000000 --- a/tests/catch/unit/printfExe/CMakeLists.txt +++ /dev/null @@ -1,5 +0,0 @@ -add_executable(printfFlags EXCLUDE_FROM_ALL printfFlags.cc) -add_executable(printfSepcifiers EXCLUDE_FROM_ALL printfSepcifiers.cc) - -add_dependencies(build_tests printfFlags) -add_dependencies(build_tests printfSepcifiers) From cdd437c9c636cd375621758635688f8d8b92ed6c Mon Sep 17 00:00:00 2001 From: Finlay Date: Mon, 11 Jul 2022 11:41:32 +0100 Subject: [PATCH 25/30] EXSWCPHIPT-113 - Unit tests for hipArray3DCreate for default and surface arrays (#2714) --- tests/catch/unit/memory/CMakeLists.txt | 2 + tests/catch/unit/memory/DriverContext.cc | 4 - tests/catch/unit/memory/DriverContext.hh | 2 - tests/catch/unit/memory/hipArray3DCreate.cc | 166 ++++++++++++++++++++ tests/catch/unit/memory/hipArrayCommon.hh | 12 +- tests/catch/unit/memory/hipMalloc3DArray.cc | 7 +- 6 files changed, 181 insertions(+), 12 deletions(-) create mode 100644 tests/catch/unit/memory/hipArray3DCreate.cc diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index c262620f53..997572ee6d 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -82,6 +82,7 @@ set(TEST_SRC hipMalloc3D.cc hipMalloc3DArray.cc hipArrayCreate.cc + hipArray3DCreate.cc hipDrvMemcpy3D.cc hipDrvMemcpy3DAsync.cc hipPointerGetAttribute.cc @@ -150,6 +151,7 @@ set(TEST_SRC hipMalloc3D.cc hipMalloc3DArray.cc hipArrayCreate.cc + hipArray3DCreate.cc hipDrvMemcpy3D.cc hipDrvMemcpy3DAsync.cc hipPointerGetAttribute.cc diff --git a/tests/catch/unit/memory/DriverContext.cc b/tests/catch/unit/memory/DriverContext.cc index 6791650567..dc27a62000 100644 --- a/tests/catch/unit/memory/DriverContext.cc +++ b/tests/catch/unit/memory/DriverContext.cc @@ -24,17 +24,13 @@ THE SOFTWARE. #include DriverContext::DriverContext() { -#if HT_NVIDIA HIP_CHECK(hipInit(0)); HIP_CHECK(hipDeviceGet(&device, 0)); HIP_CHECK(hipDevicePrimaryCtxRetain(&ctx, device)); HIP_CHECK(hipCtxPushCurrent(ctx)); -#endif } DriverContext::~DriverContext() { -#if HT_NVIDIA HIP_CHECK(hipCtxPopCurrent(&ctx)); HIP_CHECK(hipDevicePrimaryCtxRelease(device)); -#endif } diff --git a/tests/catch/unit/memory/DriverContext.hh b/tests/catch/unit/memory/DriverContext.hh index 76afe28f44..5593c512d4 100644 --- a/tests/catch/unit/memory/DriverContext.hh +++ b/tests/catch/unit/memory/DriverContext.hh @@ -26,10 +26,8 @@ THE SOFTWARE. class DriverContext { private: -#if HT_NVIDIA hipCtx_t ctx; hipDevice_t device; -#endif public: DriverContext(); diff --git a/tests/catch/unit/memory/hipArray3DCreate.cc b/tests/catch/unit/memory/hipArray3DCreate.cc new file mode 100644 index 0000000000..47e70e66c5 --- /dev/null +++ b/tests/catch/unit/memory/hipArray3DCreate.cc @@ -0,0 +1,166 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "DriverContext.hh" +#include "hipArrayCommon.hh" + +namespace { +void checkArrayIsExpected(const hiparray array, const HIP_ARRAY3D_DESCRIPTOR& expected_desc) { +// hipArray3DGetDescriptor doesn't currently exist (EXSWCPHIPT-87) +#if HT_AMD + std::ignore = array; + std::ignore = expected_desc; +#else + CUDA_ARRAY3D_DESCRIPTOR queried_desc; + cuArray3DGetDescriptor(&queried_desc, array); + + REQUIRE(queried_desc.Width == expected_desc.Width); + REQUIRE(queried_desc.Height == expected_desc.Height); + REQUIRE(queried_desc.Depth == expected_desc.Depth); + REQUIRE(queried_desc.Format == expected_desc.Format); + REQUIRE(queried_desc.NumChannels == expected_desc.NumChannels); + REQUIRE(queried_desc.Flags == expected_desc.Flags); +#endif +} +} // namespace + +TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_happy", "", char, uchar2, uint2, int4, short4, float, + float2, float4) { + using vec_info = vector_info; + DriverContext ctx; + + hiparray array; + HIP_ARRAY3D_DESCRIPTOR desc{}; + desc.Format = vec_info::format; + desc.NumChannels = vec_info::size; +#if HT_AMD + desc.Flags = 0; +#else + desc.Flags = GENERATE(0, CUDA_ARRAY3D_SURFACE_LDST); +#endif + + constexpr size_t size = 64; + + std::vector extents{ + {size, 0, 0}, // 1D array + {size, size, 0}, // 2D array + {size, size, size} // 3D array + }; + + for (auto& extent : extents) { + desc.Width = extent.width; + desc.Height = extent.height; + desc.Depth = extent.depth; + + CAPTURE(desc.Width, desc.Height, desc.Depth); + + HIP_CHECK(hipArray3DCreate(&array, &desc)); + checkArrayIsExpected(array, desc); + HIP_CHECK(hipArrayDestroy(array)); + } +} + +TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_MaxTexture", "", int, uint4, short, ushort2, + unsigned char, float, float4) { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-97"); + return; +#endif + + using vec_info = vector_info; + DriverContext ctx; + + hiparray array; + HIP_ARRAY3D_DESCRIPTOR desc{}; + desc.Format = vec_info::format; + desc.NumChannels = vec_info::size; +#if HT_AMD + desc.Flags = 0; +#else + desc.Flags = GENERATE(0, CUDA_ARRAY3D_SURFACE_LDST); + if (desc.Flags == CUDA_ARRAY3D_SURFACE_LDST) { + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-58"); + return; + } +#endif + CAPTURE(desc.Flags); + + const Sizes sizes(desc.Flags); + CAPTURE(sizes.max1D, sizes.max2D, sizes.max3D); + + const size_t s = 64; + SECTION("Happy") { + // stored in a vector so some values can be ifdef'd out + std::vector extentsToTest{ + make_hipExtent(sizes.max1D, 0, 0), // 1D max + make_hipExtent(sizes.max2D[0], s, 0), // 2D max width + make_hipExtent(s, sizes.max2D[1], 0), // 2D max height + make_hipExtent(sizes.max2D[0], sizes.max2D[1], 0), // 2D max + make_hipExtent(sizes.max3D[0], s, s), // 3D max width + make_hipExtent(s, sizes.max3D[1], s), // 3D max height + make_hipExtent(s, s, sizes.max3D[2]), // 3D max depth + make_hipExtent(s, sizes.max3D[1], sizes.max3D[2]), // 3D max height and depth + make_hipExtent(sizes.max3D[0], s, sizes.max3D[2]), // 3D max width and depth + make_hipExtent(sizes.max3D[0], sizes.max3D[1], s), // 3D max width and height + make_hipExtent(sizes.max3D[0], sizes.max3D[1], sizes.max3D[2]) // 3D max + }; + const auto extent = + GENERATE_COPY(from_range(std::begin(extentsToTest), std::end(extentsToTest))); + + desc.Width = extent.width; + desc.Height = extent.height; + desc.Depth = extent.depth; + + CAPTURE(desc.Width, desc.Height, desc.Depth); + + auto maxArrayCreateError = hipArray3DCreate(&array, &desc); + // this can try to alloc many GB of memory, so out of memory is acceptable + if (maxArrayCreateError == hipErrorOutOfMemory) return; + HIP_CHECK(maxArrayCreateError); + checkArrayIsExpected(array, desc); + HIP_CHECK(hipArrayDestroy(array)); + } + SECTION("Negative") { + std::vector extentsToTest { + make_hipExtent(sizes.max1D + 1, 0, 0), // 1D max + make_hipExtent(sizes.max2D[0] + 1, s, 0), // 2D max width + make_hipExtent(s, sizes.max2D[1] + 1, 0), // 2D max height + make_hipExtent(sizes.max2D[0] + 1, sizes.max2D[1] + 1, 0), // 2D max + make_hipExtent(sizes.max3D[0] + 1, s, s), // 3D max width + make_hipExtent(s, sizes.max3D[1] + 1, s), // 3D max height +#if !HT_NVIDIA // leads to hipSuccess on NVIDIA + make_hipExtent(s, s, sizes.max3D[2] + 1), // 3D max depth +#endif + make_hipExtent(s, sizes.max3D[1] + 1, sizes.max3D[2] + 1), // 3D max height and depth + make_hipExtent(sizes.max3D[0] + 1, s, sizes.max3D[2] + 1), // 3D max width and depth + make_hipExtent(sizes.max3D[0] + 1, sizes.max3D[1] + 1, s), // 3D max width and height + make_hipExtent(sizes.max3D[0] + 1, sizes.max3D[1] + 1, sizes.max3D[2] + 1) // 3D max + }; + const auto extent = + GENERATE_COPY(from_range(std::begin(extentsToTest), std::end(extentsToTest))); + + desc.Width = extent.width; + desc.Height = extent.height; + desc.Depth = extent.depth; + + CAPTURE(desc.Width, desc.Height, desc.Depth); + + HIP_CHECK_ERROR(hipArray3DCreate(&array, &desc), hipErrorInvalidValue); + } +} diff --git a/tests/catch/unit/memory/hipArrayCommon.hh b/tests/catch/unit/memory/hipArrayCommon.hh index c9823806c5..b51022eb56 100644 --- a/tests/catch/unit/memory/hipArrayCommon.hh +++ b/tests/catch/unit/memory/hipArrayCommon.hh @@ -131,8 +131,16 @@ struct Sizes { Sizes(unsigned int flag) { int device; HIP_CHECK(hipGetDevice(&device)); + static_assert( + hipArrayDefault == 0, + "hipArrayDefault is assumed to be equivalent to 0 for the following switch statment"); +#if HT_NVIDIA + static_assert(hipArraySurfaceLoadStore == CUDA_ARRAY3D_SURFACE_LDST, + "hipArraySurface is assumed to be equivalent to CUDA_ARRAY3D_SURFACE_LDST for " + "the following switch statment"); +#endif switch (flag) { - case hipArrayDefault: { + case hipArrayDefault: { // 0 hipDeviceProp_t prop; HIP_CHECK(hipGetDeviceProperties(&prop, device)); max1D = prop.maxTexture1D; @@ -140,7 +148,7 @@ struct Sizes { max3D = {prop.maxTexture3D[0], prop.maxTexture3D[1], prop.maxTexture3D[2]}; return; } - case hipArraySurfaceLoadStore: { + case hipArraySurfaceLoadStore: { // CUDA_ARRAY3D_SURFACE_LDST int value; HIP_CHECK(hipDeviceGetAttribute(&value, hipDeviceAttributeMaxSurface1D, device)); max1D = value; diff --git a/tests/catch/unit/memory/hipMalloc3DArray.cc b/tests/catch/unit/memory/hipMalloc3DArray.cc index 253913345e..be41740242 100644 --- a/tests/catch/unit/memory/hipMalloc3DArray.cc +++ b/tests/catch/unit/memory/hipMalloc3DArray.cc @@ -145,6 +145,7 @@ TEST_CASE("Unit_hipMalloc3DArray_MultiThread") { } } +namespace { void checkArrayIsExpected(hipArray_t array, const hipChannelFormatDesc& expected_desc, const hipExtent& expected_extent, const unsigned int expected_flags) { // hipArrayGetInfo doesn't currently exist (EXSWCPHIPT-87) @@ -172,6 +173,7 @@ void checkArrayIsExpected(hipArray_t array, const hipChannelFormatDesc& expected REQUIRE(expected_flags == queried_flags); #endif } +} TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_happy", "", char, uchar2, uint2, int4, short4, float, float2, float4) { @@ -187,19 +189,16 @@ TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_happy", "", char, uchar2, uint2, int4, SECTION("1D Array") { extent = make_hipExtent(size, 0, 0); - HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); } SECTION("2D Array") { extent = make_hipExtent(size, size, 0); - HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); } SECTION("3D Array") { extent = make_hipExtent(size, size, size); - HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); } + HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); checkArrayIsExpected(array, desc, extent, flags); - HIP_CHECK(hipFreeArray(array)); } From 87f1090e9539ac4fb7d1e3975002d89e0095ae48 Mon Sep 17 00:00:00 2001 From: agunashe <86270081+agunashe@users.noreply.github.com> Date: Tue, 12 Jul 2022 09:24:32 -0700 Subject: [PATCH 26/30] SWDEV-327563 - Win:setenv/unsetenv build failure. develop to staging (#2795) * SWDEV-327563 - Win:setenv/unsetenv build failure. develop to staging * SWDEV-327563 - Win:setenv/unsetenv build failure2. develop to staging * SWDEV-327563 - Win:setenv/unsetenv build failure. develop to staging Linux build failure --- tests/catch/unit/device/getDeviceCount_exe.cc | 39 ++++++++++++++----- 1 file changed, 30 insertions(+), 9 deletions(-) diff --git a/tests/catch/unit/device/getDeviceCount_exe.cc b/tests/catch/unit/device/getDeviceCount_exe.cc index 1da5f7c3a9..bf9951b6d4 100644 --- a/tests/catch/unit/device/getDeviceCount_exe.cc +++ b/tests/catch/unit/device/getDeviceCount_exe.cc @@ -22,6 +22,27 @@ THE SOFTWARE. #include #include +#include + +bool UNSETENV(std::string var) { + int result = -1; + #ifdef __unix__ + result = unsetenv(var.c_str()); + #else + result = _putenv((var + '=').c_str()); + #endif + return (result == 0) ? true: false; +} + +bool SETENV(std::string var, std::string value, int overwrite) { + int result = -1; + #ifdef __unix__ + result = setenv(var.c_str(), value.c_str(), overwrite); + #else + result = _putenv((var + '=' + value).c_str()); + #endif + return (result == 0) ? true: false; +} // Expects 1 command line arg, which is the Device Visible String int main(int argc, char** argv) { @@ -37,18 +58,18 @@ int main(int argc, char** argv) { // disable visible_devices env from shell #ifdef __HIP_PLATFORM_NVCC__ - unsetenv("CUDA_VISIBLE_DEVICES"); - setenv("CUDA_VISIBLE_DEVICES", argv[1], 1); + UNSETENV("CUDA_VISIBLE_DEVICES"); + SETENV("CUDA_VISIBLE_DEVICES", argv[1], 1); auto init_res = hipInit(0); if (hipSuccess != init_res) { std::cerr << "CUDA INIT API returned : " << hipGetErrorString(init_res) << std::endl; return -1; } #else - unsetenv("ROCR_VISIBLE_DEVICES"); - unsetenv("HIP_VISIBLE_DEVICES"); - setenv("ROCR_VISIBLE_DEVICES", argv[1], 1); - setenv("HIP_VISIBLE_DEVICES", argv[1], 1); + UNSETENV("ROCR_VISIBLE_DEVICES"); + UNSETENV("HIP_VISIBLE_DEVICES"); + SETENV("ROCR_VISIBLE_DEVICES", argv[1], 1); + SETENV("HIP_VISIBLE_DEVICES", argv[1], 1); #endif int count = 0; @@ -59,10 +80,10 @@ int main(int argc, char** argv) { } #ifdef __HIP_PLATFORM_NVCC__ - unsetenv("CUDA_VISIBLE_DEVICES"); + UNSETENV("CUDA_VISIBLE_DEVICES"); #else - unsetenv("ROCR_VISIBLE_DEVICES"); - unsetenv("HIP_VISIBLE_DEVICES"); + UNSETENV("ROCR_VISIBLE_DEVICES"); + UNSETENV("HIP_VISIBLE_DEVICES"); #endif return count; } \ No newline at end of file From 83dda013c23aff9975e481eb414d68996d15d3c4 Mon Sep 17 00:00:00 2001 From: agunashe <86270081+agunashe@users.noreply.github.com> Date: Tue, 12 Jul 2022 10:08:45 -0700 Subject: [PATCH 27/30] SWDEV-327563 - Win,Linx: Disable failing tests (#2796) Develop to staging --- .../hipTestMain/config/config_amd_linux.json | 3 ++- .../hipTestMain/config/config_amd_windows.json | 15 ++++++++++++++- 2 files changed, 16 insertions(+), 2 deletions(-) diff --git a/tests/catch/hipTestMain/config/config_amd_linux.json b/tests/catch/hipTestMain/config/config_amd_linux.json index fc2c1331ed..91c89ae7dc 100644 --- a/tests/catch/hipTestMain/config/config_amd_linux.json +++ b/tests/catch/hipTestMain/config/config_amd_linux.json @@ -3,7 +3,8 @@ [ "# Following test is related to ticket EXSWCPHIPT-41", "Unit_hipStreamGetPriority_happy", - "Unit_hipStreamPerThread_DeviceReset_1" + "Unit_hipStreamPerThread_DeviceReset_1", + "Unit_hipGraphAddKernelNode_Negative" ] } diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 3c12aee9f9..75e764ac2f 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -32,6 +32,19 @@ "Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory", "Unit_hipGraphMemcpyNodeSetParams_Functional", "Unit_hipMalloc3D_ValidatePitch", - "Unit_hipArrayCreate_happy" + "Unit_hipArrayCreate_happy", + "Unit_hipGraphAddKernelNode_Negative", + "Unit_hipHostRegister_Negative - int", + "Unit_hipHostRegister_Negative - float", + "Unit_hipHostRegister_Negative - double", + "Unit_hipMemAllocPitch_ValidatePitch", + "Unit_hipArrayCreate_happy - int", + "Unit_hipArrayCreate_happy - int4", + "Unit_hipArrayCreate_happy - short2", + "Unit_hipArrayCreate_happy - char", + "Unit_hipArrayCreate_happy - char4", + "Unit_hipArrayCreate_happy - float", + "Unit_hipArrayCreate_happy - float2", + "Unit_hipArrayCreate_happy - float4" ] } From 444c0cd157638350537dcb5e9841cf2a199d6d47 Mon Sep 17 00:00:00 2001 From: agunashe <86270081+agunashe@users.noreply.github.com> Date: Tue, 12 Jul 2022 14:28:50 -0700 Subject: [PATCH 28/30] SWDEV-327563 - hipHostRegister fix - develop to staging (#2794) --- tests/catch/unit/memory/hipHostRegister.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/catch/unit/memory/hipHostRegister.cc b/tests/catch/unit/memory/hipHostRegister.cc index c2fa6ed009..417ab0857a 100644 --- a/tests/catch/unit/memory/hipHostRegister.cc +++ b/tests/catch/unit/memory/hipHostRegister.cc @@ -187,7 +187,7 @@ TEMPLATE_TEST_CASE("Unit_hipHostRegister_Negative", "", int, float, double) { REQUIRE(devMemAvail > 0); REQUIRE(hostMemFree > 0); - size_t memFree = std::min(devMemFree, hostMemFree); // which is the limiter cpu or gpu + size_t memFree = (std::min)(devMemFree, hostMemFree); // which is the limiter cpu or gpu SECTION("hipHostRegister Negative Test - invalid memory size") { HIP_CHECK_ERROR(hipHostRegister(hostPtr, memFree, 0), hipErrorInvalidValue); From daceb2e1a1608e25f8bec631cbf88c40fbded27f Mon Sep 17 00:00:00 2001 From: Finlay Date: Wed, 13 Jul 2022 06:16:26 +0100 Subject: [PATCH 29/30] EXSWCPHIPT-134 - Negative tests for hipMalloc3DArray and hipArray3DCreate (#2784) --- tests/catch/unit/memory/hipArray3DCreate.cc | 157 +++++++++++- tests/catch/unit/memory/hipArrayCommon.hh | 43 ++++ tests/catch/unit/memory/hipArrayCreate.cc | 42 +--- tests/catch/unit/memory/hipMalloc3DArray.cc | 262 +++++++++++++++----- tests/catch/unit/memory/hipMallocArray.cc | 32 +-- 5 files changed, 413 insertions(+), 123 deletions(-) diff --git a/tests/catch/unit/memory/hipArray3DCreate.cc b/tests/catch/unit/memory/hipArray3DCreate.cc index 47e70e66c5..faa3822b9b 100644 --- a/tests/catch/unit/memory/hipArray3DCreate.cc +++ b/tests/catch/unit/memory/hipArray3DCreate.cc @@ -17,8 +17,10 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +#include #include "DriverContext.hh" #include "hipArrayCommon.hh" +#include "hip_test_common.hh" namespace { void checkArrayIsExpected(const hiparray array, const HIP_ARRAY3D_DESCRIPTOR& expected_desc) { @@ -38,6 +40,11 @@ void checkArrayIsExpected(const hiparray array, const HIP_ARRAY3D_DESCRIPTOR& ex REQUIRE(queried_desc.Flags == expected_desc.Flags); #endif } + +void testInvalidDescription(HIP_ARRAY3D_DESCRIPTOR desc) { + hiparray array; + HIP_CHECK_ERROR(hipArray3DCreate(&array, &desc), hipErrorInvalidValue); +} } // namespace TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_happy", "", char, uchar2, uint2, int4, short4, float, @@ -45,7 +52,6 @@ TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_happy", "", char, uchar2, uint2, int4, using vec_info = vector_info; DriverContext ctx; - hiparray array; HIP_ARRAY3D_DESCRIPTOR desc{}; desc.Format = vec_info::format; desc.NumChannels = vec_info::size; @@ -70,6 +76,7 @@ TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_happy", "", char, uchar2, uint2, int4, CAPTURE(desc.Width, desc.Height, desc.Depth); + hiparray array; HIP_CHECK(hipArray3DCreate(&array, &desc)); checkArrayIsExpected(array, desc); HIP_CHECK(hipArrayDestroy(array)); @@ -164,3 +171,151 @@ TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_MaxTexture", "", int, uint4, short, us HIP_CHECK_ERROR(hipArray3DCreate(&array, &desc), hipErrorInvalidValue); } } + +#if HT_NVIDIA +constexpr std::array validFlags{ + 0, + hipArraySurfaceLoadStore, + hipArrayLayered, + hipArrayLayered | hipArraySurfaceLoadStore, + hipArrayCubemap, + hipArrayCubemap | hipArrayLayered, + hipArrayCubemap | hipArraySurfaceLoadStore, + hipArrayCubemap | hipArrayLayered | hipArraySurfaceLoadStore, + hipArrayTextureGather}; +#else +constexpr std::array validFlags{ + 0, hipArrayCubemap, hipArrayCubemap | hipArrayLayered, + hipArrayCubemap | hipArraySurfaceLoadStore, + hipArrayCubemap | hipArrayLayered | hipArraySurfaceLoadStore}; +#endif + +constexpr HIP_ARRAY3D_DESCRIPTOR defaultDescriptor(unsigned int flags, size_t size) { + HIP_ARRAY3D_DESCRIPTOR desc{}; + desc.Format = HIP_AD_FORMAT_FLOAT; + desc.NumChannels = 4; + desc.Flags = flags; + desc.Width = size; + desc.Height = size; + desc.Depth = size; + +#if HT_NVIDIA + if (flags == CUDA_ARRAY3D_TEXTURE_GATHER) { + desc.Depth = 0; + } +#endif + return desc; +} + +// Providing the array pointer as nullptr should return an error +TEST_CASE("Unit_hipArray3DCreate_Negative_NullArrayPtr") { + auto desc = defaultDescriptor(0, 64); + + DriverContext ctx; + HIP_CHECK_ERROR(hipArray3DCreate(nullptr, &desc), hipErrorInvalidValue); +} + +// Providing the description pointer as nullptr should return an error +TEST_CASE("Unit_hipArray3DCreate_Negative_NullDescPtr") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-130"); + return; +#endif + + DriverContext ctx; + hiparray array; + HIP_CHECK_ERROR(hipArray3DCreate(&array, nullptr), hipErrorInvalidValue); +} + + +// Zero width arrays are not allowed +TEST_CASE("Unit_hipArray3DCreate_Negative_ZeroWidth") { + DriverContext ctx; + + unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + auto desc = defaultDescriptor(flags, 6); + desc.Width = 0; + CAPTURE(desc.Flags); + + testInvalidDescription(desc); +} + +// Zero height arrays are only allowed for 1D arrays and layered arrays +TEST_CASE("Unit_hipArray3DCreate_Negative_ZeroHeight") { + DriverContext ctx; + + unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + auto desc = defaultDescriptor(flags, 6); +#if HT_NVIDIA + std::array exceptions{CUDA_ARRAY3D_LAYERED, + CUDA_ARRAY3D_LAYERED | CUDA_ARRAY3D_SURFACE_LDST}; +#else + std::array exceptions{}; +#endif + desc.Height = 0; + + if (std::find(std::begin(exceptions), std::end(exceptions), desc.Flags) == std::end(exceptions)) { + // flag is not in list of exceptions + testInvalidDescription(desc); + } +} + +// Arrays must be created with a valid data format +TEST_CASE("Unit_hipArray3DCreate_Negative_InvalidFormat") { + DriverContext ctx; + + unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + auto desc = defaultDescriptor(flags, 6); + + desc.Format = static_cast(0xDEADBEEF); + REQUIRE(std::find(std::begin(driverFormats), std::end(driverFormats), desc.Format) == + std::end(driverFormats)); + + testInvalidDescription(desc); +} + +// An array must have either 1,2, or 4 channels +TEST_CASE("Unit_hipArray3DCreate_Negative_NumChannels") { + DriverContext ctx; + unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + auto desc = defaultDescriptor(flags, 6); + desc.NumChannels = GENERATE(0, 3, 5); + + testInvalidDescription(desc); +} + +// Using invalid flags should result in an error +TEST_CASE("Unit_hipArray3DCreate_Negative_InvalidFlags") { + DriverContext ctx; + + // FIXME: use the same flags for both tests when the values exist for hip +#if HT_NVIDIA + unsigned int flags = + GENERATE(0xDEADBEEF, CUDA_ARRAY3D_TEXTURE_GATHER | CUDA_ARRAY3D_SURFACE_LDST, + CUDA_ARRAY3D_TEXTURE_GATHER | CUDA_ARRAY3D_CUBEMAP, + CUDA_ARRAY3D_TEXTURE_GATHER | CUDA_ARRAY3D_SURFACE_LDST | CUDA_ARRAY3D_CUBEMAP); +#else + unsigned int flags = 0xDEADBEEF; +#endif + + CAPTURE(flags); + + auto desc = defaultDescriptor(flags, 6); + + + REQUIRE(std::find(std::begin(validFlags), std::end(validFlags), desc.Flags) == + std::end(validFlags)); + + testInvalidDescription(desc); +} + + +// hipArray3DCreate should handle the max numeric value gracefully. +TEST_CASE("Unit_hipArray3DCreate_Negative_NumericLimit") { + DriverContext ctx; + + unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + auto desc = defaultDescriptor(flags, std::numeric_limits::max()); + + testInvalidDescription(desc); +} diff --git a/tests/catch/unit/memory/hipArrayCommon.hh b/tests/catch/unit/memory/hipArrayCommon.hh index b51022eb56..8502ac218d 100644 --- a/tests/catch/unit/memory/hipArrayCommon.hh +++ b/tests/catch/unit/memory/hipArrayCommon.hh @@ -166,3 +166,46 @@ struct Sizes { } } }; + +inline const char* channelFormatString(hipChannelFormatKind formatKind) noexcept { + switch (formatKind) { + case hipChannelFormatKindFloat: + return "float"; + case hipChannelFormatKindSigned: + return "signed"; + case hipChannelFormatKindUnsigned: + return "unsigned"; + default: + return "error"; + } +} + +// All the possible formats for channel data in an array. +static const std::vector driverFormats{ + HIP_AD_FORMAT_UNSIGNED_INT8, HIP_AD_FORMAT_UNSIGNED_INT16, HIP_AD_FORMAT_UNSIGNED_INT32, + HIP_AD_FORMAT_SIGNED_INT8, HIP_AD_FORMAT_SIGNED_INT16, HIP_AD_FORMAT_SIGNED_INT32, + HIP_AD_FORMAT_HALF, HIP_AD_FORMAT_FLOAT}; + +// Helpful for printing errors +inline const char* formatToString(hipArray_Format f) { + switch (f) { + case HIP_AD_FORMAT_UNSIGNED_INT8: + return "Unsigned Int 8"; + case HIP_AD_FORMAT_UNSIGNED_INT16: + return "Unsigned Int 16"; + case HIP_AD_FORMAT_UNSIGNED_INT32: + return "Unsigned Int 32"; + case HIP_AD_FORMAT_SIGNED_INT8: + return "Signed Int 8"; + case HIP_AD_FORMAT_SIGNED_INT16: + return "Signed Int 16"; + case HIP_AD_FORMAT_SIGNED_INT32: + return "Signed Int 32"; + case HIP_AD_FORMAT_HALF: + return "Float 16"; + case HIP_AD_FORMAT_FLOAT: + return "Float 32"; + default: + return "not found"; + } +} diff --git a/tests/catch/unit/memory/hipArrayCreate.cc b/tests/catch/unit/memory/hipArrayCreate.cc index d55bdae255..836c0d1c45 100644 --- a/tests/catch/unit/memory/hipArrayCreate.cc +++ b/tests/catch/unit/memory/hipArrayCreate.cc @@ -109,36 +109,6 @@ TEST_CASE("Unit_hipArrayCreate_MultiThread") { } -// All the possible formats for channel data in an array. -static const std::vector formats{ - HIP_AD_FORMAT_UNSIGNED_INT8, HIP_AD_FORMAT_UNSIGNED_INT16, HIP_AD_FORMAT_UNSIGNED_INT32, - HIP_AD_FORMAT_SIGNED_INT8, HIP_AD_FORMAT_SIGNED_INT16, HIP_AD_FORMAT_SIGNED_INT32, - HIP_AD_FORMAT_HALF, HIP_AD_FORMAT_FLOAT}; - -// Helpful for printing errors -const char* formatToString(hipArray_Format f) { - switch (f) { - case HIP_AD_FORMAT_UNSIGNED_INT8: - return "Unsigned Int 8"; - case HIP_AD_FORMAT_UNSIGNED_INT16: - return "Unsigned Int 16"; - case HIP_AD_FORMAT_UNSIGNED_INT32: - return "Unsigned Int 32"; - case HIP_AD_FORMAT_SIGNED_INT8: - return "Signed Int 8"; - case HIP_AD_FORMAT_SIGNED_INT16: - return "Signed Int 16"; - case HIP_AD_FORMAT_SIGNED_INT32: - return "Signed Int 32"; - case HIP_AD_FORMAT_HALF: - return "Float 16"; - case HIP_AD_FORMAT_FLOAT: - return "Float 32"; - default: - return "not found"; - } -} - // Tests ///////////////////////////////////////// #if HT_AMD @@ -338,7 +308,7 @@ TEMPLATE_TEST_CASE("Unit_hipArrayCreate_maxTexture", "", uint, int, int4, ushort TEST_CASE("Unit_hipArrayCreate_ZeroWidth") { DriverContext ctx; HIP_ARRAY_DESCRIPTOR desc; - desc.Format = formats[0]; + desc.Format = driverFormats[0]; desc.NumChannels = 4; desc.Width = 0; desc.Height = GENERATE(0, 1024); @@ -351,13 +321,13 @@ TEST_CASE("Unit_hipArrayCreate_ZeroWidth") { // HipArrayCreate will return an error when nullptr is used as the array argument TEST_CASE("Unit_hipArrayCreate_Nullptr") { #if HT_AMD - HipTest::HIP_SKIP_TEST("Probably EXSWCPHIPT-45"); + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-130"); return; #endif DriverContext ctx; SECTION("Null array") { HIP_ARRAY_DESCRIPTOR desc; - desc.Format = formats[0]; + desc.Format = driverFormats[0]; desc.NumChannels = 4; desc.Width = 1024; desc.Height = 1024; @@ -374,7 +344,7 @@ TEST_CASE("Unit_hipArrayCreate_Nullptr") { TEST_CASE("Unit_hipArrayCreate_BadNumberChannelElement") { DriverContext ctx; HIP_ARRAY_DESCRIPTOR desc; - desc.Format = GENERATE(from_range(std::begin(formats), std::end(formats))); + desc.Format = GENERATE(from_range(std::begin(driverFormats), std::end(driverFormats))); desc.NumChannels = GENERATE(-1, 0, 3, 5, 8); desc.Width = 1024; desc.Height = GENERATE(0, 1024); @@ -393,9 +363,9 @@ TEST_CASE("Unit_hipArrayCreate_BadChannelFormat") { // create a bad format desc.Format = - std::accumulate(std::begin(formats), std::end(formats), formats[0], + std::accumulate(std::begin(driverFormats), std::end(driverFormats), driverFormats[0], [](auto i, auto f) { return static_cast(i + f); }); - for (auto&& format : formats) { + for (auto&& format : driverFormats) { REQUIRE(desc.Format != format); } diff --git a/tests/catch/unit/memory/hipMalloc3DArray.cc b/tests/catch/unit/memory/hipMalloc3DArray.cc index be41740242..1a8ec2f0d0 100644 --- a/tests/catch/unit/memory/hipMalloc3DArray.cc +++ b/tests/catch/unit/memory/hipMalloc3DArray.cc @@ -33,6 +33,7 @@ static constexpr auto ARRAY_SIZE{4}; static constexpr auto BIG_ARRAY_SIZE{100}; static constexpr auto ARRAY_LOOP{100}; + /* * This API verifies memory allocations for small and * bigger chunks of data. @@ -67,53 +68,6 @@ static void Malloc3DArray_DiffSizes(int gpu) { REQUIRE_THREAD(pavail == avail); } -/* - * Verifies the negative scenarios of hipMalloc3DArray API - */ -TEST_CASE("Unit_hipMalloc3DArray_Negative") { - constexpr int width{ARRAY_SIZE}, height{ARRAY_SIZE}, depth{ARRAY_SIZE}; - hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); - hipArray* arr; -#if HT_NVIDIA - SECTION("NullPointer to Array") { - REQUIRE(hipMalloc3DArray(nullptr, &channelDesc, make_hipExtent(width, height, depth), - hipArrayDefault) != hipSuccess); - } - - SECTION("NullPointer to Channel Descriptor") { - REQUIRE(hipMalloc3DArray(&arr, nullptr, make_hipExtent(width, height, depth), - hipArrayDefault) != hipSuccess); - } -#endif - SECTION("Width 0 in hipExtent") { - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(0, height, width), - hipArrayDefault) != hipSuccess); - } - - SECTION("Height 0 in hipExtent") { - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, 0, width), - hipArrayDefault) != hipSuccess); - } - - SECTION("Invalid Flag") { - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), 100) != - hipSuccess); - } - - SECTION("Width,Height & Depth 0 in hipExtent") { - REQUIRE(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(0, 0, 0), hipArrayDefault) != - hipSuccess); - } - - SECTION("Max int values to extent") { - REQUIRE(hipMalloc3DArray( - &arr, &channelDesc, - make_hipExtent(std::numeric_limits::max(), std::numeric_limits::max(), - std::numeric_limits::max()), - hipArrayDefault) != hipSuccess); - } -} - TEST_CASE("Unit_hipMalloc3DArray_DiffSizes") { Malloc3DArray_DiffSizes(0); HIP_CHECK_THREAD_FINALIZE(); @@ -173,7 +127,7 @@ void checkArrayIsExpected(hipArray_t array, const hipChannelFormatDesc& expected REQUIRE(expected_flags == queried_flags); #endif } -} +} // namespace TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_happy", "", char, uchar2, uint2, int4, short4, float, float2, float4) { @@ -187,15 +141,9 @@ TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_happy", "", char, uchar2, uint2, int4, constexpr size_t size = 64; hipExtent extent; - SECTION("1D Array") { - extent = make_hipExtent(size, 0, 0); - } - SECTION("2D Array") { - extent = make_hipExtent(size, size, 0); - } - SECTION("3D Array") { - extent = make_hipExtent(size, size, size); - } + SECTION("1D Array") { extent = make_hipExtent(size, 0, 0); } + SECTION("2D Array") { extent = make_hipExtent(size, size, 0); } + SECTION("3D Array") { extent = make_hipExtent(size, size, size); } HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); checkArrayIsExpected(array, desc, extent, flags); @@ -272,3 +220,203 @@ TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_MaxTexture", "", int, uint4, short, us HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, extent, flag), hipErrorInvalidValue); } } + + +#if HT_AMD +constexpr std::array validFlags{hipArrayDefault}; +#else +constexpr std::array validFlags{ + hipArrayDefault, + hipArrayDefault | hipArraySurfaceLoadStore, + hipArrayLayered, + hipArrayLayered | hipArraySurfaceLoadStore, + hipArrayCubemap, + hipArrayCubemap | hipArrayLayered, + hipArrayCubemap | hipArraySurfaceLoadStore, + hipArrayCubemap | hipArrayLayered | hipArraySurfaceLoadStore, + hipArrayTextureGather}; +#endif + +hipExtent makeExtent(unsigned int flag, size_t s) { + if (flag == hipArrayTextureGather) { + return make_hipExtent(s, s, 0); + } + return make_hipExtent(s, s, s); +} + + +// Providing the array pointer as nullptr should return an error +TEST_CASE("Unit_hipMalloc3DArray_Negative_NullArrayPtr") { + hipChannelFormatDesc desc = hipCreateChannelDesc(); + constexpr size_t s = 6; + + const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + HIP_CHECK_ERROR(hipMalloc3DArray(nullptr, &desc, makeExtent(flag, s), flag), + hipErrorInvalidValue); +} + +// Providing the description pointer as nullptr should return an error +TEST_CASE("Unit_hipMalloc3DArray_Negative_NullDescPtr") { + constexpr size_t s = 6; // 6 to keep cubemap happy + hipArray_t array; + + const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + + HIP_CHECK_ERROR(hipMalloc3DArray(&array, nullptr, makeExtent(flag, s), flag), + hipErrorInvalidValue); +} + +// Zero width arrays are not allowed +TEST_CASE("Unit_hipMalloc3DArray_Negative_ZeroWidth") { + constexpr size_t s = 6; // 6 to keep cubemap happy + hipArray_t array; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + + const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + + HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, make_hipExtent(0, s, s), flag), + hipErrorInvalidValue); +} + +// Zero height arrays are only allowed for 1D arrays and layered arrays +TEST_CASE("Unit_hipMalloc3DArray_Negative_ZeroHeight") { + constexpr size_t s = 6; // 6 to keep cubemap happy + hipArray_t array; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + std::array exceptions{hipArrayLayered, + hipArrayLayered | hipArraySurfaceLoadStore}; + + const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + + if (std::find(std::begin(exceptions), std::end(exceptions), flag) == std::end(exceptions)) { + // flag is not in list of exceptions + HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, make_hipExtent(s, 0, s), flag), + hipErrorInvalidValue); + } +} + +TEST_CASE("Unit_hipMalloc3DArray_Negative_InvalidFlags") { + constexpr size_t s = 6; // 6 to keep cubemap happy + hipArray_t array; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + +#if HT_AMD + const unsigned int flag = 0xDEADBEEF; +#else + const unsigned int flag = + GENERATE(0xDEADBEEF, hipArrayTextureGather | hipArraySurfaceLoadStore, + hipArrayTextureGather | hipArrayCubemap, + hipArrayTextureGather | hipArraySurfaceLoadStore | hipArrayCubemap); +#endif + + CAPTURE(flag); + + REQUIRE(std::find(std::begin(validFlags), std::end(validFlags), flag) == std::end(validFlags)); + + HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, makeExtent(flag, s), flag), hipErrorInvalidValue); +} + +void testInvalidDescription(hipChannelFormatDesc desc){ + constexpr size_t s = 6; // 6 to keep cubemap happy + hipArray_t array; + +#if HT_NVIDIA + hipError_t expectedError = hipErrorUnknown; +#else + hipError_t expectedError = hipErrorInvalidValue; +#endif + + const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, makeExtent(flag, s), flag), expectedError); +} + +TEST_CASE("Unit_hipMalloc3DArray_Negative_InvalidFormat") { + hipChannelFormatDesc desc = hipCreateChannelDesc(); + desc.f = GENERATE(hipChannelFormatKindNone, 0xBEEF); + testInvalidDescription(desc); +} + +TEST_CASE("Unit_hipMalloc3DArray_Negative_BadChannelLayout") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-129"); + return; +#endif + + const int bits = GENERATE(8, 16, 32); + const hipChannelFormatKind formatKind = + GENERATE(hipChannelFormatKindSigned, hipChannelFormatKindUnsigned, hipChannelFormatKindFloat); + if (bits == 8 && formatKind == hipChannelFormatKindFloat) return; + + + hipChannelFormatDesc desc = GENERATE_COPY(hipCreateChannelDesc(bits, bits, bits, 0, formatKind), + hipCreateChannelDesc(0, bits, bits, 0, formatKind), + hipCreateChannelDesc(0, bits, bits, bits, formatKind), + hipCreateChannelDesc(bits, 0, bits, 0, formatKind), + hipCreateChannelDesc(bits, bits, 0, bits, formatKind), + hipCreateChannelDesc(0, 0, bits, 0, formatKind), + hipCreateChannelDesc(0, 0, bits, bits, formatKind)); + + INFO("kind: " << channelFormatString(formatKind)); + INFO("x: " << desc.x << ", y: " << desc.y << ", z: " << desc.z << ", w: " << desc.w); + + testInvalidDescription(desc); +} + +TEST_CASE("Unit_hipMalloc3DArray_Negative_8BitFloat") { + hipChannelFormatDesc desc = GENERATE(hipCreateChannelDesc(8, 0, 0, 0, hipChannelFormatKindFloat), + hipCreateChannelDesc(8, 8, 0, 0, hipChannelFormatKindFloat), + hipCreateChannelDesc(8, 8, 8, 8, hipChannelFormatKindFloat)); + + testInvalidDescription(desc); +} + +TEST_CASE("Unit_hipMalloc3DArray_Negative_DifferentChannelSizes") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-129"); + return; +#endif + + const int bitsX = GENERATE(8, 16, 32); + const int bitsY = GENERATE(8, 16, 32); + const int bitsZ = GENERATE(8, 16, 32); + const int bitsW = GENERATE(8, 16, 32); + if (bitsX == bitsY && bitsY == bitsZ && bitsZ == bitsW) return; // skip when they are equal + + const hipChannelFormatKind channelFormat = + GENERATE(hipChannelFormatKindSigned, hipChannelFormatKindUnsigned, hipChannelFormatKindFloat); + + if (channelFormat == hipChannelFormatKindFloat && + (bitsX == 8 || bitsY == 8 || bitsZ == 8 || bitsW == 8)) + return; // 8 bit floats aren't allowed + + hipChannelFormatDesc desc = hipCreateChannelDesc(bitsX, bitsY, bitsZ, bitsW, channelFormat); + + INFO("format: " << channelFormatString(channelFormat) << ", x bits: " << bitsX + << ", y bits: " << bitsY << ", z bits: " << bitsZ << ", w bits: " << bitsW); + + + testInvalidDescription(desc); +} + +TEST_CASE("Unit_hipMalloc3DArray_Negative_BadChannelSize") { + const int badBits = GENERATE(-1, 0, 10, 100); + const hipChannelFormatKind formatKind = + GENERATE(hipChannelFormatKindSigned, hipChannelFormatKindUnsigned, hipChannelFormatKindFloat); + hipChannelFormatDesc desc = hipCreateChannelDesc(badBits, badBits, badBits, badBits, formatKind); + + INFO("Number of bits: " << badBits); + + testInvalidDescription(desc); +} + + +// hipMalloc3DArray should handle the max numeric value gracefully. +TEST_CASE("Unit_hipMalloc3DArray_Negative_NumericLimit") { + hipArray_t arrayPtr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + + size_t size = std::numeric_limits::max(); + const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); + HIP_CHECK_ERROR(hipMalloc3DArray(&arrayPtr, &desc, makeExtent(flag, size), flag), + hipErrorInvalidValue); +} diff --git a/tests/catch/unit/memory/hipMallocArray.cc b/tests/catch/unit/memory/hipMallocArray.cc index 820e917313..f1f6ebc9d2 100644 --- a/tests/catch/unit/memory/hipMallocArray.cc +++ b/tests/catch/unit/memory/hipMallocArray.cc @@ -135,20 +135,6 @@ template size_t getAllocSize(const size_t width, const size_t heigh return sizeof(T) * width * (height ? height : 1); } - -const char* channelFormatString(hipChannelFormatKind formatKind) noexcept { - switch (formatKind) { - case hipChannelFormatKindFloat: - return "float"; - case hipChannelFormatKindSigned: - return "signed"; - case hipChannelFormatKindUnsigned: - return "unsigned"; - default: - return "error"; - } -} - // Tests ///////////////////////////////////////// // Test the default array by generating a texture from it then reading from that texture. @@ -409,7 +395,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_happy", "", uint, int, int4, ushort, sho // pointer to the array in device memory hipArray_t arrayPtr{}; size_t width = 1024; - size_t height; + size_t height{}; SECTION("hipArrayDefault") { height = GENERATE(0, 1024); @@ -509,7 +495,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_MaxTexture_Default", "", uint, int4, ush // Arrays with channels of different size are not allowed. TEST_CASE("Unit_hipMallocArray_Negative_DifferentChannelSizes") { #if HT_AMD - HipTest::HIP_SKIP_TEST("EXSWCPHIPT-59"); + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-129"); return; #endif const int bitsX = GENERATE(8, 16, 32); @@ -564,10 +550,6 @@ TEST_CASE("Unit_hipMallocArray_Negative_ZeroWidth") { // Providing the array pointer as nullptr should return an error TEST_CASE("Unit_hipMallocArray_Negative_NullArrayPtr") { -#if HT_AMD - HipTest::HIP_SKIP_TEST("EXSWCPHIPT-45"); - return; -#endif hipChannelFormatDesc desc = hipCreateChannelDesc(); HIP_CHECK_ERROR(hipMallocArray(nullptr, &desc, 1024, 0, hipArrayDefault), hipErrorInvalidValue); @@ -575,10 +557,6 @@ TEST_CASE("Unit_hipMallocArray_Negative_NullArrayPtr") { // Providing the desc pointer as nullptr should return an error TEST_CASE("Unit_hipMallocArray_Negative_NullDescPtr") { -#if HT_AMD - HipTest::HIP_SKIP_TEST("EXSWCPHIPT-83"); - return; -#endif hipArray_t arrayPtr; HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, nullptr, 1024, 0, hipArrayDefault), hipErrorInvalidValue); @@ -586,10 +564,6 @@ TEST_CASE("Unit_hipMallocArray_Negative_NullDescPtr") { // Inappropriate but related flags should still return an error TEST_CASE("Unit_hipMallocArray_Negative_BadFlags") { -#if HT_AMD - HipTest::HIP_SKIP_TEST("EXSWCPHIPT-72"); - return; -#endif hipChannelFormatDesc desc = hipCreateChannelDesc(); hipArray_t arrayPtr; @@ -687,7 +661,7 @@ TEST_CASE("Unit_hipMallocArray_Negative_3ChannelElement") { // The bit channel description should not allow any channels after a zero channel TEST_CASE("Unit_hipMallocArray_Negative_ChannelAfterZeroChannel") { #if HT_AMD - HipTest::HIP_SKIP_TEST("EXSWCPHIPT-59"); + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-129"); return; #endif const int bits = GENERATE(8, 16, 32); From 9e84a11f23fc7cc477d6c3b16197a809ab6a7384 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 14 Jul 2022 11:25:25 +0530 Subject: [PATCH 30/30] Disable Unit_hipMemVmm_Basic on windows for now --- tests/catch/hipTestMain/config/config_amd_windows.json | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 75e764ac2f..547d92a4bf 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -45,6 +45,7 @@ "Unit_hipArrayCreate_happy - char4", "Unit_hipArrayCreate_happy - float", "Unit_hipArrayCreate_happy - float2", - "Unit_hipArrayCreate_happy - float4" + "Unit_hipArrayCreate_happy - float4", + "Unit_hipMemVmm_Basic" ] }