AI/rocm-systems
2
0
Fork 0
Κώδικας Ζητήματα Pull Requests Δράσεις Πακέτα Έργα Κυκλοφορίες Wiki Δραστηριότητα

SWDEV-474989 - Fix texture filter test issues

Περιήγηση Στο Πηγαίο 
Change-Id: I9b647aff79ba92f07b2ca18eac49d58ae63fa859


[ROCm/hip-tests commit: 8a9ed34f6a]
Αυτή η υποβολή περιλαμβάνεται σε:
taosang2
2024-07-30 09:35:04 -04:00
υποβλήθηκε από Tao Sang
γονέας 2253b82046
υποβολή 43b7d702ff
38 αρχεία άλλαξαν με 488 προσθήκες και 645 διαγραφές
Λήψη Αρχείου Patch Λήψη Αρχείου Diff Ανάπτυξη όλων των αρχείων Σύμπτυξη όλων των αρχείων
projects/hip-tests/catch/hipTestMain/config/config_amd_windows
+3 -244
Προβολή Αρχείου
@@ -13,7 +13,6 @@
"Unit_hipMemPoolTrimTo_Multithreaded",
"Unit_hipMemPoolSetGetAccess_Positive_MultipleGPU",
"Unit_hipMalloc_CoherentTst",
"Unit_hipTextureMipmapObj2D_Check",
"Unit_hipGraphAddHostNode_ClonedGraphwithHostNode",
"Unit_hipEventIpc",
"Unit_hipMalloc3D_Negative",
@@ -211,8 +210,8 @@
"Unit_hipHostMalloc_AllocateUseMoreThanAvailGPUMemory",
"=== SWDEV-432250:Below tests failed in stress test on 10/11/23 ===",
"Unit_hipVectorTypes_test_on_device",
"Unit_Layered1DTexture_Check_DeviceBufferToFromLayered1DArray - ushort4",
"Unit_Layered2DTexture_Check_DeviceBufferToFromLayered2DArray - float4",
"=== Below test is disabled due to defect EXSWHTEC-347 ===",
"Unit_hipPointerSetAttribute_Positive_SyncMemops",
"NOTE: The following test is disabled due to defect - EXSWHTEC-241",
"Unit_hipFuncGetAttributes_Negative_Parameters",
"NOTE: The following test is disabled due to defect - EXSWHTEC-242",
@@ -290,12 +289,6 @@
"Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - double",
"=== SWDEV-435667: Below tests failing randomly in stress test on 08/12/23 ===",
"Unit_hipMemPoolSetAccess_Negative_Parameters",
"Unit_hipMallocMipmappedArray_Negative_Parameters",
"Unit_hipFreeMipmappedArray_Negative_Parameters",
"Unit_hipGetMipmappedArrayLevel_Negative_Parameters",
"Unit_hipMipmappedArrayCreate_Negative_Parameters",
"Unit_hipMipmappedArrayDestroy_Negative_Parameters",
"Unit_hipMipmappedArrayGetLevel_Negative_Parameters",
"SWDEV-438524: Below tests taking long time to run in stress test on 15/12/23 ===",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - int",
"Unit_Coalesced_Group_Shfl_Up_Positive_Basic - unsigned int",
@@ -364,241 +357,7 @@
"Unit_hipDrvGraphAddMemsetNode_hipMalloc3D_1D",
"Unit_hipDrvGraphAddMemsetNode_hipMalloc_1D",
"Unit_hipDrvGraphAddMemsetNode_hipMallocManaged",
"Unit_tex1Dfetch_Positive_ReadModeElementType - char",
"Unit_tex1Dfetch_Positive_ReadModeElementType - unsigned char",
"Unit_tex1Dfetch_Positive_ReadModeElementType - short",
"Unit_tex1Dfetch_Positive_ReadModeElementType - unsigned short",
"Unit_tex1Dfetch_Positive_ReadModeElementType - int",
"Unit_tex1Dfetch_Positive_ReadModeElementType - unsigned int",
"Unit_tex1Dfetch_Positive_ReadModeElementType - float",
"Unit_tex1Dfetch_Positive_ReadModeNormalizedFloat - char",
"Unit_tex1Dfetch_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_tex1Dfetch_Positive_ReadModeNormalizedFloat - short",
"Unit_tex1Dfetch_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_tex1D_Positive_ReadModeNormalizedFloat - char",
"Unit_tex1D_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_tex1D_Positive_ReadModeNormalizedFloat - short",
"Unit_tex1D_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_tex1DLayered_Positive_ReadModeNormalizedFloat - char",
"Unit_tex1DLayered_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_tex1DLayered_Positive_ReadModeNormalizedFloat - short",
"Unit_tex1DLayered_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_tex1DGrad_Positive_ReadModeElementType - char",
"Unit_tex1DGrad_Positive_ReadModeElementType - unsigned char",
"Unit_tex1DGrad_Positive_ReadModeElementType - short",
"Unit_tex1DGrad_Positive_ReadModeElementType - unsigned short",
"Unit_tex1DGrad_Positive_ReadModeElementType - int",
"Unit_tex1DGrad_Positive_ReadModeElementType - unsigned int",
"Unit_tex1DGrad_Positive_ReadModeElementType - float",
"Unit_tex1DGrad_Positive_ReadModeNormalizedFloat - char",
"Unit_tex1DGrad_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_tex1DGrad_Positive_ReadModeNormalizedFloat - short",
"Unit_tex1DGrad_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_tex1DLayeredGrad_Positive_ReadModeElementType - char",
"Unit_tex1DLayeredGrad_Positive_ReadModeElementType - unsigned char",
"Unit_tex1DLayeredGrad_Positive_ReadModeElementType - short",
"Unit_tex1DLayeredGrad_Positive_ReadModeElementType - unsigned short",
"Unit_tex1DLayeredGrad_Positive_ReadModeElementType - int",
"Unit_tex1DLayeredGrad_Positive_ReadModeElementType - unsigned int",
"Unit_tex1DLayeredGrad_Positive_ReadModeElementType - float",
"Unit_tex1DLayeredGrad_Positive_ReadModeNormalizedFloat - char",
"Unit_tex1DLayeredGrad_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_tex1DLayeredGrad_Positive_ReadModeNormalizedFloat - short",
"Unit_tex1DLayeredGrad_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_tex1DLayeredLod_Positive_ReadModeElementType - char",
"Unit_tex1DLayeredLod_Positive_ReadModeElementType - unsigned char",
"Unit_tex1DLayeredLod_Positive_ReadModeElementType - short",
"Unit_tex1DLayeredLod_Positive_ReadModeElementType - unsigned short",
"Unit_tex1DLayeredLod_Positive_ReadModeElementType - int",
"Unit_tex1DLayeredLod_Positive_ReadModeElementType - unsigned int",
"Unit_tex1DLayeredLod_Positive_ReadModeElementType - float",
"Unit_tex1DLayeredLod_Positive_ReadModeNormalizedFloat - char",
"Unit_tex1DLayeredLod_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_tex1DLayeredLod_Positive_ReadModeNormalizedFloat - short",
"Unit_tex1DLayeredLod_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_tex1DLod_Positive_ReadModeElementType - char",
"Unit_tex1DLod_Positive_ReadModeElementType - unsigned char",
"Unit_tex1DLod_Positive_ReadModeElementType - short",
"Unit_tex1DLod_Positive_ReadModeElementType - unsigned short",
"Unit_tex1DLod_Positive_ReadModeElementType - int",
"Unit_tex1DLod_Positive_ReadModeElementType - unsigned int",
"Unit_tex1DLod_Positive_ReadModeElementType - float",
"Unit_tex1DLod_Positive_ReadModeNormalizedFloat - char",
"Unit_tex1DLod_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_tex1DLod_Positive_ReadModeNormalizedFloat - short",
"Unit_tex1DLod_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_tex3D_Positive_ReadModeElementType - char",
"Unit_tex3D_Positive_ReadModeElementType - unsigned char",
"Unit_tex3D_Positive_ReadModeElementType - short",
"Unit_tex3D_Positive_ReadModeElementType - unsigned short",
"Unit_tex3D_Positive_ReadModeElementType - int",
"Unit_tex3D_Positive_ReadModeElementType - unsigned int",
"Unit_tex3D_Positive_ReadModeElementType - float",
"Unit_tex3D_Positive_ReadModeNormalizedFloat - char",
"Unit_tex3D_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_tex3D_Positive_ReadModeNormalizedFloat - short",
"Unit_tex3D_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_tex3DLod_Positive_ReadModeElementType - char",
"Unit_tex3DLod_Positive_ReadModeElementType - unsigned char",
"Unit_tex3DLod_Positive_ReadModeElementType - short",
"Unit_tex3DLod_Positive_ReadModeElementType - unsigned short",
"Unit_tex3DLod_Positive_ReadModeElementType - int",
"Unit_tex3DLod_Positive_ReadModeElementType - unsigned int",
"Unit_tex3DLod_Positive_ReadModeElementType - float",
"Unit_tex3DLod_Positive_ReadModeNormalizedFloat - char",
"Unit_tex3DLod_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_tex3DLod_Positive_ReadModeNormalizedFloat - short",
"Unit_tex3DLod_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_tex3DGrad_Positive_ReadModeElementType - char",
"Unit_tex3DGrad_Positive_ReadModeElementType - unsigned char",
"Unit_tex3DGrad_Positive_ReadModeElementType - short",
"Unit_tex3DGrad_Positive_ReadModeElementType - unsigned short",
"Unit_tex3DGrad_Positive_ReadModeElementType - int",
"Unit_tex3DGrad_Positive_ReadModeElementType - unsigned int",
"Unit_tex3DGrad_Positive_ReadModeElementType - float",
"Unit_tex3DGrad_Positive_ReadModeNormalizedFloat - char",
"Unit_tex3DGrad_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_tex3DGrad_Positive_ReadModeNormalizedFloat - short",
"Unit_tex3DGrad_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_texCubemap_Positive_ReadModeElementType - char",
"Unit_texCubemap_Positive_ReadModeElementType - unsigned char",
"Unit_texCubemap_Positive_ReadModeElementType - short",
"Unit_texCubemap_Positive_ReadModeElementType - unsigned short",
"Unit_texCubemap_Positive_ReadModeElementType - int",
"Unit_texCubemap_Positive_ReadModeElementType - unsigned int",
"Unit_texCubemap_Positive_ReadModeElementType - float",
"Unit_texCubemap_Positive_ReadModeNormalizedFloat - char",
"Unit_texCubemap_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_texCubemap_Positive_ReadModeNormalizedFloat - short",
"Unit_texCubemap_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_texCubemapLod_Positive_ReadModeElementType - char",
"Unit_texCubemapLod_Positive_ReadModeElementType - unsigned char",
"Unit_texCubemapLod_Positive_ReadModeElementType - short",
"Unit_texCubemapLod_Positive_ReadModeElementType - unsigned short",
"Unit_texCubemapLod_Positive_ReadModeElementType - int",
"Unit_texCubemapLod_Positive_ReadModeElementType - unsigned int",
"Unit_texCubemapLod_Positive_ReadModeElementType - float",
"Unit_texCubemapLod_Positive_ReadModeNormalizedFloat - char",
"Unit_texCubemapLod_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_texCubemapLod_Positive_ReadModeNormalizedFloat - short",
"Unit_texCubemapLod_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_texCubemapGrad_Positive_ReadModeElementType - char",
"Unit_texCubemapGrad_Positive_ReadModeElementType - unsigned char",
"Unit_texCubemapGrad_Positive_ReadModeElementType - short",
"Unit_texCubemapGrad_Positive_ReadModeElementType - unsigned short",
"Unit_texCubemapGrad_Positive_ReadModeElementType - int",
"Unit_texCubemapGrad_Positive_ReadModeElementType - unsigned int",
"Unit_texCubemapGrad_Positive_ReadModeElementType - float",
"Unit_texCubemapGrad_Positive_ReadModeNormalizedFloat - char",
"Unit_texCubemapGrad_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_texCubemapGrad_Positive_ReadModeNormalizedFloat - short",
"Unit_texCubemapGrad_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_texCubemapLayered_Positive_ReadModeElementType - char",
"Unit_texCubemapLayered_Positive_ReadModeElementType - unsigned char",
"Unit_texCubemapLayered_Positive_ReadModeElementType - short",
"Unit_texCubemapLayered_Positive_ReadModeElementType - unsigned short",
"Unit_texCubemapLayered_Positive_ReadModeElementType - int",
"Unit_texCubemapLayered_Positive_ReadModeElementType - unsigned int",
"Unit_texCubemapLayered_Positive_ReadModeElementType - float",
"Unit_texCubemapLayered_Positive_ReadModeNormalizedFloat - char",
"Unit_texCubemapLayered_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_texCubemapLayered_Positive_ReadModeNormalizedFloat - short",
"Unit_texCubemapLayered_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_texCubemapLayeredLod_Positive_ReadModeElementType - char",
"Unit_texCubemapLayeredLod_Positive_ReadModeElementType - unsigned char",
"Unit_texCubemapLayeredLod_Positive_ReadModeElementType - short",
"Unit_texCubemapLayeredLod_Positive_ReadModeElementType - unsigned short",
"Unit_texCubemapLayeredLod_Positive_ReadModeElementType - int",
"Unit_texCubemapLayeredLod_Positive_ReadModeElementType - unsigned int",
"Unit_texCubemapLayeredLod_Positive_ReadModeElementType - float",
"Unit_texCubemapLayeredLod_Positive_ReadModeNormalizedFloat - char",
"Unit_texCubemapLayeredLod_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_texCubemapLayeredLod_Positive_ReadModeNormalizedFloat - short",
"Unit_texCubemapLayeredLod_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_texCubemapLayeredGrad_Positive_ReadModeElementType - char",
"Unit_texCubemapLayeredGrad_Positive_ReadModeElementType - unsigned char",
"Unit_texCubemapLayeredGrad_Positive_ReadModeElementType - short",
"Unit_texCubemapLayeredGrad_Positive_ReadModeElementType - unsigned short",
"Unit_texCubemapLayeredGrad_Positive_ReadModeElementType - int",
"Unit_texCubemapLayeredGrad_Positive_ReadModeElementType - unsigned int",
"Unit_texCubemapLayeredGrad_Positive_ReadModeElementType - float",
"Unit_texCubemapLayeredGrad_Positive_ReadModeNormalizedFloat - char",
"Unit_texCubemapLayeredGrad_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_texCubemapLayeredGrad_Positive_ReadModeNormalizedFloat - short",
"Unit_texCubemapLayeredGrad_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_tex2Dgather_Positive_ReadModeElementType - char",
"Unit_tex2Dgather_Positive_ReadModeElementType - unsigned char",
"Unit_tex2Dgather_Positive_ReadModeElementType - short",
"Unit_tex2Dgather_Positive_ReadModeElementType - unsigned short",
"Unit_tex2Dgather_Positive_ReadModeElementType - int",
"Unit_tex2Dgather_Positive_ReadModeElementType - unsigned int",
"Unit_tex2Dgather_Positive_ReadModeElementType - float",
"Unit_tex2D_Positive_ReadModeElementType - char",
"Unit_tex2D_Positive_ReadModeElementType - unsigned char",
"Unit_tex2D_Positive_ReadModeElementType - short",
"Unit_tex2D_Positive_ReadModeElementType - unsigned short",
"Unit_tex2D_Positive_ReadModeElementType - int",
"Unit_tex2D_Positive_ReadModeElementType - unsigned int",
"Unit_tex2D_Positive_ReadModeElementType - float",
"Unit_tex2D_Positive_ReadModeNormalizedFloat - char",
"Unit_tex2D_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_tex2D_Positive_ReadModeNormalizedFloat - short",
"Unit_tex2D_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_tex2DLayered_Positive_ReadModeElementType - char",
"Unit_tex2DLayered_Positive_ReadModeElementType - unsigned char",
"Unit_tex2DLayered_Positive_ReadModeElementType - short",
"Unit_tex2DLayered_Positive_ReadModeElementType - unsigned short",
"Unit_tex2DLayered_Positive_ReadModeElementType - int",
"Unit_tex2DLayered_Positive_ReadModeElementType - unsigned int",
"Unit_tex2DLayered_Positive_ReadModeElementType - float",
"Unit_tex2DLayered_Positive_ReadModeNormalizedFloat - char",
"Unit_tex2DLayered_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_tex2DLayered_Positive_ReadModeNormalizedFloat - short",
"Unit_tex2DLayered_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_tex2DGrad_Positive_ReadModeElementType - char",
"Unit_tex2DGrad_Positive_ReadModeElementType - unsigned char",
"Unit_tex2DGrad_Positive_ReadModeElementType - short",
"Unit_tex2DGrad_Positive_ReadModeElementType - unsigned short",
"Unit_tex2DGrad_Positive_ReadModeElementType - int",
"Unit_tex2DGrad_Positive_ReadModeElementType - unsigned int",
"Unit_tex2DGrad_Positive_ReadModeElementType - float",
"Unit_tex2DGrad_Positive_ReadModeNormalizedFloat - char",
"Unit_tex2DGrad_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_tex2DGrad_Positive_ReadModeNormalizedFloat - short",
"Unit_tex2DGrad_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_tex2DLayeredGrad_Positive_ReadModeElementType - char",
"Unit_tex2DLayeredGrad_Positive_ReadModeElementType - unsigned char",
"Unit_tex2DLayeredGrad_Positive_ReadModeElementType - short",
"Unit_tex2DLayeredGrad_Positive_ReadModeElementType - unsigned short",
"Unit_tex2DLayeredGrad_Positive_ReadModeElementType - int",
"Unit_tex2DLayeredGrad_Positive_ReadModeElementType - unsigned int",
"Unit_tex2DLayeredGrad_Positive_ReadModeElementType - float",
"Unit_tex2DLayeredGrad_Positive_ReadModeNormalizedFloat - char",
"Unit_tex2DLayeredGrad_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_tex2DLayeredGrad_Positive_ReadModeNormalizedFloat - short",
"Unit_tex2DLayeredGrad_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_tex2DLod_Positive_ReadModeElementType - char",
"Unit_tex2DLod_Positive_ReadModeElementType - unsigned char",
"Unit_tex2DLod_Positive_ReadModeElementType - short",
"Unit_tex2DLod_Positive_ReadModeElementType - unsigned short",
"Unit_tex2DLod_Positive_ReadModeElementType - int",
"Unit_tex2DLod_Positive_ReadModeElementType - unsigned int",
"Unit_tex2DLod_Positive_ReadModeElementType - float",
"Unit_tex2DLod_Positive_ReadModeNormalizedFloat - char",
"Unit_tex2DLod_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_tex2DLod_Positive_ReadModeNormalizedFloat - short",
"Unit_tex2DLod_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_tex2DLayeredLod_Positive_ReadModeElementType - char",
"Unit_tex2DLayeredLod_Positive_ReadModeElementType - unsigned char",
"Unit_tex2DLayeredLod_Positive_ReadModeElementType - short",
"Unit_tex2DLayeredLod_Positive_ReadModeElementType - unsigned short",
"Unit_tex2DLayeredLod_Positive_ReadModeElementType - int",
"Unit_tex2DLayeredLod_Positive_ReadModeElementType - unsigned int",
"Unit_tex2DLayeredLod_Positive_ReadModeElementType - float",
"Unit_tex2DLayeredLod_Positive_ReadModeNormalizedFloat - char",
"Unit_tex2DLayeredLod_Positive_ReadModeNormalizedFloat - unsigned char",
"Unit_tex2DLayeredLod_Positive_ReadModeNormalizedFloat - short",
"Unit_tex2DLayeredLod_Positive_ReadModeNormalizedFloat - unsigned short",
"Unit_hipModuleLaunchKernel_Negative_Parameters",
"Unit_hipExtModuleLaunchKernel_Negative_Parameters",
"Unit_hipLaunchKernel_Negative_Parameters",
"Unit_Kernel_Launch_bounds_Negative_OutOfBounds",
projects/hip-tests/catch/hipTestMain/config/config_nvidia_windows.json
+20 -1
Προβολή Αρχείου
@@ -39,6 +39,25 @@
"Performance_hipMemsetD32Async",
"Unit_hipMemcpyParam2D_Positive_Synchronization_Behavior",
"Unit_hipMemcpy_Positive_Synchronization_Behavior",
"Unit_hipMemMapArrayAsync_Positive_Basic"
"Unit_hipMemMapArrayAsync_Positive_Basic",
"=== SWDEV-475987 : Disable tests to merge hipother change 12/08/2024 ===",
"Unit_hipMalloc3DArray_Negative_InvalidFormat",
"Unit_hipMalloc3DArray_Negative_BadChannelLayout",
"Unit_hipMalloc3DArray_Negative_8BitFloat",
"Unit_hipMalloc3DArray_Negative_DifferentChannelSizes",
"Unit_hipMalloc3DArray_Negative_BadChannelSize",
"Unit_hipMallocMipmappedArray_Negative_InvalidFormat",
"Unit_hipMallocMipmappedArray_Negative_BadChannelLayout",
"Unit_hipMallocMipmappedArray_Negative_8BitFloat",
"Unit_hipMallocMipmappedArray_Negative_DifferentChannelSizes",
"Unit_hipMallocMipmappedArray_Negative_BadChannelSize",
"Unit_hipMallocArray_Negative_DifferentChannelSizes",
"Unit_hipMallocArray_Negative_8bitFloat - float",
"Unit_hipMallocArray_Negative_8bitFloat - float2",
"Unit_hipMallocArray_Negative_8bitFloat - float4",
"Unit_hipMallocArray_Negative_BadNumberOfBits",
"Unit_hipMallocArray_Negative_3ChannelElement",
"Unit_hipMallocArray_Negative_ChannelAfterZeroChannel",
"Unit_hipMallocArray_Negative_InvalidChannelFormat"
]
}
projects/hip-tests/catch/include/hip_texture_helper.hh
+1 -1
Προβολή Αρχείου
@@ -148,7 +148,7 @@ hipFabs(const T &t) {
}
template<typename T, hipTextureFilterMode fMode = hipFilterModePoint, bool sRGB = false>
bool hipTextureSamplingVerify(T outputData, T expected) {
bool hipTextureSamplingVerify(const T &outputData, const T &expected) {
bool testResult = false;
if (fMode == hipFilterModePoint && !sRGB) {
testResult = outputData == expected;
projects/hip-tests/catch/unit/texture/CMakeLists.txt
+7 -23
Προβολή Αρχείου
@@ -52,55 +52,34 @@ set(TEST_SRC
hipTexRefSetMaxAnisotropy.cc
)
# Mipmap APIs are not supported on Linux
if(WIN32)
# tests not for gfx90a+
set(NOT_FOR_gfx90a_AND_ABOVE_TEST
tex1D.cc
tex1DGrad.cc
tex1DLod.cc
tex1DLayeredLod.cc
tex1DLayeredGrad.cc
tex1Dfetch.cc
tex1DLayered.cc
tex2D.cc
tex2DLod.cc
tex2DGrad.cc
tex2DLayeredLod.cc
tex2DLayeredGrad.cc
tex2Dgather.cc
tex2DLayered.cc
tex3D.cc
tex3DGrad.cc
tex3DLod.cc
texCubemap.cc
texCubemapGrad.cc
texCubemapLod.cc
texCubemapLayered.cc
texCubemapLayeredGrad.cc
texCubemapLayeredLod.cc
)
endif()
set(gfx90a_AND_ABOVE_TARGETS gfx90a gfx940 gfx941 gfx942)
set(gfx90a_AND_ABOVE_TARGETS gfx90a gfx940 gfx941 gfx942 gfx950)
function(CheckRejectedArchs OFFLOAD_ARCH_STR_LOCAL)
set(ARCH_CHECK -1 PARENT_SCOPE)
set(NOT_GFX90a -1)
set(GFX90a -1)
string(REGEX MATCHALL "--offload-arch=gfx[0-9a-z]+" OFFLOAD_ARCH_LIST ${OFFLOAD_ARCH_STR_LOCAL})
foreach(OFFLOAD_ARCH IN LISTS OFFLOAD_ARCH_LIST)
string(REGEX MATCHALL "--offload-arch=(gfx[0-9a-z]+)" matches ${OFFLOAD_ARCH})
if (CMAKE_MATCH_COUNT EQUAL 1)
if (CMAKE_MATCH_1 IN_LIST gfx90a_AND_ABOVE_TARGETS)
set(GFX90a 1)
else()
set(NOT_GFX90a 1)
set(ARCH_CHECK 1 PARENT_SCOPE)
endif() # CMAKE_MATCH_1
endif() # CMAKE_MATCH_COUNT
endforeach() # OFFLOAD_ARCH_LIST
if (${NOT_GFX90a} EQUAL -1 AND ${GFX90a} EQUAL 1)
set(ARCH_CHECK 1 PARENT_SCOPE)
endif()
endfunction() # CheckAcceptedArchs
add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/tex_ref_get_module.code
@@ -143,6 +122,11 @@ set(TEST_SRC
hipTextureMipmapObj1D.cc
hipTextureMipmapObj2D.cc
hipTextureMipmapObj3D.cc
tex1DLod.cc
tex1DLayeredLod.cc
tex2DLod.cc
tex2DLayeredLod.cc
tex3DLod.cc
)
endif()
projects/hip-tests/catch/unit/texture/hipFreeMipmappedArray.cc
+2
Προβολή Αρχείου
@@ -52,6 +52,8 @@ TEST_CASE("Unit_hipFreeMipmappedArray_Negative_Parameters") {
}
SECTION("double free") {
INFO("Double free cheching isn't supported. Skipped.");
return;
hipMipmappedArray_t array;
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
hipExtent extent = make_hipExtent(4, 4, 6);
projects/hip-tests/catch/unit/texture/hipGetMipmappedArrayLevel.cc
+2 -2
Προβολή Αρχείου
@@ -50,7 +50,7 @@ TEST_CASE("Unit_hipGetMipmappedArrayLevel_Negative_Parameters") {
hipMipmappedArray_t array;
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
hipExtent extent = make_hipExtent(4, 4, 6);
unsigned int levels = 4;
unsigned int levels = 1 + std::log2(extent.depth);
HIP_CHECK(hipMallocMipmappedArray(&array, &desc, extent, levels, 0));
@@ -65,7 +65,7 @@ TEST_CASE("Unit_hipGetMipmappedArrayLevel_Negative_Parameters") {
}
SECTION("level index is greater than number of levels") {
HIP_CHECK_ERROR(hipGetMipmappedArrayLevel(&levelArray, array, 4), hipErrorInvalidValue);
HIP_CHECK_ERROR(hipGetMipmappedArrayLevel(&levelArray, array, levels), hipErrorInvalidValue);
}
HIP_CHECK(hipFreeMipmappedArray(array));
projects/hip-tests/catch/unit/texture/hipMallocMipmappedArray.cc
+1 -1
Προβολή Αρχείου
@@ -50,7 +50,7 @@ TEST_CASE("Unit_hipMallocMipmappedArray_Negative_Parameters") {
hipMipmappedArray_t array;
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
hipExtent extent = make_hipExtent(4, 4, 6);
unsigned int levels = 4;
unsigned int levels = 1 + std::log2(extent.depth);
SECTION("mipmappedArray is nullptr") {
HIP_CHECK_ERROR(hipMallocMipmappedArray(nullptr, &desc, extent, levels, 0),
projects/hip-tests/catch/unit/texture/hipMipmappedArrayCreate.cc
+1 -1
Προβολή Αρχείου
@@ -59,7 +59,7 @@ TEST_CASE("Unit_hipMipmappedArrayCreate_Negative_Parameters") {
desc.Depth = 6;
desc.Flags = 0;
unsigned int levels = 4;
unsigned int levels = 1 + std::log2(desc.Depth);
HIP_CHECK(hipFree(0));
projects/hip-tests/catch/unit/texture/hipMipmappedArrayDestroy.cc
+3 -1
Προβολή Αρχείου
@@ -55,6 +55,8 @@ TEST_CASE("Unit_hipMipmappedArrayDestroy_Negative_Parameters") {
}
SECTION("double free") {
INFO("Double free cheching isn't supported. Skipped.");
return;
hipmipmappedArray array;
HIP_ARRAY3D_DESCRIPTOR desc = {};
@@ -66,7 +68,7 @@ TEST_CASE("Unit_hipMipmappedArrayDestroy_Negative_Parameters") {
desc.Depth = 6;
desc.Flags = 0;
unsigned int levels = 4;
unsigned int levels = 1 + std::log2(desc.Depth);
HIP_CHECK(hipMipmappedArrayCreate(&array, &desc, levels));
projects/hip-tests/catch/unit/texture/hipMipmappedArrayGetLevel.cc
+2 -2
Προβολή Αρχείου
@@ -59,7 +59,7 @@ TEST_CASE("Unit_hipMipmappedArrayGetLevel_Negative_Parameters") {
desc.Depth = 6;
desc.Flags = 0;
unsigned int levels = 4;
unsigned int levels = 1 + std::log2(desc.Depth);
HIP_CHECK(hipFree(0));
HIP_CHECK(hipMipmappedArrayCreate(&array, &desc, levels));
@@ -75,7 +75,7 @@ TEST_CASE("Unit_hipMipmappedArrayGetLevel_Negative_Parameters") {
}
SECTION("level index is greater than number of levels") {
HIP_CHECK_ERROR(hipMipmappedArrayGetLevel(&levelArray, array, 4), hipErrorInvalidValue);
HIP_CHECK_ERROR(hipMipmappedArrayGetLevel(&levelArray, array, levels), hipErrorInvalidValue);
}
HIP_CHECK(hipMipmappedArrayDestroy(array));
projects/hip-tests/catch/unit/texture/hipTexObjPitch.cc
+5 -4
Προβολή Αρχείου
@@ -57,15 +57,16 @@ static __global__ void texture2dCopyKernel(hipTextureObject_t texObj,
* - Textures supported on device
* - HIP_VERSION >= 5.2
*/
TEMPLATE_TEST_CASE("Unit_hipTexObjPitch_texture2D", "", float, int,
unsigned char, int16_t, char, unsigned int) {
TEMPLATE_TEST_CASE("Unit_hipTexObjPitch_texture2D", "", char, unsigned char, short,
unsigned short, int, unsigned int, float) {
CHECK_IMAGE_SUPPORT
#if HT_NVIDIA
(void)hipGetLastError(); // Prevent negative tests affecting this
#endif
#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/kernels.hh
+49 -3
Προβολή Αρχείου
@@ -21,6 +21,7 @@ THE SOFTWARE.
*/
#pragma once
#pragma clang diagnostic ignored "-Wunused-parameter"
#include <hip/hip_runtime_api.h>
#include <hip/hip_cooperative_groups.h>
@@ -35,70 +36,82 @@ __host__ __device__ inline float GetCoordinate(size_t iteration, size_t N, size_
template <typename TexelType>
__global__ void tex1DfetchKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid = cg::this_grid().thread_rank();
if (tid >= N) return;
out[tid] = tex1D<TexelType>(tex_obj, tid);
out[tid] = tex1Dfetch<TexelType>(tex_obj, tid);
#endif
}
template <typename TexelType>
__global__ void tex1DKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj,
size_t width, size_t num_subdivisions, bool normalized_coords) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid = cg::this_grid().thread_rank();
if (tid >= N) return;
float x = GetCoordinate(tid, N, width, num_subdivisions, normalized_coords);
out[tid] = tex1D<TexelType>(tex_obj, x);
#endif
}
template <typename TexelType>
__global__ void tex1DLodKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj,
size_t width, size_t num_subdivisions, bool normalized_coords,
float level_of_detail) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid = cg::this_grid().thread_rank();
if (tid >= N) return;
float x = GetCoordinate(tid, N, width, num_subdivisions, normalized_coords);
out[tid] = tex1DLod<TexelType>(tex_obj, x, level_of_detail);
#endif
}
template <typename TexelType>
__global__ void tex1DLayeredLodKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj,
size_t width, size_t num_subdivisions, bool normalized_coords,
int layer, float level_of_detail) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid = cg::this_grid().thread_rank();
if (tid >= N) return;
float x = GetCoordinate(tid, N, width, num_subdivisions, normalized_coords);
out[tid] = tex1DLayeredLod<TexelType>(tex_obj, x, layer, level_of_detail);
#endif
}
template <typename TexelType>
__global__ void tex1DGradKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj,
size_t width, size_t num_subdivisions, bool normalized_coords,
float dx, float dy) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid = cg::this_grid().thread_rank();
if (tid >= N) return;
float x = GetCoordinate(tid, N, width, num_subdivisions, normalized_coords);
out[tid] = tex1DGrad<TexelType>(tex_obj, x, dx, dy);
#endif
}
template <typename TexelType>
__global__ void tex1DLayeredGradKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj,
size_t width, size_t num_subdivisions,
bool normalized_coords, int layer, float dx, float dy) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid = cg::this_grid().thread_rank();
if (tid >= N) return;
float x = GetCoordinate(tid, N, width, num_subdivisions, normalized_coords);
out[tid] = tex1DLayeredGrad<TexelType>(tex_obj, x, layer, dx, dy);
#endif
}
template <typename TexelType>
__global__ void tex2DgatherKernel(TexelType* const out, int comp, size_t N_x, size_t N_y,
hipTextureObject_t tex_obj, size_t width, size_t height,
size_t num_subdivisions, bool normalized_coords) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
if (tid_x >= N_x) return;
@@ -109,12 +122,14 @@ __global__ void tex2DgatherKernel(TexelType* const out, int comp, size_t N_x, si
float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords);
out[tid_y * N_x + tid_x] = tex2Dgather<TexelType>(tex_obj, x, y, comp);
#endif
}
template <typename TexelType>
__global__ void tex2DKernel(TexelType* const out, size_t N_x, size_t N_y,
hipTextureObject_t tex_obj, size_t width, size_t height,
size_t num_subdivisions, bool normalized_coords) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
if (tid_x >= N_x) return;
@@ -125,6 +140,7 @@ __global__ void tex2DKernel(TexelType* const out, size_t N_x, size_t N_y,
float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords);
out[tid_y * N_x + tid_x] = tex2D<TexelType>(tex_obj, x, y);
#endif
}
template <typename TexelType>
@@ -132,6 +148,7 @@ __global__ void tex2DGradKernel(TexelType* const out, size_t N_x, size_t N_y,
hipTextureObject_t tex_obj, size_t width, size_t height,
size_t num_subdivisions, bool normalized_coords, float2 dx,
float2 dy) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
if (tid_x >= N_x) return;
@@ -142,6 +159,7 @@ __global__ void tex2DGradKernel(TexelType* const out, size_t N_x, size_t N_y,
float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords);
out[tid_y * N_x + tid_x] = tex2DGrad<TexelType>(tex_obj, x, y, dx, dy);
#endif
}
template <typename TexelType>
@@ -149,6 +167,7 @@ __global__ void tex2DLayeredGradKernel(TexelType* const out, size_t N_x, size_t
hipTextureObject_t tex_obj, size_t width, size_t height,
size_t num_subdivisions, bool normalized_coords, float layer,
float2 dx, float2 dy) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
if (tid_x >= N_x) return;
@@ -159,12 +178,14 @@ __global__ void tex2DLayeredGradKernel(TexelType* const out, size_t N_x, size_t
float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords);
out[tid_y * N_x + tid_x] = tex2DLayeredGrad<TexelType>(tex_obj, x, y, layer, dx, dy);
#endif
}
template <typename TexelType>
__global__ void tex2DLodKernel(TexelType* const out, size_t N_x, size_t N_y,
hipTextureObject_t tex_obj, size_t width, size_t height,
size_t num_subdivisions, bool normalized_coords, float level) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
if (tid_x >= N_x) return;
@@ -175,6 +196,7 @@ __global__ void tex2DLodKernel(TexelType* const out, size_t N_x, size_t N_y,
float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords);
out[tid_y * N_x + tid_x] = tex2DLod<TexelType>(tex_obj, x, y, level);
#endif
}
template <typename TexelType>
@@ -182,6 +204,7 @@ __global__ void tex2DLayeredLodKernel(TexelType* const out, size_t N_x, size_t N
hipTextureObject_t tex_obj, size_t width, size_t height,
size_t num_subdivisions, bool normalized_coords, int layer,
float level) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
if (tid_x >= N_x) return;
@@ -192,12 +215,14 @@ __global__ void tex2DLayeredLodKernel(TexelType* const out, size_t N_x, size_t N
float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords);
out[tid_y * N_x + tid_x] = tex2DLayeredLod<TexelType>(tex_obj, x, y, layer, level);
#endif
}
template <typename TexelType>
__global__ void tex3DKernel(TexelType* const out, size_t N_x, size_t N_y, size_t N_z,
hipTextureObject_t tex_obj, size_t width, size_t height, size_t depth,
size_t num_subdivisions, bool normalized_coords) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
if (tid_x >= N_x) return;
@@ -212,6 +237,7 @@ __global__ void tex3DKernel(TexelType* const out, size_t N_x, size_t N_y, size_t
float z = GetCoordinate(tid_z, N_z, depth, num_subdivisions, normalized_coords);
out[tid_z * N_x * N_y + tid_y * N_x + tid_x] = tex3D<TexelType>(tex_obj, x, y, z);
#endif
}
template <typename TexelType>
@@ -219,6 +245,7 @@ __global__ void tex3DLodKernel(TexelType* const out, size_t N_x, size_t N_y, siz
hipTextureObject_t tex_obj, size_t width, size_t height,
size_t depth, size_t num_subdivisions, bool normalized_coords,
float level) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
if (tid_x >= N_x) return;
@@ -233,6 +260,7 @@ __global__ void tex3DLodKernel(TexelType* const out, size_t N_x, size_t N_y, siz
float z = GetCoordinate(tid_z, N_z, depth, num_subdivisions, normalized_coords);
out[tid_z * N_x * N_y + tid_y * N_x + tid_x] = tex3DLod<TexelType>(tex_obj, x, y, z, level);
#endif
}
template <typename TexelType>
@@ -240,6 +268,7 @@ __global__ void tex3DGradKernel(TexelType* const out, size_t N_x, size_t N_y, si
hipTextureObject_t tex_obj, size_t width, size_t height,
size_t depth, size_t num_subdivisions, bool normalized_coords,
float4 dx, float4 dy) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
if (tid_x >= N_x) return;
@@ -254,12 +283,14 @@ __global__ void tex3DGradKernel(TexelType* const out, size_t N_x, size_t N_y, si
float z = GetCoordinate(tid_z, N_z, depth, num_subdivisions, normalized_coords);
out[tid_z * N_x * N_y + tid_y * N_x + tid_x] = tex3DGrad<TexelType>(tex_obj, x, y, z, dx, dy);
#endif
}
template <typename TexelType>
__global__ void texCubemapKernel(TexelType* const out, size_t N_x, size_t N_y, size_t N_z,
hipTextureObject_t tex_obj, size_t width, size_t height,
size_t depth, size_t num_subdivisions, bool normalized_coords) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
if (tid_x >= N_x) return;
@@ -274,6 +305,7 @@ __global__ void texCubemapKernel(TexelType* const out, size_t N_x, size_t N_y, s
float z = GetCoordinate(tid_z, N_z, depth, num_subdivisions, normalized_coords);
out[tid_z * N_x * N_y + tid_y * N_x + tid_x] = texCubemap<TexelType>(tex_obj, x, y, z);
#endif
}
template <typename TexelType>
@@ -281,6 +313,7 @@ __global__ void texCubemapLodKernel(TexelType* const out, size_t N_x, size_t N_y
hipTextureObject_t tex_obj, size_t width, size_t height,
size_t depth, size_t num_subdivisions, bool normalized_coords,
float level) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
if (tid_x >= N_x) return;
@@ -295,6 +328,7 @@ __global__ void texCubemapLodKernel(TexelType* const out, size_t N_x, size_t N_y
float z = GetCoordinate(tid_z, N_z, depth, num_subdivisions, normalized_coords);
out[tid_z * N_x * N_y + tid_y * N_x + tid_x] = texCubemapLod<TexelType>(tex_obj, x, y, z, level);
#endif
}
template <typename TexelType>
@@ -302,6 +336,7 @@ __global__ void texCubemapGradKernel(TexelType* const out, size_t N_x, size_t N_
hipTextureObject_t tex_obj, size_t width, size_t height,
size_t depth, size_t num_subdivisions, bool normalized_coords,
float4 dx, float4 dy) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
if (tid_x >= N_x) return;
@@ -317,23 +352,27 @@ __global__ void texCubemapGradKernel(TexelType* const out, size_t N_x, size_t N_
out[tid_z * N_x * N_y + tid_y * N_x + tid_x] =
texCubemapGrad<TexelType>(tex_obj, x, y, z, dx, dy);
#endif
}
template <typename TexelType>
__global__ void tex1DLayeredKernel(TexelType* const out, size_t N, hipTextureObject_t tex_obj,
size_t width, size_t num_subdivisions, bool normalized_coords,
size_t layer) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid = cg::this_grid().thread_rank();
if (tid >= N) return;
float x = GetCoordinate(tid, N, width, num_subdivisions, normalized_coords);
out[tid] = tex1DLayered<TexelType>(tex_obj, x, layer);
#endif
}
template <typename TexelType>
__global__ void tex2DLayeredKernel(TexelType* const out, size_t N_x, size_t N_y,
hipTextureObject_t tex_obj, size_t width, size_t height,
size_t num_subdivisions, bool normalized_coords, size_t layer) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
if (tid_x >= N_x) return;
@@ -344,6 +383,7 @@ __global__ void tex2DLayeredKernel(TexelType* const out, size_t N_x, size_t N_y,
float y = GetCoordinate(tid_y, N_y, height, num_subdivisions, normalized_coords);
out[tid_y * N_x + tid_x] = tex2DLayered<TexelType>(tex_obj, x, y, layer);
#endif
}
template <typename TexelType>
@@ -351,6 +391,7 @@ __global__ void texCubemapLayeredKernel(TexelType* const out, size_t N_x, size_t
hipTextureObject_t tex_obj, size_t width, size_t height,
size_t depth, size_t num_subdivisions,
bool normalized_coords, size_t layer) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
if (tid_x >= N_x) return;
@@ -366,6 +407,7 @@ __global__ void texCubemapLayeredKernel(TexelType* const out, size_t N_x, size_t
out[tid_z * N_x * N_y + tid_y * N_x + tid_x] =
texCubemapLayered<TexelType>(tex_obj, x, y, z, layer);
#endif
}
template <typename TexelType>
@@ -373,6 +415,7 @@ __global__ void texCubemapLayeredLodKernel(TexelType* const out, size_t N_x, siz
hipTextureObject_t tex_obj, size_t width, size_t height,
size_t depth, size_t num_subdivisions,
bool normalized_coords, size_t layer, float level) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
if (tid_x >= N_x) return;
@@ -388,6 +431,7 @@ __global__ void texCubemapLayeredLodKernel(TexelType* const out, size_t N_x, siz
out[tid_z * N_x * N_y + tid_y * N_x + tid_x] =
texCubemapLayeredLod<TexelType>(tex_obj, x, y, z, layer, level);
#endif
}
template <typename TexelType>
@@ -396,6 +440,7 @@ __global__ void texCubemapLayeredGradKernel(TexelType* const out, size_t N_x, si
size_t height, size_t depth, size_t num_subdivisions,
bool normalized_coords, size_t layer, float4 dx,
float4 dy) {
#if !__HIP_NO_IMAGE_SUPPORT
const auto tid_x = blockIdx.x * blockDim.x + threadIdx.x;
if (tid_x >= N_x) return;
@@ -411,4 +456,5 @@ __global__ void texCubemapLayeredGradKernel(TexelType* const out, size_t N_x, si
out[tid_z * N_x * N_y + tid_y * N_x + tid_x] =
texCubemapLayeredGrad<TexelType>(tex_obj, x, y, z, layer, dx, dy);
}
#endif
}
projects/hip-tests/catch/unit/texture/test_fixture.hh
+24 -6
Προβολή Αρχείου
@@ -24,10 +24,10 @@ THE SOFTWARE.
#include <hip_test_common.hh>
#include <resource_guards.hh>
#include "texture_reference.hh"
#include "utils.hh"
#include "vec4.hh"
#include "texture_reference.hh"
#include "hip_texture_helper.hh"
template <typename TestType> struct TextureTestParams {
hipExtent extent;
@@ -64,7 +64,8 @@ template <typename TestType> struct TextureTestParams {
bool Layered() const { return layers > 1; }
void GenerateTextureDesc(decltype(hipReadModeElementType) read_mode = hipReadModeElementType) {
void GenerateTextureDesc(decltype(hipReadModeElementType) read_mode = hipReadModeElementType,
bool mipmap = false) {
constexpr bool is_floating_point = std::is_floating_point_v<TestType>;
memset(&tex_desc, 0, sizeof(tex_desc));
@@ -75,7 +76,10 @@ template <typename TestType> struct TextureTestParams {
tex_desc.filterMode = GENERATE(hipFilterModePoint, hipFilterModeLinear);
}
tex_desc.normalizedCoords = GENERATE(false, true);
tex_desc.normalizedCoords = true;
if (!mipmap) { // mipMap requires normalizedCoords = true
tex_desc.normalizedCoords = GENERATE(false, true);
}
auto address_mode_x = hipAddressModeClamp;
auto address_mode_y = address_mode_x;
@@ -116,7 +120,7 @@ struct TextureTestFixture {
hipResourceDesc res_desc;
LinearAllocGuard<VecType> host_alloc;
TextureReference<VecType> tex_h;
TextureReference<VecType, normalized_read> tex_h;
ArrayAllocGuardType<VecType> tex_alloc_d;
TextureGuard tex;
LinearAllocGuard<OutType> out_alloc_d;
@@ -168,4 +172,18 @@ struct TextureTestFixture {
hipMemcpyDeviceToHost));
HIP_CHECK(hipDeviceSynchronize());
}
};
template <typename ValType> bool Verify(const ValType& devValue, const ValType& hostValue) {
bool match = false;
if (params.tex_desc.filterMode == hipFilterModeLinear)
match = hipTextureSamplingVerify<ValType, hipFilterModeLinear>(devValue, hostValue);
else
match = hipTextureSamplingVerify<ValType, hipFilterModePoint>(devValue, hostValue);
if (!match) {
WARN((match ? "Matched: " : "Mismatched: ")
<< " GPU output : " << getString(devValue) << " CPU expected: " << getString(hostValue)
<< "\n");
}
return match;
}
};
projects/hip-tests/catch/unit/texture/tex1D.cc
+3 -10
Προβολή Αρχείου
@@ -77,10 +77,7 @@ TEMPLATE_TEST_CASE("Unit_tex1D_Positive_ReadModeElementType", "", char, unsigned
INFO("x: " << std::fixed << std::setprecision(16) << x);
auto ref_val = fixture.tex_h.Tex1D(x, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
@@ -129,12 +126,8 @@ TEMPLATE_TEST_CASE("Unit_tex1D_Positive_ReadModeNormalizedFloat", "", char, unsi
INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
INFO("x: " << std::fixed << std::setprecision(16) << x);
auto ref_val =
Vec4Map<TestType>(fixture.tex_h.Tex1D(x, params.tex_desc), NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
auto ref_val = fixture.tex_h.Tex1D(x, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
projects/hip-tests/catch/unit/texture/tex1DGrad.cc
+5 -12
Προβολή Αρχείου
@@ -57,7 +57,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeElementType", "", char, unsi
params.num_subdivisions = 4;
params.GenerateTextureDesc();
TextureTestFixture<TestType, false, true> fixture{params};
TextureTestFixture<TestType, false, false> fixture{params};
const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
tex1DGradKernel<vec4<TestType>><<<num_blocks, num_threads>>>(
@@ -77,10 +77,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeElementType", "", char, unsi
INFO("x: " << std::fixed << std::setprecision(16) << x);
auto ref_val = fixture.tex_h.Tex1D(x, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
@@ -110,7 +107,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeNormalizedFloat", "", char,
params.num_subdivisions = 4;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
TextureTestFixture<TestType, true, true> fixture{params};
TextureTestFixture<TestType, true, false> fixture{params};
const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
tex1DGradKernel<vec4<float>><<<num_blocks, num_threads>>>(
@@ -129,12 +126,8 @@ TEMPLATE_TEST_CASE("Unit_tex1DGrad_Positive_ReadModeNormalizedFloat", "", char,
INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
INFO("x: " << std::fixed << std::setprecision(16) << x);
auto ref_val =
Vec4Map<TestType>(fixture.tex_h.Tex1D(x, params.tex_desc), NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
auto ref_val = fixture.tex_h.Tex1D(x, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
projects/hip-tests/catch/unit/texture/tex1DLayered.cc
+4 -10
Προβολή Αρχείου
@@ -81,10 +81,8 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayered_Positive_ReadModeElementType", "", char, u
INFO("x: " << std::fixed << std::setprecision(16) << x);
const auto ref_val = fixture.tex_h.Tex1DLayered(x, layer, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
}
@@ -138,12 +136,8 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayered_Positive_ReadModeNormalizedFloat", "", cha
INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
INFO("x: " << std::fixed << std::setprecision(16) << x);
auto ref_val = Vec4Map<TestType>(fixture.tex_h.Tex1DLayered(x, layer, params.tex_desc),
NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
auto ref_val = fixture.tex_h.Tex1DLayered(x, layer, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
}
projects/hip-tests/catch/unit/texture/tex1DLayeredGrad.cc
+5 -12
Προβολή Αρχείου
@@ -58,7 +58,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredGrad_Positive_ReadModeElementType", "", cha
params.num_subdivisions = 4;
params.GenerateTextureDesc();
TextureTestFixture<TestType, false, true> fixture{params};
TextureTestFixture<TestType, false, false> fixture{params};
const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
@@ -81,10 +81,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredGrad_Positive_ReadModeElementType", "", cha
INFO("x: " << std::fixed << std::setprecision(16) << x);
auto ref_val = fixture.tex_h.Tex1DLayered(x, layer, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
}
@@ -116,7 +113,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredGrad_Positive_ReadModeNormalizedFloat", "",
params.num_subdivisions = 4;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
TextureTestFixture<TestType, true, true> fixture{params};
TextureTestFixture<TestType, true, false> fixture{params};
const auto [num_threads, num_blocks] = GetLaunchConfig(1024, params.NumItersX());
@@ -139,12 +136,8 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredGrad_Positive_ReadModeNormalizedFloat", "",
INFO("Filter mode: " << FilteringModeToString(params.tex_desc.filterMode));
INFO("x: " << std::fixed << std::setprecision(16) << x);
auto ref_val = Vec4Map<TestType>(fixture.tex_h.Tex1DLayered(x, layer, params.tex_desc),
NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
auto ref_val = fixture.tex_h.Tex1DLayered(x, layer, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
}
projects/hip-tests/catch/unit/texture/tex1DLayeredLod.cc
+5 -12
Προβολή Αρχείου
@@ -56,7 +56,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredLod_Positive_ReadModeElementType", "", char
params.extent = make_hipExtent(1024, 0, 0);
params.layers = 2;
params.num_subdivisions = 4;
params.GenerateTextureDesc();
params.GenerateTextureDesc(hipReadModeElementType, true);
TextureTestFixture<TestType, false, true> fixture{params};
@@ -81,10 +81,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredLod_Positive_ReadModeElementType", "", char
INFO("x: " << std::fixed << std::setprecision(16) << x);
auto ref_val = fixture.tex_h.Tex1DLayered(x, layer, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
}
@@ -114,7 +111,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredLod_Positive_ReadModeNormalizedFloat", "",
params.extent = make_hipExtent(1024, 0, 0);
params.layers = 2;
params.num_subdivisions = 4;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
params.GenerateTextureDesc(hipReadModeNormalizedFloat, true);
TextureTestFixture<TestType, true, true> fixture{params};
@@ -138,12 +135,8 @@ TEMPLATE_TEST_CASE("Unit_tex1DLayeredLod_Positive_ReadModeNormalizedFloat", "",
INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
INFO("x: " << std::fixed << std::setprecision(16) << x);
auto ref_val = Vec4Map<TestType>(fixture.tex_h.Tex1DLayered(x, layer, params.tex_desc),
NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
auto ref_val = fixture.tex_h.Tex1DLayered(x, layer, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
}
projects/hip-tests/catch/unit/texture/tex1DLod.cc
+5 -12
Προβολή Αρχείου
@@ -55,7 +55,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeElementType", "", char, unsig
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(1024, 0, 0);
params.num_subdivisions = 4;
params.GenerateTextureDesc();
params.GenerateTextureDesc(hipReadModeElementType, true);
TextureTestFixture<TestType, false, true> fixture{params};
@@ -77,10 +77,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeElementType", "", char, unsig
INFO("x: " << std::fixed << std::setprecision(16) << x);
auto ref_val = fixture.tex_h.Tex1D(x, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
@@ -108,7 +105,7 @@ TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeNormalizedFloat", "", char, u
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(1024, 0, 0);
params.num_subdivisions = 4;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
params.GenerateTextureDesc(hipReadModeNormalizedFloat, true);
TextureTestFixture<TestType, true, true> fixture{params};
@@ -129,12 +126,8 @@ TEMPLATE_TEST_CASE("Unit_tex1DLod_Positive_ReadModeNormalizedFloat", "", char, u
INFO("Address mode: " << AddressModeToString(params.tex_desc.addressMode[0]));
INFO("x: " << std::fixed << std::setprecision(16) << x);
auto ref_val =
Vec4Map<TestType>(fixture.tex_h.Tex1D(x, params.tex_desc), NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
auto ref_val = fixture.tex_h.Tex1D(x, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
projects/hip-tests/catch/unit/texture/tex1Dfetch.cc
+8 -13
Προβολή Αρχείου
@@ -24,7 +24,7 @@ THE SOFTWARE.
#include <hip_test_common.hh>
#include <resource_guards.hh>
#include "test_fixture.hh"
#include "kernels.hh"
#include "utils.hh"
#include "vec4.hh"
@@ -91,10 +91,7 @@ TEMPLATE_TEST_CASE("Unit_tex1Dfetch_Positive_ReadModeElementType", "", char, uns
for (auto i = 0u; i < out_alloc_h.size(); ++i) {
INFO("Index: " << i);
const auto ref_val = tex_h[i];
REQUIRE(ref_val.x == out_alloc_h[i].x);
REQUIRE(ref_val.y == out_alloc_h[i].y);
REQUIRE(ref_val.z == out_alloc_h[i].z);
REQUIRE(ref_val.w == out_alloc_h[i].w);
REQUIRE(out_alloc_h[i] == ref_val);
}
}
@@ -135,11 +132,11 @@ TEMPLATE_TEST_CASE("Unit_tex1Dfetch_Positive_ReadModeNormalizedFloat", "", char,
hipTextureDesc tex_desc;
memset(&tex_desc, 0, sizeof(tex_desc));
tex_desc.filterMode = hipFilterModePoint;
tex_desc.readMode = hipReadModeElementType;
tex_desc.readMode = hipReadModeNormalizedFloat;
tex_desc.normalizedCoords = false;
tex_desc.addressMode[0] = hipAddressModeClamp;
LinearAllocGuard<vec4<float>> out_alloc_d(LinearAllocs::hipMalloc, alloc_size);
LinearAllocGuard<vec4<float>> out_alloc_d(LinearAllocs::hipMalloc, tex_h.size() * sizeof(vec4<float>));
TextureGuard tex(&res_desc, &tex_desc);
const auto num_threads = std::min<size_t>(1024, tex_h.size());
@@ -148,16 +145,14 @@ TEMPLATE_TEST_CASE("Unit_tex1Dfetch_Positive_ReadModeNormalizedFloat", "", char,
<<<num_blocks, num_threads>>>(out_alloc_d.ptr(), tex_h.size(), tex.object());
std::vector<vec4<float>> out_alloc_h(tex_h.size());
HIP_CHECK(hipMemcpy(out_alloc_h.data(), out_alloc_d.ptr(), alloc_size, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(out_alloc_h.data(), out_alloc_d.ptr(), tex_h.size() * sizeof(vec4<float>),
hipMemcpyDeviceToHost));
HIP_CHECK(hipDeviceSynchronize());
for (auto i = 0u; i < out_alloc_h.size(); ++i) {
INFO("Index: " << i);
const auto ref_val = Vec4Map<TestType>(tex_h[i], NormalizeInteger<TestType>);
REQUIRE(ref_val.x == out_alloc_h[i].x);
REQUIRE(ref_val.y == out_alloc_h[i].y);
REQUIRE(ref_val.z == out_alloc_h[i].z);
REQUIRE(ref_val.w == out_alloc_h[i].w);
const auto ref_val = Vec4Map(tex_h[i]);
REQUIRE(out_alloc_h[i] == ref_val);
}
}
projects/hip-tests/catch/unit/texture/tex2D.cc
+12 -12
Προβολή Αρχείου
@@ -56,6 +56,10 @@ TEMPLATE_TEST_CASE("Unit_tex2D_Positive_ReadModeElementType", "", char, unsigned
params.extent = make_hipExtent(16, 4, 0);
params.num_subdivisions = 4;
params.GenerateTextureDesc();
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1]) {
INFO("Different address modes on X, Y aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType> fixture{params};
@@ -94,10 +98,7 @@ TEMPLATE_TEST_CASE("Unit_tex2D_Positive_ReadModeElementType", "", char, unsigned
INFO("y: " << std::fixed << std::setprecision(16) << y);
const auto ref_val = fixture.tex_h.Tex2D(x, y, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
@@ -126,6 +127,10 @@ TEMPLATE_TEST_CASE("Unit_tex2D_Positive_ReadModeNormalizedFloat", "", char, unsi
params.extent = make_hipExtent(16, 4, 0);
params.num_subdivisions = 4;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1]) {
INFO("Different address modes on X, Y aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, true> fixture{params};
@@ -149,7 +154,7 @@ TEMPLATE_TEST_CASE("Unit_tex2D_Positive_ReadModeNormalizedFloat", "", char, unsi
for (auto i = 0u; i < params.NumItersX() * params.NumItersY(); ++i) {
float x = i % params.NumItersX();
float y = i / params.NumItersY();
float y = i / params.NumItersX();
x = GetCoordinate(x, params.NumItersX(), params.Width(), params.num_subdivisions,
params.tex_desc.normalizedCoords);
@@ -162,13 +167,8 @@ TEMPLATE_TEST_CASE("Unit_tex2D_Positive_ReadModeNormalizedFloat", "", char, unsi
INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1]));
INFO("x: " << std::fixed << std::setprecision(16) << x);
INFO("y: " << std::fixed << std::setprecision(16) << y);
auto ref_val =
Vec4Map<TestType>(fixture.tex_h.Tex2D(x, y, params.tex_desc), NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
auto ref_val = fixture.tex_h.Tex2D(x, y, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
projects/hip-tests/catch/unit/texture/tex2DGrad.cc
+13 -14
Προβολή Αρχείου
@@ -56,8 +56,11 @@ TEMPLATE_TEST_CASE("Unit_tex2DGrad_Positive_ReadModeElementType", "", char, unsi
params.extent = make_hipExtent(16, 4, 0);
params.num_subdivisions = 4;
params.GenerateTextureDesc();
TextureTestFixture<TestType, false, true> fixture{params};
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1]) {
INFO("Different address modes on X, Y aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, false, false> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX());
const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(32, params.NumItersY());
@@ -95,10 +98,7 @@ TEMPLATE_TEST_CASE("Unit_tex2DGrad_Positive_ReadModeElementType", "", char, unsi
INFO("y: " << std::fixed << std::setprecision(16) << y);
const auto ref_val = fixture.tex_h.Tex2D(x, y, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
@@ -127,8 +127,11 @@ TEMPLATE_TEST_CASE("Unit_tex2DGrad_Positive_ReadModeNormalizedFloat", "", char,
params.extent = make_hipExtent(16, 4, 0);
params.num_subdivisions = 4;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
TextureTestFixture<TestType, true, true> fixture{params};
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1]) {
INFO("Different address modes on X, Y aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, true, false> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX());
const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(32, params.NumItersY());
@@ -165,12 +168,8 @@ TEMPLATE_TEST_CASE("Unit_tex2DGrad_Positive_ReadModeNormalizedFloat", "", char,
INFO("x: " << std::fixed << std::setprecision(16) << x);
INFO("y: " << std::fixed << std::setprecision(16) << y);
auto ref_val =
Vec4Map<TestType>(fixture.tex_h.Tex2D(x, y, params.tex_desc), NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
auto ref_val = fixture.tex_h.Tex2D(x, y, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
projects/hip-tests/catch/unit/texture/tex2DLayered.cc
+11 -12
Προβολή Αρχείου
@@ -57,7 +57,10 @@ TEMPLATE_TEST_CASE("Unit_tex2DLayered_Positive_ReadModeElementType", "", char, u
params.layers = 2;
params.num_subdivisions = 4;
params.GenerateTextureDesc();
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1]) {
INFO("Different address modes on X, Y aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX());
@@ -98,10 +101,7 @@ TEMPLATE_TEST_CASE("Unit_tex2DLayered_Positive_ReadModeElementType", "", char, u
INFO("y: " << std::fixed << std::setprecision(16) << y);
const auto ref_val = fixture.tex_h.Tex2DLayered(x, y, layer, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
}
@@ -132,7 +132,10 @@ TEMPLATE_TEST_CASE("Unit_tex2DLayered_Positive_ReadModeNormalizedFloat", "", cha
params.layers = 2;
params.num_subdivisions = 4;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1]) {
INFO("Different address modes on X, Y aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, true> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX());
@@ -172,12 +175,8 @@ TEMPLATE_TEST_CASE("Unit_tex2DLayered_Positive_ReadModeNormalizedFloat", "", cha
INFO("x: " << std::fixed << std::setprecision(16) << x);
INFO("y: " << std::fixed << std::setprecision(16) << y);
auto ref_val = Vec4Map<TestType>(fixture.tex_h.Tex2DLayered(x, y, layer, params.tex_desc),
NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
auto ref_val = fixture.tex_h.Tex2DLayered(x, y, layer, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
}
projects/hip-tests/catch/unit/texture/tex2DLayeredGrad.cc
+13 -14
Προβολή Αρχείου
@@ -57,8 +57,11 @@ TEMPLATE_TEST_CASE("Unit_tex2DLayeredGrad_Positive_ReadModeElementType", "", cha
params.layers = 2;
params.num_subdivisions = 4;
params.GenerateTextureDesc();
TextureTestFixture<TestType, false, true> fixture{params};
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1]) {
INFO("Different address modes on X, Y aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, false, false> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX());
const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(32, params.NumItersY());
@@ -98,10 +101,7 @@ TEMPLATE_TEST_CASE("Unit_tex2DLayeredGrad_Positive_ReadModeElementType", "", cha
INFO("y: " << std::fixed << std::setprecision(16) << y);
const auto ref_val = fixture.tex_h.Tex2DLayered(x, y, layer, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
}
@@ -132,8 +132,11 @@ TEMPLATE_TEST_CASE("Unit_tex2DLayeredGrad_Positive_ReadModeNormalizedFloat", "",
params.layers = 2;
params.num_subdivisions = 4;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
TextureTestFixture<TestType, true, true> fixture{params};
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1]) {
INFO("Different address modes on X, Y aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, true, false> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX());
const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(32, params.NumItersY());
@@ -172,12 +175,8 @@ TEMPLATE_TEST_CASE("Unit_tex2DLayeredGrad_Positive_ReadModeNormalizedFloat", "",
INFO("x: " << std::fixed << std::setprecision(16) << x);
INFO("y: " << std::fixed << std::setprecision(16) << y);
auto ref_val = Vec4Map<TestType>(fixture.tex_h.Tex2DLayered(x, y, layer, params.tex_desc),
NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
auto ref_val = fixture.tex_h.Tex2DLayered(x, y, layer, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
}
projects/hip-tests/catch/unit/texture/tex2DLayeredLod.cc
+13 -14
Προβολή Αρχείου
@@ -56,8 +56,11 @@ TEMPLATE_TEST_CASE("Unit_tex2DLayeredLod_Positive_ReadModeElementType", "", char
params.extent = make_hipExtent(16, 4, 0);
params.layers = 2;
params.num_subdivisions = 4;
params.GenerateTextureDesc();
params.GenerateTextureDesc(hipReadModeElementType, true);
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1]) {
INFO("Different address modes on X, Y aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, false, true> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX());
@@ -98,10 +101,7 @@ TEMPLATE_TEST_CASE("Unit_tex2DLayeredLod_Positive_ReadModeElementType", "", char
INFO("y: " << std::fixed << std::setprecision(16) << y);
const auto ref_val = fixture.tex_h.Tex2DLayered(x, y, layer, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
}
@@ -131,8 +131,11 @@ TEMPLATE_TEST_CASE("Unit_tex2DLayeredLod_Positive_ReadModeNormalizedFloat", "",
params.extent = make_hipExtent(16, 4, 0);
params.layers = 2;
params.num_subdivisions = 4;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
params.GenerateTextureDesc(hipReadModeNormalizedFloat, true);
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1]) {
INFO("Different address modes on X, Y aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, true, true> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX());
@@ -172,12 +175,8 @@ TEMPLATE_TEST_CASE("Unit_tex2DLayeredLod_Positive_ReadModeNormalizedFloat", "",
INFO("x: " << std::fixed << std::setprecision(16) << x);
INFO("y: " << std::fixed << std::setprecision(16) << y);
auto ref_val = Vec4Map<TestType>(fixture.tex_h.Tex2DLayered(x, y, layer, params.tex_desc),
NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
auto ref_val = fixture.tex_h.Tex2DLayered(x, y, layer, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
}
projects/hip-tests/catch/unit/texture/tex2DLod.cc
+13 -14
Προβολή Αρχείου
@@ -55,8 +55,11 @@ TEMPLATE_TEST_CASE("Unit_tex2DLod_Positive_ReadModeElementType", "", char, unsig
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(16, 4, 0);
params.num_subdivisions = 4;
params.GenerateTextureDesc();
params.GenerateTextureDesc(hipReadModeElementType, true);
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1]) {
INFO("Different address modes on X, Y aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, false, true> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX());
@@ -95,10 +98,7 @@ TEMPLATE_TEST_CASE("Unit_tex2DLod_Positive_ReadModeElementType", "", char, unsig
INFO("y: " << std::fixed << std::setprecision(16) << y);
const auto ref_val = fixture.tex_h.Tex2D(x, y, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
@@ -126,8 +126,11 @@ TEMPLATE_TEST_CASE("Unit_tex2DLod_Positive_ReadModeNormalizedFloat", "", char, u
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(16, 4, 0);
params.num_subdivisions = 4;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
params.GenerateTextureDesc(hipReadModeNormalizedFloat, true);
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1]) {
INFO("Different address modes on X, Y aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, true, true> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX());
@@ -165,12 +168,8 @@ TEMPLATE_TEST_CASE("Unit_tex2DLod_Positive_ReadModeNormalizedFloat", "", char, u
INFO("x: " << std::fixed << std::setprecision(16) << x);
INFO("y: " << std::fixed << std::setprecision(16) << y);
auto ref_val =
Vec4Map<TestType>(fixture.tex_h.Tex2D(x, y, params.tex_desc), NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
auto ref_val = fixture.tex_h.Tex2D(x, y, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
projects/hip-tests/catch/unit/texture/tex2Dgather.cc
+7 -5
Προβολή Αρχείου
@@ -56,7 +56,11 @@ TEMPLATE_TEST_CASE("Unit_tex2Dgather_Positive_ReadModeElementType", "", char, un
params.extent = make_hipExtent(16, 4, 0);
params.num_subdivisions = 4;
params.GenerateTextureDesc();
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1]) {
INFO("Different address modes on X, Y aren't supported. Skipped.");
return;
}
if (params.tex_desc.filterMode == hipFilterModeLinear) return;
TextureTestFixture<TestType> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(32, params.NumItersX());
@@ -94,12 +98,10 @@ TEMPLATE_TEST_CASE("Unit_tex2Dgather_Positive_ReadModeElementType", "", char, un
INFO("Address mode Y: " << AddressModeToString(params.tex_desc.addressMode[1]));
INFO("x: " << std::fixed << std::setprecision(16) << x);
INFO("y: " << std::fixed << std::setprecision(16) << y);
INFO("comp: " << comp);
const auto ref_val = fixture.tex_h.Tex2DGather(x, y, comp, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
projects/hip-tests/catch/unit/texture/tex3D.cc
+16 -14
Προβολή Αρχείου
@@ -51,12 +51,18 @@ THE SOFTWARE.
TEMPLATE_TEST_CASE("Unit_tex3D_Positive_ReadModeElementType", "", char, unsigned char, short,
unsigned short, int, unsigned int, float) {
CHECK_IMAGE_SUPPORT;
#if HT_NVIDIA
(void)hipGetLastError(); // Prevent negative tests affecting this
#endif
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(2, 4, 2);
params.num_subdivisions = 2;
params.GenerateTextureDesc();
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1] ||
params.tex_desc.addressMode[0] != params.tex_desc.addressMode[2]) {
INFO("Different address modes on X, Y, Z aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
@@ -104,10 +110,7 @@ TEMPLATE_TEST_CASE("Unit_tex3D_Positive_ReadModeElementType", "", char, unsigned
INFO("z: " << std::fixed << std::setprecision(16) << z);
const auto ref_val = fixture.tex_h.Tex3D(x, y, z, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
@@ -136,7 +139,11 @@ TEMPLATE_TEST_CASE("Unit_tex3D_Positive_ReadModeNormalizedFloat", "", char, unsi
params.extent = make_hipExtent(2, 2, 2);
params.num_subdivisions = 2;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1] ||
params.tex_desc.addressMode[0] != params.tex_desc.addressMode[2]) {
INFO("Different address modes on X, Y, Z aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, true> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
@@ -183,13 +190,8 @@ TEMPLATE_TEST_CASE("Unit_tex3D_Positive_ReadModeNormalizedFloat", "", char, unsi
INFO("y: " << std::fixed << std::setprecision(16) << y);
INFO("z: " << std::fixed << std::setprecision(16) << z);
auto ref_val = Vec4Map<TestType>(fixture.tex_h.Tex3D(x, y, z, params.tex_desc),
NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
auto ref_val = fixture.tex_h.Tex3D(x, y, z, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
projects/hip-tests/catch/unit/texture/tex3DGrad.cc
+15 -14
Προβολή Αρχείου
@@ -56,8 +56,12 @@ TEMPLATE_TEST_CASE("Unit_tex3DGrad_Positive_ReadModeElementType", "", char, unsi
params.extent = make_hipExtent(2, 2, 2);
params.num_subdivisions = 2;
params.GenerateTextureDesc();
TextureTestFixture<TestType, false, true> fixture{params};
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1] ||
params.tex_desc.addressMode[0] != params.tex_desc.addressMode[2]) {
INFO("Different address modes on X, Y, Z aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, false, false> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY());
@@ -104,10 +108,7 @@ TEMPLATE_TEST_CASE("Unit_tex3DGrad_Positive_ReadModeElementType", "", char, unsi
INFO("z: " << std::fixed << std::setprecision(16) << z);
const auto ref_val = fixture.tex_h.Tex3D(x, y, z, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
@@ -136,8 +137,12 @@ TEMPLATE_TEST_CASE("Unit_tex3DGrad_Positive_ReadModeNormalizedFloat", "", char,
params.extent = make_hipExtent(2, 2, 2);
params.num_subdivisions = 2;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
TextureTestFixture<TestType, true, true> fixture{params};
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1] ||
params.tex_desc.addressMode[0] != params.tex_desc.addressMode[2]) {
INFO("Different address modes on X, Y, Z aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, true, false> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
const auto [num_threads_y, num_blocks_y] = GetLaunchConfig(10, params.NumItersY());
@@ -183,12 +188,8 @@ TEMPLATE_TEST_CASE("Unit_tex3DGrad_Positive_ReadModeNormalizedFloat", "", char,
INFO("y: " << std::fixed << std::setprecision(16) << y);
INFO("z: " << std::fixed << std::setprecision(16) << z);
auto ref_val = Vec4Map<TestType>(fixture.tex_h.Tex3D(x, y, z, params.tex_desc),
NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
auto ref_val = fixture.tex_h.Tex3D(x, y, z, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
projects/hip-tests/catch/unit/texture/tex3DLod.cc
+15 -14
Προβολή Αρχείου
@@ -55,8 +55,12 @@ TEMPLATE_TEST_CASE("Unit_tex3DLod_Positive_ReadModeElementType", "", char, unsig
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(2, 2, 2);
params.num_subdivisions = 2;
params.GenerateTextureDesc();
params.GenerateTextureDesc(hipReadModeElementType, true);
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1] ||
params.tex_desc.addressMode[0] != params.tex_desc.addressMode[2]) {
INFO("Different address modes on X, Y, Z aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, false, true> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
@@ -104,10 +108,7 @@ TEMPLATE_TEST_CASE("Unit_tex3DLod_Positive_ReadModeElementType", "", char, unsig
INFO("z: " << std::fixed << std::setprecision(16) << z);
const auto ref_val = fixture.tex_h.Tex3D(x, y, z, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
@@ -135,8 +136,12 @@ TEMPLATE_TEST_CASE("Unit_tex3DLod_Positive_ReadModeNormalizedFloat", "", char, u
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(2, 2, 2);
params.num_subdivisions = 2;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
params.GenerateTextureDesc(hipReadModeNormalizedFloat, true);
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1] ||
params.tex_desc.addressMode[0] != params.tex_desc.addressMode[2]) {
INFO("Different address modes on X, Y, Z aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, true, true> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
@@ -183,12 +188,8 @@ TEMPLATE_TEST_CASE("Unit_tex3DLod_Positive_ReadModeNormalizedFloat", "", char, u
INFO("y: " << std::fixed << std::setprecision(16) << y);
INFO("z: " << std::fixed << std::setprecision(16) << z);
auto ref_val = Vec4Map<TestType>(fixture.tex_h.Tex3D(x, y, z, params.tex_desc),
NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[i].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[i].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[i].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[i].w);
auto ref_val = fixture.tex_h.Tex3D(x, y, z, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[i], ref_val));
}
}
projects/hip-tests/catch/unit/texture/texCubemap.cc
+17 -14
Προβολή Αρχείου
@@ -51,13 +51,18 @@ THE SOFTWARE.
TEMPLATE_TEST_CASE("Unit_texCubemap_Positive_ReadModeElementType", "", char, unsigned char, short,
unsigned short, int, unsigned int, float) {
CHECK_IMAGE_SUPPORT;
INFO("texCubemap isn't supported. Skipped.");
return;
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(2, 2, 6);
params.num_subdivisions = 4;
params.cubemap = true;
params.GenerateTextureDesc();
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1] ||
params.tex_desc.addressMode[0] != params.tex_desc.addressMode[2]) {
INFO("Different address modes on X, Y, Z aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
@@ -107,10 +112,7 @@ TEMPLATE_TEST_CASE("Unit_texCubemap_Positive_ReadModeElementType", "", char, uns
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
const auto ref_val = fixture.tex_h.TexCubemap(x, y, z, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[index], ref_val));
}
}
}
@@ -136,13 +138,18 @@ TEMPLATE_TEST_CASE("Unit_texCubemap_Positive_ReadModeElementType", "", char, uns
TEMPLATE_TEST_CASE("Unit_texCubemap_Positive_ReadModeNormalizedFloat", "", char, unsigned char,
short, unsigned short) {
CHECK_IMAGE_SUPPORT;
INFO("texCubemap isn't supported. Skipped.");
return;
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(2, 2, 6);
params.num_subdivisions = 4;
params.cubemap = true;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1] ||
params.tex_desc.addressMode[0] != params.tex_desc.addressMode[2]) {
INFO("Different address modes on X, Y, Z aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, true> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
@@ -191,12 +198,8 @@ TEMPLATE_TEST_CASE("Unit_texCubemap_Positive_ReadModeNormalizedFloat", "", char,
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
auto ref_val = Vec4Map<TestType>(fixture.tex_h.TexCubemap(x, y, z, params.tex_desc),
NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
auto ref_val = fixture.tex_h.TexCubemap(x, y, z, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[index], ref_val));
}
}
}
projects/hip-tests/catch/unit/texture/texCubemapGrad.cc
+17 -14
Προβολή Αρχείου
@@ -51,13 +51,18 @@ THE SOFTWARE.
TEMPLATE_TEST_CASE("Unit_texCubemapGrad_Positive_ReadModeElementType", "", char, unsigned char,
short, unsigned short, int, unsigned int, float) {
CHECK_IMAGE_SUPPORT;
INFO("texCubemap isn't supported. Skipped.");
return;
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(2, 2, 6);
params.num_subdivisions = 4;
params.cubemap = true;
params.GenerateTextureDesc();
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1] ||
params.tex_desc.addressMode[0] != params.tex_desc.addressMode[2]) {
INFO("Different address modes on X, Y, Z aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, false, true> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
@@ -107,10 +112,7 @@ TEMPLATE_TEST_CASE("Unit_texCubemapGrad_Positive_ReadModeElementType", "", char,
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
const auto ref_val = fixture.tex_h.TexCubemap(x, y, z, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[index], ref_val));
}
}
}
@@ -136,13 +138,18 @@ TEMPLATE_TEST_CASE("Unit_texCubemapGrad_Positive_ReadModeElementType", "", char,
TEMPLATE_TEST_CASE("Unit_texCubemapGrad_Positive_ReadModeNormalizedFloat", "", char, unsigned char,
short, unsigned short) {
CHECK_IMAGE_SUPPORT;
INFO("texCubemap isn't supported. Skipped.");
return;
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(2, 2, 6);
params.num_subdivisions = 4;
params.cubemap = true;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1] ||
params.tex_desc.addressMode[0] != params.tex_desc.addressMode[2]) {
INFO("Different address modes on X, Y, Z aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, true, true> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
@@ -191,12 +198,8 @@ TEMPLATE_TEST_CASE("Unit_texCubemapGrad_Positive_ReadModeNormalizedFloat", "", c
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
auto ref_val = Vec4Map<TestType>(fixture.tex_h.TexCubemap(x, y, z, params.tex_desc),
NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
auto ref_val = fixture.tex_h.TexCubemap(x, y, z, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[index], ref_val));
}
}
}
projects/hip-tests/catch/unit/texture/texCubemapLayered.cc
+17 -14
Προβολή Αρχείου
@@ -51,14 +51,19 @@ THE SOFTWARE.
TEMPLATE_TEST_CASE("Unit_texCubemapLayered_Positive_ReadModeElementType", "", char, unsigned char,
short, unsigned short, int, unsigned int, float) {
CHECK_IMAGE_SUPPORT;
INFO("texCubemap isn't supported. Skipped.");
return;
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(2, 2, 6);
params.num_subdivisions = 4;
params.layers = 1;
params.cubemap = true;
params.GenerateTextureDesc();
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1] ||
params.tex_desc.addressMode[0] != params.tex_desc.addressMode[2]) {
INFO("Different address modes on X, Y, Z aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
@@ -109,10 +114,7 @@ TEMPLATE_TEST_CASE("Unit_texCubemapLayered_Positive_ReadModeElementType", "", ch
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
const auto ref_val = fixture.tex_h.TexCubemap(x, y, z, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[index], ref_val));
}
}
}
@@ -139,14 +141,19 @@ TEMPLATE_TEST_CASE("Unit_texCubemapLayered_Positive_ReadModeElementType", "", ch
TEMPLATE_TEST_CASE("Unit_texCubemapLayered_Positive_ReadModeNormalizedFloat", "", char,
unsigned char, short, unsigned short) {
CHECK_IMAGE_SUPPORT;
INFO("texCubemap isn't supported. Skipped.");
return;
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(2, 2, 6);
params.num_subdivisions = 4;
params.layers = 1;
params.cubemap = true;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1] ||
params.tex_desc.addressMode[0] != params.tex_desc.addressMode[2]) {
INFO("Different address modes on X, Y, Z aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, true> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
@@ -196,12 +203,8 @@ TEMPLATE_TEST_CASE("Unit_texCubemapLayered_Positive_ReadModeNormalizedFloat", ""
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
auto ref_val = Vec4Map<TestType>(fixture.tex_h.TexCubemap(x, y, z, params.tex_desc),
NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
auto ref_val = fixture.tex_h.TexCubemap(x, y, z, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[index], ref_val));
}
}
}
projects/hip-tests/catch/unit/texture/texCubemapLayeredGrad.cc
+17 -14
Προβολή Αρχείου
@@ -51,14 +51,19 @@ THE SOFTWARE.
TEMPLATE_TEST_CASE("Unit_texCubemapLayeredGrad_Positive_ReadModeElementType", "", char,
unsigned char, short, unsigned short, int, unsigned int, float) {
CHECK_IMAGE_SUPPORT;
INFO("texCubemap isn't supported. Skipped.");
return;
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(2, 2, 6);
params.num_subdivisions = 4;
params.layers = 1;
params.cubemap = true;
params.GenerateTextureDesc();
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1] ||
params.tex_desc.addressMode[0] != params.tex_desc.addressMode[2]) {
INFO("Different address modes on X, Y, Z aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, false, true> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
@@ -110,10 +115,7 @@ TEMPLATE_TEST_CASE("Unit_texCubemapLayeredGrad_Positive_ReadModeElementType", ""
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
const auto ref_val = fixture.tex_h.TexCubemap(x, y, z, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[index], ref_val));
}
}
}
@@ -140,14 +142,19 @@ TEMPLATE_TEST_CASE("Unit_texCubemapLayeredGrad_Positive_ReadModeElementType", ""
TEMPLATE_TEST_CASE("Unit_texCubemapLayeredGrad_Positive_ReadModeNormalizedFloat", "", char,
unsigned char, short, unsigned short) {
CHECK_IMAGE_SUPPORT;
INFO("texCubemap isn't supported. Skipped.");
return;
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(2, 2, 6);
params.num_subdivisions = 4;
params.layers = 1;
params.cubemap = true;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1] ||
params.tex_desc.addressMode[0] != params.tex_desc.addressMode[2]) {
INFO("Different address modes on X, Y, Z aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, true, true> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
@@ -198,12 +205,8 @@ TEMPLATE_TEST_CASE("Unit_texCubemapLayeredGrad_Positive_ReadModeNormalizedFloat"
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
auto ref_val = Vec4Map<TestType>(fixture.tex_h.TexCubemap(x, y, z, params.tex_desc),
NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
auto ref_val = fixture.tex_h.TexCubemap(x, y, z, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[index], ref_val));
}
}
}
projects/hip-tests/catch/unit/texture/texCubemapLayeredLod.cc
+18 -15
Προβολή Αρχείου
@@ -51,14 +51,19 @@ THE SOFTWARE.
TEMPLATE_TEST_CASE("Unit_texCubemapLayeredLod_Positive_ReadModeElementType", "", char,
unsigned char, short, unsigned short, int, unsigned int, float) {
CHECK_IMAGE_SUPPORT;
INFO("texCubemap isn't supported. Skipped.");
return;
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(2, 2, 6);
params.num_subdivisions = 4;
params.layers = 1;
params.cubemap = true;
params.GenerateTextureDesc();
params.GenerateTextureDesc(hipReadModeElementType, true);
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1] ||
params.tex_desc.addressMode[0] != params.tex_desc.addressMode[2]) {
INFO("Different address modes on X, Y, Z aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, false, true> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
@@ -110,10 +115,7 @@ TEMPLATE_TEST_CASE("Unit_texCubemapLayeredLod_Positive_ReadModeElementType", "",
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
const auto ref_val = fixture.tex_h.TexCubemap(x, y, z, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[index], ref_val));
}
}
}
@@ -140,14 +142,19 @@ TEMPLATE_TEST_CASE("Unit_texCubemapLayeredLod_Positive_ReadModeElementType", "",
TEMPLATE_TEST_CASE("Unit_texCubemapLayeredLod_Positive_ReadModeNormalizedFloat", "", char,
unsigned char, short, unsigned short) {
CHECK_IMAGE_SUPPORT;
INFO("texCubemap isn't supported. Skipped.");
return;
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(2, 2, 6);
params.num_subdivisions = 4;
params.layers = 1;
params.cubemap = true;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1] ||
params.tex_desc.addressMode[0] != params.tex_desc.addressMode[2]) {
INFO("Different address modes on X, Y, Z aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, true, true> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
@@ -198,12 +205,8 @@ TEMPLATE_TEST_CASE("Unit_texCubemapLayeredLod_Positive_ReadModeNormalizedFloat",
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
auto ref_val = Vec4Map<TestType>(fixture.tex_h.TexCubemap(x, y, z, params.tex_desc),
NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
auto ref_val = fixture.tex_h.TexCubemap(x, y, z, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[index], ref_val));
}
}
}
projects/hip-tests/catch/unit/texture/texCubemapLod.cc
+19 -15
Προβολή Αρχείου
@@ -51,13 +51,18 @@ THE SOFTWARE.
TEMPLATE_TEST_CASE("Unit_texCubemapLod_Positive_ReadModeElementType", "", char, unsigned char,
short, unsigned short, int, unsigned int, float) {
CHECK_IMAGE_SUPPORT;
INFO("texCubemap isn't supported. Skipped.");
return;
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(2, 2, 6);
params.num_subdivisions = 4;
params.cubemap = true;
params.GenerateTextureDesc();
params.GenerateTextureDesc(hipReadModeElementType, true);
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1] ||
params.tex_desc.addressMode[0] != params.tex_desc.addressMode[2]) {
INFO("Different address modes on X, Y, Z aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, false, true> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
@@ -107,10 +112,8 @@ TEMPLATE_TEST_CASE("Unit_texCubemapLod_Positive_ReadModeElementType", "", char,
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
const auto ref_val = fixture.tex_h.TexCubemap(x, y, z, params.tex_desc);
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
REQUIRE(fixture.Verify(fixture.out_alloc_h[index], ref_val));
}
}
}
@@ -136,13 +139,18 @@ TEMPLATE_TEST_CASE("Unit_texCubemapLod_Positive_ReadModeElementType", "", char,
TEMPLATE_TEST_CASE("Unit_texCubemapLod_Positive_ReadModeNormalizedFloat", "", char, unsigned char,
short, unsigned short) {
CHECK_IMAGE_SUPPORT;
INFO("texCubemap isn't supported. Skipped.");
return;
TextureTestParams<TestType> params = {};
params.extent = make_hipExtent(2, 2, 6);
params.num_subdivisions = 4;
params.cubemap = true;
params.GenerateTextureDesc(hipReadModeNormalizedFloat);
if (params.tex_desc.addressMode[0] != params.tex_desc.addressMode[1] ||
params.tex_desc.addressMode[0] != params.tex_desc.addressMode[2]) {
INFO("Different address modes on X, Y, Z aren't supported. Skipped.");
return;
}
TextureTestFixture<TestType, true, true> fixture{params};
const auto [num_threads_x, num_blocks_x] = GetLaunchConfig(10, params.NumItersX());
@@ -191,12 +199,8 @@ TEMPLATE_TEST_CASE("Unit_texCubemapLod_Positive_ReadModeNormalizedFloat", "", ch
auto index = k * params.NumItersX() * params.NumItersY() + j * params.NumItersX() + i;
auto ref_val = Vec4Map<TestType>(fixture.tex_h.TexCubemap(x, y, z, params.tex_desc),
NormalizeInteger<TestType>);
REQUIRE(ref_val.x == fixture.out_alloc_h[index].x);
REQUIRE(ref_val.y == fixture.out_alloc_h[index].y);
REQUIRE(ref_val.z == fixture.out_alloc_h[index].z);
REQUIRE(ref_val.w == fixture.out_alloc_h[index].w);
auto ref_val = fixture.tex_h.TexCubemap(x, y, z, params.tex_desc);
REQUIRE(fixture.Verify(fixture.out_alloc_h[index], ref_val));
}
}
}
projects/hip-tests/catch/unit/texture/texture_reference.hh
+93 -54
Προβολή Αρχείου
@@ -23,35 +23,40 @@ THE SOFTWARE.
#pragma once
#include <cmath>
#include "fixed_point.hh"
#if defined(_WIN64)
typedef __int64 ssize_t;
#endif // _WIN64
template <typename TexelType> class TextureReference {
template <typename TexelType, bool normalized_read>
class TextureReference {
using valType = decltype(TexelType::x);
static constexpr bool supportFilter = normalized_read || std::is_floating_point<valType>::value;
public:
TextureReference(TexelType* alloc, hipExtent extent, size_t layers)
: alloc_{alloc}, extent_{extent}, layers_{layers} {}
: alloc_(alloc), extent_{extent}, layers_{layers} {}
TexelType Tex1D(float x, const hipTextureDesc& tex_desc) const {
auto Tex1D(float x, const hipTextureDesc& tex_desc) const {
return Tex1DLayered(x, 0, tex_desc);
}
TexelType Tex2DGather(float x, float y, int comp, const hipTextureDesc& tex_desc) const {
auto Tex2DGather(float x, float y, int comp, const hipTextureDesc& tex_desc) const {
x = tex_desc.normalizedCoords ? x * extent_.width : x;
y = tex_desc.normalizedCoords ? y * extent_.height : y;
#if HT_AMD
const auto [i, alpha] = GetLinearFilteringParams(x - 0.5f);
const auto [j, beta] = GetLinearFilteringParams(y - 0.5f);
#else
const auto [i, alpha] = GetLinearFilteringParams(x);
const auto [j, beta] = GetLinearFilteringParams(y);
#endif
const auto T_i0j0 = Sample(i, j, 0, tex_desc.addressMode);
const auto T_i1j0 = Sample(i + 1.0f, j, 0, tex_desc.addressMode);
const auto T_i0j1 = Sample(i, j + 1.0f, 0, tex_desc.addressMode);
const auto T_i1j1 = Sample(i + 1.0f, j + 1.0f, 0, tex_desc.addressMode);
const auto IndexVec4 = [](auto vec, int comp) {
const auto IndexVec4 = [](auto& vec, int comp) {
switch (comp) {
case 0:
return vec.x;
@@ -66,33 +71,36 @@ template <typename TexelType> class TextureReference {
}
};
TexelType texel;
texel.x = IndexVec4(T_i0j1, comp);
texel.y = IndexVec4(T_i1j1, comp);
texel.z = IndexVec4(T_i1j0, comp);
texel.w = IndexVec4(T_i0j0, comp);
decltype(T_i0j0) texel {
IndexVec4(T_i0j1, comp),
IndexVec4(T_i1j1, comp),
IndexVec4(T_i1j0, comp),
IndexVec4(T_i0j0, comp)};
return texel;
}
TexelType Tex2D(float x, float y, const hipTextureDesc& tex_desc) const {
auto Tex2D(float x, float y, const hipTextureDesc& tex_desc) const {
return Tex2DLayered(x, y, 0, tex_desc);
}
TexelType Tex3D(float x, float y, float z, const hipTextureDesc& tex_desc) const {
auto Tex3D(float x, float y, float z, const hipTextureDesc& tex_desc) const {
x = tex_desc.normalizedCoords ? x * extent_.width : x;
y = tex_desc.normalizedCoords ? y * extent_.height : y;
z = tex_desc.normalizedCoords ? z * extent_.depth : z;
if (tex_desc.filterMode == hipFilterModePoint) {
return Sample(floorf(x), floorf(y), floorf(z), tex_desc.addressMode);
} else if (tex_desc.filterMode == hipFilterModeLinear) {
return LinearFiltering(x, y, z, tex_desc.addressMode);
if constexpr (supportFilter) {
return LinearFiltering(x, y, z, tex_desc.addressMode);
} else {
throw std::invalid_argument("hipFilterModeLinear not supported");
}
} else {
throw std::invalid_argument("Invalid hipFilterMode value");
}
}
TexelType TexCubemap(float x, float y, float z, const hipTextureDesc& tex_desc) const {
auto TexCubemap(float x, float y, float z, const hipTextureDesc& tex_desc) const {
x = tex_desc.normalizedCoords ? x * extent_.width : x;
y = tex_desc.normalizedCoords ? y * extent_.height : y;
z = tex_desc.normalizedCoords ? z * extent_.depth : z;
@@ -144,30 +152,42 @@ template <typename TexelType> class TextureReference {
if (tex_desc.filterMode == hipFilterModePoint) {
return Sample(roundf(coord1), roundf(coord2), face, tex_desc.addressMode);
} else if (tex_desc.filterMode == hipFilterModeLinear) {
return LinearFiltering(coord1, coord2, face, tex_desc.addressMode);
if constexpr (supportFilter) {
return LinearFiltering(coord1, coord2, face, tex_desc.addressMode);
} else {
throw std::invalid_argument("hipFilterModeLinear not supported");
}
} else {
throw std::invalid_argument("Invalid hipFilterMode value");
}
}
TexelType Tex1DLayered(float x, int layer, const hipTextureDesc& tex_desc) const {
auto Tex1DLayered(float x, int layer, const hipTextureDesc& tex_desc) const {
x = tex_desc.normalizedCoords ? x * extent_.width : x;
if (tex_desc.filterMode == hipFilterModePoint) {
return Sample(floorf(x), layer, tex_desc.addressMode);
} else if (tex_desc.filterMode == hipFilterModeLinear) {
return LinearFiltering(x, layer, tex_desc.addressMode);
if constexpr (supportFilter) {
return LinearFiltering(x, layer, tex_desc.addressMode);
} else {
throw std::invalid_argument("hipFilterModeLinear not supported");
}
} else {
throw std::invalid_argument("Invalid hipFilterMode value");
}
}
TexelType Tex2DLayered(float x, float y, int layer, const hipTextureDesc& tex_desc) const {
auto Tex2DLayered(float x, float y, int layer, const hipTextureDesc& tex_desc) const {
x = tex_desc.normalizedCoords ? x * extent_.width : x;
y = tex_desc.normalizedCoords ? y * extent_.height : y;
if (tex_desc.filterMode == hipFilterModePoint) {
return Sample(floorf(x), floorf(y), layer, tex_desc.addressMode);
} else if (tex_desc.filterMode == hipFilterModeLinear) {
return LinearFiltering(x, y, layer, tex_desc.addressMode);
if constexpr (supportFilter) {
return LinearFiltering(x, y, layer, tex_desc.addressMode);
} else {
throw std::invalid_argument("hipFilterModeLinear not supported");
}
} else {
throw std::invalid_argument("Invalid hipFilterMode value");
}
@@ -188,15 +208,13 @@ template <typename TexelType> class TextureReference {
const hipExtent extent_;
const size_t layers_;
template <typename T> TexelType Vec4Sum(T arg) const { return Vec4Add(arg, Zero()); }
template <typename T, typename... Ts> TexelType Vec4Sum(T arg, Ts... args) const {
return Vec4Add(arg, Vec4Sum(args...));
TexelType Zero() const {
TexelType ret {0, 0, 0, 0};
return ret;
}
TexelType Zero() const {
TexelType ret;
memset(&ret, 0, sizeof(ret));
vec4<float> Zerof() const {
vec4<float> ret {0., 0., 0., 0.};
return ret;
}
@@ -217,41 +235,63 @@ template <typename TexelType> class TextureReference {
}
}
TexelType Sample(float x, int layer, const hipTextureAddressMode* address_mode) const {
auto Sample(float x, int layer, const hipTextureAddressMode* address_mode) const {
x = ApplyAddressMode(x, extent_.width, address_mode[0]);
if (std::isnan(x)) {
return Zero();
if constexpr (normalized_read) {
if (std::isnan(x)) {
return Zerof();
}
return Vec4Map(ptr(layer)[static_cast<size_t>(x)]);
} else {
if (std::isnan(x)) {
return Zero();
}
return ptr(layer)[static_cast<size_t>(x)];
}
return ptr(layer)[static_cast<size_t>(x)];
}
TexelType Sample(float x, float y, int layer, const hipTextureAddressMode* address_mode) const {
auto Sample(float x, float y, int layer, const hipTextureAddressMode* address_mode) const {
x = ApplyAddressMode(x, extent_.width, address_mode[0]);
y = ApplyAddressMode(y, extent_.height, address_mode[1]);
if (std::isnan(x) || std::isnan(y)) {
return Zero();
if constexpr (normalized_read) {
if (std::isnan(x) || std::isnan(y)) {
return Zerof();
}
return Vec4Map(
ptr(layer)[static_cast<size_t>(y) * extent_.width + static_cast<size_t>(x)]);
} else {
if (std::isnan(x) || std::isnan(y)) {
return Zero();
}
return ptr(layer)[static_cast<size_t>(y) * extent_.width + static_cast<size_t>(x)];
}
return ptr(layer)[static_cast<size_t>(y) * extent_.width + static_cast<size_t>(x)];
}
TexelType Sample(float x, float y, float z, const hipTextureAddressMode* address_mode) const {
auto Sample(float x, float y, float z, const hipTextureAddressMode* address_mode) const {
x = ApplyAddressMode(x, extent_.width, address_mode[0]);
y = ApplyAddressMode(y, extent_.height, address_mode[1]);
z = ApplyAddressMode(z, extent_.depth, address_mode[2]);
if (std::isnan(x) || std::isnan(y) || std::isnan(z)) {
return Zero();
if constexpr (normalized_read) {
if (std::isnan(x) || std::isnan(y) || std::isnan(z)) {
return Zerof();
}
return Vec4Map(
ptr(0)[static_cast<size_t>(z) * extent_.width * extent_.height +
static_cast<size_t>(y) * extent_.width + static_cast<size_t>(x)]);
} else {
if (std::isnan(x) || std::isnan(y) || std::isnan(z)) {
return Zero();
}
return ptr(0)[static_cast<size_t>(z) * extent_.width * extent_.height +
static_cast<size_t>(y) * extent_.width + static_cast<size_t>(x)];
}
return ptr(0)[static_cast<size_t>(z) * extent_.width * extent_.height +
static_cast<size_t>(y) * extent_.width + static_cast<size_t>(x)];
}
TexelType LinearFiltering(float x, int layer, const hipTextureAddressMode* address_mode) const {
// LinearFiltering won't be called when valType isn't float or normalized_read is false
auto LinearFiltering(float x, int layer, const hipTextureAddressMode* address_mode) const {
const auto [i, alpha] = GetLinearFilteringParams(x);
const auto T_i0 = Sample(i, layer, address_mode);
@@ -259,11 +299,10 @@ template <typename TexelType> class TextureReference {
const auto term_i0 = Vec4Scale((1.0f - alpha), T_i0);
const auto term_i1 = Vec4Scale(alpha, T_i1);
return Vec4Sum(term_i0, term_i1);
return term_i0 + term_i1;
}
TexelType LinearFiltering(float x, float y, int layer,
auto LinearFiltering(float x, float y, int layer,
const hipTextureAddressMode* address_mode) const {
const auto [i, alpha] = GetLinearFilteringParams(x);
const auto [j, beta] = GetLinearFilteringParams(y);
@@ -278,10 +317,10 @@ template <typename TexelType> class TextureReference {
const auto term_i0j1 = Vec4Scale((1.0f - alpha) * beta, T_i0j1);
const auto term_i1j1 = Vec4Scale(alpha * beta, T_i1j1);
return Vec4Sum(term_i0j0, term_i1j0, term_i0j1, term_i1j1);
return term_i0j0 + term_i1j0 + term_i0j1 + term_i1j1;
}
TexelType LinearFiltering(float x, float y, float z,
auto LinearFiltering(float x, float y, float z,
const hipTextureAddressMode* address_mode) const {
const auto [i, alpha] = GetLinearFilteringParams(x);
const auto [j, beta] = GetLinearFilteringParams(y);
@@ -305,8 +344,8 @@ template <typename TexelType> class TextureReference {
const auto term_i0j1k1 = Vec4Scale((1.0f - alpha) * beta * gamma, T_i0j1k1);
const auto term_i1j1k1 = Vec4Scale(alpha * beta * gamma, T_i1j1k1);
return Vec4Sum(term_i0j0k0, term_i1j0k0, term_i0j1k0, term_i1j1k0, term_i0j0k1, term_i1j0k1,
term_i0j1k1, term_i1j1k1);
return term_i0j0k0 + term_i1j0k0 + term_i0j1k0 + term_i1j1k0 + term_i0j0k1 + term_i1j0k1 +
term_i0j1k1 + term_i1j1k1;
}
float ApplyClamp(float coord, size_t dim) const {
projects/hip-tests/catch/unit/texture/vec4.hh
+7 -8
Προβολή Αρχείου
@@ -75,13 +75,12 @@ template <typename T> inline void MakeVec4(const T x, const T y, const T z, cons
return vec;
}
template <typename T, typename F> inline auto Vec4Map(const vec4<T>& vec, F f) {
vec4<decltype(f(vec.x))> ret;
ret.x = f(vec.x);
ret.y = f(vec.y);
ret.z = f(vec.z);
ret.w = f(vec.w);
template <typename T> inline vec4<float> Vec4Map(const T& vec) {
vec4<float> ret;
ret.x = NormalizeInteger(vec.x);
ret.y = NormalizeInteger(vec.y);
ret.z = NormalizeInteger(vec.z);
ret.w = NormalizeInteger(vec.w);
return ret;
}
@@ -103,4 +102,4 @@ template <typename T> inline __host__ __device__ auto Vec4Add(const T& vec1, con
ret.w = vec1.w + vec2.w;
return ret;
}
}