diff --git a/projects/hip-tests/catch/unit/memory/CMakeLists.txt b/projects/hip-tests/catch/unit/memory/CMakeLists.txt index af72155549..1028010aa5 100644 --- a/projects/hip-tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip-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/projects/hip-tests/catch/unit/memory/DriverContext.cc b/projects/hip-tests/catch/unit/memory/DriverContext.cc new file mode 100644 index 0000000000..6791650567 --- /dev/null +++ b/projects/hip-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/projects/hip-tests/catch/unit/memory/DriverContext.hh b/projects/hip-tests/catch/unit/memory/DriverContext.hh new file mode 100644 index 0000000000..76afe28f44 --- /dev/null +++ b/projects/hip-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/projects/hip-tests/catch/unit/memory/hipArrayCommon.hh b/projects/hip-tests/catch/unit/memory/hipArrayCommon.hh new file mode 100644 index 0000000000..1c6100a6e9 --- /dev/null +++ b/projects/hip-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/projects/hip-tests/catch/unit/memory/hipArrayCreate.cc b/projects/hip-tests/catch/unit/memory/hipArrayCreate.cc index 3bbf6cbd54..9ef637821e 100644 --- a/projects/hip-tests/catch/unit/memory/hipArrayCreate.cc +++ b/projects/hip-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/projects/hip-tests/catch/unit/memory/hipMallocArray.cc b/projects/hip-tests/catch/unit/memory/hipMallocArray.cc index 22e1ce9864..8859aadb45 100644 --- a/projects/hip-tests/catch/unit/memory/hipMallocArray.cc +++ b/projects/hip-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));