Author: Joseph Huber Date: 2026-08-07T21:08:18-05:00 New Revision: 756e26ab79917987f678bb9a81a470c4c61516f4
URL: https://github.com/llvm/llvm-project/commit/756e26ab79917987f678bb9a81a470c4c61516f4 DIFF: https://github.com/llvm/llvm-project/commit/756e26ab79917987f678bb9a81a470c4c61516f4.diff LOG: [libclc] Add math smoke test for cos special values (#214786) Summary: Basic math version of unit tests used to give reference for how these can test the known special values. SImilar to LLVM libc's smoke math tests on the GPU. More can be added later, Added: libclc/test/conformance/math/cos.cl libclc/test/conformance/workgroups/work_group_reduce.cl Modified: libclc/test/conformance/conformance.h libclc/test/lit.cfg.py Removed: libclc/test/conformance/work_group_reduce.cl ################################################################################ diff --git a/libclc/test/conformance/conformance.h b/libclc/test/conformance/conformance.h index 6c98681b29e86..d5f7f40e27daa 100644 --- a/libclc/test/conformance/conformance.h +++ b/libclc/test/conformance/conformance.h @@ -29,4 +29,61 @@ #define CHECK_GT(LHS, RHS) CHECK((LHS) > (RHS)) #define CHECK_GE(LHS, RHS) CHECK((LHS) >= (RHS)) +// Representable values between X and Y. Does not handle fractional ULP values +// and rounds up instead, not fully accurate for all cases. +static inline uint __clc_test_ulp_f32(float x, float y) { + uint a = as_uint(x), b = as_uint(y); + a = (a >> 31) ? 0x80000000U - (a & 0x7fffffffU) : 0x80000000U + a; + b = (b >> 31) ? 0x80000000U - (b & 0x7fffffffU) : 0x80000000U + b; + return a > b ? a - b : b - a; +} + +// The reference must be of the result's own type, or it converts and the +// distance is measured on the wrong grid. +#define CHECK_ULP_F32(GOT, REF, ULPS) \ + do { \ + _Static_assert(__builtin_types_compatible_p(__typeof__(GOT), float) && \ + __builtin_types_compatible_p(__typeof__(REF), float), \ + "CHECK_ULP_F32 needs float operands"); \ + CHECK(__clc_test_ulp_f32((GOT), (REF)) <= (ULPS)); \ + } while (0) + +#ifdef __opencl_c_fp64 +static inline ulong __clc_test_ulp_f64(double x, double y) { + ulong a = as_ulong(x), b = as_ulong(y); + a = (a >> 63) ? 0x8000000000000000UL - (a & 0x7fffffffffffffffUL) + : 0x8000000000000000UL + a; + b = (b >> 63) ? 0x8000000000000000UL - (b & 0x7fffffffffffffffUL) + : 0x8000000000000000UL + b; + return a > b ? a - b : b - a; +} + +#define CHECK_ULP_F64(GOT, REF, ULPS) \ + do { \ + _Static_assert(__builtin_types_compatible_p(__typeof__(GOT), double) && \ + __builtin_types_compatible_p(__typeof__(REF), double), \ + "CHECK_ULP_F64 needs double operands"); \ + CHECK(__clc_test_ulp_f64((GOT), (REF)) <= (ULPS)); \ + } while (0) +#endif + +#ifdef cl_khr_fp16 +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +static inline ushort __clc_test_ulp_f16(half x, half y) { + ushort a = as_ushort(x), b = as_ushort(y); + a = (a >> 15) ? 0x8000U - (a & 0x7fffU) : 0x8000U + a; + b = (b >> 15) ? 0x8000U - (b & 0x7fffU) : 0x8000U + b; + return a > b ? a - b : b - a; +} + +#define CHECK_ULP_F16(GOT, REF, ULPS) \ + do { \ + _Static_assert(__builtin_types_compatible_p(__typeof__(GOT), half) && \ + __builtin_types_compatible_p(__typeof__(REF), half), \ + "CHECK_ULP_F16 needs half operands"); \ + CHECK(__clc_test_ulp_f16((GOT), (REF)) <= (ULPS)); \ + } while (0) +#endif + #endif // LIBCLC_TEST_CONFORMANCE_H diff --git a/libclc/test/conformance/math/cos.cl b/libclc/test/conformance/math/cos.cl new file mode 100644 index 0000000000000..650e8991e7b03 --- /dev/null +++ b/libclc/test/conformance/math/cos.cl @@ -0,0 +1,88 @@ +// RUN: %libclc-compile +// RUN: %libclc-run --kernel cos_f32 --threads-x 1 %t +// RUN: %libclc-run --kernel cos_f64 --threads-x 1 %t +// RUN: %libclc-run --kernel cos_f16 --threads-x 1 %t + +#include "conformance.h" + +// OpenCL C v3.0.19, Sec. 7.4, Table 65: cos() is allowed 4 ULP. The arguments +// cover the argument reduction paths. +__kernel void cos_f32(void) { + CHECK_ULP_F32(cos(TEST_INPUT(float, 1.0f)), 0x1.14a28p-1f, 4); + + // Nearest float to pi/2, where the result all but cancels. + CHECK_ULP_F32(cos(TEST_INPUT(float, 0x1.921fb6p+0f)), -0x1.777a5cp-25f, 4); + + // Nearest float to pi. + CHECK_ULP_F32(cos(TEST_INPUT(float, 0x1.921fb6p+1f)), -0x1p+0f, 4); + + CHECK_ULP_F32(cos(TEST_INPUT(float, 100.0f)), 0x1.b981dcp-1f, 4); + + // Needs more bits of 2/pi than a double holds. + CHECK_ULP_F32(cos(TEST_INPUT(float, 0x1.0p+64f)), -0x1.ffdb8p-1f, 4); + CHECK_ULP_F32(cos(TEST_INPUT(float, 0x1.fffffep+127f)), 0x1.b4bf2cp-1f, 4); + + // Not correctly rounded, so the bound is exercised. + CHECK_ULP_F32(cos(TEST_INPUT(float, 0x1.0c158cp+0f)), 0x1.fffe96p-2f, 4); + + CHECK_EQ(cos(TEST_INPUT(float, 0.0f)), 1.0f); + CHECK_EQ(cos(TEST_INPUT(float, -0.0f)), 1.0f); + CHECK(isnan(cos(TEST_INPUT(float, INFINITY)))); + CHECK(isnan(cos(TEST_INPUT(float, -INFINITY)))); + CHECK(isnan(cos(TEST_INPUT(float, NAN)))); +} + +// OpenCL C v3.0.19, Sec. 7.4, Table 68: cos() is allowed 4 ULP. +__kernel void cos_f64(void) { +#ifdef __opencl_c_fp64 + CHECK_ULP_F64(cos(TEST_INPUT(double, 1.0)), 0x1.14a280fb5068cp-1, 4); + + // Nearest double to pi/2, then to pi. + CHECK_ULP_F64(cos(TEST_INPUT(double, 0x1.921fb54442d18p+0)), + 0x1.1a62633145c07p-54, 4); + CHECK_ULP_F64(cos(TEST_INPUT(double, 0x1.921fb54442d18p+1)), -0x1p+0, 4); + + CHECK_ULP_F64(cos(TEST_INPUT(double, 100.0)), 0x1.b981dbf665fdfp-1, 4); + + CHECK_ULP_F64(cos(TEST_INPUT(double, 0x1.0p+64)), -0x1.ffdb7fa3fe34dp-1, 4); + CHECK_ULP_F64(cos(TEST_INPUT(double, 0x1.fffffffffffffp+1023)), + -0x1.fffe62ecfab75p-1, 4); + + // Not correctly rounded, so the bound is exercised. + CHECK_ULP_F64(cos(TEST_INPUT(double, 0x1.0cf7ef9db22d1p+9)), + -0x1.7f59309b88661p-1, 4); + + CHECK_EQ(cos(TEST_INPUT(double, 0.0)), 1.0); + CHECK_EQ(cos(TEST_INPUT(double, -0.0)), 1.0); + CHECK(isnan(cos(TEST_INPUT(double, (double)INFINITY)))); + CHECK(isnan(cos(TEST_INPUT(double, -(double)INFINITY)))); + CHECK(isnan(cos(TEST_INPUT(double, (double)NAN)))); +#endif +} + +// OpenCL C v3.0.19, Sec. 7.4, Table 69 (Full Profile): half cos() is allowed +// 2 ULP. +__kernel void cos_f16(void) { +#ifdef cl_khr_fp16 + CHECK_ULP_F16(cos(TEST_INPUT(half, 1.0h)), 0x1.14cp-1h, 2); + + // Nearest half to pi/2, then to pi. + CHECK_ULP_F16(cos(TEST_INPUT(half, 0x1.92p+0h)), 0x1.fb4p-12h, 2); + CHECK_ULP_F16(cos(TEST_INPUT(half, 0x1.92p+1h)), -0x1p+0h, 2); + + // Widest argument needing reduction. + CHECK_ULP_F16(cos(TEST_INPUT(half, 0x1.ffcp+15h)), -0x1.c3cp-3h, 2); + + // Not correctly rounded, so the bound is exercised. + CHECK_ULP_F16(cos(TEST_INPUT(half, 0x1.8acp-4h)), 0x1.fd8p-1h, 2); + + // Smallest subnormal, where the result rounds to one. + CHECK_ULP_F16(cos(TEST_INPUT(half, 0x1.0p-24h)), 0x1p+0h, 2); + + CHECK_EQ(cos(TEST_INPUT(half, 0.0h)), 1.0h); + CHECK_EQ(cos(TEST_INPUT(half, -0.0h)), 1.0h); + CHECK(isnan(cos(TEST_INPUT(half, (half)INFINITY)))); + CHECK(isnan(cos(TEST_INPUT(half, (half)-INFINITY)))); + CHECK(isnan(cos(TEST_INPUT(half, (half)NAN)))); +#endif +} diff --git a/libclc/test/conformance/work_group_reduce.cl b/libclc/test/conformance/workgroups/work_group_reduce.cl similarity index 100% rename from libclc/test/conformance/work_group_reduce.cl rename to libclc/test/conformance/workgroups/work_group_reduce.cl diff --git a/libclc/test/lit.cfg.py b/libclc/test/lit.cfg.py index d6c49dbfe5878..0e592f736aff7 100644 --- a/libclc/test/lit.cfg.py +++ b/libclc/test/lit.cfg.py @@ -83,10 +83,12 @@ def calculate_arch_features(arch_string): if test_arch and offload_libdir and os.path.isfile(path): clang = os.path.join(config.llvm_tools_dir, "clang") loader = os.path.join(config.llvm_tools_dir, "llvm-gpu-loader") + # Shared headers are included by name from any depth below conformance/. + conformance_dir = os.path.join(config.test_source_root, "conformance") compile_cmd = ( f"{clang} --target={config.libclc_target} " f"-march={test_arch} -cl-std=CL3.0 -nogpulib " - f"--libclc-lib=:{path} %s -o %t" + f"--libclc-lib=:{path} -I{conformance_dir} %s -o %t" ) config.environment["LD_LIBRARY_PATH"] = os.pathsep.join( [offload_libdir, config.environment.get("LD_LIBRARY_PATH", "")] _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
