From 08f750571da58c463c59faa58ff0f04ac41f11d9 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 23 May 2018 14:43:47 +0530 Subject: [PATCH 01/14] Optimize memcpy2D kernel use [ROCm/hip commit: 9a76d5b94c86258a12c81fe5e3db6e9908391968] --- projects/hip/src/hip_memory.cpp | 16 ++++++++++++---- 1 file changed, 12 insertions(+), 4 deletions(-) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index fc60dffe86..b76f7f0524 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -1495,10 +1495,18 @@ __global__ void hip_copy2d_n(T* dst, const T* src, size_t width, size_t height, size_t idx = blockIdx.x * blockDim.x + threadIdx.x; size_t idy = blockIdx.y * blockDim.y + threadIdx.y; - if((idx < width) && (idy < height)){ + size_t floorWidth = (width/sizeof(T)); + if((idx < floorWidth)){ T *dstPtr = (T *)((uint8_t*) dst + idy * destPitch); T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch); dstPtr[idx] = srcPtr[idx]; + } else { + size_t bytesToCopy = width - floorWidth; + uint8_t *dstPtr = (uint8_t *) ((uint8_t*) dst + idy * destPitch); + uint8_t *srcPtr = (uint8_t *) ((uint8_t*) src + idy * srcPitch); + for(int i =0 ; i < bytesToCopy ; i++) { + dstPtr[idx+i]= srcPtr[idx+i]; + } } } } // namespace @@ -1516,7 +1524,7 @@ void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t sizeBytes) { template void ihipMemcpy2dKernel(hipStream_t stream, T* dst, const T* src, size_t width, size_t height, size_t destPitch, size_t srcPitch) { size_t threadsPerBlock = 16; - uint32_t grid_dim_x = clamp_integer( (width+(threadsPerBlock-1)) / threadsPerBlock, 1, UINT32_MAX); + uint32_t grid_dim_x = clamp_integer( ((width/sizeof(T))+(threadsPerBlock-1)) / threadsPerBlock, 1, UINT32_MAX); uint32_t grid_dim_y = clamp_integer( (height+(threadsPerBlock-1)) / threadsPerBlock, 1, UINT32_MAX); hipLaunchKernelGGL(hip_copy2d_n, dim3(grid_dim_x,grid_dim_y), dim3(threadsPerBlock,threadsPerBlock), 0u, stream, dst, src, width, height, destPitch, srcPitch); @@ -1622,7 +1630,7 @@ hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch stream->locked_copySync((unsigned char*)dst + i * dpitch, (unsigned char*)src + i * spitch, width, kind); } else { - ihipMemcpy2dKernel (stream, static_cast (dst), static_cast (src), width, height, dpitch, spitch); + ihipMemcpy2dKernel (stream, static_cast (dst), static_cast (src), width, height, dpitch, spitch); stream->locked_wait(); } } catch (ihipException& ex) { @@ -1661,7 +1669,7 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp e = hip_internal::memcpyAsync((unsigned char*)dst + i * dpitch, (unsigned char*)src + i * spitch, width, kind, stream); } else{ - ihipMemcpy2dKernel (stream, static_cast (dst), static_cast (src), width, height, dpitch, spitch); + ihipMemcpy2dKernel (stream, static_cast (dst), static_cast (src), width, height, dpitch, spitch); } } catch (ihipException& ex) { e = ex._code; From fb1425959e3e93bd3425b6f310a5d7a73366ef3e Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 24 May 2018 08:27:24 +0530 Subject: [PATCH 02/14] Correct remaining bytes in copy 2d kernel [ROCm/hip commit: dc179e0c33a6b5c7061d20c2151b22c001c0b279] --- projects/hip/src/hip_memory.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index b76f7f0524..a1bc1416a0 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -1501,7 +1501,7 @@ __global__ void hip_copy2d_n(T* dst, const T* src, size_t width, size_t height, T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch); dstPtr[idx] = srcPtr[idx]; } else { - size_t bytesToCopy = width - floorWidth; + size_t bytesToCopy = width - (floorWidth * sizeof(T)); uint8_t *dstPtr = (uint8_t *) ((uint8_t*) dst + idy * destPitch); uint8_t *srcPtr = (uint8_t *) ((uint8_t*) src + idy * srcPitch); for(int i =0 ; i < bytesToCopy ; i++) { From fb745baa7e674fe209a60d7bd6f5c335324f3ee9 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 24 May 2018 17:00:12 +0530 Subject: [PATCH 03/14] Fix memcpy2d kernel dims [ROCm/hip commit: 981e56a68f897e7dbd5c8c768be29dd05006eaaf] --- projects/hip/src/hip_memory.cpp | 26 ++++++++++++++------------ 1 file changed, 14 insertions(+), 12 deletions(-) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index a1bc1416a0..2f0c793033 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -1496,17 +1496,19 @@ __global__ void hip_copy2d_n(T* dst, const T* src, size_t width, size_t height, size_t idx = blockIdx.x * blockDim.x + threadIdx.x; size_t idy = blockIdx.y * blockDim.y + threadIdx.y; size_t floorWidth = (width/sizeof(T)); - if((idx < floorWidth)){ - T *dstPtr = (T *)((uint8_t*) dst + idy * destPitch); - T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch); - dstPtr[idx] = srcPtr[idx]; - } else { - size_t bytesToCopy = width - (floorWidth * sizeof(T)); - uint8_t *dstPtr = (uint8_t *) ((uint8_t*) dst + idy * destPitch); - uint8_t *srcPtr = (uint8_t *) ((uint8_t*) src + idy * srcPitch); - for(int i =0 ; i < bytesToCopy ; i++) { - dstPtr[idx+i]= srcPtr[idx+i]; - } + if((idx < width) && (idy < height)) { + if((idx < floorWidth)){ + T *dstPtr = (T *)((uint8_t*) dst + idy * destPitch); + T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch); + dstPtr[idx] = srcPtr[idx]; + } else { + size_t bytesToCopy = width - (floorWidth * sizeof(T)); + uint8_t *dstPtr = (uint8_t *) ((uint8_t*) dst + idy * destPitch); + uint8_t *srcPtr = (uint8_t *) ((uint8_t*) src + idy * srcPitch); + for(int i =0 ; i < bytesToCopy ; i++) { + dstPtr[idx+i]= srcPtr[idx+i]; + } + } } } } // namespace @@ -1524,7 +1526,7 @@ void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t sizeBytes) { template void ihipMemcpy2dKernel(hipStream_t stream, T* dst, const T* src, size_t width, size_t height, size_t destPitch, size_t srcPitch) { size_t threadsPerBlock = 16; - uint32_t grid_dim_x = clamp_integer( ((width/sizeof(T))+(threadsPerBlock-1)) / threadsPerBlock, 1, UINT32_MAX); + uint32_t grid_dim_x = clamp_integer( (width+(threadsPerBlock*sizeof(T)-1)) / (threadsPerBlock*sizeof(T)), 1, UINT32_MAX); uint32_t grid_dim_y = clamp_integer( (height+(threadsPerBlock-1)) / threadsPerBlock, 1, UINT32_MAX); hipLaunchKernelGGL(hip_copy2d_n, dim3(grid_dim_x,grid_dim_y), dim3(threadsPerBlock,threadsPerBlock), 0u, stream, dst, src, width, height, destPitch, srcPitch); From 35169c519116d24d1917273894f769243bdaf518 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 24 May 2018 23:30:27 +0530 Subject: [PATCH 04/14] Clean up and fix remaining bytes copy [ROCm/hip commit: 4ff059d641c796c9a9e2732745a53c86eb400b3a] --- projects/hip/src/hip_memory.cpp | 22 +++++++++------------- 1 file changed, 9 insertions(+), 13 deletions(-) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 2f0c793033..fa68d7cfb1 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -1496,19 +1496,15 @@ __global__ void hip_copy2d_n(T* dst, const T* src, size_t width, size_t height, size_t idx = blockIdx.x * blockDim.x + threadIdx.x; size_t idy = blockIdx.y * blockDim.y + threadIdx.y; size_t floorWidth = (width/sizeof(T)); - if((idx < width) && (idy < height)) { - if((idx < floorWidth)){ - T *dstPtr = (T *)((uint8_t*) dst + idy * destPitch); - T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch); - dstPtr[idx] = srcPtr[idx]; - } else { - size_t bytesToCopy = width - (floorWidth * sizeof(T)); - uint8_t *dstPtr = (uint8_t *) ((uint8_t*) dst + idy * destPitch); - uint8_t *srcPtr = (uint8_t *) ((uint8_t*) src + idy * srcPitch); - for(int i =0 ; i < bytesToCopy ; i++) { - dstPtr[idx+i]= srcPtr[idx+i]; - } - } + T *dstPtr = (T *)((uint8_t*) dst + idy * destPitch); + T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch); + if((idx < floorWidth) && (idy < height)){ + dstPtr[idx] = srcPtr[idx]; + } else if((idx < width) && (idy < height)){ + size_t bytesToCopy = width - (floorWidth * sizeof(T)); + dstPtr += floorWidth; + srcPtr += floorWidth; + __builtin_memcpy(reinterpret_cast(dstPtr), reinterpret_cast(srcPtr),bytesToCopy); } } } // namespace From 4021f68f648392cff10070ca6be795f47fc20df3 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 24 May 2018 23:51:52 +0530 Subject: [PATCH 05/14] Use 64x4 grid dims [ROCm/hip commit: d8cb47242bed6190b509a1e092721f148f9102e3] --- projects/hip/src/hip_memory.cpp | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index fa68d7cfb1..9b6758ddf4 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -1521,10 +1521,11 @@ void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t sizeBytes) { template void ihipMemcpy2dKernel(hipStream_t stream, T* dst, const T* src, size_t width, size_t height, size_t destPitch, size_t srcPitch) { - size_t threadsPerBlock = 16; - uint32_t grid_dim_x = clamp_integer( (width+(threadsPerBlock*sizeof(T)-1)) / (threadsPerBlock*sizeof(T)), 1, UINT32_MAX); - uint32_t grid_dim_y = clamp_integer( (height+(threadsPerBlock-1)) / threadsPerBlock, 1, UINT32_MAX); - hipLaunchKernelGGL(hip_copy2d_n, dim3(grid_dim_x,grid_dim_y), dim3(threadsPerBlock,threadsPerBlock), 0u, stream, dst, src, + size_t threadsPerBlock_x = 64; + size_t threadsPerBlock_y = 4; + uint32_t grid_dim_x = clamp_integer( (width+(threadsPerBlock_x*sizeof(T)-1)) / (threadsPerBlock_x*sizeof(T)), 1, UINT32_MAX); + uint32_t grid_dim_y = clamp_integer( (height+(threadsPerBlock_y-1)) / threadsPerBlock_y, 1, UINT32_MAX); + hipLaunchKernelGGL(hip_copy2d_n, dim3(grid_dim_x,grid_dim_y), dim3(threadsPerBlock_x,threadsPerBlock_y), 0u, stream, dst, src, width, height, destPitch, srcPitch); } From 2708f3325d2f8d1d4dc918eed8b70232b68f97ed Mon Sep 17 00:00:00 2001 From: Jorghi12 Date: Sat, 26 May 2018 00:40:14 -0400 Subject: [PATCH 06/14] Adding double/long int signatures for abs Adding overloads for abs that are found in cuda's math_functions. [ROCm/hip commit: 4383d6c6debdb19563d02147517dae6215e35157] --- projects/hip/src/math_functions.cpp | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/projects/hip/src/math_functions.cpp b/projects/hip/src/math_functions.cpp index dedc40f2ae..bd0116080d 100644 --- a/projects/hip/src/math_functions.cpp +++ b/projects/hip/src/math_functions.cpp @@ -59,6 +59,12 @@ __device__ int abs(int x) { __device__ long long abs(long long x) { return x >= 0 ? x : -x; } +__device__ double abs(double x) { + return x >= 0 ? x : -x; +} +__device__ long int abs(long int x) { + return x >= 0 ? x : -x; +} __device__ float fabsf(float x) { return hc::precise_math::fabsf(x); } __device__ float fdimf(float x, float y) { return hc::precise_math::fdimf(x, y); } __device__ float fdividef(float x, float y) { return x / y; } From a82eb2dfa2d7f4584ed063b9d558f72c071d0499 Mon Sep 17 00:00:00 2001 From: Jorghi12 Date: Sat, 26 May 2018 00:41:24 -0400 Subject: [PATCH 07/14] Adding double/long int signatures for abs Adding overloads for abs that are found in cuda's math_functions. [ROCm/hip commit: 4f6904b5c7fa7150efbaf2859467b47808f1f3fc] --- projects/hip/include/hip/hcc_detail/math_functions.h | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/hip/include/hip/hcc_detail/math_functions.h b/projects/hip/include/hip/hcc_detail/math_functions.h index e717df07c1..e7312f32f2 100644 --- a/projects/hip/include/hip/hcc_detail/math_functions.h +++ b/projects/hip/include/hip/hcc_detail/math_functions.h @@ -58,6 +58,8 @@ __device__ float expf(float x); __device__ float expm1f(float x); __device__ int abs(int x); __device__ long long abs(long long x); +__device__ double abs(double x); +__device__ long int abs(long int x); __device__ float fabsf(float x); __device__ float fdimf(float x, float y); __device__ float fdividef(float x, float y); From c46563addc9caf2c676c0454cadc79edad28f14f Mon Sep 17 00:00:00 2001 From: Jorghi12 Date: Sat, 26 May 2018 16:21:14 -0400 Subject: [PATCH 08/14] Update math_functions.cpp CUDA also has a function named labs. [ROCm/hip commit: ec2edb2c927e86944a0c5e077fa634357ef3d1d4] --- projects/hip/src/math_functions.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/projects/hip/src/math_functions.cpp b/projects/hip/src/math_functions.cpp index bd0116080d..982073275b 100644 --- a/projects/hip/src/math_functions.cpp +++ b/projects/hip/src/math_functions.cpp @@ -65,6 +65,9 @@ __device__ double abs(double x) { __device__ long int abs(long int x) { return x >= 0 ? x : -x; } +__device__ long long int labs(long long int x) { + return x >= 0 ? x : -x; +} __device__ float fabsf(float x) { return hc::precise_math::fabsf(x); } __device__ float fdimf(float x, float y) { return hc::precise_math::fdimf(x, y); } __device__ float fdividef(float x, float y) { return x / y; } From 50ccc9fc63b0db2270880a76beb94e28b4342972 Mon Sep 17 00:00:00 2001 From: Jorghi12 Date: Sat, 26 May 2018 16:22:10 -0400 Subject: [PATCH 09/14] Update math_functions.h CUDA also has a function named labs. [ROCm/hip commit: 6e2b9d054794540a83e74d33326a3eda7c039eca] --- projects/hip/include/hip/hcc_detail/math_functions.h | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/hip/include/hip/hcc_detail/math_functions.h b/projects/hip/include/hip/hcc_detail/math_functions.h index e7312f32f2..ea3bd97571 100644 --- a/projects/hip/include/hip/hcc_detail/math_functions.h +++ b/projects/hip/include/hip/hcc_detail/math_functions.h @@ -60,6 +60,7 @@ __device__ int abs(int x); __device__ long long abs(long long x); __device__ double abs(double x); __device__ long int abs(long int x); +__device__ long long int labs(long long int x); __device__ float fabsf(float x); __device__ float fdimf(float x, float y); __device__ float fdividef(float x, float y); From 24b35fc052f816d3b124d82a33c3a454e77f932b Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 28 May 2018 15:02:06 +0530 Subject: [PATCH 10/14] Add 1d texture types for NVCC path [ROCm/hip commit: 024f77ce6197cce3e041e7fbb0ee1e916436ff14] --- projects/hip/include/hip/nvcc_detail/hip_runtime_api.h | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index bd3ffc1cc1..e00405cc71 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -155,6 +155,8 @@ typedef struct cudaArray* hipArray_const_t; typedef cudaTextureObject_t hipTextureObject_t; typedef cudaSurfaceObject_t hipSurfaceObject_t; +#define hipTextureType1D cudaTextureType1D +#define hipTextureType1DLayered cudaTextureType1DLayered #define hipTextureType2D cudaTextureType2D #define hipTextureType3D cudaTextureType3D #define hipDeviceMapHost cudaDeviceMapHost From d5039612f8f6c7a7c9c1ab5764284e6e835f1ec1 Mon Sep 17 00:00:00 2001 From: lthakur Date: Tue, 29 May 2018 14:08:01 +0530 Subject: [PATCH 11/14] HIP test case for 1D texture fetch (#424) [ROCm/hip commit: a05ac35ab120eee2049fa17ea4e278875d36e2f5] --- .../hip/tests/src/texture/hipBindTexObj1D.cpp | 100 ++++++++++++++++ .../tests/src/texture/tex1Dfetch_linear.cpp | 112 ++++++++++++++++++ 2 files changed, 212 insertions(+) create mode 100644 projects/hip/tests/src/texture/hipBindTexObj1D.cpp create mode 100644 projects/hip/tests/src/texture/tex1Dfetch_linear.cpp diff --git a/projects/hip/tests/src/texture/hipBindTexObj1D.cpp b/projects/hip/tests/src/texture/hipBindTexObj1D.cpp new file mode 100644 index 0000000000..2bcf78b1ca --- /dev/null +++ b/projects/hip/tests/src/texture/hipBindTexObj1D.cpp @@ -0,0 +1,100 @@ +/* +Copyright (c) 2015-2017 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. +*/ + + +/* HIT_START + * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * RUN: %t + * HIT_END + */ + +#include "hip/hip_runtime.h" +#include "hip/hip_runtime_api.h" +#include "test_common.h" +#include +#include + +#define N 512 +using namespace std; + +texture tex; + +bool testResult = true; + +__global__ void kernel(int *out) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + out[x] = tex1Dfetch(tex, x); +} + +void runTest(void); + +int main(int argc, char **argv) { + runTest(); + + if (testResult) { + passed(); + } else { + exit(EXIT_FAILURE); + } +} + +void runTest() { + string out; + int *tex_buf; + int val[N], i, output[N]; + size_t size = 0; + + for (i = 0; i < N; i++) { + val[i] = i; + output[i] = 0; + } + hipChannelFormatDesc chan_desc = + hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindUnsigned); + + hipMalloc(&tex_buf, N * sizeof(int)); + + hipMemcpy(tex_buf, val, N * sizeof(int), hipMemcpyHostToDevice); + + tex.addressMode[0] = hipAddressModeWrap; + tex.filterMode = hipFilterModeLinear; + tex.normalized = true; + + hipBindTexture(&size, &tex, (void *)tex_buf, &chan_desc, N * sizeof(int)); + + dim3 dimBlock(64, 1, 1); + dim3 dimGrid(N / dimBlock.x, 1, 1); + + hipLaunchKernelGGL(kernel, dim3(dimGrid), dim3(dimBlock), 0, 0, output); + + hipDeviceSynchronize(); + + hipMemcpy(output, tex_buf, N * sizeof(int), hipMemcpyDeviceToHost); + + for (i = 0; i < N; i++) { + if (output[i] != val[i]) { + testResult = false; + return; + } + } + hipUnbindTexture(&tex); + hipFree(tex_buf); +} diff --git a/projects/hip/tests/src/texture/tex1Dfetch_linear.cpp b/projects/hip/tests/src/texture/tex1Dfetch_linear.cpp new file mode 100644 index 0000000000..4f755873f7 --- /dev/null +++ b/projects/hip/tests/src/texture/tex1Dfetch_linear.cpp @@ -0,0 +1,112 @@ +/* +Copyright (c) 2015-2017 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. +*/ + +/*HIT_START + * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * RUN: %t + * HIT_END + */ + +#include "hip/hip_runtime.h" +#include "hip/hip_runtime_api.h" +#include "test_common.h" +#include +#include + +#define N 512 +using namespace std; + +bool testResult = true; + +__global__ void tex1d_kernel(float *val, hipTextureObject_t obj) { + int k = blockIdx.x * blockDim.x + threadIdx.x; + val[k] = tex1Dfetch(obj, k); +} + +void runTest(void); + +int main(int argc, char **argv) { + runTest(); + + if (testResult) { + passed(); + } else { + exit(EXIT_FAILURE); + } +} + +void runTest() { + + // Allocating the required buffer on gpu device + float *tex_buf, *tex_buf_check; + float val[N], output[N]; + int i; + for (i = 0; i < N; i++) + val[i] = (i + 1) * (i + 1); + hipMalloc(&tex_buf, N * sizeof(float)); + + hipMalloc(&tex_buf_check, N * sizeof(float)); + + hipMemcpy(tex_buf, val, N * sizeof(float), hipMemcpyHostToDevice); + + hipMemset(tex_buf_check, 0, N * sizeof(float)); + hipResourceDesc res_lin; + + memset(&res_lin, 0, sizeof(res_lin)); + + res_lin.resType = hipResourceTypeLinear; + res_lin.res.linear.devPtr = tex_buf; + res_lin.res.linear.desc.f = hipChannelFormatKindFloat; + res_lin.res.linear.desc.x = 32; + res_lin.res.linear.sizeInBytes = N * sizeof(float); + + hipTextureDesc tex_desc; + memset(&tex_desc, 0, sizeof(tex_desc)); + tex_desc.readMode = hipReadModeElementType; + + // Creating texture object + + hipTextureObject_t tex_obj = 0; + + hipCreateTextureObject(&tex_obj, &res_lin, &tex_desc, NULL); + + dim3 dimBlock(64, 1, 1); + dim3 dimGrid(N / dimBlock.x, 1, 1); + + for (i = 0; i < N; i++) + output[i] = 0; + + hipLaunchKernelGGL(tex1d_kernel, dim3(dimGrid), dim3(dimBlock), 0, 0, + tex_buf_check, tex_obj); + hipDeviceSynchronize(); + + hipMemcpy(output, tex_buf_check, N * sizeof(float), hipMemcpyDeviceToHost); + + for (i = 0; i < N; i++) + if (output[i] != val[i]) { + testResult = false; + } + + hipDestroyTextureObject(tex_obj); + hipFree(tex_buf); + hipFree(tex_buf_check); +} From 56de4b168d1a5f2543413f78b7deb7a7a9518e41 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 29 May 2018 16:03:37 +0530 Subject: [PATCH 12/14] Fix tex2D tests for result value [ROCm/hip commit: ca156449e21379ffe26da667d6f3b71b33295c79] --- .../hip/tests/src/texture/hipTextureObj2D.cpp | 216 +++++++++--------- .../hip/tests/src/texture/hipTextureRef2D.cpp | 186 +++++++-------- 2 files changed, 201 insertions(+), 201 deletions(-) diff --git a/projects/hip/tests/src/texture/hipTextureObj2D.cpp b/projects/hip/tests/src/texture/hipTextureObj2D.cpp index 1bf51bc2cb..504b632612 100644 --- a/projects/hip/tests/src/texture/hipTextureObj2D.cpp +++ b/projects/hip/tests/src/texture/hipTextureObj2D.cpp @@ -1,108 +1,108 @@ -/* HIT_START - * BUILD: %t %s ../test_common.cpp - * RUN: %t - * HIT_END - */ -#include -#include -#include - -#include -#include "test_common.h" - -bool testResult = true; - -__global__ void tex2DKernel(float* outputData, hipTextureObject_t textureObject, int width, - int height) { - int x = blockIdx.x * blockDim.x + threadIdx.x; - int y = blockIdx.y * blockDim.y + threadIdx.y; - outputData[y * width + x] = tex2D(textureObject, x, y); -} - -void runTest(int argc, char** argv); - -int main(int argc, char** argv) { - runTest(argc, argv); - - if (testResult) { - passed(); - } else { - exit(EXIT_FAILURE); - } -} - -void runTest(int argc, char** argv) { - unsigned int width = 256; - unsigned int height = 256; - unsigned int size = width * height * sizeof(float); - float* hData = (float*)malloc(size); - memset(hData, 0, size); - for (int i = 0; i < height; i++) { - for (int j = 0; j < width; j++) { - hData[i * width + j] = i * width + j; - } - } - printf("hData: "); - for (int i = 0; i < 64; i++) { - printf("%f ", hData[i]); - } - printf("\n"); - - hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat); - hipArray* hipArray; - hipMallocArray(&hipArray, &channelDesc, width, height); - - hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice); - - struct hipResourceDesc resDesc; - memset(&resDesc, 0, sizeof(resDesc)); - resDesc.resType = hipResourceTypeArray; - resDesc.res.array.array = hipArray; - - // Specify texture object parameters - struct hipTextureDesc texDesc; - memset(&texDesc, 0, sizeof(texDesc)); - texDesc.addressMode[0] = hipAddressModeWrap; - texDesc.addressMode[1] = hipAddressModeWrap; - texDesc.filterMode = hipFilterModePoint; - texDesc.readMode = hipReadModeElementType; - texDesc.normalizedCoords = 0; - - // Create texture object - hipTextureObject_t textureObject = 0; - hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL); - - float* dData = NULL; - hipMalloc((void**)&dData, size); - - dim3 dimBlock(16, 16, 1); - dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1); - - hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, textureObject, - width, height); - - hipDeviceSynchronize(); - - float* hOutputData = (float*)malloc(size); - memset(hOutputData, 0, size); - hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost); - - printf("dData: "); - for (int i = 0; i < 64; i++) { - printf("%f ", hOutputData[i]); - } - printf("\n"); - for (int i = 0; i < height; i++) { - for (int j = 0; j < width; j++) { - if (hData[i * width + j] != hOutputData[i * width + j]) { - printf("Difference [ %d %d ]:%f ----%f\n", i, j, hData[i * width + j], - hOutputData[i * width + j]); - testResult = false; - break; - } - } - } - hipDestroyTextureObject(textureObject); - hipFree(dData); - hipFreeArray(hipArray); -} +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * RUN: %t + * HIT_END + */ +#include +#include +#include + +#include +#include "test_common.h" + +__global__ void tex2DKernel(float* outputData, hipTextureObject_t textureObject, int width, + int height) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + outputData[y * width + x] = tex2D(textureObject, x, y); +} + +int runTest(int argc, char** argv); + +int main(int argc, char** argv) { + int testResult = runTest(argc, argv); + + if (testResult) { + passed(); + } else { + exit(EXIT_FAILURE); + } +} + +void runTest(int argc, char** argv) { + int testResult = 1; + unsigned int width = 256; + unsigned int height = 256; + unsigned int size = width * height * sizeof(float); + float* hData = (float*)malloc(size); + memset(hData, 0, size); + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + hData[i * width + j] = i * width + j; + } + } + printf("hData: "); + for (int i = 0; i < 64; i++) { + printf("%f ", hData[i]); + } + printf("\n"); + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat); + hipArray* hipArray; + hipMallocArray(&hipArray, &channelDesc, width, height); + + hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice); + + struct hipResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = hipResourceTypeArray; + resDesc.res.array.array = hipArray; + + // Specify texture object parameters + struct hipTextureDesc texDesc; + memset(&texDesc, 0, sizeof(texDesc)); + texDesc.addressMode[0] = hipAddressModeWrap; + texDesc.addressMode[1] = hipAddressModeWrap; + texDesc.filterMode = hipFilterModePoint; + texDesc.readMode = hipReadModeElementType; + texDesc.normalizedCoords = 0; + + // Create texture object + hipTextureObject_t textureObject = 0; + hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL); + + float* dData = NULL; + hipMalloc((void**)&dData, size); + + dim3 dimBlock(16, 16, 1); + dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1); + + hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, textureObject, + width, height); + + hipDeviceSynchronize(); + + float* hOutputData = (float*)malloc(size); + memset(hOutputData, 0, size); + hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost); + + printf("dData: "); + for (int i = 0; i < 64; i++) { + printf("%f ", hOutputData[i]); + } + printf("\n"); + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + if (hData[i * width + j] != hOutputData[i * width + j]) { + printf("Difference [ %d %d ]:%f ----%f\n", i, j, hData[i * width + j], + hOutputData[i * width + j]); + testResult = 0; + break; + } + } + } + hipDestroyTextureObject(textureObject); + hipFree(dData); + hipFreeArray(hipArray); + return testResult; +} diff --git a/projects/hip/tests/src/texture/hipTextureRef2D.cpp b/projects/hip/tests/src/texture/hipTextureRef2D.cpp index c4c0b9e2fe..b912f789a7 100644 --- a/projects/hip/tests/src/texture/hipTextureRef2D.cpp +++ b/projects/hip/tests/src/texture/hipTextureRef2D.cpp @@ -1,93 +1,93 @@ -/* HIT_START - * BUILD: %t %s ../test_common.cpp - * RUN: %t - * HIT_END - */ -#include -#include -#include - -#include -#include "test_common.h" -texture tex; - -bool testResult = true; - -__global__ void tex2DKernel(float* outputData, - int width, int height) { - int x = blockIdx.x * blockDim.x + threadIdx.x; - int y = blockIdx.y * blockDim.y + threadIdx.y; - outputData[y * width + x] = tex2D(tex, x, y); -} - -void runTest(int argc, char** argv); - -int main(int argc, char** argv) { - runTest(argc, argv); - if (testResult) { - passed(); - } else { - exit(EXIT_FAILURE); - } -} - -void runTest(int argc, char** argv) { - unsigned int width = 256; - unsigned int height = 256; - unsigned int size = width * height * sizeof(float); - float* hData = (float*)malloc(size); - memset(hData, 0, size); - for (int i = 0; i < height; i++) { - for (int j = 0; j < width; j++) { - hData[i * width + j] = i * width + j; - } - } - printf("hData: "); - for (int i = 0; i < 64; i++) { - printf("%f ", hData[i]); - } - printf("\n"); - - hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat); - hipArray* hipArray; - hipMallocArray(&hipArray, &channelDesc, width, height); - - hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice); - - tex.addressMode[0] = hipAddressModeWrap; - tex.addressMode[1] = hipAddressModeWrap; - tex.filterMode = hipFilterModePoint; - tex.normalized = 0; - - hipBindTextureToArray(tex, hipArray, channelDesc); - - float* dData = NULL; - hipMalloc((void**)&dData, size); - - dim3 dimBlock(16, 16, 1); - dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1); - hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, width, height); - hipDeviceSynchronize(); - - float* hOutputData = (float*)malloc(size); - memset(hOutputData, 0, size); - hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost); - - printf("dData: "); - for (int i = 0; i < 64; i++) { - printf("%f ", hOutputData[i]); - } - printf("\n"); - for (int i = 0; i < height; i++) { - for (int j = 0; j < width; j++) { - if (hData[i * width + j] != hOutputData[i * width + j]) { - printf("Difference [ %d %d ]:%f ----%f\n", i, j, hData[i * width + j], - hOutputData[i * width + j]); - testResult = false; - break; - } - } - } - hipFree(dData); - hipFreeArray(hipArray); -} +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * RUN: %t + * HIT_END + */ +#include +#include +#include + +#include +#include "test_common.h" +texture tex; + +__global__ void tex2DKernel(float* outputData, + int width, int height) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + outputData[y * width + x] = tex2D(tex, x, y); +} + +int runTest(int argc, char** argv); + +int main(int argc, char** argv) { + int testResult = runTest(argc, argv); + if (testResult) { + passed(); + } else { + exit(EXIT_FAILURE); + } +} + +int runTest(int argc, char** argv) { + int testResult = 1; + unsigned int width = 256; + unsigned int height = 256; + unsigned int size = width * height * sizeof(float); + float* hData = (float*)malloc(size); + memset(hData, 0, size); + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + hData[i * width + j] = i * width + j; + } + } + printf("hData: "); + for (int i = 0; i < 64; i++) { + printf("%f ", hData[i]); + } + printf("\n"); + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat); + hipArray* hipArray; + hipMallocArray(&hipArray, &channelDesc, width, height); + + hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice); + + tex.addressMode[0] = hipAddressModeWrap; + tex.addressMode[1] = hipAddressModeWrap; + tex.filterMode = hipFilterModePoint; + tex.normalized = 0; + + hipBindTextureToArray(tex, hipArray, channelDesc); + + float* dData = NULL; + hipMalloc((void**)&dData, size); + + dim3 dimBlock(16, 16, 1); + dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1); + hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, width, height); + hipDeviceSynchronize(); + + float* hOutputData = (float*)malloc(size); + memset(hOutputData, 0, size); + hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost); + + printf("dData: "); + for (int i = 0; i < 64; i++) { + printf("%f ", hOutputData[i]); + } + printf("\n"); + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + if (hData[i * width + j] != hOutputData[i * width + j]) { + printf("Difference [ %d %d ]:%f ----%f\n", i, j, hData[i * width + j], + hOutputData[i * width + j]); + testResult = 0; + break; + } + } + } + hipFree(dData); + hipFreeArray(hipArray); + return testResult; +} From ecf51641e3a03c17b5e2e71593411119457d2444 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 29 May 2018 16:18:14 +0530 Subject: [PATCH 13/14] Fix runTest return type [ROCm/hip commit: 7f7a7041a5b70b430023bab31ca42b0eeb5da5c7] --- projects/hip/tests/src/texture/hipTextureObj2D.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/src/texture/hipTextureObj2D.cpp b/projects/hip/tests/src/texture/hipTextureObj2D.cpp index 504b632612..e214295989 100644 --- a/projects/hip/tests/src/texture/hipTextureObj2D.cpp +++ b/projects/hip/tests/src/texture/hipTextureObj2D.cpp @@ -29,7 +29,7 @@ int main(int argc, char** argv) { } } -void runTest(int argc, char** argv) { +int runTest(int argc, char** argv) { int testResult = 1; unsigned int width = 256; unsigned int height = 256; From 6dc1a919ace20dddb7a847a9e1ef763c19fadda2 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 30 May 2018 10:59:07 +0530 Subject: [PATCH 14/14] Fix hipBindTexture on NVCC path [ROCm/hip commit: d67b614b22a3bb2d70ba67e76a48d918519b49bb] --- projects/hip/include/hip/nvcc_detail/hip_runtime_api.h | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index e00405cc71..d54a8de99c 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -1163,8 +1163,8 @@ inline static hipError_t hipBindTexture(size_t* offset, const struct texture -inline static hipError_t hipBindTexture(size_t* offset, struct texture* tex, - const void* devPtr, const struct hipChannelFormatDesc* desc, +inline static hipError_t hipBindTexture(size_t* offset, struct texture& tex, + const void* devPtr, const struct hipChannelFormatDesc& desc, size_t size = UINT_MAX) { return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size)); } @@ -1174,6 +1174,11 @@ inline static hipError_t hipUnbindTexture(struct texture* tex) return hipCUDAErrorTohipError(cudaUnbindTexture(tex)); } +inline static hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* devPtr, + const hipChannelFormatDesc* desc, size_t size = UINT_MAX){ + return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size)); +} + template inline static hipError_t hipBindTextureToArray(struct texture& tex, hipArray_const_t array,