KhronosGroup / OpenCL-CTS

The OpenCL Conformance Tests
Apache License 2.0
184 stars 195 forks source link

test_printf producing false expected results #2009

Closed coldav closed 2 months ago

coldav commented 2 months ago

in the oneAPI Construction Kit, our overnight testing on 10th shows repeated fails on the test_printf test:

15)testing printf("%+#21.15E",789456123.0) verifyOutputBuffer failed with kernel: __kernel void test3(void) { printf("%+#21.15E\n",789456123.0); }

expected: +7.894561280000000E+08 got: +7.894561230000000E+08

This was run on x86 in our "host" target. I believe the 123 part is actually correct, so the test seems wrong.

https://github.com/KhronosGroup/OpenCL-CTS/pull/1988 maybe related as it was committed around the same time.

To show with the oneAPI Construction Kit.

git clone git@github.com:codeplaysoftware/oneapi-construction-kit.git
cd oneapi-construction-kit
cmake -GNinja -DCA_MUX_TARGETS_TO_ENABLE="host" -DCA_ENABLE_API=cl       -DCA_LLVM_INSTALL_DIR=<path_to_17_or_18_llvm_install>t -Bbuild
ninja -c build libOpenCL.so
svenvh commented 2 months ago

@coldav could you tell which subtest this is? This should be printed before the line starting with 0).

coldav commented 2 months ago

@coldav could you tell which subtest this is? This should be printed before the line starting with 0).

float...

coldav commented 2 months ago

With a little bit more digging if we disable double support we do get the 128 answer. We think it's acceptable to do the conversion to double though so both the 123 and 128 are correct. As an aside: Some platforms show /home/harald/OpenCL-CTS/test_conformance/printf/test_printf.cpp: In function ‘_cl_program* {anonymous}::makeMixedFormatPrintfProgram(_cl_kernel**, cl_context, cl_device_id, unsigned int, unsigned int, const std::string&)’: /home/harald/OpenCL-CTS/test_conformance/printf/test_printf.cpp:300:14: error: ‘uint8_t’ is not a member of ‘std’; did you mean ‘wint_t’? 300 | std::uint8_t is_int = genrand_int32(gMTdata) % 2; | ^~~~~~~ | wint_t /home/harald/OpenCL-CTS/test_conformance/printf/test_printf.cpp:303:32: error: ‘is_int’ was not declared in this scope; did you mean ‘u_int’? 303 | set_round(deviceRound, is_int != 0 ? kint : kfloat); | ^~~~~~ | u_int ninja: build stopped: subcommand failed. so there is likely a missing include.

svenvh commented 2 months ago

Thanks, I figured this had to be related to 789456123.0 not being exactly representable in fp32. I'm wondering if we should update the test to use a constant that's representable in both fp32 and fp64 instead; the obvious candidate would be 789456128.0.

I'm not sure if this is a recent regression though; the float subtests haven't been modified recently I believe.

so there is likely a missing include.

That seems like a recent regression indeed. I can put up a fix unless you already have a patch for it. update: to be fixed by #2011

coldav commented 2 months ago

Sorry, I meant to reply. Thanks for the patch. This should stay open for the original problem though.

bashbaug commented 2 months ago

I'm wondering if we should update the test to use a constant that's representable in both fp32 and fp64 instead [...]

Can we make this change, for now at least? We're encountering the same problem on our devices that support fp64.

dneto0 commented 2 months ago

I get the same failure on Linux x86-64 Gcc 13.2.0 with Clvk (top of tree), and NVIDIA Quadro P1000 Driver 525.147.5.0 (2204418368)

for completeness, subtests 14 and 15 have the same failure for the same value.

14)testing printf("%-#20.15e",789456123.0)
verifyOutputBuffer failed with kernel: 
__kernel void test3(void)
{
   printf("%-#20.15e\n",789456123.0);
}

 expected: 7.894561280000000e+08
 got:      7.894561230000000e+08
15)testing printf("%+#21.15E",789456123.0)
verifyOutputBuffer failed with kernel: 
__kernel void test3(void)
{
   printf("%+#21.15E\n",789456123.0);
}

 expected: +7.894561280000000E+08
 got:      +7.894561230000000E+08
dneto0 commented 2 months ago

I'm wondering if we should update the test to use a constant that's representable in both fp32 and fp64 instead [...]

Can we make this change, for now at least? We're encountering the same problem on our devices that support fp64.

I posted #2015 which does that.

bashbaug commented 2 months ago

Thanks, I approved #2015.

We're not meeting for the next couple of weeks, and this should be non-contentious, so as soon as somebody else signs off on it we should merge it.

Kerilk commented 2 months ago

@bashbaug I approved as well

bashbaug commented 2 months ago

Thanks! @coldav let us know if this doesn't fix your problem for some reason.