Skip to content
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.

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

backport unreachable #2852

Open
wants to merge 43 commits into
base: main
Choose a base branch
from
Open
Show file tree
Hide file tree
Changes from 5 commits
Commits
Show all changes
43 commits
Select commit Hold shift + click to select a range
fdc6c36
backport and fix `unreachable`
davebayer Nov 18, 2024
0246df1
remove `_CCCL_UNREACHABLE`
davebayer Nov 18, 2024
0d83892
update tests
davebayer Nov 18, 2024
f6f2574
define cccl library `unreachable` version
davebayer Nov 18, 2024
94101f2
use `_CCCL_NORETURN` instead of standard `[[noreturn]]`
davebayer Nov 18, 2024
4f0c516
use `_CUDA_VSTD` within libcu++
davebayer Nov 18, 2024
a8e3f60
fix for cuda compilation with NVHPC
davebayer Nov 18, 2024
bed27be
fix compilation for host/device unreachable
davebayer Nov 18, 2024
86e4ce8
separate host and device implementations
davebayer Nov 18, 2024
f7affcf
fall back to the original implementation
davebayer Nov 19, 2024
3208ba5
Merge branch 'main' into backport_and_fix_unreachable
davebayer Nov 19, 2024
913ef98
[pre-commit.ci] auto code formatting
pre-commit-ci[bot] Nov 19, 2024
4dfdb61
fix tests
davebayer Nov 20, 2024
7bed5c3
Merge branch 'main' into backport_and_fix_unreachable
davebayer Nov 20, 2024
897086a
Move implementation of `_LIBCUDACXX_TEMPLATE` to CCCL (#2832)
miscco Nov 21, 2024
786d442
Try to work around issue with NVHPC in conjunction of older CTK versi…
miscco Nov 21, 2024
6b5fa22
Refactoring (#2905)
bernhardmgruber Nov 21, 2024
44f0331
add "`interface`" to `_CCCL_PUSH_MACROS` (#2919)
ericniebler Nov 21, 2024
c9a6e6a
Replace inconsistent Doxygen macros with `_CCCL_DOXYGEN_INVOKED` (#2921)
ericniebler Nov 21, 2024
a50019d
implement C++26 `std::span::at` (#2924)
davebayer Nov 22, 2024
d0f5bd2
move msvc compiler macros to new version (#2885)
davebayer Nov 22, 2024
5b57a4c
Reorganize PTX tests to match generator (#2930)
bernhardmgruber Nov 22, 2024
96e8199
Reorganize PTX docs to match generator (#2929)
bernhardmgruber Nov 22, 2024
92a22f2
Improve build instructions for libcu++ (#2881)
miscco Nov 22, 2024
14484a6
Reorganize PTX headers to match generator (#2925)
bernhardmgruber Nov 22, 2024
e897f1f
implement C++26 `std::span`'s constructor from `std::initializer_list…
davebayer Nov 22, 2024
3b23083
Add tuple protocol to `cuda::std::complex` from C++26 (#2882)
davebayer Nov 22, 2024
0a0773f
Move implementation of `_LIBCUDACXX_TEMPLATE` to CCCL (#2832)
miscco Nov 21, 2024
3911b25
Replace inconsistent Doxygen macros with `_CCCL_DOXYGEN_INVOKED` (#2921)
ericniebler Nov 21, 2024
b760c7f
move msvc compiler macros to new version (#2885)
davebayer Nov 22, 2024
8ee0b18
implement C++26 `std::span`'s constructor from `std::initializer_list…
davebayer Nov 22, 2024
ef9ae72
Add missing qualifier for cuda namespace (#2940)
bernhardmgruber Nov 23, 2024
5e4d6e0
Try to fix a clang warning: (#2941)
bernhardmgruber Nov 23, 2024
6dbeb46
minor consistency improvements in concepts macros (#2928)
ericniebler Nov 24, 2024
2a47168
Drop some of the mdspan fold implementation (#2949)
miscco Nov 25, 2024
a4477a4
[STF] Implement CUDASTF_DOT_TIMING for the ctx.cuda_kernel construct …
caugonnet Nov 25, 2024
09db35f
Avoid potential null dereference in `annotated_ptr` (#2951)
miscco Nov 25, 2024
8ca5380
make compiler version comparison utility generic (#2952)
davebayer Nov 25, 2024
3c6fcd2
Add SM100 descriptor to target (#2954)
miscco Nov 25, 2024
aededb9
Regenerate `cuda::ptx` headers/docs and run format (#2937)
bernhardmgruber Nov 25, 2024
4d53204
Regenerate PTX test (#2953)
bernhardmgruber Nov 25, 2024
bdff820
Do not include extended floating point headers if they are not needed…
miscco Nov 25, 2024
651bbcf
[CUDAX] Add copy_bytes and fill_bytes overloads for mdspan (#2932)
pciolkosz Nov 25, 2024
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
6 changes: 4 additions & 2 deletions cub/cub/agent/agent_sub_warp_merge_sort.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -45,6 +45,8 @@

#include <thrust/system/cuda/detail/core/util.h>

#include <cuda/std/utility>

#include <nv/target>

CUB_NAMESPACE_BEGIN
Expand Down Expand Up @@ -120,7 +122,7 @@ class AgentSubWarpSort
{
return lhs < rhs;
}
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}

#if defined(__CUDA_FP16_TYPES_EXIST__)
Expand All @@ -135,7 +137,7 @@ class AgentSubWarpSort
{
NV_IF_TARGET(NV_PROVIDES_SM_53, (return __hlt(lhs, rhs);), (return __half2float(lhs) < __half2float(rhs);));
}
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}
#endif // __CUDA_FP16_TYPES_EXIST__
};
Expand Down
3 changes: 2 additions & 1 deletion cub/cub/thread/thread_reduce.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -53,6 +53,7 @@
#include <cuda/std/cassert> // assert
#include <cuda/std/cstdint> // uint16_t
#include <cuda/std/functional> // cuda::std::plus
#include <cuda/std/utility> // unreachable

CUB_NAMESPACE_BEGIN

Expand Down Expand Up @@ -522,7 +523,7 @@ ThreadReduceSimd(const Input& input, ReductionOp reduction_op) -> ::cuda::std::r
}
return unsafe_bitcast<UnpackedType>(result)[0];
}
_CCCL_UNREACHABLE(); // nvcc 11.x warning workaround (never reached)
::cuda::std::unreachable(); // nvcc 11.x warning workaround (never reached)
}

} // namespace internal
Expand Down
6 changes: 3 additions & 3 deletions cub/cub/util_device.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -730,7 +730,7 @@ private:
template <int, typename FunctorT>
CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t invoke_static(FunctorT&, ::cuda::std::false_type)
{
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}

template <typename FunctorT>
Expand All @@ -742,7 +742,7 @@ private:
template <typename FunctorT>
CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t do_invoke(FunctorT&, ::cuda::std::false_type)
{
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}
};

Expand Down Expand Up @@ -773,7 +773,7 @@ private:
template <int, typename FunctorT>
CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t invoke_static(FunctorT&, ::cuda::std::false_type)
{
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}
};

Expand Down
3 changes: 2 additions & 1 deletion cub/test/catch2_test_device_histogram.cu
Original file line number Diff line number Diff line change
Expand Up @@ -33,6 +33,7 @@
#include <cuda/std/array>
#include <cuda/std/bit>
#include <cuda/std/type_traits>
#include <cuda/std/utility>

#include <algorithm>
#include <limits>
Expand Down Expand Up @@ -305,7 +306,7 @@ void test_even_and_range(LevelT max_level, int max_level_count, OffsetT width, O
{
return static_cast<int>((sample - min) * fp_scales[channel]);
}
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
};
auto h_histogram = compute_reference_result<Channels, CounterT>(
h_samples, sample_to_bin_index, num_levels, width, height, row_pitch);
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -17,6 +17,7 @@
#include <cuda/std/__type_traits/type_list.h>
#include <cuda/std/__utility/declval.h>
#include <cuda/std/__utility/integer_sequence.h>
#include <cuda/std/__utility/unreachable.h>
#include <cuda/std/span>
#include <cuda/std/tuple>

Expand Down Expand Up @@ -79,7 +80,7 @@ _CCCL_NODISCARD _CUDAX_API constexpr auto fool_compiler(const dimensions<T, Exte
{
return ex;
}
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}

template <typename QueryLevel, typename Hierarchy>
Expand Down Expand Up @@ -118,7 +119,7 @@ struct get_level_helper
{
return (*this)(levels...);
}
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}
};
} // namespace detail
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -12,6 +12,7 @@
#define _CUDAX__HIERARCHY_HIERARCHY_LEVELS

#include <cuda/std/__type_traits/type_list.h>
#include <cuda/std/__utility/unreachable.h>

#include <cuda/experimental/__hierarchy/dimensions.cuh>

Expand Down Expand Up @@ -285,7 +286,7 @@ template <typename Unit, typename Level>
return dims_product<typename Level::product_type>(
extents_impl<SplitLevel, Level>(), extents_impl<Unit, SplitLevel>());
}
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}

template <typename Unit, typename Level>
Expand All @@ -302,7 +303,7 @@ template <typename Unit, typename Level>
dims_product<typename Level::product_type>(index_impl<SplitLevel, Level>(), extents_impl<Unit, SplitLevel>()),
index_impl<Unit, SplitLevel>());
}
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}
} // namespace detail

Expand Down
1 change: 0 additions & 1 deletion libcudacxx/include/cuda/__cccl_config
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,6 @@
#include <cuda/std/__cccl/rtti.h> // IWYU pragma: export
#include <cuda/std/__cccl/sequence_access.h> // IWYU pragma: export
#include <cuda/std/__cccl/system_header.h> // IWYU pragma: export
#include <cuda/std/__cccl/unreachable.h> // IWYU pragma: export
#include <cuda/std/__cccl/version.h> // IWYU pragma: export
#include <cuda/std/__cccl/visibility.h> // IWYU pragma: export

Expand Down
3 changes: 2 additions & 1 deletion libcudacxx/include/cuda/__memcpy_async/memcpy_async_tx.h
Original file line number Diff line number Diff line change
Expand Up @@ -34,6 +34,7 @@
# include <cuda/__ptx/ptx_helper_functions.h>
# include <cuda/std/__atomic/scopes.h>
# include <cuda/std/__type_traits/is_trivially_copyable.h>
# include <cuda/std/__utility/unreachable.h>
# include <cuda/std/cstdint>

_LIBCUDACXX_BEGIN_NAMESPACE_CUDA_DEVICE
Expand Down Expand Up @@ -75,7 +76,7 @@ _CCCL_DEVICE inline async_contract_fulfillment memcpy_async_tx(
// or from shared to remote cluster dsmem. To copy to remote
// dsmem, we need to arrive on a cluster-scoped barrier, which
// is not yet implemented. So we trap in this case as well.
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}),
(__cuda_ptx_memcpy_async_tx_is_not_supported_before_SM_90__();));

Expand Down
13 changes: 7 additions & 6 deletions libcudacxx/include/cuda/__memcpy_async/memcpy_completion.h
Original file line number Diff line number Diff line change
Expand Up @@ -29,6 +29,7 @@
#include <cuda/__memcpy_async/is_local_smem_barrier.h>
#include <cuda/__memcpy_async/try_get_barrier_handle.h>
#include <cuda/std/__atomic/scopes.h>
#include <cuda/std/__utility/unreachable.h>
#include <cuda/std/cstdint>

#if defined(_CCCL_CUDA_COMPILER)
Expand Down Expand Up @@ -84,7 +85,7 @@ struct __memcpy_completion_impl
// This completion mechanism should not be used with a shared
// memory barrier. Or at least, we do not currently envision
// bulk group to be used with shared memory barriers.
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
case __completion_mechanism::__mbarrier_complete_tx:
#if __cccl_ptx_isa >= 800
// Pre-sm90, the mbarrier_complete_tx completion mechanism is not available.
Expand All @@ -100,7 +101,7 @@ struct __memcpy_completion_impl
return async_contract_fulfillment::none;
default:
// Get rid of "control reaches end of non-void function":
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}
}

Expand Down Expand Up @@ -129,16 +130,16 @@ struct __memcpy_completion_impl
return async_contract_fulfillment::async;
case __completion_mechanism::__mbarrier_complete_tx:
// Non-smem barriers do not have an mbarrier_complete_tx mechanism..
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
case __completion_mechanism::__async_bulk_group:
// This completion mechanism is currently not expected to be used with barriers.
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
case __completion_mechanism::__sync:
// sync: In this case, we do not need to do anything.
return async_contract_fulfillment::none;
default:
// Get rid of "control reaches end of non-void function":
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}
}

