Movatterモバイル変換

[0]ホーム
URL:

Skip to content

Navigation Menu

Sign in
Appearance settings

Search code, repositories, users, issues, pull requests...

Provide feedback

We read every piece of feedback, and take your input very seriously.

Saved searches

Use saved searches to filter your results more quickly

 Sign up
Appearance settings
This repository was archived by the owner on Mar 21, 2024. It is now read-only.
/libcudacxxPublic archive

Commit26d57af

Browse files
committed
Define internal macros for__host__ and__device__ to isolate against include order issues
1 parent070697d commit26d57af

File tree

31 files changed

+863
-878
lines changed

31 files changed

+863
-878
lines changed

‎.upstream-tests/test/std/containers/views/mdspan/foo_customizations.hpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,7 @@ namespace Foo {
77
template<classT>
88
structfoo_ptr {
99
T* data;
10-
__MDSPAN_HOST_DEVICE
10+
_LIBCUDACXX_HOST_DEVICE
1111
constexprfoo_ptr(T* ptr):data(ptr) {}
1212
};
1313

@@ -75,7 +75,7 @@ class layout_foo::mapping {
7575
__MDSPAN_INLINE_FUNCTION_DEFAULTEDconstexprmapping()noexcept = default;
7676
__MDSPAN_INLINE_FUNCTION_DEFAULTEDconstexprmapping(mappingconst&)noexcept = default;
7777

78-
__MDSPAN_HOST_DEVICE
78+
_LIBCUDACXX_HOST_DEVICE
7979
constexprmapping(extents_typeconst& __exts)noexcept
8080
:__extents(__exts)
8181
{ }

‎codegen/codegen.cpp‎

Lines changed: 12 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -99,10 +99,10 @@ int main() {
9999
};
100100

101101
for(auto& s : scopes) {
102-
out <<"static inline__device__ void __cuda_membar_" << s.first <<"() { asm volatile(\"membar" << membar_scopes[s.first] <<";\":::\"memory\"); }\n";
102+
out <<"static inline_LIBCUDACXX_DEVICE void __cuda_membar_" << s.first <<"() { asm volatile(\"membar" << membar_scopes[s.first] <<";\":::\"memory\"); }\n";
103103
for(auto& sem : fence_semantics)
104-
out <<"static inline__device__ void" <<fencename(sem.first, s.first) <<"() { asm volatile(\"fence" << sem.second << s.second <<";\":::\"memory\"); }\n";
105-
out <<"static inline__device__ void __atomic_thread_fence_cuda(int __memorder," <<scopenametag(s.first) <<") {\n";
104+
out <<"static inline_LIBCUDACXX_DEVICE void" <<fencename(sem.first, s.first) <<"() { asm volatile(\"fence" << sem.second << s.second <<";\":::\"memory\"); }\n";
105+
out <<"static inline_LIBCUDACXX_DEVICE void __atomic_thread_fence_cuda(int __memorder," <<scopenametag(s.first) <<") {\n";
106106
out <<" NV_DISPATCH_TARGET(\n";
107107
out <<" NV_PROVIDES_SM_70, (\n";
108108
out <<" switch (__memorder) {\n";
@@ -131,7 +131,7 @@ int main() {
131131
for(auto& sz : ld_sizes) {
132132
for(auto& sem : ld_semantics) {
133133
out <<"template<class _CUDA_A, class _CUDA_B>";
134-
out <<"static inline__device__ void __cuda_load_" << sem.first <<"_" << sz <<"_" << s.first <<"(_CUDA_A __ptr, _CUDA_B& __dst) {";
134+
out <<"static inline_LIBCUDACXX_DEVICE void __cuda_load_" << sem.first <<"_" << sz <<"_" << s.first <<"(_CUDA_A __ptr, _CUDA_B& __dst) {";
135135
if(ld_as_atom)
136136
out <<"asm volatile(\"atom.add" << (sem.first =="volatile" ?"" : sem.second.c_str()) << s.second <<".u" << sz <<" %0, [%1], 0;\" :";
137137
else
@@ -141,7 +141,7 @@ int main() {
141141
}
142142
for(auto& cv: cv_qualifier) {
143143
out <<"template<class _Type, typename _CUDA_VSTD::enable_if<sizeof(_Type)==" << sz/8 <<", int>::type = 0>\n";
144-
out <<"__device__ void __atomic_load_cuda(const" << cv <<"_Type *__ptr, _Type *__ret, int __memorder," <<scopenametag(s.first) <<") {\n";
144+
out <<"_LIBCUDACXX_DEVICE void __atomic_load_cuda(const" << cv <<"_Type *__ptr, _Type *__ret, int __memorder," <<scopenametag(s.first) <<") {\n";
145145
out <<" uint" << sz <<"_t __tmp = 0;\n";
146146
out <<" NV_DISPATCH_TARGET(\n";
147147
out <<" NV_PROVIDES_SM_70, (\n";
@@ -170,14 +170,14 @@ int main() {
170170
for(auto& sz : st_sizes) {
171171
for(auto& sem : st_semantics) {
172172
out <<"template<class _CUDA_A, class _CUDA_B>";
173-
out <<"static inline__device__ void __cuda_store_" << sem.first <<"_" << sz <<"_" << s.first <<"(_CUDA_A __ptr, _CUDA_B __src) {";
173+
out <<"static inline_LIBCUDACXX_DEVICE void __cuda_store_" << sem.first <<"_" << sz <<"_" << s.first <<"(_CUDA_A __ptr, _CUDA_B __src) {";
174174
out <<"asm volatile(\"st" << sem.second << (sem.first =="volatile" ?"" : s.second.c_str()) <<".b" << sz <<" [%0], %1;\" ::";
175175
out <<"\"l\"(__ptr),\"" <<registers("b", sz) <<"\"(__src)";
176176
out <<" :\"memory\"); }\n";
177177
}
178178
for(auto& cv: cv_qualifier) {
179179
out <<"template<class _Type, typename cuda::std::enable_if<sizeof(_Type)==" << sz/8 <<", int>::type = 0>\n";
180-
out <<"__device__ void __atomic_store_cuda(" << cv <<"_Type *__ptr, _Type *__val, int __memorder," <<scopenametag(s.first) <<") {\n";
180+
out <<"_LIBCUDACXX_DEVICE void __atomic_store_cuda(" << cv <<"_Type *__ptr, _Type *__val, int __memorder," <<scopenametag(s.first) <<") {\n";
181181
out <<" uint" << sz <<"_t __tmp = 0;\n";
182182
out <<" memcpy(&__tmp, __val," << sz/8 <<");\n";
183183
out <<" NV_DISPATCH_TARGET(\n";
@@ -215,7 +215,7 @@ int main() {
215215
out <<"template<class _CUDA_A, class _CUDA_B, class _CUDA_C, class _CUDA_D>";
216216
else
217217
out <<"template<class _CUDA_A, class _CUDA_B, class _CUDA_C>";
218-
out <<"static inline__device__ void __cuda_" << rmw.first <<"_" << sem.first <<"_" << type.first << sz <<"_" << s.first <<"(";
218+
out <<"static inline_LIBCUDACXX_DEVICE void __cuda_" << rmw.first <<"_" << sem.first <<"_" << type.first << sz <<"_" << s.first <<"(";
219219
if(rmw.first =="compare_exchange")
220220
out <<"_CUDA_A __ptr, _CUDA_B& __dst, _CUDA_C __cmp, _CUDA_D __op";
221221
else
@@ -238,7 +238,7 @@ int main() {
238238
for(auto& cv: cv_qualifier) {
239239
if(rmw.first =="compare_exchange") {
240240
out <<"template<class _Type, typename cuda::std::enable_if<sizeof(_Type)==" << sz/8 <<", int>::type = 0>\n";
241-
out <<"__device__ bool __atomic_compare_exchange_cuda(" << cv <<"_Type *__ptr, _Type *__expected, const _Type *__desired, bool, int __success_memorder, int __failure_memorder," <<scopenametag(s.first) <<") {\n";
241+
out <<"_LIBCUDACXX_DEVICE bool __atomic_compare_exchange_cuda(" << cv <<"_Type *__ptr, _Type *__expected, const _Type *__desired, bool, int __success_memorder, int __failure_memorder," <<scopenametag(s.first) <<") {\n";
242242
out <<" uint" << sz <<"_t __tmp = 0, __old = 0, __old_tmp;\n";
243243
out <<" memcpy(&__tmp, __desired," << sz/8 <<");\n";
244244
out <<" memcpy(&__old, __expected," << sz/8 <<");\n";
@@ -276,7 +276,7 @@ int main() {
276276
out <<"template<class _Type, typename cuda::std::enable_if<sizeof(_Type)==" << sz/8;
277277
if(rmw.first =="exchange") {
278278
out <<", int>::type = 0>\n";
279-
out <<"__device__ void __atomic_exchange_cuda(" << cv <<"_Type *__ptr, _Type *__val, _Type *__ret, int __memorder," <<scopenametag(s.first) <<") {\n";
279+
out <<"_LIBCUDACXX_DEVICE void __atomic_exchange_cuda(" << cv <<"_Type *__ptr, _Type *__val, _Type *__ret, int __memorder," <<scopenametag(s.first) <<") {\n";
280280
out <<" uint" << sz <<"_t __tmp = 0;\n";
281281
out <<" memcpy(&__tmp, __val," << sz/8 <<");\n";
282282
}
@@ -293,7 +293,7 @@ int main() {
293293
out <<" && cuda::std::is_integral<_Type>::value, int>::type = 0>\n";
294294
else
295295
out <<", int>::type = 0>\n";
296-
out <<"__device__ _Type __atomic_" << rmw.first <<"_cuda(" << cv <<"_Type *__ptr, _Type __val, int __memorder," <<scopenametag(s.first) <<") {\n";
296+
out <<"_LIBCUDACXX_DEVICE _Type __atomic_" << rmw.first <<"_cuda(" << cv <<"_Type *__ptr, _Type __val, int __memorder," <<scopenametag(s.first) <<") {\n";
297297
out <<" _Type __ret;\n";
298298
if(type.first =="f" && sz ==32)
299299
out <<" float";
@@ -345,7 +345,7 @@ int main() {
345345
std::vector<std::string> addsub{"add","sub" };
346346
for(auto& op : addsub) {
347347
out <<"template<class _Type>\n";
348-
out <<"__device__ _Type* __atomic_fetch_" << op <<"_cuda(_Type *" << cv <<"*__ptr, ptrdiff_t __val, int __memorder," <<scopenametag(s.first) <<") {\n";
348+
out <<"_LIBCUDACXX_DEVICE _Type* __atomic_fetch_" << op <<"_cuda(_Type *" << cv <<"*__ptr, ptrdiff_t __val, int __memorder," <<scopenametag(s.first) <<") {\n";
349349
out <<" _Type* __ret;\n";
350350
out <<" uint64_t __tmp = 0;\n";
351351
out <<" memcpy(&__tmp, &__val, 8);\n";

0 commit comments

Comments
 (0)
[8]ページ先頭
©2009-2025 Movatter.jp