From 3d29c352b93b59e9ca6565af599eebd5c686deac Mon Sep 17 00:00:00 2001
From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com>
Date: Mon, 19 Jun 2023 22:10:06 +0530
Subject: [PATCH] SWDEV-397570 - Move tex2d call under IMAGE_SUPPORT guards.
 (#294)

Change-Id: Id7239f026dfb463c8c6d82b9656ac5b6716956fc
---
 .../2_Cookbook/11_texture_driver/tex2dKernel.cpp | 16 ++++++++++++++++
 1 file changed, 16 insertions(+)

diff --git a/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp b/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp
index a1d3985de5..d20b1d0864 100644
--- a/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp
+++ b/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp
@@ -32,49 +32,65 @@ texture<int4, hipTextureType2D, hipReadModeElementType> texInt4;
 texture<float4, hipTextureType2D, hipReadModeElementType> texFloat4;
 
 extern "C" __global__ void tex2dKernelChar(char* outputData, int width, int height) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
     int x = blockIdx.x * blockDim.x + threadIdx.x;
     int y = blockIdx.y * blockDim.y + threadIdx.y;
     outputData[y * width + x] = tex2D(texChar, x, y);
+#endif
 }
 
 extern "C" __global__ void tex2dKernelShort(short* outputData, int width, int height) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
     int x = blockIdx.x * blockDim.x + threadIdx.x;
     int y = blockIdx.y * blockDim.y + threadIdx.y;
     outputData[y * width + x] = tex2D(texShort, x, y);
+#endif
 }
 
 extern "C" __global__ void tex2dKernelInt(int* outputData, int width, int height) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
     int x = blockIdx.x * blockDim.x + threadIdx.x;
     int y = blockIdx.y * blockDim.y + threadIdx.y;
     outputData[y * width + x] = tex2D(texInt, x, y);
+#endif
 }
 
 extern "C" __global__ void tex2dKernelFloat(float* outputData, int width, int height) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
     int x = blockIdx.x * blockDim.x + threadIdx.x;
     int y = blockIdx.y * blockDim.y + threadIdx.y;
     outputData[y * width + x] = tex2D(texFloat, x, y);
+#endif
 }
 
 extern "C" __global__ void tex2dKernelChar4(char4* outputData, int width, int height) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
     int x = blockIdx.x * blockDim.x + threadIdx.x;
     int y = blockIdx.y * blockDim.y + threadIdx.y;
     outputData[y * width + x] = tex2D(texChar4, x, y);
+#endif
 }
 
 extern "C" __global__ void tex2dKernelShort4(short4* outputData, int width, int height) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
     int x = blockIdx.x * blockDim.x + threadIdx.x;
     int y = blockIdx.y * blockDim.y + threadIdx.y;
     outputData[y * width + x] = tex2D(texShort4, x, y);
+#endif
 }
 
 extern "C" __global__ void tex2dKernelInt4(int4* outputData, int width, int height) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
     int x = blockIdx.x * blockDim.x + threadIdx.x;
     int y = blockIdx.y * blockDim.y + threadIdx.y;
     outputData[y * width + x] = tex2D(texInt4, x, y);
+#endif
 }
 
 extern "C" __global__ void tex2dKernelFloat4(float4* outputData, int width, int height) {
+#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
     int x = blockIdx.x * blockDim.x + threadIdx.x;
     int y = blockIdx.y * blockDim.y + threadIdx.y;
     outputData[y * width + x] = tex2D(texFloat4, x, y);
+#endif
 }