From ccaea193b22d1b12c7d306c32303891d872e171a Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 30 Jul 2019 02:56:47 +0530 Subject: [PATCH] [docs]Fix texture reference APIs usage part --- docs/markdown/hip_porting_guide.md | 34 +----------------------------- 1 file changed, 1 insertion(+), 33 deletions(-) diff --git a/docs/markdown/hip_porting_guide.md b/docs/markdown/hip_porting_guide.md index 90854c3a46..6915945d4e 100644 --- a/docs/markdown/hip_porting_guide.md +++ b/docs/markdown/hip_porting_guide.md @@ -485,39 +485,7 @@ AMD compilers currently load all data into both the L1 and L2 caches, so __ldg i We recommend the following for functional portability: - For programs that use textures only to benefit from improved caching, use the __ldg instruction -- Programs that use texture object APIs, work well on HIP -- For program that use texture reference APIs, use conditional compilation (see [Identify HIP Target Platform](#identify-hip-target-platform)) - - For the `__HIP_PLATFORM_HCC__` path, pass an additional argument to the kernel and in texture fetch API inside kernel as shown below:- - -``` -texture tex; - -__global__ void tex2DKernel(float* outputData, -#ifdef __HIP_PLATFORM_HCC__ - hipTextureObject_t textureObject, -#endif - int width, - int height) -{ - int x = blockIdx.x*blockDim.x + threadIdx.x; - int y = blockIdx.y*blockDim.y + threadIdx.y; -#ifdef __HIP_PLATFORM_HCC__ - outputData[y*width + x] = tex2D(tex, textureObject, x, y); -#else - outputData[y*width + x] = tex2D(tex, x, y); -#endif -} - -// Host code: -void myFunc () -{ - // ... - -#ifdef __HIP_PLATFORM_HCC__ - hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, tex.textureObject, width, height); -#else - hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, width, height); -#endif +- Programs that use texture object and reference APIs, work well on HIP ```