Skip to content

Commit 1618acb

Browse files
Add checked C2H CUDA buffer generators
1 parent d09fb7f commit 1618acb

11 files changed

Lines changed: 762 additions & 30 deletions

c2h/generators.cu

Lines changed: 49 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -3,20 +3,28 @@
33

44
#include <cub/device/device_copy.cuh>
55

6+
#include <thrust/detail/config/device_system.h>
7+
#include <thrust/detail/raw_pointer_cast.h>
68
#include <thrust/for_each.h>
79
#include <thrust/iterator/counting_iterator.h>
810
#include <thrust/iterator/transform_iterator.h>
911
#include <thrust/tabulate.h>
12+
#include <thrust/version.h>
1013

1114
#include <cuda/iterator>
1215
#include <cuda/std/optional>
16+
#include <cuda/std/span>
17+
#include <cuda/stream>
18+
19+
#include <cstddef>
20+
#include <cstdint>
21+
#include <limits>
1322

1423
#include <c2h/bfloat16.cuh>
1524
#include <c2h/custom_type.h>
1625
#include <c2h/detail/generators.cuh>
1726
#include <c2h/device_policy.h>
1827
#include <c2h/extended_types.h>
19-
#include <c2h/generators.h>
2028
#include <c2h/half.cuh>
2129
#include <c2h/vector.h>
2230

@@ -65,31 +73,46 @@ public:
6573

6674
float* prepare_random_generator(seed_t seed, std::size_t num_items)
6775
{
68-
m_distribution.resize(num_items);
76+
return prepare_random_generator(::cuda::stream_ref{::cudaStream_t{}}, seed, num_items);
77+
}
78+
79+
float* prepare_random_generator(::cuda::stream_ref stream, seed_t seed, std::size_t num_items)
80+
{
81+
resize_distribution(num_items);
6982

7083
#if C2H_HAS_CURAND
7184
curandSetPseudoRandomGeneratorSeed(m_gen, seed.get());
72-
#else
85+
#else // C2H_HAS_CURAND
7386
m_gen.seed(seed.get());
74-
#endif
87+
#endif // C2H_HAS_CURAND
7588

76-
generate();
89+
generate(stream);
7790

7891
return thrust::raw_pointer_cast(m_distribution.data());
7992
}
8093

8194
// re-fills the currently held distribution vector with new random values
82-
void generate()
95+
void generate(::cuda::stream_ref stream)
8396
{
8497
#if C2H_HAS_CURAND
98+
curandSetStream(m_gen, stream.get());
8599
curandGenerateUniform(m_gen, thrust::raw_pointer_cast(m_distribution.data()), m_distribution.size());
86100
#else
87-
thrust::tabulate(device_policy, m_distribution.begin(), m_distribution.end(), i_to_rnd_t{m_gen});
101+
thrust::tabulate(device_policy.on(stream.get()), m_distribution.begin(), m_distribution.end(), i_to_rnd_t{m_gen});
88102
m_gen.discard(m_distribution.size());
89103
#endif
90104
}
91105

92106
private:
107+
void resize_distribution(std::size_t num_items)
108+
{
109+
#if THRUST_VERSION >= 300100
110+
m_distribution.resize(num_items, thrust::no_init);
111+
#else // THRUST_VERSION >= 300100
112+
m_distribution.resize(num_items);
113+
#endif // THRUST_VERSION >= 300100
114+
}
115+
93116
#if C2H_HAS_CURAND
94117
curandGenerator_t
95118
#else
@@ -113,6 +136,11 @@ float* prepare_random_data(seed_t seed, std::size_t num_items)
113136
return generator.value().prepare_random_generator(seed, num_items);
114137
}
115138

139+
float* prepare_random_data(::cuda::stream_ref stream, seed_t seed, std::size_t num_items)
140+
{
141+
return generator.value().prepare_random_generator(stream, seed, num_items);
142+
}
143+
116144
void cleanup_generator()
117145
{
118146
_CCCL_VERIFY(generator.has_value(), "");
@@ -136,6 +164,18 @@ struct random_to_custom_t
136164
};
137165

