From cc134f7c580c2ef2b15733eb8dcf306d98c651ff Mon Sep 17 00:00:00 2001 From: Vladislav Sytchenko Date: Fri, 13 Mar 2020 17:35:25 -0400 Subject: [PATCH] Rework hipNormalizedFloatValueTex test This is currently so buggy that it causes a runtime crash on Nvidia platfrom... Disable the new version for hcc and vdi, header fixes are required for it to pass. Currently tex1D returns a char, when the actual sampled pixel value is a float, so the hi 3 bytes get lost. Change-Id: I8222a4d8d1d8b101eb43f3f8dfbe4818f885f8ea --- .../texture/hipNormalizedFloatValueTex.cpp | 112 ++++++++++-------- 1 file changed, 65 insertions(+), 47 deletions(-) diff --git a/tests/src/texture/hipNormalizedFloatValueTex.cpp b/tests/src/texture/hipNormalizedFloatValueTex.cpp index 04660ad70b..b4aa3e9c05 100644 --- a/tests/src/texture/hipNormalizedFloatValueTex.cpp +++ b/tests/src/texture/hipNormalizedFloatValueTex.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc hcc vdi * TEST: %t * HIT_END */ @@ -30,73 +30,92 @@ THE SOFTWARE. #define SIZE 10 static float getNormalizedValue(const float value, - const enum hipArray_Format texFormat) { - switch (texFormat) { - case HIP_AD_FORMAT_SIGNED_INT8: - return (value / SCHAR_MAX); - case HIP_AD_FORMAT_UNSIGNED_INT8: - return (value / UCHAR_MAX); - case HIP_AD_FORMAT_SIGNED_INT16: - return (value / SHRT_MAX); - case HIP_AD_FORMAT_UNSIGNED_INT16: - return (value / USHRT_MAX); - default: - return value; - } + const hipChannelFormatDesc& desc) { + if ((desc.x == 8) && (desc.f == hipChannelFormatKindSigned)) + return (value / SCHAR_MAX); + if ((desc.x == 8) && (desc.f == hipChannelFormatKindUnsigned)) + return (value / UCHAR_MAX); + if ((desc.x == 16) && (desc.f == hipChannelFormatKindSigned)) + return (value / SHRT_MAX); + if ((desc.x == 16) && (desc.f == hipChannelFormatKindUnsigned)) + return (value / USHRT_MAX); + return value; } #if __HIP__ __hip_pinned_shadow__ #endif -texture textureNormalizedVal_1D; +texture texc; +#if __HIP__ +__hip_pinned_shadow__ +#endif +texture texuc; + +#if __HIP__ +__hip_pinned_shadow__ +#endif +texture texs; + +#if __HIP__ +__hip_pinned_shadow__ +#endif +texture texus; + + +template __global__ void normalizedValTextureTest(unsigned int numElements, float* pDst) { unsigned int elementID = hipThreadIdx_x; if(elementID >= numElements) - return; + return; float coord =(float) elementID/numElements; - pDst[elementID] = tex1D(textureNormalizedVal_1D, coord); + if(std::is_same::value) + pDst[elementID] = tex1D(texc, coord); + else if(std::is_same::value) + pDst[elementID] = tex1D(texuc, coord); + else if(std::is_same::value) + pDst[elementID] = tex1D(texs, coord); + else if(std::is_same::value) + pDst[elementID] = tex1D(texus, coord); } template -bool textureTest(enum hipArray_Format texFormat) +bool textureTest(texture *tex) { + hipChannelFormatDesc desc = hipCreateChannelDesc(); + hipArray_t dData; + HIPCHECK(hipMallocArray(&dData, &desc, SIZE, 1, hipArrayDefault)); + T hData[] = {65, 66, 67, 68, 69, 70, 71, 72, 73, 74}; - T *dData = NULL; - HIPCHECK(hipMalloc((void **) &dData, sizeof(T)*SIZE)); - HIPCHECK(hipMemcpyHtoD((hipDeviceptr_t)dData, hData, sizeof(T)*SIZE)); - textureReference* texRef = &textureNormalizedVal_1D; - HIPCHECK(hipTexRefSetAddressMode(texRef, 0, hipAddressModeClamp)); - HIPCHECK(hipTexRefSetAddressMode(texRef, 1, hipAddressModeClamp)); - HIPCHECK(hipTexRefSetFilterMode(texRef, hipFilterModePoint)); - HIPCHECK(hipTexRefSetFlags(texRef, HIP_TRSF_NORMALIZED_COORDINATES)); - HIPCHECK(hipTexRefSetFormat(texRef, texFormat, 1)); - - size_t offSet = 0; - HIPCHECK(hipTexRefSetAddress(&offSet, texRef, (hipDeviceptr_t)dData, sizeof(T)*SIZE)); - - bool testResult = true; + HIPCHECK(hipMemcpy2DToArray(dData, 0, 0, hData, sizeof(T)*SIZE, sizeof(T)*SIZE, 1, hipMemcpyHostToDevice)); + + tex->normalized = true; + tex->channelDesc = desc; + HIPCHECK(hipBindTextureToArray(tex, dData, &desc)); + float *dOutputData = NULL; HIPCHECK(hipMalloc((void **) &dOutputData, sizeof(float)*SIZE)); - - hipLaunchKernelGGL(HIP_KERNEL_NAME(normalizedValTextureTest), dim3(1,1,1), dim3(SIZE,1,1), 0, 0, SIZE, dOutputData); + + hipLaunchKernelGGL(normalizedValTextureTest, dim3(1,1,1), dim3(SIZE,1,1), 0, 0, SIZE, dOutputData); float *hOutputData = new float[SIZE]; - HIPCHECK(hipMemcpyDtoH(hOutputData, (hipDeviceptr_t)dOutputData, (sizeof(float)*SIZE))); - + HIPCHECK(hipMemcpy(hOutputData, dOutputData, (sizeof(float)*SIZE), hipMemcpyDeviceToHost)); + + bool testResult = true; for(int i = 0; i < SIZE; i++) { - float expected = getNormalizedValue(float(hData[i]), texFormat); + float expected = getNormalizedValue(float(hData[i]), desc); if(expected != hOutputData[i]) { - printf("mismatch at index:%d for texType:%d output:%f\n",i,texFormat,hOutputData[i]); + printf("mismatch at index:%d output:%f expected:%f\n",i,hOutputData[i],expected); testResult = false; - break; + break; } } - hipFree(dData); - hipFree(dOutputData); + + HIPCHECK(hipFreeArray(dData)); + HIPCHECK(hipFree(dOutputData)); delete [] hOutputData; return testResult; } @@ -113,12 +132,11 @@ int main(int argc, char** argv) std::cout << "Arch - AMD GPU :: " << props.gcnArch << std::endl; #endif - status &= textureTest (HIP_AD_FORMAT_SIGNED_INT8); - status &= textureTest (HIP_AD_FORMAT_UNSIGNED_INT8); - status &= textureTest (HIP_AD_FORMAT_SIGNED_INT16); - status &= textureTest(HIP_AD_FORMAT_UNSIGNED_INT16); - status &= textureTest (HIP_AD_FORMAT_FLOAT); - + status &= textureTest (&texc); + status &= textureTest (&texuc); + status &= textureTest (&texs); + status &= textureTest(&texus); + if(status){ passed(); }