From eebf627237e6439b1de3a9534afb83a19e4617b8 Mon Sep 17 00:00:00 2001 From: Finlay Date: Wed, 25 May 2022 11:43:18 +0100 Subject: [PATCH] hipMallocArray tests for default array and surface load store (#2675) [ROCm/hip-tests commit: 205107aa3cfdfc8d619e7e785b289c3fb6023d93] --- .../catch/include/hip_test_common.hh | 9 + .../catch/unit/memory/hipMallocArray.cc | 323 +++++++++++++++--- 2 files changed, 285 insertions(+), 47 deletions(-) diff --git a/projects/hip-tests/catch/include/hip_test_common.hh b/projects/hip-tests/catch/include/hip_test_common.hh index 4566600909..496f202b8c 100644 --- a/projects/hip-tests/catch/include/hip_test_common.hh +++ b/projects/hip-tests/catch/include/hip_test_common.hh @@ -163,7 +163,16 @@ inline bool isImageSupported() { return imageSupport != 0; } +/** + * Causes the test to stop and be skipped at runtime. + * reason: Message describing the reason the test has been skipped. + */ +static inline void HIP_SKIP_TEST(char const* const reason) noexcept { + // ctest is setup to parse for "HIP_SKIP_THIS_TEST", at which point it will skip the test. + std::cout << "Skipping test. Reason: " << reason << '\n' << "HIP_SKIP_THIS_TEST" << std::endl; } +} + // This must be called in the beginning of image test app's main() to indicate whether image // is supported. diff --git a/projects/hip-tests/catch/unit/memory/hipMallocArray.cc b/projects/hip-tests/catch/unit/memory/hipMallocArray.cc index 3e521bfcb5..a73566b40d 100644 --- a/projects/hip-tests/catch/unit/memory/hipMallocArray.cc +++ b/projects/hip-tests/catch/unit/memory/hipMallocArray.cc @@ -25,7 +25,6 @@ hipMallocArray API test scenarios 4. Multithreaded scenario */ - #include static constexpr auto NUM_W{4}; @@ -53,20 +52,16 @@ static void MallocArray_DiffSizes(int gpu) { std::vector array_size; array_size.push_back(NUM_W); array_size.push_back(BIGNUM_W); - for (auto &size : array_size) { + for (auto& size : array_size) { hipArray* A_d[ARRAY_LOOP]; size_t tot, avail, ptot, pavail; hipChannelFormatDesc desc = hipCreateChannelDesc(); HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); for (int i = 0; i < ARRAY_LOOP; i++) { if (size == NUM_W) { - HIP_CHECK(hipMallocArray(&A_d[i], &desc, - NUM_W, NUM_H, - hipArrayDefault)); + HIP_CHECK(hipMallocArray(&A_d[i], &desc, NUM_W, NUM_H, hipArrayDefault)); } else { - HIP_CHECK(hipMallocArray(&A_d[i], &desc, - BIGNUM_W, BIGNUM_H, - hipArrayDefault)); + HIP_CHECK(hipMallocArray(&A_d[i], &desc, BIGNUM_W, BIGNUM_H, hipArrayDefault)); } } for (int i = 0; i < ARRAY_LOOP; i++) { @@ -79,10 +74,6 @@ static void MallocArray_DiffSizes(int gpu) { } } -static void MallocArrayThreadFunc(int gpu) { - MallocArray_DiffSizes(gpu); -} - /* * This testcase verifies the negative scenarios of * hipMallocArray API @@ -92,56 +83,31 @@ TEST_CASE("Unit_hipMallocArray_Negative") { hipChannelFormatDesc desc = hipCreateChannelDesc(); #if HT_NVIDIA SECTION("NullPointer to Array") { - REQUIRE(hipMallocArray(nullptr, &desc, - NUM_W, NUM_H, hipArrayDefault) != hipSuccess); + REQUIRE(hipMallocArray(nullptr, &desc, NUM_W, NUM_H, hipArrayDefault) != hipSuccess); } SECTION("NullPointer to Channel Descriptor") { - REQUIRE(hipMallocArray(&A_d, nullptr, - NUM_W, NUM_H, hipArrayDefault) != hipSuccess); + REQUIRE(hipMallocArray(&A_d, nullptr, NUM_W, NUM_H, hipArrayDefault) != hipSuccess); } #endif SECTION("Width 0 in hipMallocArray") { - REQUIRE(hipMallocArray(&A_d, &desc, - 0, NUM_H, hipArrayDefault) != hipSuccess); + REQUIRE(hipMallocArray(&A_d, &desc, 0, NUM_H, hipArrayDefault) != hipSuccess); } SECTION("Height 0 in hipMallocArray") { - REQUIRE(hipMallocArray(&A_d, &desc, - NUM_W, 0, hipArrayDefault) == hipSuccess); + REQUIRE(hipMallocArray(&A_d, &desc, NUM_W, 0, hipArrayDefault) == hipSuccess); } - SECTION("Invalid Flag") { - REQUIRE(hipMallocArray(&A_d, &desc, - NUM_W, NUM_H, 100) != hipSuccess); - } + SECTION("Invalid Flag") { REQUIRE(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, 100) != hipSuccess); } SECTION("Max int values") { - REQUIRE(hipMallocArray(&A_d, &desc, - std::numeric_limits::max(), - std::numeric_limits::max(), - hipArrayDefault) != hipSuccess); + REQUIRE(hipMallocArray(&A_d, &desc, std::numeric_limits::max(), + std::numeric_limits::max(), hipArrayDefault) != hipSuccess); } } -/* - * This testcase verifies the basic scenario of - * hipMallocArray API for different datatypes - * of size 10 - */ -TEMPLATE_TEST_CASE("Unit_hipMallocArray_Basic", - "", int, unsigned int, float) { - hipArray* A_d; - hipChannelFormatDesc desc = hipCreateChannelDesc(); - REQUIRE(hipMallocArray(&A_d, &desc, - NUM_W, NUM_H, - hipArrayDefault) == hipSuccess); - HIP_CHECK(hipFreeArray(A_d)); -} -TEST_CASE("Unit_hipMallocArray_DiffSizes") { - MallocArray_DiffSizes(0); -} +TEST_CASE("Unit_hipMallocArray_DiffSizes") { MallocArray_DiffSizes(0); } /* @@ -156,10 +122,10 @@ TEST_CASE("Unit_hipMallocArray_MultiThread") { size_t tot, avail, ptot, pavail; HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); for (int i = 0; i < devCnt; i++) { - threadlist.push_back(std::thread(MallocArrayThreadFunc, i)); + threadlist.push_back(std::thread(MallocArray_DiffSizes, i)); } - for (auto &t : threadlist) { + for (auto& t : threadlist) { t.join(); } HIP_CHECK(hipMemGetInfo(&avail, &tot)); @@ -170,3 +136,266 @@ 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); + for (size_t i = 0; i < vector_info::size; ++i) { + as[i] = as[i] + static_cast(1); + } +} + +// read from a surface and write to another +template __global__ void incSurface(hipSurfaceObject_t surf, size_t height) { + // Calculate surface coordinates + unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; + unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; + if (height == 0) { + T data; + surf1Dread(&data, surf, x * sizeof(T)); + addOne(&data); // change the value to show that write works + surf1Dwrite(data, surf, x * sizeof(T)); + } else { + T data; + surf2Dread(&data, surf, x * sizeof(T), y); + addOne(&data); // change the value to show that write works + surf2Dwrite(data, surf, x * sizeof(T), y); + } +} + +// Helpers /////////////////////////////////////// + +template size_t getAllocSize(const size_t width, const size_t height) noexcept { + 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); +} + +// Tests ///////////////////////////////////////// + +// Test the default array by generating a texture from it then reading from that texture. +// Textures are read-only so write to the array then copy from the texture into normal device memory +template +void testArrayAsTexture(hipArray_t arrayPtr, const size_t width, const size_t height) { + using scalar_type = typename vector_info::type; + constexpr auto vec_size = vector_info::size; + + const auto h = height ? height : 1; + const size_t pitch = width * sizeof(T); // no padding + const auto size = pitch * h; + + // create an array to initialize the hip array, then later use it to hold the result + std::vector hostData(width * h * vec_size); + + // Setup backing array + // assign ascending values to the data array to show indexing is working. + std::iota(std::begin(hostData), std::end(hostData), 0); + HIP_CHECK( + hipMemcpy2DToArray(arrayPtr, 0, 0, hostData.data(), pitch, pitch, h, hipMemcpyHostToDevice)); + + + // create texture + hipTextureObject_t textObj{}; + hipResourceDesc resDesc{}; + memset(&resDesc, 0, sizeof(hipResourceDesc)); + // enum to store how to resDesc.res union is being used + resDesc.resType = hipResourceTypeArray; + resDesc.res.array.array = arrayPtr; + + hipTextureDesc textDesc{}; + memset(&textDesc, 0, sizeof(hipTextureDesc)); + textDesc.filterMode = + hipFilterModePoint; // use the actual values in the texture, not normalized data + textDesc.readMode = hipReadModeElementType; // don't convert the data to floats + textDesc.normalizedCoords = 1; // use normalized coordinates (0.0-1.0) + + HIP_CHECK(hipCreateTextureObject(&textObj, &resDesc, &textDesc, 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(hipDestroyTextureObject(textObj)); + HIP_CHECK(hipFree(device_data)); +} + +// Test the an array created with the SurfaceLoadStore flag by generating a surface and reading from +// it and writing to it. +template +void testArrayAsSurface(hipArray_t arrayPtr, const size_t width, const size_t height) { + using scalar_type = typename vector_info::type; + constexpr auto vec_size = vector_info::size; + + const auto h = height ? height : 1; + const size_t pitch = width * sizeof(T); // no padding + const auto size = pitch * h; + + std::vector hostData(width * h * vec_size); + + // Setup backing array + // assign ascending values to the data array to show indexing is working. + std::iota(std::begin(hostData), std::end(hostData), 0); + HIP_CHECK( + hipMemcpy2DToArray(arrayPtr, 0, 0, hostData.data(), pitch, pitch, h, hipMemcpyHostToDevice)); + + + // create surface + hipSurfaceObject_t surfObj{}; + hipResourceDesc resDesc; + memset(&resDesc, 0, sizeof(hipResourceDesc)); + resDesc.resType = hipResourceTypeArray; + + resDesc.res.array.array = arrayPtr; + HIP_CHECK(hipCreateSurfaceObject(&surfObj, &resDesc)); + + + // run kernel + T* device_data{}; + HIP_CHECK(hipMalloc(&device_data, size)); + // This will increment the values of the surface, so this is undone later + incSurface<<>>(surfObj, height); + 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(hipMemcpy2DFromArray(hostData.data(), pitch, arrayPtr, 0, 0, pitch, h, + hipMemcpyDeviceToHost)); + + + // undo the increment + std::for_each(std::begin(hostData), std::end(hostData), + [](scalar_type& x) { x -= static_cast(1); }); + checkDataIsAscending(hostData); + + // clean up + HIP_CHECK(hipDestroySurfaceObject(surfObj)); + 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, + char4, float, float2, float4) { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-62"); +#endif + + hipChannelFormatDesc desc = hipCreateChannelDesc(); + + size_t init_free = getFreeMem(); + + // pointer to the array in device memory + hipArray_t arrayPtr{}; + size_t width = 1024; + size_t height = GENERATE(0, 1024); + + SECTION("hipArrayDefault") { + INFO("flag is hipArrayDefault"); + INFO("height: " << height); + + HIP_CHECK(hipMallocArray(&arrayPtr, &desc, width, height, hipArrayDefault)); + testArrayAsTexture(arrayPtr, width, height); + } +#if HT_NVIDIA // surfaces and texture gather not supported on AMD + SECTION("hipArraySurfaceLoadStore") { + INFO("flag is hipArraySurfaceLoadStore"); + INFO("height: " << height); + + HIP_CHECK(hipMallocArray(&arrayPtr, &desc, width, height, hipArraySurfaceLoadStore)); + testArrayAsSurface(arrayPtr, width, height); + } +#endif + + size_t final_free = getFreeMem(); + + const size_t alloc_size = getAllocSize(width, height); + // alloc will be chunked, so this is not exact + REQUIRE(init_free - final_free >= alloc_size); + + HIP_CHECK(hipFreeArray(arrayPtr)); +}