CHIP-SPV / chipStar

chipStar is a tool for compiling and running HIP/CUDA on SPIR-V via OpenCL or Level Zero APIs.
Other
185 stars 30 forks source link

HIP/tests/catch/unit/deviceLib/{Single,Double}PrecisionMathFunctions tests are not very useful #625

Open franz opened 12 months ago

franz commented 12 months ago
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.

pjaaskel commented 12 months ago

Should this be reported to the AMD HIP upstream?

pvelesko commented 12 months ago

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?

franz commented 11 months ago

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.