Expand All @@ -158,7 +159,7 @@ struct __memcpy_completion_impl
return async_contract_fulfillment::none;
default:
// Get rid of "control reaches end of non-void function":
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}
}
};
Expand Down
52 changes: 0 additions & 52 deletions libcudacxx/include/cuda/std/__cccl/unreachable.h

This file was deleted.

5 changes: 3 additions & 2 deletions libcudacxx/include/cuda/std/__exception/terminate.h
Original file line number Diff line number Diff line change
Expand Up @@ -22,6 +22,7 @@
# pragma system_header
#endif // no system header

#include <cuda/std/__utility/unreachable.h>
#include <cuda/std/cstdlib> // ::exit

_CCCL_DIAG_PUSH
Expand All @@ -32,7 +33,7 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD_NOVERSION // purposefully not using versioning n
_CCCL_NORETURN _LIBCUDACXX_HIDE_FROM_ABI void __cccl_terminate() noexcept
{
NV_IF_ELSE_TARGET(NV_IS_HOST, (::exit(-1);), (__trap();))
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}

#if 0 // Expose once atomic is universally available
Expand All @@ -59,7 +60,7 @@ _LIBCUDACXX_HIDE_FROM_ABI terminate_handler get_terminate() noexcept
_CCCL_NORETURN _LIBCUDACXX_HIDE_FROM_ABI void terminate() noexcept
{
__cccl_terminate();
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}