138166
void gen_custom_type_state(
167+
seed_t seed,
168+
char* d_out,
169+
custom_type_state_t min,
170+
custom_type_state_t max,
171+
std::size_t elements,
172+
std::size_t element_size)
173+
{
174+
gen_custom_type_state(::cuda::stream_ref{::cudaStream_t{}}, seed, d_out, min, max, elements, element_size);
175+
}
176+
177+
void gen_custom_type_state(
178+
::cuda::stream_ref stream,
139179
seed_t seed,
140180
char* d_out,
141181
custom_type_state_t /* min */,
@@ -144,8 +184,8 @@ void gen_custom_type_state(
144184
std::size_t element_size)
145185
{
146186
// FIXME(bgruber): implement min/max handling for custom_type_state_t
147-
float* d_in = prepare_random_data(seed, elements * 2);
148-
thrust::for_each(device_policy,
187+
float* d_in = prepare_random_data(stream, seed, elements * 2);
188+
thrust::for_each(device_policy.on(stream.get()),
149189
thrust::counting_iterator<std::size_t>{0},
150190
thrust::counting_iterator<std::size_t>{elements},
151191
random_to_custom_t{d_in, d_out, element_size});

c2h/generators_gen_values.cu

Lines changed: 14 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -4,20 +4,28 @@
44
#include <thrust/tabulate.h>
55
#include <thrust/transform.h>
66

7+
#include <cuda/stream>
8+
79
#include <c2h/bfloat16.cuh>
810
#include <c2h/detail/generators.cuh>
911
#include <c2h/device_policy.h>
1012
#include <c2h/extended_types.h>
11-
#include <c2h/generators.h>
1213
#include <c2h/half.cuh>
1314

1415
namespace c2h::detail
1516
{
1617
template <typename T>
1718
void gen_values_between(seed_t seed, ::cuda::std::span<T> data, T min, T max)
1819
{
19-
const auto* dist = prepare_random_data(seed, data.size());
20-
thrust::transform(device_policy, dist, dist + data.size(), data.begin(), random_to_item_t<T>(min, max));
20+
gen_values_between(::cuda::stream_ref{::cudaStream_t{}}, seed, data, min, max);
21+
}
22+
23+
template <typename T>
24+
void gen_values_between(::cuda::stream_ref stream, seed_t seed, ::cuda::std::span<T> data, T min, T max)
25+
{
26+
const auto* dist = prepare_random_data(stream, seed, data.size());
27+
thrust::transform(
28+
device_policy.on(stream.get()), dist, dist + data.size(), data.begin(), random_to_item_t<T>(min, max));
2129
}
2230

2331
template <typename T>
@@ -39,8 +47,9 @@ void gen_values_cyclic(modulo_t mod, ::cuda::std::span<T> data)
3947
thrust::tabulate(device_policy, data.begin(), data.end(), counter_to_cyclic_item_t<T>{mod.get()});
4048
}
4149

42-
#define INSTANTIATE_RND(TYPE) \
43-
template void gen_values_between<TYPE>(seed_t, ::cuda::std::span<TYPE> data, TYPE min, TYPE max)
50+
#define INSTANTIATE_RND(TYPE) \
51+
template void gen_values_between<TYPE>(seed_t, ::cuda::std::span<TYPE> data, TYPE min, TYPE max); \
52+
template void gen_values_between<TYPE>(::cuda::stream_ref, seed_t, ::cuda::std::span<TYPE> data, TYPE min, TYPE max)
4453
#define INSTANTIATE_MOD(TYPE) template void gen_values_cyclic<TYPE>(modulo_t, ::cuda::std::span<TYPE> data)
4554

4655
#define INSTANTIATE(TYPE) \

c2h/generators_vector.cu

Lines changed: 20 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -1,16 +1,20 @@
11
// SPDX-FileCopyrightText: Copyright (c) 2011-2022, NVIDIA CORPORATION. All rights reserved.
22
// SPDX-License-Identifier: BSD-3-Clause
33

4+
#include <thrust/detail/config/device_system.h>
5+
#include <thrust/for_each.h>
6+
#include <thrust/iterator/counting_iterator.h>
47
#include <thrust/tabulate.h>
58

69
#include <cuda/std/cstddef>
710
#include <cuda/std/cstdint>
11+
#include <cuda/std/span>
12+
#include <cuda/stream>
813

914
#include <c2h/detail/generators.cuh>
1015
#include <c2h/device_policy.h>
1116
#include <c2h/extended_types.h>
1217
#include <c2h/fill_striped.h>
13-
#include <c2h/generators.h>
1418
#include <c2h/vector.h>
1519

1620
namespace c2h::detail
@@ -49,14 +53,21 @@ struct random_to_vec_item_t
4953
};
5054

5155
#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
52-
# define VEC_SPECIALIZATION(T) \
53-
template <> \
54-
void gen_values_between(seed_t seed, ::cuda::std::span<T> data, T min, T max) \
55-
{ \
56-
const auto* dist = prepare_random_data(seed, data.size()); \
57-
auto op = random_to_vec_item_t<T, ::cuda::std::tuple_size_v<T>>{min, max, dist, data.data()}; \
58-
thrust::for_each( \
59-
device_policy, thrust::counting_iterator<size_t>{0}, thrust::counting_iterator<size_t>{data.size()}, op); \
56+
# define VEC_SPECIALIZATION(T) \
57+
template <> \
58+
void gen_values_between(::cuda::stream_ref stream, seed_t seed, ::cuda::std::span<T> data, T min, T max) \
59+
{ \
60+
const auto* dist = prepare_random_data(stream, seed, data.size()); \
61+
auto op = random_to_vec_item_t<T, ::cuda::std::tuple_size_v<T>>{min, max, dist, data.data()}; \
62+
thrust::for_each(device_policy.on(stream.get()), \
63+
thrust::counting_iterator<size_t>{0}, \
64+
thrust::counting_iterator<size_t>{data.size()}, \
65+
op); \
66+
} \
67+
template <> \
68+
void gen_values_between(seed_t seed, ::cuda::std::span<T> data, T min, T max) \
69+
{ \
70+
gen_values_between(::cuda::stream_ref{::cudaStream_t{}}, seed, data, min, max); \
6071
}
6172

6273
VEC_SPECIALIZATION(char2);
Lines changed: 134 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,134 @@
1+
// SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. All rights reserved.
2+
// SPDX-License-Identifier: BSD-3
3+
4+
#pragma once
5+
6+
#include <cuda/std/detail/__config>
7+
8+
#if _CCCL_HAS_CTK() && !_CCCL_COMPILER(NVRTC)
9+
# include <cuda/__algorithm/copy.h>
10+
# include <cuda/buffer>
11+
# include <cuda/devices>
12+
# include <cuda/std/limits>
13+
# include <cuda/std/utility>
14+
# include <cuda/stream>
15+
#endif // _CCCL_HAS_CTK() && !_CCCL_COMPILER(NVRTC)
16+
17+
#include <cstddef>
18+
19+
#include <c2h/checked_memory_resource.cuh>
20+
#include <c2h/generator_common.h>
21+
22+
namespace c2h
23+
{
24+
#if _CCCL_HAS_CTK() && !_CCCL_COMPILER(NVRTC)
25+
namespace detail
26+
{
27+
template <typename T>
28+
[[nodiscard]] cuda::host_buffer<T> device_buffer_to_host_buffer(
29+
cuda::stream_ref stream, cuda::device_ref device, const cuda::device_buffer<T>& d_items, std::size_t num_items)
30+
{
31+
// Scope `device` for host allocation checks and possible default-stream copies.
32+
// Non-default stream/device agreement is part of the public helper contract.
33+
const ::c2h::detail::scoped_current_device device_scope{device.get()};
34+
35+
auto h_items = ::c2h::make_host_buffer<T>(stream, device, num_items, cuda::no_init);
36+
cuda::copy_bytes(stream, d_items.first(num_items), h_items);
37+
stream.sync();
38+
39+
return h_items;
40+
}
41+
42+
template <typename T>
43+
void gen_into_device_buffer(cuda::stream_ref stream, seed_t seed, cuda::device_buffer<T>& d_items, T min, T max)
44+
{
45+
::c2h::detail::gen_values_between(stream, seed, d_items.first(d_items.size()), min, max);
46+
}
47+
48+
template <template <typename> class... Ps>
49+
void gen_into_device_buffer(
50+
cuda::stream_ref stream,
51+
seed_t seed,
52+
cuda::device_buffer<custom_type_t<Ps...>>& d_items,
53+
custom_type_t<Ps...> min,
54+
custom_type_t<Ps...> max)
55+
{
56+
::c2h::detail::gen_custom_type_state(
57+
stream, seed, reinterpret_cast<char*>(d_items.data()), min, max, d_items.size(), sizeof(custom_type_t<Ps...>));
58+
}
59+
} // namespace detail
60+
61+
// `size` is the number of generated items shared by both buffers. The owning
62+
// buffers may contain additional capacity that is not part of the generated sequence.
63+
template <typename T>
64+
struct sized_device_host_buffers
65+
{
66+
cuda::device_buffer<T> d_items;
67+
cuda::host_buffer<T> h_items;
68+
std::size_t size;
69+
};
70+
71+
/**
72+
* @brief Generates random data with the existing c2h device generator and returns it in device memory.
73+
*
74+
* @pre If `stream` is non-default, it must have been created for `device`.
75+
*/
76+
template <typename T>
77+
[[nodiscard]] cuda::device_buffer<T> gen_device_buffer(
78+
cuda::stream_ref stream,
79+
cuda::device_ref device,
80+
seed_t seed,
81+
std::size_t num_items,
82+
T min = ::cuda::std::numeric_limits<T>::lowest(),
83+
T max = ::cuda::std::numeric_limits<T>::max())
84+
{
85+
// Scope `device` for generator storage backed by current-device allocation.
86+
const ::c2h::detail::scoped_current_device device_scope{device.get()};
87+
88+
auto d_items = ::c2h::make_device_buffer<T>(stream, device, num_items, cuda::no_init);
89+
::c2h::detail::gen_into_device_buffer(stream, seed, d_items, min, max);
90+
91+
return d_items;
92+
}
93+
94+
/**
95+
* @brief Generates random data with the existing c2h device generator and returns device and host buffers.
96+
*
97+
* @pre If `stream` is non-default, it must have been created for `device`.
98+
*/
99+
template <typename T>
100+
[[nodiscard]] sized_device_host_buffers<T> gen_buffers(
101+
cuda::stream_ref stream,
102+
cuda::device_ref device,
103+
seed_t seed,
104+
std::size_t num_items,
105+
T min = ::cuda::std::numeric_limits<T>::lowest(),
106+
T max = ::cuda::std::numeric_limits<T>::max())
107+
{
108+
auto d_items = ::c2h::gen_device_buffer<T>(stream, device, seed, num_items, min, max);
109+
110+
const auto items_count = d_items.size();
111+
auto h_items = ::c2h::detail::device_buffer_to_host_buffer(stream, device, d_items, items_count);
112+
113+
return {::cuda::std::move(d_items), ::cuda::std::move(h_items), items_count};
114+
}
115+
116+
/**
117+
* @brief Generates random data with the existing c2h device generator and returns it in host pageable memory.
118+
*
119+
* @pre If `stream` is non-default, it must have been created for `device`.
120+
*/
121+
template <typename T>
122+
[[nodiscard]] cuda::host_buffer<T> gen_host_buffer(
123+
cuda::stream_ref stream,
124+
cuda::device_ref device,
125+
seed_t seed,
126+
std::size_t num_items,
127+
T min = ::cuda::std::numeric_limits<T>::lowest(),
128+
T max = ::cuda::std::numeric_limits<T>::max())
129+
{
130+
auto buffers = ::c2h::gen_buffers<T>(stream, device, seed, num_items, min, max);
131+
return ::cuda::std::move(buffers.h_items);
132+
}
133+
#endif // _CCCL_HAS_CTK() && !_CCCL_COMPILER(NVRTC)
134+
} // namespace c2h

c2h/include/c2h/checked_allocator.cuh

Lines changed: 14 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -3,9 +3,10 @@
33

44
#pragma once
55

6+
#include <cuda/std/detail/__config>
7+
68
#include <thrust/device_allocator.h>
79
#include <thrust/mr/new.h>
8-
#include <thrust/system/cuda/memory.h>
910
#include <thrust/system/cuda/memory_resource.h>
1011
#include <thrust/system/cuda/pointer.h>
1112

@@ -54,18 +55,24 @@ public:
5455

5556
struct checked_host_memory_resource final : public THRUST_NS_QUALIFIER::mr::new_delete_resource_base
5657
{
57-
void* do_allocate(std::size_t bytes, std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) final
58+
[[nodiscard]] _CCCL_HOST_API void*
59+
do_allocate(std::size_t bytes, std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) final
5860
{
5961
// Some systems with integrated host/device memory have issues with allocating more memory
6062
// than is available. Check the amount of free memory before attempting to allocate on
6163
// integrated systems.
6264
int device = 0;
63-
CubDebugExit(cudaGetDevice(&device));
64-
cudaDeviceProp prop;
65-
CubDebugExit(cudaGetDeviceProperties(&prop, device));
66-
if (prop.integrated)
65+
if (cudaGetDevice(&device) != cudaSuccess)
66+
{
67+
throw std::bad_alloc{};
68+
}
69+
70+
// Validate allocation-size arithmetic before delegating to new_delete_resource_base.
71+
const std::size_t allocation_size = detail::checked_host_allocation_size(bytes, alignment);
72+
73+
if (detail::is_integrated_device(device))
6774
{
68-
auto status = detail::check_free_device_memory(bytes + alignment + sizeof(std::size_t));
75+
const auto status = detail::check_free_device_memory(allocation_size);
6976
if (status != cudaSuccess)
7077
{
7178
throw std::bad_alloc{};

0 commit comments

Comments
 (0)