chipStar
chipStar copied to clipboard
HIP/tests/catch/unit/deviceLib/{Single,Double}PrecisionMathFunctions tests are not very useful
grep GENERATE HIP/tests/catch/unit/deviceLib/SinglePrecisionMathFunctions/*
acosf.cc:GENERATE_KERNEL_FLOAT(acosf, acosf(1.0f));
acoshf.cc:GENERATE_KERNEL_FLOAT(acoshf, acoshf(1.0f));
asinf.cc:GENERATE_KERNEL_FLOAT(asinf, asinf(1.0f));
asinhf.cc:GENERATE_KERNEL_FLOAT(asinhf, asinhf(1.0f));
atan2f.cc:GENERATE_KERNEL_FLOAT(atan2f, atan2f(1.0f, 1.0f));
atanf.cc:GENERATE_KERNEL_FLOAT(atanf, atanf(1.0f));
atanhf.cc:GENERATE_KERNEL_FLOAT(atanhf, atanhf(1.0f));
....
Almost all of the tests call a function with a constant argument, and do not use the function's result:
#define GENERATE_KERNEL_DOUBLE(FUNCNAME, FUNC) \
__global__ void testKernel_##FUNCNAME(double* a) { \
FUNC; \
} \
When chipStar is compiled with anything but -O0, this gets optimized to ret void. If chipStar is run with -O0, the code is generated, but is then optimized away by the OpenCL/Level0 driver at runtime.
Should this be reported to the AMD HIP upstream?
I wrote these. The purpose of these tests was to make sure that we have all the implementations for the device library. At the time, HIP provided devicelib tests where there was a single function calling all the device-side functions in a single kernel. If the kernel failed to JIT, you wouldn't know which function you were missing. Since we're always testing -O0 these don't get optimized away and we would encounter a JIT failure. I guess these could be improved so that they don't get optimized away?
The purpose of these tests was to make sure that we have all the implementations for the device library.... you wouldn't know which function you were missing.
makes sense
I guess these could be improved so that they don't get optimized away?
Yep this should be very easy. I can take a look next week.