AI/rocm-systems
2
0
Rozštěpit 0
Zdrojový kód Úkoly Pull requesty Akce Balíčky Projekty Vydání Wiki Aktivita

Rework hipNormalizedFloatValueTex test

Procházet zdrojové kódy 
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<char, hipTextureType1D, hipReadModeNormalizedFloat> returns a char, when the actual sampled pixel value is a float, so the hi 3 bytes get lost.

Change-Id: I8222a4d8d1d8b101eb43f3f8dfbe4818f885f8ea
Tento commit je obsažen v:
Vladislav Sytchenko
2020-03-13 17:35:25 -04:00
odevzdal Vladislav Sytchenko
rodič 9343eee07c
revize cc134f7c58
1 změnil soubory, kde provedl 65 přidání a 47 odebrání
Stáhněte soubor záplaty Stáhněte rozdílový soubor Rozbalit všechny soubory Sbalit všechny soubory
tests/src/texture/hipNormalizedFloatValueTex.cpp
+65 -47
Zobrazit soubor
@@ -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<float, hipTextureType1D, hipReadModeElementType> textureNormalizedVal_1D;
texture<char, hipTextureType1D, hipReadModeNormalizedFloat> texc;
#if __HIP__
__hip_pinned_shadow__
#endif
texture<unsigned char, hipTextureType1D, hipReadModeNormalizedFloat> texuc;
#if __HIP__
__hip_pinned_shadow__
#endif
texture<short, hipTextureType1D, hipReadModeNormalizedFloat> texs;
#if __HIP__
__hip_pinned_shadow__
#endif
texture<unsigned short, hipTextureType1D, hipReadModeNormalizedFloat> texus;
template<typename T>
__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<T, char>::value)
pDst[elementID] = tex1D(texc, coord);
else if(std::is_same<T, unsigned char>::value)
pDst[elementID] = tex1D(texuc, coord);
else if(std::is_same<T, short>::value)
pDst[elementID] = tex1D(texs, coord);
else if(std::is_same<T, unsigned short>::value)
pDst[elementID] = tex1D(texus, coord);
}
template<typename T>
bool textureTest(enum hipArray_Format texFormat)
bool textureTest(texture<T, hipTextureType1D, hipReadModeNormalizedFloat> *tex)
{
hipChannelFormatDesc desc = hipCreateChannelDesc<T>();
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<T>, 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<char> (HIP_AD_FORMAT_SIGNED_INT8);
status &= textureTest<unsigned char> (HIP_AD_FORMAT_UNSIGNED_INT8);
status &= textureTest<short> (HIP_AD_FORMAT_SIGNED_INT16);
status &= textureTest<unsigned short>(HIP_AD_FORMAT_UNSIGNED_INT16);
status &= textureTest<float> (HIP_AD_FORMAT_FLOAT);
status &= textureTest<char> (&texc);
status &= textureTest<unsigned char> (&texuc);
status &= textureTest<short> (&texs);
status &= textureTest<unsigned short>(&texus);
if(status){
passed();
}