From a5522a55efdef54444afee2bbb3e9cda1bb20b24 Mon Sep 17 00:00:00 2001 From: Vladana Stojiljkovic Date: Mon, 12 Aug 2024 15:19:04 +0200 Subject: [PATCH] SWDEV-478767 - Implement tests for hipTexRefSetAddress2D Change-Id: Id5b946c216097c606a1e4b839514f68d351adc19 [ROCm/hip-tests commit: b24e61671aebfeed4ef0fb42c6cc071c962e35ba] --- .../catch/unit/texture/CMakeLists.txt | 10 ++ .../unit/texture/hipTexRefSetAddress2D.cc | 124 ++++++++++++++++++ .../catch/unit/texture/tex_ref_get_module.cc | 28 ++++ 3 files changed, 162 insertions(+) create mode 100644 projects/hip-tests/catch/unit/texture/hipTexRefSetAddress2D.cc create mode 100644 projects/hip-tests/catch/unit/texture/tex_ref_get_module.cc diff --git a/projects/hip-tests/catch/unit/texture/CMakeLists.txt b/projects/hip-tests/catch/unit/texture/CMakeLists.txt index 6c70b1ed93..ea947c3ed2 100644 --- a/projects/hip-tests/catch/unit/texture/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/texture/CMakeLists.txt @@ -44,6 +44,7 @@ set(TEST_SRC hipTextureObj2DCheckSRGBModes.cc hipTexObjectTests.cc hipTextureObjectTests.cc + hipTexRefSetAddress2D.cc ) # tests not for gfx90a+ @@ -89,6 +90,13 @@ function(CheckRejectedArchs OFFLOAD_ARCH_STR_LOCAL) endforeach() # OFFLOAD_ARCH_LIST endfunction() # CheckAcceptedArchs +add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/tex_ref_get_module.code + COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} --std=c++17 ${CMAKE_CURRENT_SOURCE_DIR}/tex_ref_get_module.cc + -o tex_ref_get_module.code + -I${HIP_PATH}/include/ --rocm-path=${ROCM_PATH} + DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/tex_ref_get_module.cc) +add_custom_target(tex_ref_get_module ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/tex_ref_get_module.code) + if(HIP_PLATFORM MATCHES "amd") if (DEFINED OFFLOAD_ARCH_STR) CheckRejectedArchs(${OFFLOAD_ARCH_STR}) @@ -127,3 +135,5 @@ endif() hip_add_exe_to_target(NAME TextureTest TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests) + +add_dependencies(TextureTest tex_ref_get_module) diff --git a/projects/hip-tests/catch/unit/texture/hipTexRefSetAddress2D.cc b/projects/hip-tests/catch/unit/texture/hipTexRefSetAddress2D.cc new file mode 100644 index 0000000000..b5db9e1b2c --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/hipTexRefSetAddress2D.cc @@ -0,0 +1,124 @@ +/* +Copyright (c) 2024 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 clang diagnostic ignored "-Wunused-parameter" +#include +#include + +#if defined(__HIP_PLATFORM_AMD__) || CUDA_VERSION < CUDA_12000 + +TEST_CASE("Unit_hipTexRefSetAddress2D_Negative_Parameters") { + CHECK_IMAGE_SUPPORT + + constexpr int width = 256; + constexpr int height = 256; + + hipCtx_t ctx; + hipDevice_t device; + + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipCtxCreate(&ctx, 0, device)); + + hipTexRef tex_ref = nullptr; + hipModule_t module = nullptr; + HIP_CHECK(hipModuleLoad(&module, "tex_ref_get_module.code")); + HIP_CHECK(hipModuleGetTexRef(&tex_ref, module, "tex")); + + int size = width * height * sizeof(float); + float* h_data = new float[size]; + for (int i = 0; i < width * height; ++i) { + h_data[i] = static_cast(i); + } + + hipDeviceptr_t d_data; + size_t dest_pitch; + HIP_CHECK(hipMemAllocPitch(&d_data, &dest_pitch, width * sizeof(float), height, sizeof(float))); + HIP_CHECK(hipMemcpy2D((void*)d_data, dest_pitch, h_data, width * sizeof(float), + width * sizeof(float), height, hipMemcpyHostToDevice)); + + HIP_ARRAY_DESCRIPTOR array_desc; + array_desc.Format = HIP_AD_FORMAT_FLOAT; + array_desc.Height = height; + array_desc.Width = width; + array_desc.NumChannels = 1; + + SECTION("Null texture") { +#if HT_AMD + HIP_CHECK_ERROR(hipTexRefSetAddress2D(nullptr, &array_desc, d_data, dest_pitch), + hipErrorInvalidValue); +#else + HIP_CHECK_ERROR(hipTexRefSetAddress2D(nullptr, &array_desc, d_data, dest_pitch), + hipErrorInvalidResourceHandle); +#endif + } + + SECTION("Null array descriptor") { + HIP_CHECK_ERROR(hipTexRefSetAddress2D(tex_ref, nullptr, d_data, dest_pitch), + hipErrorInvalidValue); + } + + free(h_data); + HIP_CHECK(hipFree((void*)d_data)); + HIP_CHECK(hipModuleUnload(module)); + HIP_CHECK(hipCtxDestroy(ctx)); +} + +TEST_CASE("Unit_hipTexRefSetAddress2D_Positive") { + CHECK_IMAGE_SUPPORT + + constexpr int width = 256; + constexpr int height = 256; + + hipCtx_t ctx; + hipDevice_t device; + + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipCtxCreate(&ctx, 0, device)); + + hipTexRef tex_ref = nullptr; + hipModule_t module = nullptr; + HIP_CHECK(hipModuleLoad(&module, "tex_ref_get_module.code")); + HIP_CHECK(hipModuleGetTexRef(&tex_ref, module, "tex")); + + int size = width * height * sizeof(float); + float* h_data = new float[size]; + for (int i = 0; i < width * height; ++i) { + h_data[i] = static_cast(i); + } + + hipDeviceptr_t d_data; + size_t dest_pitch; + HIP_CHECK(hipMemAllocPitch(&d_data, &dest_pitch, width * sizeof(float), height, sizeof(float))); + HIP_CHECK(hipMemcpy2D((void*)d_data, dest_pitch, h_data, width * sizeof(float), + width * sizeof(float), height, hipMemcpyHostToDevice)); + + HIP_ARRAY_DESCRIPTOR array_desc; + array_desc.Format = HIP_AD_FORMAT_FLOAT; + array_desc.Height = height; + array_desc.Width = width; + array_desc.NumChannels = 1; + HIP_CHECK(hipTexRefSetAddress2D(tex_ref, &array_desc, d_data, dest_pitch)); + + free(h_data); + HIP_CHECK(hipFree((void*)d_data)); + HIP_CHECK(hipModuleUnload(module)); + HIP_CHECK(hipCtxDestroy(ctx)); +} + +#endif diff --git a/projects/hip-tests/catch/unit/texture/tex_ref_get_module.cc b/projects/hip-tests/catch/unit/texture/tex_ref_get_module.cc new file mode 100644 index 0000000000..af7e542f75 --- /dev/null +++ b/projects/hip-tests/catch/unit/texture/tex_ref_get_module.cc @@ -0,0 +1,28 @@ +/* +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 + +#if defined(__HIP_PLATFORM_AMD__) || CUDA_VERSION < CUDA_12000 + +texture tex; + +#endif // defined(__HIP_PLATFORM_AMD__) || CUDA_VERSION < CUDA_12000