AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-397570 - Move tex2d call under IMAGE_SUPPORT guards. (#294)

Browse Source 
Change-Id: Id7239f026dfb463c8c6d82b9656ac5b6716956fc
This commit is contained in:
ROCm CI Service Account
2023-06-19 22:10:06 +05:30
committed by GitHub
parent af4826551c
commit 3d29c352b9
1 changed files with 16 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp
+16
View File
@@ -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
}