_LIBCUDACXX_END_NAMESPACE_STD_NOVERSION
Expand Down
3 changes: 2 additions & 1 deletion libcudacxx/include/cuda/std/__iterator/advance.h
Original file line number Diff line number Diff line change
Expand Up @@ -28,6 +28,7 @@
#include <cuda/std/__iterator/iterator_traits.h>
#include <cuda/std/__utility/convert_to_integral.h>
#include <cuda/std/__utility/move.h>
#include <cuda/std/__utility/unreachable.h>

_LIBCUDACXX_BEGIN_NAMESPACE_STD

Expand Down Expand Up @@ -226,7 +227,7 @@ struct __fn
}
return __n;
}
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}
};
_LIBCUDACXX_END_NAMESPACE_CPO
Expand Down
5 changes: 3 additions & 2 deletions libcudacxx/include/cuda/std/__iterator/distance.h
Original file line number Diff line number Diff line change
Expand Up @@ -29,6 +29,7 @@
#include <cuda/std/__ranges/size.h>
#include <cuda/std/__type_traits/decay.h>
#include <cuda/std/__type_traits/remove_cvref.h>
#include <cuda/std/__utility/unreachable.h>

_LIBCUDACXX_BEGIN_NAMESPACE_STD

Expand Down Expand Up @@ -92,7 +93,7 @@ struct __fn
{
return __last - decay_t<_Ip>(__first);
}
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}

