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

SWDEV-405157 - Rewrite sample 11_texture_driver to use texture objects (#348)

Browse Source 
Change-Id: I107bfc06fabd62f43e6665b8b038226fe2154fc5
This commit is contained in:
ROCm CI Service Account
2023-07-08 20:55:09 +05:30
committed by GitHub
parent a461ae2fc9
commit c60a07656f
2 changed files with 57 additions and 54 deletions
Download Patch File Download Diff File Expand all files Collapse all files
samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp
+16 -26
View File
@@ -21,76 +21,66 @@ THE SOFTWARE.
*/
#include "hip/hip_runtime.h"
texture<char, hipTextureType2D, hipReadModeElementType> texChar;
texture<short, hipTextureType2D, hipReadModeElementType> texShort;
texture<int, hipTextureType2D, hipReadModeElementType> texInt;
texture<float, hipTextureType2D, hipReadModeElementType> texFloat;
texture<char4, hipTextureType2D, hipReadModeElementType> texChar4;
texture<short4, hipTextureType2D, hipReadModeElementType> texShort4;
texture<int4, hipTextureType2D, hipReadModeElementType> texInt4;
texture<float4, hipTextureType2D, hipReadModeElementType> texFloat4;
extern "C" __global__ void tex2dKernelChar(char* outputData, int width, int height) {
extern "C" __global__ void tex2dKernelChar(char* outputData,hipTextureObject_t texObj, 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);
outputData[y * width + x] = tex2D<char>(texObj, x, y);
#endif
}
extern "C" __global__ void tex2dKernelShort(short* outputData, int width, int height) {
extern "C" __global__ void tex2dKernelShort(short* outputData,hipTextureObject_t texObj, 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);
outputData[y * width + x] = tex2D<short>(texObj, x, y);
#endif
}
extern "C" __global__ void tex2dKernelInt(int* outputData, int width, int height) {
extern "C" __global__ void tex2dKernelInt(int* outputData,hipTextureObject_t texObj ,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);
outputData[y * width + x] = tex2D<int>(texObj, x, y);
#endif
}
extern "C" __global__ void tex2dKernelFloat(float* outputData, int width, int height) {
extern "C" __global__ void tex2dKernelFloat(float* outputData,hipTextureObject_t texObj, 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);
outputData[y * width + x] = tex2D<float>(texObj, x, y);
#endif
}
extern "C" __global__ void tex2dKernelChar4(char4* outputData, int width, int height) {
extern "C" __global__ void tex2dKernelChar4(char4* outputData,hipTextureObject_t texObj, 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);
outputData[y * width + x] = tex2D<char4>(texObj, x, y);
#endif
}
extern "C" __global__ void tex2dKernelShort4(short4* outputData, int width, int height) {
extern "C" __global__ void tex2dKernelShort4(short4* outputData,hipTextureObject_t texObj, 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);
outputData[y * width + x] = tex2D<short4>(texObj, x, y);
#endif
}
extern "C" __global__ void tex2dKernelInt4(int4* outputData, int width, int height) {
extern "C" __global__ void tex2dKernelInt4(int4* outputData,hipTextureObject_t texObj, 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);
outputData[y * width + x] = tex2D<int4>(texObj, x, y);
#endif
}
extern "C" __global__ void tex2dKernelFloat4(float4* outputData, int width, int height) {
extern "C" __global__ void tex2dKernelFloat4(float4* outputData,hipTextureObject_t texObj, 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);
outputData[y * width + x] = tex2D<float4>(texObj, x, y);
#endif
}