AI/rocm-systems
2
0
Forka 0
Codice Problemi Pull Request Actions Pacchetti Progetti Rilasci Wiki Attività

EXSWHTEC-224 - Test cases ID clean up and documentation for Texture Management (#92)

Sfoglia il codice sorgente 
- Update hipBindTextureToMipmappedArray.cc
- Fix build error for undeclared identifier 'ret'
- Disable tests which fail in external CI
This commit is contained in:
milos-mozetic
2023-08-14 17:15:35 +02:00
committato da GitHub
parent 9c24abc77b
commit 7c7884a2d9
19 ha cambiato i file con 1321 aggiunte e 106 eliminazioni
Scarica il file Patch Scarica il file Diff Expand all files Collapse all files
catch/unit/texture/hipTextureObj3DCheckModes.cc
+64 -62
Vedi File
@@ -22,32 +22,38 @@ THE SOFTWARE.
#include <hip_test_checkers.hh>
#include <hip_texture_helper.hh>
/**
* @addtogroup hipCreateTextureObject hipCreateTextureObject
* @{
* @ingroup TextureTest
*/
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) {
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
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
outputData[z * width * height + y * width + x] = tex3D<float>(textureObject,
normalizedCoords ? (x + offsetX) / width : x + offsetX,
normalizedCoords ? (y + offsetY) / height : y + offsetY,
normalizedCoords ? (z + offsetZ) / depth : z + offsetZ);
outputData[z * width * height + y * width + x] =
tex3D<float>(textureObject, normalizedCoords ? (x + offsetX) / width : x + offsetX,
normalizedCoords ? (y + offsetY) / height : y + offsetY,
normalizedCoords ? (z + offsetZ) / depth : z + offsetZ);
#endif
}
template<hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool normalizedCoords>
static void runTest(const int width, const int height, const int depth, const float offsetX, const float offsetY,
const float offsetZ) {
//printf("%s(addressMode=%d, filterMode=%d, normalizedCoords=%d, width=%d, height=%d, depth=%d, offsetX=%f, offsetY=%f, offsetZ=%f)\n",
// __FUNCTION__, addressMode, filterMode, normalizedCoords, width, height,
// depth, offsetX, offsetY, offsetZ);
template <hipTextureAddressMode addressMode, hipTextureFilterMode filterMode, bool normalizedCoords>
static void runTest(const int width, const int height, const int depth, const float offsetX,
const float offsetY, const float offsetZ) {
// printf("%s(addressMode=%d, filterMode=%d, normalizedCoords=%d, width=%d, height=%d, depth=%d,
// offsetX=%f, offsetY=%f, offsetZ=%f)\n",
// __FUNCTION__, addressMode, filterMode, normalizedCoords, width, height,
// depth, offsetX, offsetY, offsetZ);
bool result = true;
unsigned int size = width * height * depth * sizeof(float);
float *hData = (float*) malloc(size);
float* hData = (float*)malloc(size);
memset(hData, 0, size);
for (int i = 0; i < depth; i++) {
@@ -61,13 +67,14 @@ static void runTest(const int width, const int height, const int depth, const fl
// Allocate array and copy image data
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
hipArray *arr;
hipArray* arr;
HIP_CHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayDefault));
HIP_CHECK(
hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayDefault));
hipMemcpy3DParms myparms;
memset(&myparms, 0, sizeof(myparms));
myparms.srcPos = make_hipPos(0,0,0);
myparms.dstPos = make_hipPos(0,0,0);
myparms.srcPos = make_hipPos(0, 0, 0);
myparms.dstPos = make_hipPos(0, 0, 0);
myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(float), width, height);
myparms.dstArray = arr;
myparms.extent = make_hipExtent(width, height, depth);
@@ -105,20 +112,20 @@ static void runTest(const int width, const int height, const int depth, const fl
return;
}
float *dData = nullptr;
HIP_CHECK(hipMalloc((void**) &dData, size));
float* dData = nullptr;
HIP_CHECK(hipMalloc((void**)&dData, size));
HIP_CHECK(hipMemset(dData, 0, size));
dim3 dimBlock(8, 8, 8); // 512 threads
dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y,
dim3 dimBlock(8, 8, 8); // 512 threads
dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y - 1) / dimBlock.y,
(depth + dimBlock.z - 1) / dimBlock.z);
hipLaunchKernelGGL(tex3DKernel<normalizedCoords>, dimGrid, dimBlock, 0, 0, dData,
textureObject, width, height, depth, offsetX, offsetY, offsetZ);
HIP_CHECK(hipGetLastError());
hipLaunchKernelGGL(tex3DKernel<normalizedCoords>, dimGrid, dimBlock, 0, 0, dData, textureObject,
width, height, depth, offsetX, offsetY, offsetZ);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipDeviceSynchronize());
float *hOutputData = (float*) malloc(size);
float* hOutputData = (float*)malloc(size);
memset(hOutputData, 0, size);
HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
@@ -130,8 +137,8 @@ static void runTest(const int width, const int height, const int depth, const fl
width, height, depth, offsetX + k, offsetY + j, offsetZ + i, hData);
if (!hipTextureSamplingVerify<float, filterMode>(hOutputData[index], expectedValue)) {
INFO("Mismatch at (" << offsetX + k << ", " << offsetY + j << ", " << offsetZ + i << "):" <<
hOutputData[index] << " expected:" << expectedValue);
INFO("Mismatch at (" << offsetX + k << ", " << offsetY + j << ", " << offsetZ + i
<< "):" << hOutputData[index] << " expected:" << expectedValue);
result = false;
goto line1;
}
@@ -145,9 +152,20 @@ line1:
HIP_CHECK(hipFreeArray(arr));
free(hData);
REQUIRE(result);
}
/**
* Test Description
* ------------------------
* - Uses different addressing and filtering modes for 3D arrays.
* Test source
* ------------------------
* - unit/texture/hipTextureObj3DCheckModes.cc
* Test requirements
* ------------------------
* - Textures supported on device
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_hipTextureObj3DCheckModes") {
CHECK_IMAGE_SUPPORT
@@ -159,58 +177,42 @@ TEST_CASE("Unit_hipTextureObj3DCheckModes") {
}
SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") {
runTest<hipAddressModeClamp, hipFilterModePoint, false>
(256, 256, 256, -3.9, 6.1, 9.5);
runTest<hipAddressModeClamp, hipFilterModePoint, false>
(256, 256, 256, 4.4, -7.0, 5.3);
runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, 256, -3.9, 6.1, 9.5);
runTest<hipAddressModeClamp, hipFilterModePoint, false>(256, 256, 256, 4.4, -7.0, 5.3);
}
SECTION("hipAddressModeBorder, hipFilterModePoint, regularCoords") {
runTest<hipAddressModeBorder, hipFilterModePoint, false>
(256, 256, 256, -8.5, 2.9, 5.8);
runTest<hipAddressModeBorder, hipFilterModePoint, false>
(256, 256, 256, 12.5, 6.7, 11.4);
runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, 256, 256, -8.5, 2.9, 5.8);
runTest<hipAddressModeBorder, hipFilterModePoint, false>(256, 256, 256, 12.5, 6.7, 11.4);
}
SECTION("hipAddressModeClamp, hipFilterModeLinear, regularCoords") {
runTest<hipAddressModeClamp, hipFilterModeLinear, false>
(256, 256, 256, -0.4, -0.4, -0.4);
runTest<hipAddressModeClamp, hipFilterModeLinear, false>
(256, 256, 256, 4, 14.6, -0.3);
runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, 256, 256, -0.4, -0.4, -0.4);
runTest<hipAddressModeClamp, hipFilterModeLinear, false>(256, 256, 256, 4, 14.6, -0.3);
}
SECTION("hipAddressModeBorder, hipFilterModeLinear, regularCoords") {
runTest<hipAddressModeBorder, hipFilterModeLinear, false>
(256, 256, 256, 6.9, 7.4, 0.4);
runTest<hipAddressModeBorder, hipFilterModeLinear, false>
(256, 256, 256, 12.5, 23.7, 0.34);
runTest<hipAddressModeBorder, hipFilterModeLinear, false>(256, 256, 256, 6.9, 7.4, 0.4);
runTest<hipAddressModeBorder, hipFilterModeLinear, false>(256, 256, 256, 12.5, 23.7, 0.34);
}
SECTION("hipAddressModeClamp, hipFilterModePoint, normalizedCoords") {
runTest<hipAddressModeClamp, hipFilterModePoint, true>
(256, 256, 256, -3, 8.9, -4);
runTest<hipAddressModeClamp, hipFilterModePoint, true>
(256, 256, 256, 4, -0.1, 8.2);
runTest<hipAddressModeClamp, hipFilterModePoint, true>(256, 256, 256, -3, 8.9, -4);
runTest<hipAddressModeClamp, hipFilterModePoint, true>(256, 256, 256, 4, -0.1, 8.2);
}
SECTION("hipAddressModeBorder, hipFilterModePoint, normalizedCoords") {
runTest<hipAddressModeBorder, hipFilterModePoint, true>
(256, 256, 256, -8.5, 15.9, 0.1);
runTest<hipAddressModeBorder, hipFilterModePoint, true>
(256, 256, 256, 12.5, -17.9, -0.35);
runTest<hipAddressModeBorder, hipFilterModePoint, true>(256, 256, 256, -8.5, 15.9, 0.1);
runTest<hipAddressModeBorder, hipFilterModePoint, true>(256, 256, 256, 12.5, -17.9, -0.35);
}
SECTION("hipAddressModeClamp, hipFilterModeLinear, normalizedCoords") {
runTest<hipAddressModeClamp, hipFilterModeLinear, true>
(256, 256, 256, -3, 5.8, 0.89);
runTest<hipAddressModeClamp, hipFilterModeLinear, true>
(256, 256, 256, 4, 9.1, 2.08);
runTest<hipAddressModeClamp, hipFilterModeLinear, true>(256, 256, 256, -3, 5.8, 0.89);
runTest<hipAddressModeClamp, hipFilterModeLinear, true>(256, 256, 256, 4, 9.1, 2.08);
}
SECTION("hipAddressModeBorder, hipFilterModeLinear, normalizedCoords") {
runTest<hipAddressModeBorder, hipFilterModeLinear, true>
(256, 256, 256, -8.5, 6.6, 3.67);
runTest<hipAddressModeBorder, hipFilterModeLinear, true>
(256, 256, 256, 12.5, 0.01, -9.9);
runTest<hipAddressModeBorder, hipFilterModeLinear, true>(256, 256, 256, -8.5, 6.6, 3.67);
runTest<hipAddressModeBorder, hipFilterModeLinear, true>(256, 256, 256, 12.5, 0.01, -9.9);
}
}