Skip to content

Commit

Permalink
Ensure that cuda::aligned_size_t is usable in a constepxr context
Browse files Browse the repository at this point in the history
Fixes [BUG]: `aligned_size_t` constructor and conversion operator are not constexpr NVIDIA#1563
  • Loading branch information
miscco committed Mar 23, 2024
1 parent 4182329 commit 8aef726
Show file tree
Hide file tree
Showing 2 changed files with 49 additions and 8 deletions.
20 changes: 12 additions & 8 deletions libcudacxx/include/cuda/std/detail/libcxx/include/__cuda/barrier.h
Original file line number Diff line number Diff line change
Expand Up @@ -46,14 +46,18 @@ _LIBCUDACXX_BEGIN_NAMESPACE_CUDA
template<thread_scope _Scope>
class pipeline;

template<_CUDA_VSTD::size_t _Alignment>
struct aligned_size_t {
static constexpr _CUDA_VSTD::size_t align = _Alignment;
_CUDA_VSTD::size_t value;
_LIBCUDACXX_INLINE_VISIBILITY
explicit aligned_size_t(size_t __s) : value(__s) { }
_LIBCUDACXX_INLINE_VISIBILITY
operator size_t() const { return value; }
template <_CUDA_VSTD::size_t _Alignment>
struct aligned_size_t
{
static constexpr _CUDA_VSTD::size_t align = _Alignment;
_CUDA_VSTD::size_t value;
_LIBCUDACXX_INLINE_VISIBILITY explicit constexpr aligned_size_t(size_t __s)
: value(__s)
{}
_LIBCUDACXX_INLINE_VISIBILITY constexpr operator size_t() const
{
return value;
}
};

// Type only used for logging purpose
Expand Down
37 changes: 37 additions & 0 deletions libcudacxx/test/libcudacxx/cuda/barrier/aligned_size_t.pass.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,37 @@
//===----------------------------------------------------------------------===//
//
// Part of libcu++, the C++ Standard Library for your entire system,
// 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.
//
//===----------------------------------------------------------------------===//

// UNSUPPORTED: libcpp-has-no-threads
// UNSUPPORTED: pre-sm-70

// <cuda/barrier>

#include <cuda/barrier>
#include <cuda/std/cassert>
#include <cuda/std/type_traits>

__host__ __device__ constexpr bool test() {
using aligned_t = cuda::aligned_size_t<1>;
static_assert(!cuda::std::is_default_constructible<aligned_t>::value, "");
static_assert(aligned_t::align == 1, "");
{
const aligned_t aligned{42};
assert(aligned.value == 42);
assert(static_cast<cuda::std::size_t>(aligned) == 42);
}
return true;
}

int main(int, char**)
{
test();
static_assert(test(), "");
return 0;
}

0 comments on commit 8aef726

Please sign in to comment.