Skip to content

Assert before terminating on throw in device code #7358

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
davebayer merged 3 commits into from
Jan 26, 2026
Merged

Assert before terminating on throw in device code #7358

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
d85659d
Assert before terminating on throwing in device code
davebayer Jan 26, 2026
1976f08
Simplify formatting
davebayer Jan 26, 2026
c68fa41
fix docs
davebayer Jan 26, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
26 changes: 13 additions & 13 deletions docs/cccl/development/macro.rst
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -345,17 +345,17 @@ In C++23, the ``if consteval`` statement (`link <https://en.cppreference.com/w/c

CUDA doesn't support exceptions in device code, however, sometimes we need to write host/device functions that use exceptions on host and ``__trap()`` on device. CCCL provides a set of macros that should be used in place of the standard C++ keywords to make the code compile in both, host and device code.

+-----------------------------+-----------------------------------------------------------------------+
| ``_CCCL_TRY`` | Replacement for the ``try`` keyword |
+-----------------------------+-----------------------------------------------------------------------+
| ``_CCCL_CATCH (X)`` | Replacement for the ``catch (/*X*/)`` statement |
+-----------------------------+-----------------------------------------------------------------------+
| ``_CCCL_CATCH_ALL`` | Replacement for the ``catch (...)`` statement |
+-----------------------------+-----------------------------------------------------------------------+
| ``_CCCL_THROW`` | Replacement for the ``throw /*arg*/`` expression |
+-----------------------------+-----------------------------------------------------------------------+
| ``_CCCL_RETHROW`` | Replacement for the plain ``throw`` expression |
+-----------------------------+-----------------------------------------------------------------------+
+-----------------------------+------------------------------------------------------------------------------------------------------------------+
| ``_CCCL_TRY`` | Replacement for the ``try`` keyword. |
+-----------------------------+------------------------------------------------------------------------------------------------------------------+
| ``_CCCL_CATCH (X)`` | Replacement for the ``catch (/*X*/)`` statement. |
+-----------------------------+------------------------------------------------------------------------------------------------------------------+
| ``_CCCL_CATCH_ALL`` | Replacement for the ``catch (...)`` statement. |
+-----------------------------+------------------------------------------------------------------------------------------------------------------+
| ``_CCCL_THROW(X, ...)`` | Replacement for the ``throw X(...)`` expression. ``X`` must be fully qualified type, without the leading ``::``. |
+-----------------------------+------------------------------------------------------------------------------------------------------------------+
| ``_CCCL_RETHROW`` | Replacement for the plain ``throw`` expression. |
+-----------------------------+------------------------------------------------------------------------------------------------------------------+

*Note*: The ``_CCCL_CATCH`` clause must always introduce a named variable, like: ``_CCCL_CATCH(const exception_type& var)``.

Expand All @@ -373,12 +373,12 @@ Example:
{
return ptr;
}
_CCCL_THROW(std::bad_alloc{}); // on device calls cuda::std::terminate()
_CCCL_THROW(std::bad_alloc); // on device calls cuda::std::terminate()
}

__host__ __device__ void do_something(int* buff)
{
_CCCL_THROW(std::runtime_error{"Something went wrong"}); // on device calls cuda::std::terminate()
_CCCL_THROW(std::runtime_error, "Something went wrong"); // on device calls cuda::std::terminate()
}

__host__ __device__ void fn(cuda::std::size_t n)
Expand Down
6 changes: 3 additions & 3 deletions libcudacxx/include/cuda/__driver/driver_api.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -951,7 +951,7 @@ __cutensormap_size_bytes(::cuda::std::size_t __num_items, ::CUtensorMapDataType
case ::CU_TENSOR_MAP_DATA_TYPE_FLOAT16:
if (__num_items > __max_size / 2)
{
_CCCL_THROW(::std::invalid_argument{"Number of items must be less than or equal to 2^64 / 2"});
_CCCL_THROW(std::invalid_argument, "Number of items must be less than or equal to 2^64 / 2");
}
return __num_items * 2;
case ::CU_TENSOR_MAP_DATA_TYPE_INT32:
Expand All @@ -962,15 +962,15 @@ __cutensormap_size_bytes(::cuda::std::size_t __num_items, ::CUtensorMapDataType
case ::CU_TENSOR_MAP_DATA_TYPE_TFLOAT32_FTZ:
if (__num_items > __max_size / 4)
{
_CCCL_THROW(::std::invalid_argument{"Number of items must be less than or equal to 2^64 / 4"});
_CCCL_THROW(std::invalid_argument, "Number of items must be less than or equal to 2^64 / 4");
}
return __num_items * 4;
case ::CU_TENSOR_MAP_DATA_TYPE_INT64:
case ::CU_TENSOR_MAP_DATA_TYPE_UINT64:
case ::CU_TENSOR_MAP_DATA_TYPE_FLOAT64:
if (__num_items > __max_size / 8)
{
_CCCL_THROW(::std::invalid_argument{"Number of items must be less than or equal to 2^64 / 8"});
_CCCL_THROW(std::invalid_argument, "Number of items must be less than or equal to 2^64 / 8");
}
return __num_items * 8;
# if _CCCL_CTK_AT_LEAST(12, 8)
Expand Down
4 changes: 2 additions & 2 deletions libcudacxx/include/cuda/__mdspan/mdspan_to_dlpack.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -190,11 +190,11 @@ __to_dlpack(const ::cuda::std::mdspan<_ElementType, _Extents, _Layout, _Accessor
{
if (::cuda::std::cmp_greater(__mdspan.extent(__i), __max_extent))
{
_CCCL_THROW(::std::invalid_argument{"Extent is too large"});
_CCCL_THROW(std::invalid_argument, "Extent is too large");
}
if (::cuda::std::cmp_greater(__mdspan.stride(__i), __max_extent))
{
_CCCL_THROW(::std::invalid_argument{"Stride is too large"});
_CCCL_THROW(std::invalid_argument, "Stride is too large");
}
__wrapper.__shape[__i] = static_cast<::cuda::std::int64_t>(__mdspan.extent(__i));
__wrapper.__strides[__i] = static_cast<::cuda::std::int64_t>(__mdspan.stride(__i));
Expand Down
Loading