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

[ROCm/hip-tests commit: c60a07656f]
This commit is contained in:
ROCm CI Service Account
2023-07-08 20:55:09 +05:30
committed by GitHub
parent a1346c114a
commit dac192e64b
2 changed files with 57 additions and 54 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/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
}
projects/hip-tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp
+41 -28
View File
@@ -65,6 +65,19 @@ static inline constexpr int rank() {
return sizeof(T) / sizeof(decltype(T::x));
}
#ifdef __HIP_PLATFORM_NVIDIA__
template <typename T,
typename std::enable_if<std::is_same<T, int4>::value ||
std::is_same<T, short4>::value ||
std::is_same<T, char4>::value ||
std::is_same<T, float4>::value>::type *t = nullptr>
static inline bool operator!=(const T& a, const T& b)
{
return (a.x != b.x) || (a.y != b.y) || (a.z != b.z) || (a.w != b.w);
}
#endif
template<typename T>
static inline T getRandom() {
double r = 0;
@@ -139,43 +152,42 @@ bool runTest(hipModule_t &module, const char *refName, const char *funcName) {
}
}
hipArray *array;
HIP_ARRAY_DESCRIPTOR desc;
desc.Format = format;
desc.NumChannels = channels;
desc.Width = width;
desc.Height = height;
HIP_CHECK(hipArrayCreate(&array, &desc));
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
hipArray_t array;
HIP_CHECK(hipMallocArray(&array, &channelDesc, width, height));
hip_Memcpy2D copyParam;
memset(&copyParam, 0, sizeof(copyParam));
copyParam.dstMemoryType = hipMemoryTypeArray;
copyParam.dstArray = array;
copyParam.srcMemoryType = hipMemoryTypeHost;
copyParam.srcHost = hData;
copyParam.srcPitch = width * sizeof(T);
copyParam.WidthInBytes = copyParam.srcPitch;
copyParam.Height = height;
HIP_CHECK(hipMemcpyParam2D(&copyParam));
const size_t spitch = width * sizeof(T);
textureReference *texref;
HIP_CHECK(hipModuleGetTexRef(&texref, module, refName));
HIP_CHECK(hipTexRefSetAddressMode(texref, 0, hipAddressModeClamp));
HIP_CHECK(hipTexRefSetAddressMode(texref, 1, hipAddressModeClamp));
HIP_CHECK(hipTexRefSetFilterMode(texref, hipFilterModePoint));
HIP_CHECK(hipTexRefSetFlags(texref, HIP_TRSF_READ_AS_INTEGER));
HIP_CHECK(hipTexRefSetFormat(texref, format, channels));
HIP_CHECK(hipTexRefSetArray(texref, array, HIP_TRSA_OVERRIDE_FORMAT));
HIP_CHECK(hipMemcpy2DToArray(array, 0, 0, hData, spitch, width * sizeof(T),
height, hipMemcpyHostToDevice));
hipResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = hipResourceTypeArray;
resDesc.res.array.array = array;
hipTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.addressMode[0] = hipAddressModeClamp;
texDesc.addressMode[1] = hipAddressModeClamp;
texDesc.filterMode = hipFilterModePoint;
texDesc.readMode = hipReadModeElementType;
texDesc.normalizedCoords = 0;
hipTextureObject_t texObj;
HIP_CHECK(hipCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr));
T *dData = NULL;
HIP_CHECK(hipMalloc((void** )&dData, size));
struct {
void *_Ad;
hipTextureObject_t _texObj;
unsigned int _Bd;
unsigned int _Cd;
} args;
args._Ad = (void*) dData;
args._texObj = texObj;
args._Bd = width;
args._Cd = height;
@@ -192,7 +204,7 @@ bool runTest(hipModule_t &module, const char *refName, const char *funcName) {
HIP_CHECK(
hipModuleLaunchKernel(Function, 16, 16, 1, temp1, temp2, 1, 0, 0, NULL,
(void** )&config));
hipDeviceSynchronize();
HIP_CHECK(hipDeviceSynchronize());
T *hOutputData = (T*) malloc(size);
memset(hOutputData, 0, size);
@@ -207,7 +219,7 @@ bool runTest(hipModule_t &module, const char *refName, const char *funcName) {
}
}
}
HIP_CHECK(hipUnbindTexture(texref));
HIP_CHECK(hipDestroyTextureObject(texObj));
HIP_CHECK(hipFree(dData));
HIP_CHECK(hipFreeArray(array));
free(hOutputData);
@@ -230,7 +242,8 @@ int main(int argc, char** argv) {
printf("Texture is not support on the device. Skipped.\n");
return 0;
}
hipInit(0);
HIP_CHECK(hipInit(0));
HIP_CHECK(hipSetDevice(0));
hipModule_t module;
HIP_CHECK(hipModuleLoad(&module, fileName));
testResult = testResult && runTest<char>(module, "texChar", "tex2dKernelChar");