AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-295298: introduce warp sync functions

Browse Source 
The following builtins from the CUDA spec are implemented:

  - __all_sync, __any_sync, __ballot_sync and __activemask
  - __match_any_sync and __match_all_sync
  - __shfl_sync, __shfl_up_sync, __shfl_down_sync, and __shfl_xor_sync

The following builtins are NOT implemented, pending support in the compiler:

  - __reduce_add_sync, __reduce_min_sync, __reduce_max_sync
  - __reduce_and_sync, __reduce_or_sync, __reduce_xor_sync

Change-Id: I07dedbbfe5449f4b5c9b040bed59f5603ccec8c3
This commit is contained in:
Sameer Sahasrabuddhe
2024-01-11 16:19:34 +05:30
committed by Rahul Garg
parent 4b3c4709ff
commit c5ab5680b4
4 changed files with 306 additions and 34 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/src/hiprtc/CMakeLists.txt
+2 -1
View File
@@ -146,8 +146,9 @@ ${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_surface_functions.h
${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_complex.h
${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_math_constants.h
${PROJECT_SOURCE_DIR}/include/hip/amd_detail/math_fwd.h
${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_device_functions.h
${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_warp_functions.h
${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_device_functions.h
${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_warp_sync_functions.h
${PROJECT_SOURCE_DIR}/include/hip/amd_detail/hip_cooperative_groups_helper.h
${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_cooperative_groups.h
${PROJECT_SOURCE_DIR}/include/hip/amd_detail/amd_hip_unsafe_atomics.h