From 5ae6d93526f600d8b6cfb627312136e10eabb57d Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 30 May 2018 21:54:51 +0530 Subject: [PATCH] Fixed texture obj 1Dfetch test --- tests/src/texture/hipTextureObj1DFetch.cpp | 99 ++++++++++++++++++ tests/src/texture/tex1Dfetch_linear.cpp | 112 --------------------- 2 files changed, 99 insertions(+), 112 deletions(-) create mode 100644 tests/src/texture/hipTextureObj1DFetch.cpp delete mode 100644 tests/src/texture/tex1Dfetch_linear.cpp diff --git a/tests/src/texture/hipTextureObj1DFetch.cpp b/tests/src/texture/hipTextureObj1DFetch.cpp new file mode 100644 index 0000000000..faf6541ce1 --- /dev/null +++ b/tests/src/texture/hipTextureObj1DFetch.cpp @@ -0,0 +1,99 @@ +/* +Copyright (c) 2015 - present 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 + * RUN: %t + * HIT_END + */ + +#include "hip/hip_runtime.h" +#include "test_common.h" + +#define N 512 + +__global__ void tex1dKernel(float *val, hipTextureObject_t obj) { + int k = blockIdx.x * blockDim.x + threadIdx.x; + if (k < N) + val[k] = tex1Dfetch(obj, k); +} + +int runTest(void); + +int main(int argc, char **argv) { + int testResult = runTest(); + if(testResult) { + passed(); + } else { + exit(EXIT_FAILURE); + } +} + +int runTest() { + int testResult = 1; + // Allocating the required buffer on gpu device + float *texBuf, *texBufOut; + float val[N], output[N]; + for (int i = 0; i < N; i++) { + val[i] = (i + 1) * (i + 1); + output[i] = 0.0; + } + HIPCHECK(hipMalloc(&texBuf, N * sizeof(float))); + HIPCHECK(hipMalloc(&texBufOut, N * sizeof(float))); + HIPCHECK(hipMemcpy(texBuf, val, N * sizeof(float), hipMemcpyHostToDevice)); + HIPCHECK(hipMemset(texBufOut, 0, N * sizeof(float))); + hipResourceDesc resDescLinear; + + memset(&resDescLinear, 0, sizeof(resDescLinear)); + resDescLinear.resType = hipResourceTypeLinear; + resDescLinear.res.linear.devPtr = texBuf; + resDescLinear.res.linear.desc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat); + resDescLinear.res.linear.sizeInBytes = N * sizeof(float); + + hipTextureDesc texDesc; + memset(&texDesc, 0, sizeof(texDesc)); + texDesc.readMode = hipReadModeElementType; + + // Creating texture object + hipTextureObject_t texObj = 0; + HIPCHECK(hipCreateTextureObject(&texObj, &resDescLinear, &texDesc, NULL)); + + dim3 dimBlock(64, 1, 1); + dim3 dimGrid(N / dimBlock.x, 1, 1); + + hipLaunchKernelGGL(tex1dKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, + texBufOut, texObj); + HIPCHECK(hipDeviceSynchronize()); + + HIPCHECK(hipMemcpy(output, texBufOut, N * sizeof(float), hipMemcpyDeviceToHost)); + + for(int i = 0; i < N; i++) + if (output[i] != val[i]) { + testResult = 0; + break; + } + + HIPCHECK(hipDestroyTextureObject(texObj)); + HIPCHECK(hipFree(texBuf)); + HIPCHECK(hipFree(texBufOut)); + return testResult; +} diff --git a/tests/src/texture/tex1Dfetch_linear.cpp b/tests/src/texture/tex1Dfetch_linear.cpp deleted file mode 100644 index 4f755873f7..0000000000 --- a/tests/src/texture/tex1Dfetch_linear.cpp +++ /dev/null @@ -1,112 +0,0 @@ -/* -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); -}