Skip to content

Commit 2dede26

Browse files
pciolkoszpciolkosz
authored
[cudax->libcu++] Move host_launch to libcu++ (#6536)
1 parent a7f2041 commit 2dede26

File tree

5 files changed

+42
-48
lines changed

5 files changed

+42
-48
lines changed

‎cudax/include/cuda/experimental/__container/async_buffer.cuh‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -27,6 +27,7 @@
2727

2828
#include <cuda/__container/heterogeneous_iterator.h>
2929
#include <cuda/__container/uninitialized_async_buffer.h>
30+
#include <cuda/__launch/host_launch.h>
3031
#include <cuda/__memory_resource/any_resource.h>
3132
#include <cuda/__memory_resource/get_memory_resource.h>
3233
#include <cuda/__memory_resource/properties.h>
@@ -48,7 +49,6 @@
4849

4950
#include <cuda/experimental/__detail/utility.cuh>
5051
#include <cuda/experimental/__execution/policy.cuh>
51-
#include <cuda/experimental/__launch/host_launch.cuh>
5252
#include <cuda/experimental/__utility/ensure_current_device.cuh>
5353

5454
#include <cuda/std/__cccl/prologue.h>
@@ -590,7 +590,7 @@ __fill_n(cuda::stream_ref __stream, _Tp* __first, ::cuda::std::size_t __count, c
590590

591591
if constexpr (_IsHostOnly)
592592
{
593-
::cuda::experimental::host_launch(
593+
::cuda::host_launch(
594594
__stream, ::cuda::std::uninitialized_fill_n<_Tp*, ::cuda::std::size_t, _Tp>, __first, __count, __value);
595595
}
596596
else

‎cudax/include/cuda/experimental/launch.cuh‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -12,7 +12,6 @@
1212
#define __CUDAX_LAUNCH___
1313

1414
#include <cuda/experimental/__launch/configuration.cuh>
15-
#include <cuda/experimental/__launch/host_launch.cuh>
1615
#include <cuda/experimental/__launch/launch.cuh>
1716
#include <cuda/experimental/__launch/param_kind.cuh>
1817
#include <cuda/experimental/__stream/device_transform.cuh>

‎cudax/test/CMakeLists.txt‎

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -66,10 +66,6 @@ foreach(cudax_target IN LISTS cudax_TARGETS)
6666
launch/launch_smoke.cu
6767
)
6868

69-
cudax_add_catch2_test(test_target host_launch ${cudax_target}
70-
launch/host_launch.cu
71-
)
72-
7369
cudax_add_catch2_test(test_target launch_configuration ${cudax_target}
7470
launch/configuration.cu
7571
)

‎cudax/include/cuda/experimental/__launch/host_launch.cuh‎ renamed to ‎libcudacxx/include/cuda/__launch/host_launch.h‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1,15 +1,15 @@
11
//===----------------------------------------------------------------------===//
22
//
3-
// Part of CUDA Experimental in CUDA C++ Core Libraries,
3+
// Part of libcu++, the C++ Standard Library for your entire system,
44
// under the Apache License v2.0 with LLVM Exceptions.
55
// See https://llvm.org/LICENSE.txt for license information.
66
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
77
// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
88
//
99
//===----------------------------------------------------------------------===//
1010

11-
#ifndef _CUDAX__LAUNCH_HOST_LAUNCH
12-
#define _CUDAX__LAUNCH_HOST_LAUNCH
11+
#ifndef _CUDA___LAUNCH_HOST_LAUNCH_H
12+
#define _CUDA___LAUNCH_HOST_LAUNCH_H
1313

1414
#include <cuda/__cccl_config>
1515

@@ -31,8 +31,8 @@
3131

3232
#include <cuda/std/__cccl/prologue.h>
3333

34-
namespace cuda::experimental
35-
{
34+
_CCCL_BEGIN_NAMESPACE_CUDA
35+
3636
template <class _Callable, class... _Args>
3737
struct __stream_callback_data
3838
{
@@ -104,8 +104,8 @@ _CCCL_HOST_API void host_launch(stream_ref __stream, ::cuda::std::reference_wrap
104104
::cuda::__driver::__launchHostFunc(
105105
__stream.get(), __host_func_launcher<_Callable>, ::cuda::std::addressof(__callable.get()));
106106
}
107-
} // namespace cuda::experimental
107+
_CCCL_END_NAMESPACE_CUDA
108108

109109
#include <cuda/std/__cccl/epilogue.h>
110110

111-
#endif // !_CUDAX__LAUNCH_HOST_LAUNCH
111+
#endif // !_CUDA___LAUNCH_HOST_LAUNCH_H

‎cudax/test/launch/host_launch.cu‎ renamed to ‎libcudacxx/test/libcudacxx/cuda/ccclrt/launch/host_launch.cu‎

Lines changed: 33 additions & 34 deletions
Original file line numberDiff line numberDiff line change
@@ -1,44 +1,43 @@
11
//===----------------------------------------------------------------------===//
22
//
3-
// Part of CUDA Experimental in CUDA C++ Core Libraries,
3+
// Part of libcu++, the C++ Standard Library for your entire system,
44
// under the Apache License v2.0 with LLVM Exceptions.
55
// See https://llvm.org/LICENSE.txt for license information.
66
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
77
// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
88
//
99
//===----------------------------------------------------------------------===//
10+
#include <cuda/__launch/host_launch.h>
11+
#include <cuda/__stream/stream.h>
1012
#include <cuda/atomic>
1113
#include <cuda/memory>
1214

13-
#include <cuda/experimental/launch.cuh>
14-
#include <cuda/experimental/stream.cuh>
15-
1615
#include <cooperative_groups.h>
1716
#include <testing.cuh>
1817

19-
void block_stream(cudax::stream_ref stream, cuda::atomic<int>& atomic)
18+
void block_stream(cuda::stream_ref stream, cuda::atomic<int>& atomic)
2019
{
2120
auto block_lambda = [&]() {
2221
while (atomic != 1)
2322
;
2423
};
25-
cudax::host_launch(stream, block_lambda);
24+
cuda::host_launch(stream, block_lambda);
2625
}
2726

28-
void unblock_and_wait_stream(cudax::stream_ref stream, cuda::atomic<int>& atomic)
27+
void unblock_and_wait_stream(cuda::stream_ref stream, cuda::atomic<int>& atomic)
2928
{
30-
CUDAX_REQUIRE(!stream.is_done());
29+
CCCLRT_REQUIRE(!stream.is_done());
3130
atomic = 1;
3231
stream.sync();
3332
atomic = 0;
3433
}
3534

36-
void launch_local_lambda(cudax::stream_ref stream, int& set, int set_to)
35+
void launch_local_lambda(cuda::stream_ref stream, int& set, int set_to)
3736
{
3837
auto lambda = [&set, set_to]() {
3938
set = set_to;
4039
};
41-
cudax::host_launch(stream, lambda);
40+
cuda::host_launch(stream, lambda);
4241
}
4342

4443
template <typename Lambda>
@@ -58,7 +57,7 @@ struct lambda_wrapper
5857
if constexpr (cuda::std::is_same_v<cuda::std::invoke_result_t<Lambda>, void*>)
5958
{
6059
// If lambda returns the address it captured, confirm this object wasn't moved
61-
CUDAX_REQUIRE(lambda() == this);
60+
CCCLRT_REQUIRE(lambda() == this);
6261
}
6362
else
6463
{
@@ -69,7 +68,7 @@ struct lambda_wrapper
6968
// Make sure we fail if const is added to this wrapper anywhere
7069
void operator()() const
7170
{
72-
CUDAX_REQUIRE(false);
71+
CCCLRT_REQUIRE(false);
7372
}
7473
};
7574

@@ -107,13 +106,13 @@ private:
107106
MoveOnlyCallable() = default;
108107
};
109108

110-
C2H_TEST("Host launch", "")
109+
C2H_CCCLRT_TEST("Host launch", "")
111110
{
112111
cuda::device_ref device{0};
113112
device.init();
114113

115114
cuda::atomic<int> atomic = 0;
116-
cudax::stream stream{device};
115+
cuda::stream stream{device};
117116
int i = 0;
118117

119118
auto set_lambda = [&](int set) {
@@ -124,35 +123,35 @@ C2H_TEST("Host launch", "")
124123
{
125124
block_stream(stream, atomic);
126125

127-
cudax::host_launch(stream, set_lambda, 2);
126+
cuda::host_launch(stream, set_lambda, 2);
128127

129128
unblock_and_wait_stream(stream, atomic);
130-
CUDAX_REQUIRE(i == 2);
129+
CCCLRT_REQUIRE(i == 2);
131130
}
132131

133132
SECTION("Can launch multiple functions")
134133
{
135134
block_stream(stream, atomic);
136135
auto check_lambda = [&]() {
137-
CUDAX_REQUIRE(i == 4);
136+
CCCLRT_REQUIRE(i == 4);
138137
};
139138

140-
cudax::host_launch(stream, set_lambda, 3);
141-
cudax::host_launch(stream, set_lambda, 4);
142-
cudax::host_launch(stream, check_lambda);
143-
cudax::host_launch(stream, set_lambda, 5);
139+
cuda::host_launch(stream, set_lambda, 3);
140+
cuda::host_launch(stream, set_lambda, 4);
141+
cuda::host_launch(stream, check_lambda);
142+
cuda::host_launch(stream, set_lambda, 5);
144143
unblock_and_wait_stream(stream, atomic);
145-
CUDAX_REQUIRE(i == 5);
144+
CCCLRT_REQUIRE(i == 5);
146145
}
147146

148147
SECTION("Non trivially copyable")
149148
{
150149
std::string s = "hello";
151150

152-
cudax::host_launch(
151+
cuda::host_launch(
153152
stream,
154153
[&](auto str_arg) {
155-
CUDAX_REQUIRE(s == str_arg);
154+
CCCLRT_REQUIRE(s == str_arg);
156155
},
157156
s);
158157
stream.sync();
@@ -164,17 +163,17 @@ C2H_TEST("Host launch", "")
164163
i = 21;
165164
});
166165

167-
cudax::host_launch(stream, wrapped_lambda);
166+
cuda::host_launch(stream, wrapped_lambda);
168167
stream.sync();
169-
CUDAX_REQUIRE(i == 21)
168+
CCCLRT_REQUIRE(i == 21);
170169
}
171170

172171
SECTION("Can launch a local function and return")
173172
{
174173
block_stream(stream, atomic);
175174
launch_local_lambda(stream, i, 42);
176175
unblock_and_wait_stream(stream, atomic);
177-
CUDAX_REQUIRE(i == 42);
176+
CCCLRT_REQUIRE(i == 42);
178177
}
179178

180179
SECTION("Launch by reference")
@@ -188,9 +187,9 @@ C2H_TEST("Host launch", "")
188187
wrapper_ptr = static_cast<void*>(&another_lambda_setter);
189188

190189
block_stream(stream, atomic);
191-
host_launch(stream, cuda::std::ref(another_lambda_setter));
190+
cuda::host_launch(stream, cuda::std::ref(another_lambda_setter));
192191
unblock_and_wait_stream(stream, atomic);
193-
CUDAX_REQUIRE(i == 84);
192+
CCCLRT_REQUIRE(i == 84);
194193
}
195194

196195
SECTION("Launch by reference with arguments")
@@ -201,9 +200,9 @@ C2H_TEST("Host launch", "")
201200
result = j;
202201
};
203202
block_stream(stream, atomic);
204-
host_launch(stream, cuda::std::ref(lambda), i);
203+
cuda::host_launch(stream, cuda::std::ref(lambda), i);
205204
unblock_and_wait_stream(stream, atomic);
206-
CUDAX_REQUIRE(result == 10);
205+
CCCLRT_REQUIRE(result == 10);
207206
}
208207

209208
SECTION("Launch by reference with arguments captured by reference")
@@ -213,14 +212,14 @@ C2H_TEST("Host launch", "")
213212
j = 10;
214213
};
215214
block_stream(stream, atomic);
216-
host_launch(stream, cuda::std::ref(lambda), cuda::std::ref(i));
215+
cuda::host_launch(stream, cuda::std::ref(lambda), cuda::std::ref(i));
217216
unblock_and_wait_stream(stream, atomic);
218-
CUDAX_REQUIRE(i == 10);
217+
CCCLRT_REQUIRE(i == 10);
219218
}
220219

221220
SECTION("Check that host_launch works with move only callables and arguments")
222221
{
223-
cudax::host_launch(stream, MoveOnlyCallable::make(), MoveOnlyArg::make());
222+
cuda::host_launch(stream, MoveOnlyCallable::make(), MoveOnlyArg::make());
224223
stream.sync();
225224
}
226225
}

0 commit comments

Comments
 (0)