NVIDIA / NVIDIA/cccl

[BUG]: SIMD complex trig functions exceed tolerance for __nv_bfloat16 on MSVC

Open
#9,631 1 comment 0 reactions 1 assignee Claimed by @fbusato View on GitHub
bug
Dominant language
C++
Stars
2.5k
Forks
486
Avg merge
2d 6h
Merged PRs (30d)
295

Description

## Describe the bug

The SIMD complex trigonometric/hyperbolic test
`std/numerics/simd/simd.complex/complex_math_trig.pass.cpp` fails a device-side
floating-point closeness assertion for the `__nv_bfloat16` element type on the
MSVC + CTK 12.9 toolchain. The `float`, `double`, and `__half` cases pass; only
`__nv_bfloat16` diverges, and it does so beyond the test's existing 50% relative
tolerance for that type (`simd_test_utils.h:126`).

The failure is toolchain-specific, not MSVC-general: it reproduces on
**CTK 12.9 / MSVC 14.44 / C++17** but **passes on CTK 13.3 / MSVC 14.50 / C++17**.
The same test passes in Linux CI (nvcc 12.9 + GCC/Clang), so the divergence is tied
to the MSVC host compiler combined with the older CTK, not nvcc 12.9 alone. (No C++20
test signal yet — those jobs are currently blocked by an unrelated build failure.)

```
simd_test_utils.h:126: block: [0,0,0], thread: [0,0,0]
Assertion `fptest_close_pct(static_cast(a), static_cast(b), 5.e-1f)` failed.
CUDA ERROR, line 124: cudaErrorAssert: device-side assert triggered
```

This was found while enabling MSVC GPU jobs in CI. The `__nv_bfloat16` case is
currently skipped on MSVC to keep the rest of the test running; this issue tracks
re-enabling it.

## How to Reproduce

1. Build/run the libcu++ lit suite with MSVC as the host compiler (CTK 12.9, MSVC
14.44, C++17), on an RTX 2080 (SM75).
2. Run `std/numerics/simd/simd.complex/complex_math_trig.pass.cpp`.
3. The `test_runtime()` path instantiates `test_type<__nv_bfloat16, 1>()` /
`test_type<__nv_bfloat16, 4>()`, whose `is_fp_close` checks against
`cuda::std::{sin,cos,tan,...}` exceed the 50% bfloat16 tolerance and trip the
device-side assert.

## Expected behavior

The `__nv_bfloat16` complex trig results should match the scalar reference within the
type's tolerance on the MSVC + CTK 12.9 toolchain, as they do on GCC/Clang and on
CTK 13.3 / MSVC 14.50, allowing the case to run there without a skip.

## System information

- OS: Windows
- Failing toolchain: CUDA 12.9, MSVC 14.44 (cl 19.44), C++17
- Passing toolchain: CUDA 13.3, MSVC 14.50, C++17
- GPU: NVIDIA GeForce RTX 2080 (SM75)

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.