AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

SWDEV-432490 - fix __HIP_NO_IMAGE_SUPPORT logic

Просмотр исходного кода 
1.Skip texture/image tests when __HIP_NO_IMAGE_SUPPORT is set to prevent inevitable silent failure
2.Simplify __HIP_NO_IMAGE_SUPPORT macros

Change-Id: I54ef12a239298c534a213edf39c05c2dad06a7f4


[ROCm/hip-tests commit: 0ee0e4ab4a]
Этот коммит содержится в:
Todd tiantuo Li
2024-01-26 03:44:44 -08:00
коммит произвёл Rakesh Roy
родитель 7df8798269
Коммит 14f6363879
26 изменённых файлов: 305 добавлений и 39 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
projects/hip-tests/catch/unit/memory/hipArrayCommon.hh
+2 -2
Просмотреть файл
@@ -33,7 +33,7 @@ constexpr size_t ChannelToRead = 1;
template <typename T>
__global__ void readFromTexture(T* output, hipTextureObject_t texObj, size_t width, size_t height,
bool textureGather) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
// Calculate normalized texture coordinates
const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
@@ -50,7 +50,7 @@ __global__ void readFromTexture(T* output, hipTextureObject_t texObj, size_t wid
output[y * width + x] = tex2D<T>(texObj, u, v);
}
}
#endif
#endif
}
template <typename T> void checkDataIsAscending(const std::vector<T>& hostData) {
projects/hip-tests/catch/unit/memory/hipArrayCreate.cc
+41 -1
Просмотреть файл
@@ -75,6 +75,11 @@ static void ArrayCreate_DiffSizes(int gpu) {
TEST_CASE("Unit_hipArrayCreate_DiffSizes") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
ArrayCreate_DiffSizes(0);
HIP_CHECK_THREAD_FINALIZE();
}
@@ -87,6 +92,11 @@ and verifies the hipArrayCreate API with small and big chunks data
TEST_CASE("Unit_hipArrayCreate_MultiThread") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
std::vector<std::thread> threadlist;
int devCnt = 0;
@@ -188,7 +198,7 @@ void testArrayAsTexture(hipArray_t array, const size_t width, const size_t heigh
std::fill(std::begin(hostData), std::end(hostData), 0);
HIP_CHECK(hipMemcpy(hostData.data(), device_data, size, hipMemcpyDeviceToHost));
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
checkDataIsAscending(hostData);
#endif
@@ -203,6 +213,11 @@ TEMPLATE_TEST_CASE("Unit_hipArrayCreate_happy", "", uint, int, int4, ushort, sho
char4, float, float2, float4) {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
using vec_info = vector_info<TestType>;
DriverContext ctx;
@@ -228,6 +243,11 @@ TEMPLATE_TEST_CASE("Unit_hipArrayCreate_maxTexture", "", uint, int, int4, ushort
uchar2, char4, float, float2, float4) {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
using vec_info = vector_info<TestType>;
DriverContext ctx;
@@ -288,6 +308,11 @@ TEMPLATE_TEST_CASE("Unit_hipArrayCreate_maxTexture", "", uint, int, int4, ushort
TEST_CASE("Unit_hipArrayCreate_ZeroWidth") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
DriverContext ctx;
HIP_ARRAY_DESCRIPTOR desc;
desc.Format = driverFormats[0];
@@ -304,6 +329,11 @@ TEST_CASE("Unit_hipArrayCreate_ZeroWidth") {
TEST_CASE("Unit_hipArrayCreate_Nullptr") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
DriverContext ctx;
SECTION("Null array") {
HIP_ARRAY_DESCRIPTOR desc;
@@ -324,6 +354,11 @@ TEST_CASE("Unit_hipArrayCreate_Nullptr") {
TEST_CASE("Unit_hipArrayCreate_BadNumberChannelElement") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
DriverContext ctx;
HIP_ARRAY_DESCRIPTOR desc;
desc.Format = GENERATE(from_range(std::begin(driverFormats), std::end(driverFormats)));
@@ -342,6 +377,11 @@ TEST_CASE("Unit_hipArrayCreate_BadNumberChannelElement") {
TEST_CASE("Unit_hipArrayCreate_BadChannelFormat") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
DriverContext ctx;
HIP_ARRAY_DESCRIPTOR desc;
projects/hip-tests/catch/unit/surface/surf1D.cc
+18 -3
Просмотреть файл
@@ -35,7 +35,7 @@ THE SOFTWARE.
template <typename T>
__global__ void surf1DKernelR(hipSurfaceObject_t surfaceObject, T* outputData, int width) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
if (x < width) {
surf1Dread(outputData + x, surfaceObject, x * sizeof(T));
@@ -45,7 +45,7 @@ __global__ void surf1DKernelR(hipSurfaceObject_t surfaceObject, T* outputData, i
template <typename T>
__global__ void surf1DKernelW(hipSurfaceObject_t surfaceObject, T* inputData, int width) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
if (x < width) {
surf1Dwrite(inputData[x], surfaceObject, x * sizeof(T));
@@ -56,7 +56,7 @@ __global__ void surf1DKernelW(hipSurfaceObject_t surfaceObject, T* inputData, in
template <typename T>
__global__ void surf1DKernelRW(hipSurfaceObject_t surfaceObject, hipSurfaceObject_t outputSurfObj,
int width) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
if (x < width) {
T data;
@@ -247,6 +247,11 @@ TEMPLATE_TEST_CASE("Unit_surf1Dread_Positive_Basic", "", char, uchar, short, ush
uint4, float4) {
CHECK_IMAGE_SUPPORT;
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
const int width = GENERATE(31, 67, 131, 263);
runTestR<TestType>(width);
}
@@ -268,6 +273,11 @@ TEMPLATE_TEST_CASE("Unit_surf1Dwrite_Positive_Basic", "", char, uchar, short, us
uint4, float4) {
CHECK_IMAGE_SUPPORT;
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
const int width = GENERATE(31, 67, 131, 263);
runTestW<TestType>(width);
}
@@ -289,6 +299,11 @@ TEMPLATE_TEST_CASE("Unit_surf1D_Positive_ReadWrite", "", char, uchar, short, ush
uint4, float4) {
CHECK_IMAGE_SUPPORT;
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
const int width = GENERATE(31, 67, 131, 263);
runTestRW<TestType>(width);
}
projects/hip-tests/catch/unit/surface/surf2D.cc
+18 -3
Просмотреть файл
@@ -38,7 +38,7 @@ THE SOFTWARE.
template <typename T>
__global__ void surf2DKernelR(hipSurfaceObject_t surfaceObject, T* outputData, int width,
int height) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x < width && y < height) {
@@ -50,7 +50,7 @@ __global__ void surf2DKernelR(hipSurfaceObject_t surfaceObject, T* outputData, i
template <typename T>
__global__ void surf2DKernelW(hipSurfaceObject_t surfaceObject, T* inputData, int width,
int height) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x < width && y < height) {
@@ -62,7 +62,7 @@ __global__ void surf2DKernelW(hipSurfaceObject_t surfaceObject, T* inputData, in
template <typename T>
__global__ void surf2DKernelRW(hipSurfaceObject_t surfaceObject, hipSurfaceObject_t outputSurfObj,
int width, int height) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x < width && y < height) {
@@ -288,6 +288,11 @@ TEMPLATE_TEST_CASE("Unit_surf2Dread_Positive_Basic", "", char, uchar, short, ush
uint4, float4) {
CHECK_IMAGE_SUPPORT;
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
const int width = GENERATE(31, 67);
const int height = GENERATE(131, 263);
runTestR<TestType>(width, height);
@@ -310,6 +315,11 @@ TEMPLATE_TEST_CASE("Unit_surf2Dwrite_Positive_Basic", "", char, uchar, short, us
uint4, float4) {
CHECK_IMAGE_SUPPORT;
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
const int width = GENERATE(31, 67);
const int height = GENERATE(131, 263);
runTestW<TestType>(width, height);
@@ -332,6 +342,11 @@ TEMPLATE_TEST_CASE("Unit_surf2D_Positive_ReadWrite", "", char, uchar, short, ush
uint4, float4) {
CHECK_IMAGE_SUPPORT;
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
const int width = GENERATE(31, 67);
const int height = GENERATE(131, 263);
runTestRW<TestType>(width, height);
projects/hip-tests/catch/unit/surface/surf3D.cc
+18 -3
Просмотреть файл
@@ -36,7 +36,7 @@ THE SOFTWARE.
template <typename T>
__global__ void surf3DKernelR(hipSurfaceObject_t surfaceObject, T* outputData, int width,
int height, int depth) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
@@ -49,7 +49,7 @@ __global__ void surf3DKernelR(hipSurfaceObject_t surfaceObject, T* outputData, i
template <typename T>
__global__ void surf3DKernelW(hipSurfaceObject_t surfaceObject, T* inputData, int width, int height,
int depth) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
@@ -62,7 +62,7 @@ __global__ void surf3DKernelW(hipSurfaceObject_t surfaceObject, T* inputData, in
template <typename T>
__global__ void surf3DKernelRW(hipSurfaceObject_t surfaceObject, hipSurfaceObject_t outputSurfObj,
int width, int height, int depth) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
@@ -334,6 +334,11 @@ TEMPLATE_TEST_CASE("Unit_surf3Dread_Positive_Basic", "", char, uchar, short, ush
uint4, float4) {
CHECK_IMAGE_SUPPORT;
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
const int width = GENERATE(31, 67);
const int height = GENERATE(131, 263);
const int depth = GENERATE(4, 11);
@@ -357,6 +362,11 @@ TEMPLATE_TEST_CASE("Unit_surf3Dwrite_Positive_Basic", "", char, uchar, short, us
uint4, float4) {
CHECK_IMAGE_SUPPORT;
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
const int width = GENERATE(31, 67);
const int height = GENERATE(131, 263);
const int depth = GENERATE(4, 11);
@@ -380,6 +390,11 @@ TEMPLATE_TEST_CASE("Unit_surf3D_Positive_ReadWrite", "", char, uchar, short, ush
uint4, float4) {
CHECK_IMAGE_SUPPORT;
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
const int width = GENERATE(31, 67);
const int height = GENERATE(131, 263);
const int depth = GENERATE(4, 11);
projects/hip-tests/catch/unit/texture/hipBindTex2DPitch.cc
+6 -1
Просмотреть файл
@@ -32,7 +32,7 @@ texture<TYPE_t, 2, hipReadModeElementType> tex;
// texture object is a kernel argument
static __global__ void texture2dCopyKernel(TYPE_t* dst) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
if ( (x < SIZE_W) && (y < SIZE_H) ) {
@@ -45,6 +45,11 @@ static __global__ void texture2dCopyKernel(TYPE_t* dst) {
TEST_CASE("Unit_hipBindTexture2D_Pitch") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
TYPE_t* B;
TYPE_t* A;
TYPE_t* devPtrB;
projects/hip-tests/catch/unit/texture/hipBindTexRef1DFetch.cc
+6 -1
Просмотреть файл
@@ -27,7 +27,7 @@ THE SOFTWARE.
texture<float, 1, hipReadModeElementType> tex;
static __global__ void kernel(float *out) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
if (x < N) {
out[x] = tex1Dfetch(tex, x);
@@ -38,6 +38,11 @@ static __global__ void kernel(float *out) {
TEST_CASE("Unit_hipBindTexture_tex1DfetchVerification") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
float *texBuf;
float val[N], output[N];
size_t offset = 0;
projects/hip-tests/catch/unit/texture/hipBindTextureToMipmappedArray.cc
+14 -3
Просмотреть файл
@@ -33,7 +33,7 @@ THE SOFTWARE.
texture<float, 2, hipReadModeElementType> texRef;
// MipMap is currently supported only on windows
#if (defined(_WIN32) && !defined(__HIP_NO_IMAGE_SUPPORT))
#if (defined(_WIN32) && !__HIP_NO_IMAGE_SUPPORT)
__global__ void tex2DKernel(float* outputData, int width, int height, float level) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
@@ -136,10 +136,16 @@ static void runMipMapTest(unsigned int width, unsigned int height, unsigned int
*/
TEST_CASE("Unit_hipTextureMipmapRef2D_Positive_Check") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
// Height Width Vector
std::vector<unsigned int> hw_vec = {2048, 1024, 512, 256, 64};
std::vector<unsigned int> mip_vec = {8, 4, 2, 1};
#if (defined(_WIN32) && !defined(__HIP_NO_IMAGE_SUPPORT))
#if (defined(_WIN32) && !__HIP_NO_IMAGE_SUPPORT)
for (auto& hw : hw_vec) {
for (auto& mip : mip_vec) {
if ((hw / static_cast<int>(pow(2, (mip * 2)))) > 0) {
@@ -175,7 +181,12 @@ TEST_CASE("Unit_hipTextureMipmapRef2D_Positive_Check") {
TEST_CASE("Unit_hipTextureMipmapRef2D_Negative_Parameters") {
CHECK_IMAGE_SUPPORT
#if (defined(_WIN32) && !defined(__HIP_NO_IMAGE_SUPPORT))
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
#if defined(_WIN32)
unsigned int width = 64;
unsigned int height = 64;
unsigned int mipmap_level = 1;
projects/hip-tests/catch/unit/texture/hipNormalizedFloatValueTex.cc
+6 -1
Просмотреть файл
@@ -46,7 +46,7 @@ texture<unsigned char, hipTextureType1D, hipReadModeNormalizedFloat> texuc;
template<typename T>
__global__ void normalizedValTextureTest(unsigned int numElements,
float* pDst) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
unsigned int elementID = threadIdx.x;
if (elementID >= numElements)
return;
@@ -149,6 +149,11 @@ static void runTest_hipTextureFilterMode() {
TEST_CASE("Unit_hipNormalizedFloatValueTex_CheckModes") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
#if HT_AMD
hipDeviceProp_t props;
HIP_CHECK(hipGetDeviceProperties(&props, 0));
projects/hip-tests/catch/unit/texture/hipSimpleTexture1DLayered.cc
+13 -1
Просмотреть файл
@@ -27,7 +27,7 @@ THE SOFTWARE.
template <typename TestType>
__global__ void simpleKernelLayered1DArray(hipTextureObject_t tex, TestType* outputData,
unsigned int width, unsigned int layer) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
outputData[layer * width + x] = tex1DLayered<TestType>(tex, x, layer);
#endif
@@ -62,6 +62,12 @@ TEMPLATE_TEST_CASE("Unit_Layered1DTexture_Check_HostBufferToFromLayered1DArray",
char2, uchar2, short2, ushort2, int2, uint2, float2,
char4, uchar4, short4, ushort4, int4, uint4, float4) {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
constexpr int SIZE = 512;
constexpr int num_layers = 5;
constexpr unsigned int width = SIZE;
@@ -205,6 +211,12 @@ TEMPLATE_TEST_CASE("Unit_Layered1DTexture_Check_DeviceBufferToFromLayered1DArray
char2, uchar2, short2, ushort2, int2, uint2, float2,
char4, uchar4, short4, ushort4, int4, uint4, float4) {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
constexpr int SIZE = 512;
constexpr int num_layers = 5;
constexpr unsigned int width = SIZE;
projects/hip-tests/catch/unit/texture/hipSimpleTexture2DLayered.cc
+13 -1
Просмотреть файл
@@ -28,7 +28,7 @@ template <typename TestType>
__global__ void simpleKernelLayered2DArray(hipTextureObject_t tex, TestType* outputData,
unsigned int width, unsigned int height,
unsigned int layer) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
outputData[layer * width * height + y * width + x] = tex2DLayered<TestType>(tex, x, y, layer);
@@ -64,6 +64,12 @@ TEMPLATE_TEST_CASE("Unit_Layered2DTexture_Check_HostBufferToFromLayered2DArray",
char2, uchar2, short2, ushort2, int2, uint2, float2,
char4, uchar4, short4, ushort4, int4, uint4, float4) {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
constexpr int SIZE = 512;
constexpr int num_layers = 5;
constexpr unsigned int width = SIZE;
@@ -210,6 +216,12 @@ TEMPLATE_TEST_CASE("Unit_Layered2DTexture_Check_DeviceBufferToFromLayered2DArray
char2, uchar2, short2, ushort2, int2, uint2, float2,
char4, uchar4, short4, ushort4, int4, uint4, float4) {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
constexpr int SIZE = 512;
constexpr int num_layers = 5;
constexpr unsigned int width = SIZE;
projects/hip-tests/catch/unit/texture/hipSimpleTexture3D.cc
+6 -1
Просмотреть файл
@@ -31,7 +31,7 @@ texture<char, hipTextureType3D, hipReadModeElementType> texc;
template <typename T>
__global__ void simpleKernel3DArray(T* outputData, int width,
int height, int depth) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
for (int i = 0; i < depth; i++) {
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
@@ -116,6 +116,11 @@ static void runSimpleTexture3D_Check(int width, int height, int depth,
TEST_CASE("Unit_hipSimpleTexture3D_Check_DataTypes") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
for ( int i = 1; i < 25; i++ ) {
runSimpleTexture3D_Check<float>(i, i, i, &texf);
runSimpleTexture3D_Check<int>(i+1, i, i, &texi);
projects/hip-tests/catch/unit/texture/hipTex1DFetchCheckModes.cc
+6 -1
Просмотреть файл
@@ -24,7 +24,7 @@ THE SOFTWARE.
#define offset 3
static __global__ void tex1dKernel(float *val, hipTextureObject_t obj) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int k = blockIdx.x * blockDim.x + threadIdx.x;
if (k < (N - offset))
val[k] = tex1Dfetch<float>(obj, k+offset);
@@ -110,6 +110,11 @@ static void runTest(hipTextureAddressMode addressMode,
TEST_CASE("Unit_tex1Dfetch_CheckModes") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
SECTION("hipAddressModeClamp AND hipFilterModePoint") {
runTest(hipAddressModeClamp, hipFilterModePoint);
}
projects/hip-tests/catch/unit/texture/hipTexObjPitch.cc
+6 -1
Просмотреть файл
@@ -37,7 +37,7 @@ THE SOFTWARE.
template <typename TYPE_t>
static __global__ void texture2dCopyKernel(hipTextureObject_t texObj,
TYPE_t* dst) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
for (int i = 0; i < SIZE_H; i++)
for (int j = 0; j < SIZE_W; j++)
dst[SIZE_W*i+j] = tex2D<TYPE_t>(texObj, j, i);
@@ -61,6 +61,11 @@ TEMPLATE_TEST_CASE("Unit_hipTexObjPitch_texture2D", "", float, int,
unsigned char, int16_t, char, unsigned int) {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
TestType* B;
TestType* A;
TestType* devPtrB;
projects/hip-tests/catch/unit/texture/hipTextureMipmapObj1D.cc
+19 -2
Просмотреть файл
@@ -33,7 +33,7 @@ static constexpr bool printLog = false; // Print log for debugging
template <typename T, hipTextureReadMode readMode>
static __global__ void populateMipmapNextLevelArray(hipSurfaceObject_t surfOut,
hipTextureObject_t texIn, unsigned int width, T* data = nullptr) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
float px = 1.0 / float(width);
@@ -61,7 +61,7 @@ static __global__ void populateMipmapNextLevelArray(hipSurfaceObject_t surfOut,
template <typename T>
__global__ void getMipmap(hipTextureObject_t texMipmap, unsigned int width,
float offsetX, float lod, T* data = nullptr) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
float px = 1.0 / float(width);
@@ -316,6 +316,11 @@ TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType",
char4, uchar4, short4, ushort4, int4, uint4, float4) {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
SECTION(
"Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModePoint, "
"hipAddressModeClamp 23") {
@@ -366,6 +371,12 @@ TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloa
char2, uchar2, short2, ushort2,
char4, uchar4, short4, ushort4) {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
SECTION(
"Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, "
"hipAddressModeClamp 23") {
@@ -438,6 +449,12 @@ TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj1D_Check - hipReadModeNormalizedFloa
TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType float only", "",
float, float1, float2, float4) {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
SECTION(
"Unit_hipTextureMipmapObj1D_Check - hipReadModeElementType, hipFilterModeLinear, "
"hipAddressModeClamp 23, 0.") {
projects/hip-tests/catch/unit/texture/hipTextureMipmapObj2D.cc
+19 -2
Просмотреть файл
@@ -40,7 +40,7 @@ static constexpr bool printLog = false; // Print log for debugging
template <typename T, hipTextureReadMode readMode>
static __global__ void populateMipmapNextLevelArray(hipSurfaceObject_t surfOut,
hipTextureObject_t texIn, unsigned int width, unsigned int height, T* data = nullptr) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
@@ -71,7 +71,7 @@ static __global__ void populateMipmapNextLevelArray(hipSurfaceObject_t surfOut,
template <typename T>
static __global__ void getMipmap(hipTextureObject_t texMipmap, unsigned int width,
unsigned int height, float offsetX, float offsetY, float lod, T* data = nullptr) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
@@ -349,6 +349,11 @@ TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType",
char4, uchar4, short4, ushort4, int4, uint4, float4) {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
SECTION(
"Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModePoint, "
"hipAddressModeClamp 23, 21") {
@@ -399,6 +404,12 @@ TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloa
char2, uchar2, short2, ushort2,
char4, uchar4, short4, ushort4) {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
SECTION(
"Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, "
"hipAddressModeClamp 23, 21") {
@@ -471,6 +482,12 @@ TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj2D_Check - hipReadModeNormalizedFloa
TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType float only", "",
float, float1, float2, float4) {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
SECTION(
"Unit_hipTextureMipmapObj2D_Check - hipReadModeElementType, hipFilterModeLinear, "
"hipAddressModeClamp 23, 17, 0., 0.") {
projects/hip-tests/catch/unit/texture/hipTextureMipmapObj3D.cc
+19 -2
Просмотреть файл
@@ -33,7 +33,7 @@ static constexpr bool printLog = false; // Print log for debugging
template <typename T, hipTextureReadMode readMode>
static __global__ void populateMipmapNextLevelArray(hipSurfaceObject_t surfOut, hipTextureObject_t texIn,
unsigned int width, unsigned int height, unsigned int depth, T* data = nullptr) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
@@ -74,7 +74,7 @@ static __global__ void getMipmap(hipTextureObject_t texMipmap, unsigned int widt
unsigned int height,
unsigned int depth, float offsetX, float offsetY, float offsetZ, float lod,
T* data = nullptr) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
@@ -367,6 +367,11 @@ TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType",
char4, uchar4, short4, ushort4, int4, uint4, float4) {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
SECTION(
"Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModePoint, "
"hipAddressModeClamp 23, 21, 47") {
@@ -417,6 +422,12 @@ TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloa
char2, uchar2, short2, ushort2,
char4, uchar4, short4, ushort4) {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
SECTION(
"Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloat, hipFilterModePoint, "
"hipAddressModeClamp 23, 21, 67") {
@@ -489,6 +500,12 @@ TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj3D_Check - hipReadModeNormalizedFloa
TEMPLATE_TEST_CASE("Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType float only", "",
float, float1, float2, float4) {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
SECTION(
"Unit_hipTextureMipmapObj3D_Check - hipReadModeElementType, hipFilterModeLinear, "
"hipAddressModeClamp 23, 17, 301, 0., 0., 0.") {
projects/hip-tests/catch/unit/texture/hipTextureObj1DCheckModes.cc
+6 -1
Просмотреть файл
@@ -31,7 +31,7 @@ THE SOFTWARE.
template<bool normalizedCoords>
__global__ void tex1DKernel(float *outputData, hipTextureObject_t textureObject,
int width, float offsetX) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
outputData[x] = tex1D<float>(textureObject, normalizedCoords ? (x + offsetX) / width : x + offsetX);
#endif
@@ -122,6 +122,11 @@ static void runTest(const int width, const float offsetX) {
TEST_CASE("Unit_hipTextureObj1DCheckModes") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") {
runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, -3);
runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 4);
projects/hip-tests/catch/unit/texture/hipTextureObj1DCheckSRGBModes.cc
+22 -2
Просмотреть файл
@@ -25,7 +25,7 @@ THE SOFTWARE.
template<bool normalizedCoords>
__global__ void tex1DRGBAKernel(float4 *outputData, hipTextureObject_t textureObject,
int width, float offsetX) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
outputData[x] = tex1D<float4>(textureObject,
normalizedCoords ? (x + offsetX) / width : x + offsetX);
@@ -33,7 +33,7 @@ __global__ void tex1DRGBAKernel(float4 *outputData, hipTextureObject_t textureOb
}
__global__ void tex1DRGBAKernelFetch(float4 *outputData, hipTextureObject_t textureObject, float offsetX) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
outputData[x] = tex1Dfetch<float4>(textureObject, int(x + offsetX));
#endif
@@ -177,6 +177,11 @@ line1:
TEST_CASE("Unit_hipTextureObj1DCheckRGBAModes - array") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
SECTION("RGBA 1D hipAddressModeClamp, hipFilterModePoint, hipResourceTypeArray, regularCoords") {
runTest<hipAddressModeClamp, hipFilterModePoint, hipResourceTypeArray, false>(255, -3.9);
runTest<hipAddressModeClamp, hipFilterModePoint, hipResourceTypeArray, false>(255, 4.4);
@@ -228,6 +233,11 @@ TEST_CASE("Unit_hipTextureObj1DCheckRGBAModes - array") {
TEST_CASE("Unit_hipTextureObj1DCheckSRGBAModes - array") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
SECTION("SRGBA 1D hipAddressModeClamp, hipFilterModePoint, hipResourceTypeArray, regularCoords") {
runTest<hipAddressModeClamp, hipFilterModePoint, hipResourceTypeArray, false, true>(255, -3.9);
runTest<hipAddressModeClamp, hipFilterModePoint, hipResourceTypeArray, false, true>(255, 4.4);
@@ -277,6 +287,11 @@ TEST_CASE("Unit_hipTextureObj1DCheckSRGBAModes - array") {
TEST_CASE("Unit_hipTextureObj1DCheckRGBAModes - buffer") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
SECTION("RGBA 1D hipAddressModeClamp, hipFilterModePoint, hipResourceTypeLinear, regularCoords") {
runTest<hipAddressModeClamp, hipFilterModePoint, hipResourceTypeLinear, false, false>(255);
}
@@ -289,6 +304,11 @@ TEST_CASE("Unit_hipTextureObj1DCheckRGBAModes - buffer") {
TEST_CASE("Unit_hipTextureObj1DCheckSRGBAModes - buffer") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
SECTION("SRGBA 1D hipAddressModeClamp, hipFilterModePoint, hipResourceTypeLinear, regularCoords") {
runTest<hipAddressModeClamp, hipFilterModePoint, hipResourceTypeLinear, false, true>(255);
}
projects/hip-tests/catch/unit/texture/hipTextureObj1DFetch.cc
+6 -1
Просмотреть файл
@@ -23,7 +23,7 @@ THE SOFTWARE.
#define N 512
static __global__ void tex1dKernel(float *val, hipTextureObject_t obj) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int k = blockIdx.x * blockDim.x + threadIdx.x;
if (k < N) {
val[k] = tex1Dfetch<float>(obj, k);
@@ -35,6 +35,11 @@ static __global__ void tex1dKernel(float *val, hipTextureObject_t obj) {
TEST_CASE("Unit_hipCreateTextureObject_tex1DfetchVerification") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
// Allocating the required buffer on gpu device
float *texBuf, *texBufOut;
float val[N], output[N];
projects/hip-tests/catch/unit/texture/hipTextureObj2D.cc
+6 -1
Просмотреть файл
@@ -28,7 +28,7 @@ THE SOFTWARE.
__global__ void tex2DKernel(float* outputData,
hipTextureObject_t textureObject, int width) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__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<float>(textureObject, x, y);
@@ -50,6 +50,11 @@ __global__ void tex2DKernel(float* outputData,
TEST_CASE("Unit_hipTextureObj2D_Check") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
constexpr int SIZE = 256;
constexpr unsigned int width = SIZE;
constexpr unsigned int height = SIZE;
projects/hip-tests/catch/unit/texture/hipTextureObj2DCheckModes.cc
+6 -1
Просмотреть файл
@@ -32,7 +32,7 @@ template<bool normalizedCoords>
__global__ void tex2DKernel(float *outputData, hipTextureObject_t textureObject,
int width, int height, float offsetX,
float offsetY) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__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<float>(textureObject,
@@ -134,6 +134,11 @@ line1:
TEST_CASE("Unit_hipTextureObj2DCheckModes") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") {
runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, -3.9, 6.1);
runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, 4.4, -7.0);
projects/hip-tests/catch/unit/texture/hipTextureObj2DCheckSRGBModes.cc
+11 -1
Просмотреть файл
@@ -26,7 +26,7 @@ template<bool normalizedCoords>
__global__ void tex2DRGBAKernel(float4 *outputData, hipTextureObject_t textureObject,
int width, int height, float offsetX,
float offsetY) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__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<float4>(textureObject,
@@ -157,6 +157,11 @@ line1:
TEST_CASE("Unit_hipTextureObj2DCheckRGBAModes") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
SECTION("RGBA 2D hipAddressModeClamp, hipFilterModePoint, regularCoords") {
runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, -3.9, 6.1);
runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, 4.4, -7.0);
@@ -202,6 +207,11 @@ TEST_CASE("Unit_hipTextureObj2DCheckRGBAModes") {
TEST_CASE("Unit_hipTextureObj2DCheckSRGBAModes") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
SECTION("SRGBA 2D hipAddressModeClamp, hipFilterModePoint, regularCoords") {
runTest<hipAddressModeClamp, hipFilterModePoint, false, true>(256, 256, -3.9, 6.1);
runTest<hipAddressModeClamp, hipFilterModePoint, false, true>(256, 256, 4.4, -7.0);
projects/hip-tests/catch/unit/texture/hipTextureObj3DCheckModes.cc
+6 -1
Просмотреть файл
@@ -34,7 +34,7 @@ bool LinearFilter3D = false;
template <bool normalizedCoords>
__global__ void tex3DKernel(float* outputData, hipTextureObject_t textureObject, int width,
int height, int depth, float offsetX, float offsetY, float offsetZ) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
@@ -173,6 +173,11 @@ line1:
TEST_CASE("Unit_hipTextureObj3DCheckModes") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
int device = 0;
hipDeviceProp_t props;
HIPCHECK(hipGetDeviceProperties(&props, device));
projects/hip-tests/catch/unit/texture/hipTextureObjFetchVector.cc
+6 -1
Просмотреть файл
@@ -27,7 +27,7 @@ THE SOFTWARE.
template <typename T>
__global__ void tex1dKernelFetch(T *val, hipTextureObject_t obj, int N) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__HIP_NO_IMAGE_SUPPORT
int k = blockIdx.x * blockDim.x + threadIdx.x;
if (k < N) {
val[k] = tex1Dfetch<T>(obj, k);
@@ -143,6 +143,11 @@ bool runTest() {
TEST_CASE("Unit_hipTextureFetch_vector") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
// test for char
runTest<char1>();
runTest<char2>();
projects/hip-tests/catch/unit/texture/hipTextureRef2D.cc
+6 -1
Просмотреть файл
@@ -26,7 +26,7 @@ THE SOFTWARE.
texture<float, 2, hipReadModeElementType> tex;
__global__ void tex2DKernel(float* outputData, int width) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
#if !__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(tex, x, y);
@@ -36,6 +36,11 @@ __global__ void tex2DKernel(float* outputData, int width) {
TEST_CASE("Unit_hipTextureRef2D_Check") {
CHECK_IMAGE_SUPPORT
#if __HIP_NO_IMAGE_SUPPORT
HipTest::HIP_SKIP_TEST("__HIP_NO_IMAGE_SUPPORT is set");
return;
#endif
constexpr int SIZE = 256;
constexpr unsigned int width = SIZE;
constexpr unsigned int height = SIZE;