-
Notifications
You must be signed in to change notification settings - Fork 135
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Make thrust::transform use cub::DeviceTransform
* Introduces address stability detection and opt-in in libcu++ * Mark BabelStream lambdas address oblivious
- Loading branch information
1 parent
27ba2e4
commit eeb0202
Showing
8 changed files
with
289 additions
and
28 deletions.
There are no files selected for viewing
64 changes: 64 additions & 0 deletions
64
libcudacxx/include/cuda/std/__type_traits/address_stability.h
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,64 @@ | ||
//===----------------------------------------------------------------------===// | ||
// | ||
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | ||
// See https://llvm.org/LICENSE.txt for license information. | ||
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. | ||
// | ||
//===----------------------------------------------------------------------===// | ||
|
||
#ifndef _LIBCUDACXX___TYPE_TRAITS_ADDRESS_STABILITY_H | ||
#define _LIBCUDACXX___TYPE_TRAITS_ADDRESS_STABILITY_H | ||
|
||
#include <cuda/std/detail/__config> | ||
|
||
#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) | ||
# pragma GCC system_header | ||
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) | ||
# pragma clang system_header | ||
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) | ||
# pragma system_header | ||
#endif // no system header | ||
|
||
#include <cuda/std/__type_traits/integral_constant.h> | ||
#include <cuda/std/__type_traits/void_t.h> | ||
#include <cuda/std/__utility/move.h> | ||
|
||
_LIBCUDACXX_BEGIN_NAMESPACE_STD | ||
|
||
// need a separate implementation trait because we SFINAE with a type parameter before the variadic pack | ||
template <typename F, typename SFINAE, typename... Args> | ||
struct __can_copy_arguments_impl : false_type | ||
{}; | ||
|
||
template <typename F, typename... Args> | ||
struct __can_copy_arguments_impl<F, void_t<decltype(F::can_copy_arguments)>, Args...> | ||
{ | ||
static constexpr bool value = F::can_copy_arguments; | ||
}; | ||
|
||
//! Trait telling whether a function object relies on the memory address of its arguments when called with the given set | ||
//! of types. The nested value is true when the addresses of the arguments do not matter and arguments can be provided | ||
//! from arbitrary copies of the respective sources. | ||
template <typename F, typename... Args> | ||
using can_copy_arguments = __can_copy_arguments_impl<F, void, Args...>; | ||
|
||
//! Wrapper for a callable to mark it as allowing copied arguments | ||
template <typename F> | ||
struct copied_arguments_allowing_wrapper : F | ||
{ | ||
using F::operator(); | ||
static constexpr bool can_copy_arguments = true; | ||
}; | ||
|
||
//! Creates a new function object from an existing one, allowing its arguments to be copies of whatever source they come | ||
//! from. This implies that the addresses of the arguments are irrelevant to the function object. | ||
template <typename F> | ||
_CCCL_HOST_DEVICE constexpr auto allow_copied_arguments(F f) -> copied_arguments_allowing_wrapper<F> | ||
{ | ||
return copied_arguments_allowing_wrapper<F>{_CUDA_VSTD::move(f)}; | ||
} | ||
|
||
_LIBCUDACXX_END_NAMESPACE_STD | ||
|
||
#endif // _LIBCUDACXX___TYPE_TRAITS_ADDRESS_STABILITY_H |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,26 @@ | ||
#include <cuda/std/__type_traits/address_stability.h> | ||
|
||
#include <unittest/unittest.h> | ||
|
||
// TODO(bgruber): move this test into libcu++ | ||
|
||
struct MyPlus | ||
{ | ||
_CCCL_HOST_DEVICE auto operator()(int a, int b) const -> int | ||
{ | ||
return a + b; | ||
} | ||
}; | ||
|
||
void TestAddressStability() | ||
{ | ||
using ::cuda::std::allow_copied_arguments; | ||
using ::cuda::std::can_copy_arguments; | ||
|
||
static_assert(!can_copy_arguments<thrust::plus<int>, int, int>::value, ""); | ||
static_assert(can_copy_arguments<decltype(allow_copied_arguments(thrust::plus<int>{})), int, int>::value, ""); | ||
|
||
static_assert(!can_copy_arguments<MyPlus, int, int>::value, ""); | ||
static_assert(can_copy_arguments<decltype(allow_copied_arguments(MyPlus{})), int, int>::value, ""); | ||
} | ||
DECLARE_UNITTEST(TestAddressStability); |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Oops, something went wrong.