diff --git a/docs/markdown/hip_porting_driver_api.md b/docs/markdown/hip_porting_driver_api.md index 0912e676cc..47cb7fb009 100644 --- a/docs/markdown/hip_porting_driver_api.md +++ b/docs/markdown/hip_porting_driver_api.md @@ -231,3 +231,45 @@ int main(){ return 0; } ``` + +## HIP Module and Texture Driver API + +HIP supports texture driver APIs however texture reference should be declared in host scope. Following code explains the use of texture reference for __HIP_PLATFORM_HCC__ platform. + +``` +// Code to generate code object + +#include "hip/hip_runtime.h" +extern texture tex; + +__global__ void tex2dKernel(hipLaunchParm lp, float* outputData, + int width, + int height) +{ + int x = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x; + int y = hipBlockIdx_y*hipBlockDim_y + hipThreadIdx_y; + outputData[y*width + x] = tex2D(tex, x, y); +} + +``` +``` +// Host code: + +texture tex; + +void myFunc () +{ + // ... + + textureReference* texref; + hipModuleGetTexRef(&texref, Module1, "tex"); + hipTexRefSetAddressMode(texref, 0, hipAddressModeWrap); + hipTexRefSetAddressMode(texref, 1, hipAddressModeWrap); + hipTexRefSetFilterMode(texref, hipFilterModePoint); + hipTexRefSetFlags(texref, 0); + hipTexRefSetFormat(texref, HIP_AD_FORMAT_FLOAT, 1); + hipTexRefSetArray(texref, array, HIP_TRSA_OVERRIDE_FORMAT); + + // ... +} +``` \ No newline at end of file diff --git a/docs/markdown/hip_porting_guide.md b/docs/markdown/hip_porting_guide.md index 12ec931f2a..aeb7b171d6 100644 --- a/docs/markdown/hip_porting_guide.md +++ b/docs/markdown/hip_porting_guide.md @@ -465,34 +465,36 @@ a performance impact. ### Textures and Cache Control ->Texture support is under-development and not yet supported by HIP. - Compute programs sometimes use textures either to access dedicated texture caches or to use the texture-sampling hardware for interpolation and clamping. The former approach uses simple point samplers with linear interpolation, essentially only reading a single point. The latter approach uses the sampler hardware to interpolate and combine multiple point samples. AMD hardware, as well as recent competing hardware, has a unified texture/L1 cache, so it no longer has a dedicated texture cache. But the nvcc path often caches global loads in the L2 cache, and some programs may benefit from explicit control of the L1 cache contents. We recommend the __ldg instruction for this purpose. -HIP currently lacks texture support; a future revision will add this capability. Also, AMD compilers currently load all data into both the L1 and L2 caches, so __ldg is treated as a no-op. +AMD compilers currently load all data into both the L1 and L2 caches, so __ldg is treated as a no-op. We recommend the following for functional portability: - For programs that use textures only to benefit from improved caching, use the __ldg instruction -- Alternatively, use conditional compilation (see [Identify HIP Target Platform](#identify-hip-target-platform)) - - For the `__HIP_PLATFORM_NVCC__` path, use the full texture path - - For the `__HIP_PLATFORM_HCC__` path, pass an additional pointer to the kernel and reference it using regular device memory-load instructions rather than texture loads. Some applications may already take this step, since it allows experimentation with caching behavior. +- 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 t_features; +texture tex; -void __global__ MyKernel(float *d_features /* pass pointer parameter, if not already available */...) -{ - // ... - -#ifdef __HIP_PLATFORM_NVCC__ - float tval = tex1Dfetch(t_features,addr); -#else - float tval = d_features[addr]; +__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: @@ -500,23 +502,15 @@ void myFunc () { // ... -#ifdef __HIP_PLATFORM_NVCC__ - cudaChannelFormatDesc chDesc0 = cudaCreateChannelDesc(); - t_features.filterMode = cudaFilterModePoint; - t_features.normalized = false; - t_features.channelDesc = chDesc0; - - cudaBindTexture(NULL, &t_features, d_features, &chDesc0, npoints*nfeatures*sizeof(float)); +#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 + ``` -Additionally, many of the Rodinia benchmarks demonstrate how to modify hipified programs so that textures are not required - search for USE_TEXTURES define in the rodinia source directory. -For example, [here - - -Cuda programs that employ sampler hardware must either wait for hcc texture support or use more-sophisticated workarounds. - ## More Tips ### HIPTRACE Mode diff --git a/tests/src/texture/hipTextureRef2D.cpp b/tests/src/texture/hipTextureRef2D.cpp index eb27b23230..ebc7a04385 100644 --- a/tests/src/texture/hipTextureRef2D.cpp +++ b/tests/src/texture/hipTextureRef2D.cpp @@ -14,7 +14,9 @@ texture tex; bool testResult = true; __global__ void tex2DKernel(float* outputData, +#ifdef __HIP_PLATFORM_HCC__ hipTextureObject_t textureObject, +#endif int width, int height) { @@ -78,7 +80,7 @@ void runTest(int argc, char **argv) #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, 0, width, height); + hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, width, height); #endif hipDeviceSynchronize();