NVIDIA / NVIDIA/cccl

[BUG]: _CCCL_VERIFY bug triggers call to _invoke_watson on Windows

Open
#10,841 2 comments 0 reactions 1 assignee Claimed by @miscco View on GitHub
needs triage
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

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.