1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
|
---
grand_parent: Extended API
parent: Shapes
---
# `cuda::aligned_size_t`
Defined in headers `<cuda/barrier>` and `<cuda/pipeline>`:
```cuda
template <cuda::std::size_t Alignment>
struct cuda::aligned_size_t {
static constexpr cuda::std::size_t align = Align;
cuda::std::size_t value;
__host__ __device__ explicit constexpr aligned_size(cuda::std::size_t size);
__host__ __device__ constexpr operator cuda::std::size_t();
};
```
The class template `cuda::aligned_size_t` is a _shape_ representing an extent
of bytes with a statically defined (address and size) alignment.
*Preconditions*:
- The _address_ of the extent of bytes must be aligned to an `Alignment` alignment boundary.
- The _size_ of the extent of bytes must be a multiple of the `Alignment`.
## Template Parameters
| `Alignment` | The address and size alignement of the byte extent. |
## Data Members
| `align` | The alignment of the byte extent. |
| `value` | The size of the byte extent. |
## Member Functions
| (constructor) | Constructs an _aligned size_. If the `size` is not a multiple of `Alignment` the behavior is undefined. |
| (destructor) [implicitly declared] | Trivial implicit destructor. |
| `operator=` [implicitly declared] | Trivial implicit copy/move assignment. |
| `operator cuda::std::size_t` | Implicit conversion to [`cuda::std::size_t`]. |
## Notes
If `Alignment` is not a [valid alignment], the behavior is undefined.
## Example
```cuda
#include <cuda/barrier>
__global__ void example_kernel(void* dst, void* src, size_t size) {
cuda::barrier<cuda::thread_scope_system> bar;
init(&bar, 1);
// Implementation cannot make assumptions about alignment.
cuda::memcpy_async(dst, src, size, bar);
// Implementation can assume that dst and src are 16-bytes aligned,
// and that size is a multiple of 16, and may optimize accordingly.
cuda::memcpy_async(dst, src, cuda::aligned_size_t<16>(size), bar);
bar.arrive_and_wait();
}
```
[See it on Godbolt](https://godbolt.org/z/PWGdfTd7d){: .btn }
[valid alignment]: https://en.cppreference.com/w/c/language/object#Alignment
[`cuda::std::size_t`]: https://en.cppreference.com/w/cpp/types/size_t
|
