Skip to content

release/20.x: cuda clang: Fix argument order for __reduce_max_sync (#132881) #134295

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Apr 11, 2025
Merged

release/20.x: cuda clang: Fix argument order for __reduce_max_sync (#132881) #134295

Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
16 changes: 8 additions & 8 deletions clang/lib/Headers/__clang_cuda_intrinsics.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -515,32 +515,32 @@ __device__ inline cuuint32_t __nvvm_get_smem_pointer(void *__ptr) {
#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 800
__device__ inline unsigned __reduce_add_sync(unsigned __mask,
unsigned __value) {
return __nvvm_redux_sync_add(__mask, __value);
return __nvvm_redux_sync_add(__value, __mask);
}
__device__ inline unsigned __reduce_min_sync(unsigned __mask,
unsigned __value) {
return __nvvm_redux_sync_umin(__mask, __value);
return __nvvm_redux_sync_umin(__value, __mask);
}
__device__ inline unsigned __reduce_max_sync(unsigned __mask,
unsigned __value) {
return __nvvm_redux_sync_umax(__mask, __value);
return __nvvm_redux_sync_umax(__value, __mask);
}
__device__ inline int __reduce_min_sync(unsigned __mask, int __value) {
return __nvvm_redux_sync_min(__mask, __value);
return __nvvm_redux_sync_min(__value, __mask);
}
__device__ inline int __reduce_max_sync(unsigned __mask, int __value) {
return __nvvm_redux_sync_max(__mask, __value);
return __nvvm_redux_sync_max(__value, __mask);
}
__device__ inline unsigned __reduce_or_sync(unsigned __mask, unsigned __value) {
return __nvvm_redux_sync_or(__mask, __value);
return __nvvm_redux_sync_or(__value, __mask);
}
__device__ inline unsigned __reduce_and_sync(unsigned __mask,
unsigned __value) {
return __nvvm_redux_sync_and(__mask, __value);
return __nvvm_redux_sync_and(__value, __mask);
}
__device__ inline unsigned __reduce_xor_sync(unsigned __mask,
unsigned __value) {
return __nvvm_redux_sync_xor(__mask, __value);
return __nvvm_redux_sync_xor(__value, __mask);
}

__device__ inline void __nv_memcpy_async_shared_global_4(void *__dst,
Expand Down
Loading