From 534376cb416b591d4d220ef0284e025d410acaac Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Sun, 5 Jan 2020 23:03:50 -0800 Subject: [PATCH] Add hipBindTexture2D on NVCC path (#1773) --- hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 6 ++++++ hipamd/tests/src/texture/hipBindTex2DPitch.cpp | 11 ----------- 2 files changed, 6 insertions(+), 11 deletions(-) diff --git a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 172bc20a96..373fd80879 100644 --- a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -1383,6 +1383,12 @@ inline static hipError_t hipBindTexture(size_t* offset, struct textureReference* return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size)); } +inline static hipError_t hipBindTexture2D(size_t* offset, struct textureReference* tex, const void* devPtr, + const hipChannelFormatDesc* desc, size_t width, size_t height, + size_t pitch) { + return hipCUDAErrorTohipError(cudaBindTexture2D(offset, tex, devPtr, desc, width, height, pitch)); +} + inline static hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) { return cudaCreateChannelDesc(x, y, z, w, hipChannelFormatKindToCudaChannelFormatKind(f)); diff --git a/hipamd/tests/src/texture/hipBindTex2DPitch.cpp b/hipamd/tests/src/texture/hipBindTex2DPitch.cpp index f463942329..905e24bce3 100644 --- a/hipamd/tests/src/texture/hipBindTex2DPitch.cpp +++ b/hipamd/tests/src/texture/hipBindTex2DPitch.cpp @@ -58,19 +58,8 @@ int main (void) SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyHostToDevice)); tex.normalized = false; - #if defined(__HIP_PLATFORM_NVCC__) - - cudaError_t status = cudaBindTexture2D(&tex_ofs, &tex, devPtrA, &tex.channelDesc, - SIZE_W, SIZE_H, devPitchA); - if (status != cudaSuccess) { - printf("%serror: '%s'(%d) at %s:%d%s\n", KRED, cudaGetErrorString(status), - status,__FILE__, __LINE__, KNRM); - failed("API returned error code."); - } - #else HIPCHECK(hipBindTexture2D(&tex_ofs, &tex, devPtrA, &tex.channelDesc, SIZE_W, SIZE_H, devPitchA)); - #endif HIPCHECK(hipMalloc((void**)&devPtrB, SIZE_W*sizeof(TYPE_t)*SIZE_H)) ; hipLaunchKernelGGL(texture2dCopyKernel, dim3(4,4,1), dim3(32,32,1), 0, 0, devPtrB);