[BUG]: _CCCL_VERIFY bug triggers call to _invoke_watson on Windows
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
### Is this a duplicate?
- [x] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this bug and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)
### Type of Bug
Compile-time Error
### Component
General CCCL
### Describe the bug
_CCCL_VERIFY has a hole where _invoke_watson is exposed to device functions. Here's the breakdown:
_CCCL_VERIFY(expr, msg) expands _CCCL_ASSERT_IMPL_DEVICE(expr, msg). That's fine.
_CCCL_ASSERT_IMPL_DEVICE expands _CCCL_ASSERT_IMPL_HOST(expr, msg). That's trouble.
https://github.com/NVIDIA/cccl/blob/main/libcudacxx/include/cuda/std/__cccl/assert.h#L106
There's an if/elif chain here. I'm building on Circle and I register as an NVHPC compiler, but NVHPC is not tested here. There's an NVRTC check, then NVCC check, then _CCCL_CUDA_COMPILATION(). That macro evaluates true. This gets defined:
```
# define _CCCL_ASSERT_IMPL_DEVICE(expression, message) _CCCL_ASSERT_IMPL_HOST(expression, message)
````
At line 69:
```
#elif __has_include() && _CCCL_OS(WINDOWS) // Windows uses _STL_VERIFY from
# include
# define _CCCL_ASSERT_IMPL_HOST(expression, message) _STL_VERIFY(expression, message)
```
_CCCL_ASSERT_IMPL_HOST is defined to _STL_VERIFY. Now we have to go into the Windows SDK:
```
#ifndef _MSVC_STL_DOOM_FUNCTION
#ifdef _MSVC_STL_USE_ABORT_AS_DOOM_FUNCTION
#define _MSVC_STL_DOOM_FUNCTION(mesg) _CSTD abort()
#else // ^^^ defined(_MSVC_STL_USE_ABORT_AS_DOOM_FUNCTION) / !defined(_MSVC_STL_USE_ABORT_AS_DOOM_FUNCTION) vvv
// TRANSITION, GH-4858: after dropping Win7 support, we can directly call __fastfail(FAST_FAIL_INVALID_ARG).
#define _MSVC_STL_DOOM_FUNCTION(mesg) ::_invoke_watson(nullptr, nullptr, nullptr, 0, 0)
#endif // ^^^ !defined(_MSVC_STL_USE_ABORT_AS_DOOM_FUNCTION) ^^^
#endif // ^^^ !defined(_MSVC_STL_DOOM_FUNCTION) ^^^
#define _STL_REPORT_ERROR(mesg) \
_RPTF0(_CRT_ASSERT, mesg); \
_MSVC_STL_DOOM_FUNCTION(mesg)
#define _STL_VERIFY(cond, mesg) \
if (!(cond)) { \
_STL_REPORT_ERROR(mesg); \
} \
_Analysis_assume_(cond)
```
_STL_VERIFY expands _MSVC_STL_DOOM_FUNCTION which calls _invoke_watson. That's the undefined device symbol.
There should be an NVHPC check in __cccl/assert.h. There isn't one, but because nvc++ doesn't build on Windows, this error hasn't be raised until now.
### How to Reproduce
```
#include
__global__ void kernel(int* p, int x) {
_CCCL_VERIFY(x > 0, "x must be positive"); // undefined external _invoke_watson
*p = x;
}
```
### Expected behavior
.
### Reproduction link
_No response_
### Operating System
Windows
### nvidia-smi output
_No response_
### NVCC version
_No response_
Contributor guide
Assessment
This issue has not been assessed yet.