diff --git a/catch/unit/texture/CMakeLists.txt b/catch/unit/texture/CMakeLists.txt index 27f7037a79..8c14e1d62b 100644 --- a/catch/unit/texture/CMakeLists.txt +++ b/catch/unit/texture/CMakeLists.txt @@ -32,7 +32,7 @@ set(TEST_SRC hipSimpleTexture1DLayered.cc hipSimpleTexture2DLayered.cc hipBindTexture.cc - hipBindTex2DPitch.cc + hipBindTexture2D.cc hipTex1DFetchCheckModes.cc hipGetChanDesc.cc hipGetTextureAlignmentOffset.cc diff --git a/catch/unit/texture/hipBindTex2DPitch.cc b/catch/unit/texture/hipBindTex2DPitch.cc deleted file mode 100644 index 79aa118c73..0000000000 --- a/catch/unit/texture/hipBindTex2DPitch.cc +++ /dev/null @@ -1,91 +0,0 @@ -/* -Copyright (c) 2023 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. -*/ - -#pragma clang diagnostic ignored "-Wunused-parameter" -#include -#include - - -#if CUDA_VERSION < CUDA_12000 - -#define SIZE_H 8 -#define SIZE_W 12 -#define TYPE_t float - -texture tex; - -// texture object is a kernel argument -static __global__ void texture2dCopyKernel(TYPE_t* dst) { -#if !__HIP_NO_IMAGE_SUPPORT - int x = threadIdx.x + blockIdx.x * blockDim.x; - int y = threadIdx.y + blockIdx.y * blockDim.y; - if ( (x < SIZE_W) && (y < SIZE_H) ) { - dst[SIZE_W*y+x] = tex2D(tex, x, y); - } -#endif -} - - -TEST_CASE("Unit_hipBindTexture2D_Pitch") { - CHECK_IMAGE_SUPPORT - -#if __HIP_NO_IMAGE_SUPPORT - HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); - return; -#endif - - TYPE_t* B; - TYPE_t* A; - TYPE_t* devPtrB; - TYPE_t* devPtrA; - - B = new TYPE_t[SIZE_H*SIZE_W]; - A = new TYPE_t[SIZE_H*SIZE_W]; - for (size_t i = 1; i <= (SIZE_H * SIZE_W); i++) { - A[i-1] = i; - } - - size_t devPitchA, tex_ofs; - HIP_CHECK(hipMallocPitch(reinterpret_cast(&devPtrA), &devPitchA, - SIZE_W*sizeof(TYPE_t), SIZE_H)); - HIP_CHECK(hipMemcpy2D(devPtrA, devPitchA, A, SIZE_W*sizeof(TYPE_t), - SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyHostToDevice)); - - tex.normalized = false; - HIP_CHECK(hipBindTexture2D(&tex_ofs, &tex, devPtrA, &tex.channelDesc, - SIZE_W, SIZE_H, devPitchA)); - HIP_CHECK(hipMalloc(reinterpret_cast(&devPtrB), - SIZE_W*sizeof(TYPE_t)*SIZE_H)); - - hipLaunchKernelGGL(texture2dCopyKernel, dim3(4, 4, 1), dim3(32, 32, 1), - 0, 0, devPtrB); - HIP_CHECK(hipGetLastError()); - HIP_CHECK(hipDeviceSynchronize()); - HIP_CHECK(hipMemcpy2D(B, SIZE_W*sizeof(TYPE_t), devPtrB, - SIZE_W*sizeof(TYPE_t), SIZE_W*sizeof(TYPE_t), - SIZE_H, hipMemcpyDeviceToHost)); - HipTest::checkArray(A, B, SIZE_H, SIZE_W); - delete []A; - delete []B; - HIP_CHECK(hipFree(devPtrA)); - HIP_CHECK(hipFree(devPtrB)); -} - -#endif // CUDA_VERSION < CUDA_12000 - diff --git a/catch/unit/texture/hipBindTexture2D.cc b/catch/unit/texture/hipBindTexture2D.cc new file mode 100644 index 0000000000..db26642222 --- /dev/null +++ b/catch/unit/texture/hipBindTexture2D.cc @@ -0,0 +1,145 @@ +/* +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 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. +*/ + +#pragma clang diagnostic ignored "-Wunused-parameter" +#include +#include + +#if defined(__HIP_PLATFORM_AMD__) || CUDA_VERSION < CUDA_12000 + +#define SIZE_H 8 +#define SIZE_W 12 + +texture tex; + +// texture object is a kernel argument +static __global__ void texture2dCopyKernel(float* dst) { +#if !__HIP_NO_IMAGE_SUPPORT + int x = threadIdx.x + blockIdx.x * blockDim.x; + int y = threadIdx.y + blockIdx.y * blockDim.y; + if ((x < SIZE_W) && (y < SIZE_H)) { + dst[SIZE_W * y + x] = tex2D(tex, x, y); + } +#endif +} + +TEST_CASE("Unit_hipBindTexture2D_Positive") { + CHECK_IMAGE_SUPPORT + float* device_ptr; + size_t device_pitch, texture_offset; + HIP_CHECK(hipMallocPitch(&device_ptr, &device_pitch, SIZE_W, SIZE_H)); + HIP_CHECK(hipBindTexture2D(&texture_offset, &tex, device_ptr, &tex.channelDesc, SIZE_W, SIZE_H, + device_pitch)); + HIP_CHECK(hipUnbindTexture(tex)); + HIP_CHECK(hipFree((void*)device_ptr)); +} + +TEST_CASE("Unit_hipBindTexture2D_Pitch") { + CHECK_IMAGE_SUPPORT + +#if __HIP_NO_IMAGE_SUPPORT + HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set"); + return; +#endif + + float* b; + float* a; + float* dev_ptr_b; + float* dev_ptr_a; + + b = new float[SIZE_H * SIZE_W]; + a = new float[SIZE_H * SIZE_W]; + for (size_t i = 1; i <= (SIZE_H * SIZE_W); i++) { + a[i - 1] = i; + } + + size_t dev_pitch_a, tex_ofs; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&dev_ptr_a), &dev_pitch_a, + SIZE_W * sizeof(float), SIZE_H)); + HIP_CHECK(hipMemcpy2D(dev_ptr_a, dev_pitch_a, a, SIZE_W * sizeof(float), SIZE_W * sizeof(float), + SIZE_H, hipMemcpyHostToDevice)); + + tex.normalized = false; + HIP_CHECK( + hipBindTexture2D(&tex_ofs, &tex, dev_ptr_a, &tex.channelDesc, SIZE_W, SIZE_H, dev_pitch_a)); + HIP_CHECK(hipMalloc(reinterpret_cast(&dev_ptr_b), SIZE_W * sizeof(float) * SIZE_H)); + + hipLaunchKernelGGL(texture2dCopyKernel, dim3(4, 4, 1), dim3(32, 32, 1), 0, 0, dev_ptr_b); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy2D(b, SIZE_W * sizeof(float), dev_ptr_b, SIZE_W * sizeof(float), + SIZE_W * sizeof(float), SIZE_H, hipMemcpyDeviceToHost)); + HipTest::checkArray(a, b, SIZE_H, SIZE_W); + delete[] a; + delete[] b; + HIP_CHECK(hipFree(dev_ptr_a)); + HIP_CHECK(hipFree(dev_ptr_b)); + HIP_CHECK(hipUnbindTexture(tex)); +} + +TEST_CASE("Unit_hipBindTexture2D_Negative") { + CHECK_IMAGE_SUPPORT + float* device_ptr; + size_t device_pitch, texture_offset; + HIP_CHECK(hipMallocPitch(&device_ptr, &device_pitch, SIZE_W, SIZE_H / 2)); + + SECTION("Texture is nullptr") { +#if HT_AMD + HIP_CHECK_ERROR(hipBindTexture2D(&texture_offset, nullptr, device_ptr, &tex.channelDesc, SIZE_W, + SIZE_H, device_pitch), + hipErrorInvalidSymbol); +#else + HIP_CHECK_ERROR(hipBindTexture2D(&texture_offset, nullptr, device_ptr, &tex.channelDesc, SIZE_W, + SIZE_H, device_pitch), + hipErrorInvalidTexture); +#endif + } + + SECTION("Device ptr is nullptr") { +#if HT_AMD + HIP_CHECK_ERROR(hipBindTexture2D(&texture_offset, &tex, nullptr, &tex.channelDesc, SIZE_W, + SIZE_H, device_pitch), + hipErrorInvalidValue); +#else + HIP_CHECK_ERROR(hipBindTexture2D(&texture_offset, &tex, nullptr, &tex.channelDesc, SIZE_W, + SIZE_H, device_pitch), + hipErrorNotFound); +#endif + } + + SECTION("Width is 0") { + HIP_CHECK_ERROR(hipBindTexture2D(&texture_offset, &tex, device_ptr, &tex.channelDesc, 0, SIZE_H, + device_pitch), + hipErrorInvalidValue); + } + + SECTION("Height is 0") { + HIP_CHECK_ERROR(hipBindTexture2D(&texture_offset, &tex, device_ptr, &tex.channelDesc, SIZE_W, 0, + device_pitch), + hipErrorInvalidValue); + } + + SECTION("Pitch is 0") { + HIP_CHECK_ERROR( + hipBindTexture2D(&texture_offset, &tex, device_ptr, &tex.channelDesc, SIZE_W, SIZE_H, 0), + hipErrorInvalidValue); + } +} + +#endif // __HIP_PLATFORM_AMD__ || CUDA_VERSION < CUDA_12000