diff --git a/projects/hip/include/hip/hcc_detail/math_functions.h b/projects/hip/include/hip/hcc_detail/math_functions.h index e717df07c1..ea3bd97571 100644 --- a/projects/hip/include/hip/hcc_detail/math_functions.h +++ b/projects/hip/include/hip/hcc_detail/math_functions.h @@ -58,6 +58,9 @@ __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__ 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); 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..d54a8de99c 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 @@ -1161,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)); } @@ -1172,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, diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index fc60dffe86..9b6758ddf4 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -1495,10 +1495,16 @@ __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)){ - T *dstPtr = (T *)((uint8_t*) dst + idy * destPitch); - T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch); + size_t floorWidth = (width/sizeof(T)); + 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 @@ -1515,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-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, + 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); } @@ -1622,7 +1629,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 +1668,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; diff --git a/projects/hip/src/math_functions.cpp b/projects/hip/src/math_functions.cpp index dedc40f2ae..982073275b 100644 --- a/projects/hip/src/math_functions.cpp +++ b/projects/hip/src/math_functions.cpp @@ -59,6 +59,15 @@ __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__ 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; } 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/hipTextureObj2D.cpp b/projects/hip/tests/src/texture/hipTextureObj2D.cpp index 1bf51bc2cb..e214295989 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); + } +} + +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); + + 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; +} 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); +}