Skip to content

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

Description

@seanbaxter

Is this a duplicate?

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(<yvals.h>) && _CCCL_OS(WINDOWS) // Windows uses _STL_VERIFY from <yvals.h>
#  include <yvals.h>
#  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 <cuda/std/__cccl/assert.h>

__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

Metadata

Metadata

Assignees

Labels

needs triageIssues that require the team's attention

Type

Projects

Status
Todo

Milestone

No milestone

Relationships

None yet

Development

No branches or pull requests

Issue actions