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

SWDEV-303204 - Disable texture/image Apis on some devices. (#2381)

Browse Source 
Modify texture tests to pass on devices that don't support texture.
Add hipDeviceAttributeImageSupport to check image support in runtime

Change-Id: Ia89c494e651a6198a24448b59a91e046a9ebea38

[ROCm/hip commit: 5c8af57480]
This commit is contained in:
TomSang
2021-12-02 01:16:13 -05:00
committed by GitHub
parent e90cc323e3
commit b87ef5aa18
10 changed files with 77 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp
+7
View File
@@ -592,6 +592,13 @@ bool testTexSingleStreamMultGPU(unsigned int numOfGPUs,
int main(int argc, char** argv) {
HipTest::parseStandardArguments(argc, argv, true);
int imageSupport = 0;
hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport,
p_gpuDevice);
if (!imageSupport) {
printf("Texture is not support on the device\n");
passed();
}
bool TestPassed = true;
if (p_tests == 0x01) {
TestPassed = testTexType<float>(HIP_AD_FORMAT_FLOAT,
projects/hip/tests/src/runtimeApi/module/tex2d_kernel.cpp
+8
View File
@@ -30,36 +30,44 @@ __device__ float deviceGlobalFloat;
extern "C" __global__ void tex2dKernelFloat(float* outputData,
int width, int height) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
if ((x < width) && (y < width)) {
outputData[y * width + x] = tex2D(ftex, 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 = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
if ((x < width) && (y < width)) {
outputData[y * width + x] = tex2D(itex, x, y);
}
#endif
}
extern "C" __global__ void tex2dKernelInt16(short* outputData,
int width, int height) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
if ((x < width) && (y < width)) {
outputData[y * width + x] = tex2D(stex, x, y);
}
#endif
}
extern "C" __global__ void tex2dKernelInt8(char* outputData,
int width, int height) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
if ((x < width) && (y < width)) {
outputData[y * width + x] = tex2D(ctex, x, y);
}
#endif
}