_LIBCUDACXX_TEMPLATE(class _Rp)
Expand All @@ -107,7 +108,7 @@ struct __fn
{
return operator()(_CUDA_VRANGES::begin(__r), _CUDA_VRANGES::end(__r));
}
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}
};
_LIBCUDACXX_END_NAMESPACE_CPO
Expand Down
3 changes: 2 additions & 1 deletion libcudacxx/include/cuda/std/__iterator/move_iterator.h
Original file line number Diff line number Diff line change
Expand Up @@ -41,6 +41,7 @@
#include <cuda/std/__type_traits/is_reference.h>
#include <cuda/std/__type_traits/remove_reference.h>
#include <cuda/std/__utility/move.h>
#include <cuda/std/__utility/unreachable.h>

_LIBCUDACXX_BEGIN_NAMESPACE_STD

Expand Down Expand Up @@ -126,7 +127,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT move_iterator
{
return input_iterator_tag{};
}
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}
# endif // !_CCCL_COMPILER_MSVC_2017
#endif // _CCCL_STD_VER >= 2017
Expand Down
3 changes: 2 additions & 1 deletion libcudacxx/include/cuda/std/__ranges/size.h
Original file line number Diff line number Diff line change
Expand Up @@ -32,6 +32,7 @@
#include <cuda/std/__type_traits/remove_cvref.h>
#include <cuda/std/__utility/auto_cast.h>
#include <cuda/std/__utility/declval.h>
#include <cuda/std/__utility/unreachable.h>
#include <cuda/std/cstddef>

_LIBCUDACXX_BEGIN_NAMESPACE_RANGES
Expand Down Expand Up @@ -191,7 +192,7 @@ struct __fn
{
return static_cast<_Signed>(_CUDA_VRANGES::size(__t));
}
_CCCL_UNREACHABLE();
::cuda::std::unreachable();
}
};
_LIBCUDACXX_END_NAMESPACE_CPO
Expand Down
Loading