From 2c3d1498d4e3b03d5a7a5aa763855dcd755ea738 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 --- hipamd/src/hip_memory.cpp | 16 ++++++++++++---- 1 file changed, 12 insertions(+), 4 deletions(-) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index fc60dffe86..b76f7f0524 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/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 96b4618d26acfb7e21537ed8b48a77feb0305209 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 --- hipamd/src/hip_memory.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index b76f7f0524..a1bc1416a0 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/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 53a7c61e9b93d74067f363de5959b82712cee5b9 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 --- hipamd/src/hip_memory.cpp | 26 ++++++++++++++------------ 1 file changed, 14 insertions(+), 12 deletions(-) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index a1bc1416a0..2f0c793033 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/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 aed26538575d4a501134a17ba59002be58a5b442 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 --- hipamd/src/hip_memory.cpp | 22 +++++++++------------- 1 file changed, 9 insertions(+), 13 deletions(-) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 2f0c793033..fa68d7cfb1 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/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 27e0566c3ac9864ead727d8311dcbc5bd585cbdd 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 --- hipamd/src/hip_memory.cpp | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index fa68d7cfb1..9b6758ddf4 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/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 13f37d550f7c238aa5b9672d67ad6fc4c61e27f2 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. --- hipamd/src/math_functions.cpp | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/hipamd/src/math_functions.cpp b/hipamd/src/math_functions.cpp index dedc40f2ae..bd0116080d 100644 --- a/hipamd/src/math_functions.cpp +++ b/hipamd/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 b127488ea9a3891ecffffa6dc6b176c4aef956c7 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. --- hipamd/include/hip/hcc_detail/math_functions.h | 2 ++ 1 file changed, 2 insertions(+) diff --git a/hipamd/include/hip/hcc_detail/math_functions.h b/hipamd/include/hip/hcc_detail/math_functions.h index e717df07c1..e7312f32f2 100644 --- a/hipamd/include/hip/hcc_detail/math_functions.h +++ b/hipamd/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 61ff40a1cf416361eda2e432d43efd4f9c72ea74 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. --- hipamd/src/math_functions.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/hipamd/src/math_functions.cpp b/hipamd/src/math_functions.cpp index bd0116080d..982073275b 100644 --- a/hipamd/src/math_functions.cpp +++ b/hipamd/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 b33876ee4d7ee4c135c3217936c9288801819ce0 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. --- hipamd/include/hip/hcc_detail/math_functions.h | 1 + 1 file changed, 1 insertion(+) diff --git a/hipamd/include/hip/hcc_detail/math_functions.h b/hipamd/include/hip/hcc_detail/math_functions.h index e7312f32f2..ea3bd97571 100644 --- a/hipamd/include/hip/hcc_detail/math_functions.h +++ b/hipamd/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 3699e74857b631cc510de9f7e7fc464a6eead951 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 --- hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 2 ++ 1 file changed, 2 insertions(+) diff --git a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index bd3ffc1cc1..e00405cc71 100644 --- a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/hipamd/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 366a0a97ca2b21f222ddb917e41702ef5428d5c7 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) --- hipamd/tests/src/texture/hipBindTexObj1D.cpp | 100 ++++++++++++++++ .../tests/src/texture/tex1Dfetch_linear.cpp | 112 ++++++++++++++++++ 2 files changed, 212 insertions(+) create mode 100644 hipamd/tests/src/texture/hipBindTexObj1D.cpp create mode 100644 hipamd/tests/src/texture/tex1Dfetch_linear.cpp diff --git a/hipamd/tests/src/texture/hipBindTexObj1D.cpp b/hipamd/tests/src/texture/hipBindTexObj1D.cpp new file mode 100644 index 0000000000..2bcf78b1ca --- /dev/null +++ b/hipamd/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/hipamd/tests/src/texture/tex1Dfetch_linear.cpp b/hipamd/tests/src/texture/tex1Dfetch_linear.cpp new file mode 100644 index 0000000000..4f755873f7 --- /dev/null +++ b/hipamd/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 28c878f266dbf04c9785e55cf50e033c60dda35a 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 --- hipamd/tests/src/texture/hipTextureObj2D.cpp | 216 +++++++++---------- hipamd/tests/src/texture/hipTextureRef2D.cpp | 186 ++++++++-------- 2 files changed, 201 insertions(+), 201 deletions(-) diff --git a/hipamd/tests/src/texture/hipTextureObj2D.cpp b/hipamd/tests/src/texture/hipTextureObj2D.cpp index 1bf51bc2cb..504b632612 100644 --- a/hipamd/tests/src/texture/hipTextureObj2D.cpp +++ b/hipamd/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/hipamd/tests/src/texture/hipTextureRef2D.cpp b/hipamd/tests/src/texture/hipTextureRef2D.cpp index c4c0b9e2fe..b912f789a7 100644 --- a/hipamd/tests/src/texture/hipTextureRef2D.cpp +++ b/hipamd/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 dabe2fda39cd815913fc79624cdf5f5a903da062 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 --- hipamd/tests/src/texture/hipTextureObj2D.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/tests/src/texture/hipTextureObj2D.cpp b/hipamd/tests/src/texture/hipTextureObj2D.cpp index 504b632612..e214295989 100644 --- a/hipamd/tests/src/texture/hipTextureObj2D.cpp +++ b/hipamd/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 d74353b1128922e1268dc1dddc3af79df88a47d5 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 --- hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index e00405cc71..d54a8de99c 100644 --- a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/hipamd/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,