-
Notifications
You must be signed in to change notification settings - Fork 42
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Co-authored-by: Binyang Li <[email protected]>
- Loading branch information
1 parent
dab19e0
commit 544ff0c
Showing
65 changed files
with
770 additions
and
398 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,25 @@ | ||
| # Copyright (c) Microsoft Corporation. | ||
| # Licensed under the MIT license. | ||
|
|
||
| set(AMD_FOUND "FALSE") | ||
|
|
||
| set(CMAKE_PREFIX_PATH "/opt/rocm;${CMAKE_PREFIX_PATH}") | ||
| # Temporal fix for rocm5.6 | ||
| set(ENV{amd_comgr_DIR} "/opt/rocm/lib/cmake/amd_comgr") | ||
| set(ENV{AMDDeviceLibs_DIR} "/opt/rocm/lib/cmake/AMDDeviceLibs") | ||
|
|
||
| find_package(hip QUIET) | ||
|
|
||
| if(NOT hip_FOUND) | ||
| return() | ||
| endif() | ||
|
|
||
| enable_language(HIP) | ||
|
|
||
| set(CHECK_SRC "${CMAKE_CURRENT_SOURCE_DIR}/cmake/check_amd_gpu.hip") | ||
|
|
||
| try_run(RUN_RESULT COMPILE_SUCCESS SOURCES ${CHECK_SRC}) | ||
|
|
||
| if(COMPILE_SUCCESS AND RUN_RESULT EQUAL 0) | ||
| set(AMD_FOUND "TRUE") | ||
| endif() |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,36 @@ | ||
| # Copyright (c) Microsoft Corporation. | ||
| # Licensed under the MIT license. | ||
|
|
||
| set(NVIDIA_FOUND "FALSE") | ||
|
|
||
| find_package(CUDAToolkit) | ||
|
|
||
| if(NOT CUDAToolkit_FOUND) | ||
| return() | ||
| endif() | ||
|
|
||
| set(CMAKE_CUDA_ARCHITECTURES "60") | ||
| if(NOT CMAKE_CUDA_COMPILER) | ||
| # In case the CUDA Toolkit directory is not in the PATH | ||
| find_program(CUDA_COMPILER | ||
| NAMES nvcc | ||
| PATHS ${CUDAToolkit_BIN_DIR}) | ||
| if(NOT CUDA_COMPILER) | ||
| message(WARNING "Could not find nvcc in ${CUDAToolkit_BIN_DIR}") | ||
| unset(CMAKE_CUDA_ARCHITECTURES) | ||
| return() | ||
| endif() | ||
| set(CMAKE_CUDA_COMPILER "${CUDA_COMPILER}") | ||
| endif() | ||
| enable_language(CUDA) | ||
|
|
||
| set(CHECK_SRC "${CMAKE_CURRENT_SOURCE_DIR}/cmake/check_nvidia_gpu.cu") | ||
|
|
||
| try_run(RUN_RESULT COMPILE_SUCCESS SOURCES ${CHECK_SRC}) | ||
|
|
||
| if(COMPILE_SUCCESS AND RUN_RESULT EQUAL 0) | ||
| set(NVIDIA_FOUND "TRUE") | ||
| else() | ||
| unset(CMAKE_CUDA_ARCHITECTURES) | ||
| unset(CMAKE_CUDA_COMPILER) | ||
| endif() |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,15 @@ | ||
| // Copyright (c) Microsoft Corporation. | ||
| // Licensed under the MIT license. | ||
|
|
||
| #include <hip/hip_runtime.h> | ||
|
|
||
| __global__ void kernel() {} | ||
|
|
||
| int main() { | ||
| int cnt; | ||
| hipError_t err = hipGetDeviceCount(&cnt); | ||
| if (err != hipSuccess || cnt == 0) { | ||
| return 1; | ||
| } | ||
| return 0; | ||
| } |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,15 @@ | ||
| // Copyright (c) Microsoft Corporation. | ||
| // Licensed under the MIT license. | ||
|
|
||
| #include <cuda_runtime.h> | ||
|
|
||
| __global__ void kernel() {} | ||
|
|
||
| int main() { | ||
| int cnt; | ||
| cudaError_t err = cudaGetDeviceCount(&cnt); | ||
| if (err != cudaSuccess || cnt == 0) { | ||
| return 1; | ||
| } | ||
| return 0; | ||
| } |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,65 @@ | ||
| // Copyright (c) Microsoft Corporation. | ||
| // Licensed under the MIT license. | ||
|
|
||
| #ifndef MSCCLPP_ATOMIC_DEVICE_HPP_ | ||
| #define MSCCLPP_ATOMIC_DEVICE_HPP_ | ||
|
|
||
| #include "device.hpp" | ||
|
|
||
| #if defined(MSCCLPP_DEVICE_CUDA) | ||
| #include <cuda/atomic> | ||
| #endif // defined(MSCCLPP_DEVICE_CUDA) | ||
|
|
||
| namespace mscclpp { | ||
|
|
||
| #if defined(MSCCLPP_DEVICE_CUDA) | ||
|
|
||
| constexpr cuda::memory_order memoryOrderRelaxed = cuda::memory_order_relaxed; | ||
| constexpr cuda::memory_order memoryOrderAcquire = cuda::memory_order_acquire; | ||
| constexpr cuda::memory_order memoryOrderRelease = cuda::memory_order_release; | ||
| constexpr cuda::memory_order memoryOrderAcqRel = cuda::memory_order_acq_rel; | ||
| constexpr cuda::memory_order memoryOrderSeqCst = cuda::memory_order_seq_cst; | ||
|
|
||
| template <typename T> | ||
| MSCCLPP_HOST_DEVICE_INLINE T atomicLoad(T* ptr, cuda::memory_order memoryOrder) { | ||
| return cuda::atomic_ref<T, cuda::thread_scope_system>{*ptr}.load(memoryOrder); | ||
| } | ||
|
|
||
| template <typename T> | ||
| MSCCLPP_HOST_DEVICE_INLINE void atomicStore(T* ptr, const T& val, cuda::memory_order memoryOrder) { | ||
| cuda::atomic_ref<T, cuda::thread_scope_system>{*ptr}.store(val, memoryOrder); | ||
| } | ||
|
|
||
| template <typename T> | ||
| MSCCLPP_HOST_DEVICE_INLINE T atomicFetchAdd(T* ptr, const T& val, cuda::memory_order memoryOrder) { | ||
| return cuda::atomic_ref<T, cuda::thread_scope_system>{*ptr}.fetch_add(val, memoryOrder); | ||
| } | ||
|
|
||
| #elif defined(MSCCLPP_DEVICE_HIP) | ||
|
|
||
| constexpr auto memoryOrderRelaxed = __ATOMIC_RELAXED; | ||
| constexpr auto memoryOrderAcquire = __ATOMIC_ACQUIRE; | ||
| constexpr auto memoryOrderRelease = __ATOMIC_RELEASE; | ||
| constexpr auto memoryOrderAcqRel = __ATOMIC_ACQ_REL; | ||
| constexpr auto memoryOrderSeqCst = __ATOMIC_SEQ_CST; | ||
|
|
||
| template <typename T> | ||
| MSCCLPP_HOST_DEVICE_INLINE T atomicLoad(const T* ptr, int memoryOrder) { | ||
| return __atomic_load_n(ptr, memoryOrder); | ||
| } | ||
|
|
||
| template <typename T> | ||
| MSCCLPP_HOST_DEVICE_INLINE void atomicStore(T* ptr, const T& val, int memoryOrder) { | ||
| __atomic_store_n(ptr, val, memoryOrder); | ||
| } | ||
|
|
||
| template <typename T> | ||
| MSCCLPP_HOST_DEVICE_INLINE T atomicFetchAdd(T* ptr, const T& val, int memoryOrder) { | ||
| return __atomic_fetch_add(ptr, val, memoryOrder); | ||
| } | ||
|
|
||
| #endif // defined(MSCCLPP_DEVICE_HIP) | ||
|
|
||
| } // namespace mscclpp | ||
|
|
||
| #endif // MSCCLPP_ATOMIC_DEVICE_HPP_ |
Oops, something went wrong.