From ef72944015f5824b89ac33dd456b5d64e65a57e3 Mon Sep 17 00:00:00 2001 From: BertanDogancay Date: Thu, 15 Feb 2024 13:58:16 -0800 Subject: [PATCH] Disable unsupported ld/st instructions [ROCm/rccl commit: 6f3310605c356a7b7acf389d91ed6585737fdd6e] --- .../rccl/src/device/network/unpack/unpack.h | 6 +- projects/rccl/src/device/op128.h | 124 ++++++++---------- projects/rccl/src/device/reduce_kernel.h | 30 ++--- 3 files changed, 68 insertions(+), 92 deletions(-) diff --git a/projects/rccl/src/device/network/unpack/unpack.h b/projects/rccl/src/device/network/unpack/unpack.h index a218e44735..74f5b789cb 100644 --- a/projects/rccl/src/device/network/unpack/unpack.h +++ b/projects/rccl/src/device/network/unpack/unpack.h @@ -21,10 +21,8 @@ inline __device__ void load64gpu(const uint64_t* ptr, uint64_t &v) { asm volatile("ld.relaxed.gpu.u64 {%0}, [%1];" : "=l"(v) : "l"(ptr)); #else -#ifdef NEED_CHECKING - asm volatile("ld.volatile.global.u64 {%0}, [%1];" - : "=l"(v) : "l"(ptr)); -#endif + // asm volatile("ld.volatile.global.u64 {%0}, [%1];" + // : "=l"(v) : "l"(ptr)); #endif } diff --git a/projects/rccl/src/device/op128.h b/projects/rccl/src/device/op128.h index e5e69e28c5..2bf7d07396 100644 --- a/projects/rccl/src/device/op128.h +++ b/projects/rccl/src/device/op128.h @@ -162,38 +162,22 @@ __device__ __forceinline__ T fromPack(typename BytePackOf::Pack pack) { // Load/store of BytePack using integral addresses. template __device__ BytePack ld_global(uintptr_t addr); +// template __device__ BytePack ld_shared(uint32_t addr); template __device__ BytePack ld_volatile_global(uintptr_t addr); -//template __device__ BytePack ld_shared(uint32_t addr); -//template __device__ BytePack ld_volatile_shared(uint32_t addr); +// template __device__ BytePack ld_volatile_shared(uint32_t addr); +// template __device__ BytePack ld_relaxed_gpu_global(uintptr_t addr); template __device__ void st_global(uintptr_t addr, BytePack value); -//template __device__ void st_shared(uint32_t addr, BytePack value); +// template __device__ void st_shared(uint32_t addr, BytePack value); +// template __device__ void st_relaxed_gpu_global(uintptr_t addr, BytePack value); template<> __device__ __forceinline__ BytePack<0> ld_global<0>(uintptr_t addr) { return {}; } +// template<> __device__ __forceinline__ BytePack<0> ld_shared<0>(uint32_t addr) { return {}; } template<> __device__ __forceinline__ BytePack<0> ld_volatile_global<0>(uintptr_t addr) { return {}; } -//template<> __device__ __forceinline__ BytePack<0> ld_shared<0>(uint32_t addr) { return {}; } -//template<> __device__ __forceinline__ BytePack<0> ld_volatile_shared<0>(uint32_t addr) { return {}; } +// template<> __device__ __forceinline__ BytePack<0> ld_volatile_shared<0>(uint32_t addr) { return {}; } +// template<> __device__ __forceinline__ BytePack<0> ld_relaxed_gpu_global<0>(uintptr_t addr) { return {}; } template<> __device__ __forceinline__ void st_global<0>(uintptr_t addr, BytePack<0> value) {} -//template<> __device__ __forceinline__ void st_shared<0>(uint32_t addr, BytePack<0> value) {} - -#ifdef NEED_CHECKING -template __device__ BytePack ld_global(uintptr_t addr); -template __device__ BytePack ld_shared(uint32_t addr); -template __device__ BytePack ld_volatile_global(uintptr_t addr); -template __device__ BytePack ld_volatile_shared(uint32_t addr); -template __device__ BytePack ld_relaxed_gpu_global(uintptr_t addr); -template __device__ void st_global(uintptr_t addr, BytePack value); -template __device__ void st_shared(uint32_t addr, BytePack value); -template __device__ void st_relaxed_gpu_global(uintptr_t addr, BytePack value); - -template<> __device__ __forceinline__ BytePack<0> ld_global<0>(uintptr_t addr) { return {}; } -template<> __device__ __forceinline__ BytePack<0> ld_shared<0>(uint32_t addr) { return {}; } -template<> __device__ __forceinline__ BytePack<0> ld_volatile_global<0>(uintptr_t addr) { return {}; } -template<> __device__ __forceinline__ BytePack<0> ld_volatile_shared<0>(uint32_t addr) { return {}; } -template<> __device__ __forceinline__ BytePack<0> ld_relaxed_gpu_global<0>(uintptr_t addr) { return {}; } -template<> __device__ __forceinline__ void st_global<0>(uintptr_t addr, BytePack<0> value) {} -template<> __device__ __forceinline__ void st_shared<0>(uint32_t addr, BytePack<0> value) {} -template<> __device__ __forceinline__ void st_relaxed_gpu_global<0>(uintptr_t addr, BytePack<0> value) {} -#endif +// template<> __device__ __forceinline__ void st_shared<0>(uint32_t addr, BytePack<0> value) {} +// template<> __device__ __forceinline__ void st_relaxed_gpu_global<0>(uintptr_t addr, BytePack<0> value) {} // Used to define implementations for above prototypes. #define DEFINE_ld_st__size_space(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, space, addr_cxx_ty, addr_reg_ty) \ @@ -219,28 +203,26 @@ template<> __device__ __forceinline__ void st_relaxed_gpu_global<0>(uintptr_t ad *((data_cxx_ty *)addr) = tmp; \ } -#ifdef NEED_CHECKING -#if __CUDA_ARCH__ >= 700 - #define PTX_relaxed_gpu "relaxed.gpu" -#else - #define PTX_relaxed_gpu "volatile" -#endif +// #if __CUDA_ARCH__ >= 700 +// #define PTX_relaxed_gpu "relaxed.gpu" +// #else +// #define PTX_relaxed_gpu "volatile" +// #endif -#define DEFINE_ld_st_gpu_relaxed__size(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty) \ - template<> \ - __device__ __forceinline__ BytePack ld_relaxed_gpu_global(uintptr_t addr) { \ - data_cxx_ty tmp; \ - asm("ld." PTX_relaxed_gpu ".global." #data_ptx_ty " %0, [%1];" : "="#data_reg_ty(tmp) : "l"(addr)); \ - BytePack ans; \ - ans.native = tmp; \ - return ans; \ - } \ - template<> \ - __device__ __forceinline__ void st_relaxed_gpu_global(uintptr_t addr, BytePack value) { \ - data_cxx_ty tmp = value.native; \ - asm volatile("st." PTX_relaxed_gpu ".global." #data_ptx_ty " [%0], %1;" :: "l"(addr), #data_reg_ty(tmp) : "memory"); \ - } -#endif +// #define DEFINE_ld_st_gpu_relaxed__size(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty) \ +// template<> \ +// __device__ __forceinline__ BytePack ld_relaxed_gpu_global(uintptr_t addr) { \ +// data_cxx_ty tmp; \ +// asm("ld." PTX_relaxed_gpu ".global." #data_ptx_ty " %0, [%1];" : "="#data_reg_ty(tmp) : "l"(addr)); \ +// BytePack ans; \ +// ans.native = tmp; \ +// return ans; \ +// } \ +// template<> \ +// __device__ __forceinline__ void st_relaxed_gpu_global(uintptr_t addr, BytePack value) { \ +// data_cxx_ty tmp = value.native; \ +// asm volatile("st." PTX_relaxed_gpu ".global." #data_ptx_ty " [%0], %1;" :: "l"(addr), #data_reg_ty(tmp) : "memory"); \ +// } #define DEFINE_ld_st__size(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty) \ DEFINE_ld_st__size_space(bytes, data_cxx_ty, data_ptx_ty, data_reg_ty, global, uintptr_t, l) @@ -282,20 +264,18 @@ DEFINE_ld_st_16__space(global, uintptr_t, l) // DEFINE_ld_st_16__space(shared, uint32_t, r) #undef DEFINE_ld_st_16 -#ifdef NEED_CHECKING -template<> -__device__ __forceinline__ BytePack<16> ld_relaxed_gpu_global<16>(uintptr_t addr) { - BytePack<16> ans; - asm("ld." PTX_relaxed_gpu ".global.v2.b64 {%0,%1}, [%2];" : "=l"(ans.u64[0]), "=l"(ans.u64[1]) : "l"(addr)); - return ans; -} -template<> -__device__ __forceinline__ void st_relaxed_gpu_global<16>(uintptr_t addr, BytePack<16> value) { - asm volatile("st." PTX_relaxed_gpu ".global.v2.b64 [%0], {%1,%2};" :: "l"(addr), "l"(value.u64[0]), "l"(value.u64[1]) : "memory"); -} - -#undef PTX_relaxed_gpu -#endif +// template<> +// __device__ __forceinline__ BytePack<16> ld_relaxed_gpu_global<16>(uintptr_t addr) { +// BytePack<16> ans; +// asm("ld." PTX_relaxed_gpu ".global.v2.b64 {%0,%1}, [%2];" : "=l"(ans.u64[0]), "=l"(ans.u64[1]) : "l"(addr)); +// return ans; +// } +// template<> +// __device__ __forceinline__ void st_relaxed_gpu_global<16>(uintptr_t addr, BytePack<16> value) { +// asm volatile("st." PTX_relaxed_gpu ".global.v2.b64 [%0], {%1,%2};" :: "l"(addr), "l"(value.u64[0]), "l"(value.u64[1]) : "memory"); +// } +// +// #undef PTX_relaxed_gpu //////////////////////////////////////////////////////////////////////////////// // Atomic load/store using c++ pointers. @@ -310,17 +290,17 @@ __device__ __forceinline__ uint64_t ld_relaxed_sys_global(uint64_t *ptr) { ans = __builtin_nontemporal_load(ptr); return ans; } -#ifdef NEED_CHECKING -__device__ __forceinline__ uint64_t ld_relaxed_gpu_global(uint64_t *ptr) { - uint64_t ans; - #if __CUDA_ARCH__ >= 700 - asm("ld.relaxed.gpu.global.u64 %0, [%1];" : "=l"(ans) : "l"(cvta_to_global(ptr))); - #else - asm("ld.volatile.global.u64 %0, [%1];" : "=l"(ans) : "l"(cvta_to_global(ptr))); - #endif - return ans; -} -#endif + +// __device__ __forceinline__ uint64_t ld_relaxed_gpu_global(uint64_t *ptr) { +// uint64_t ans; +// #if __CUDA_ARCH__ >= 700 +// asm("ld.relaxed.gpu.global.u64 %0, [%1];" : "=l"(ans) : "l"(cvta_to_global(ptr))); +// #else +// asm("ld.volatile.global.u64 %0, [%1];" : "=l"(ans) : "l"(cvta_to_global(ptr))); +// #endif +// return ans; +// } + __device__ __forceinline__ uint64_t ld_acquire_sys_global(uint64_t *ptr) { uint64_t ans; ans = __atomic_load_n(ptr ,__ATOMIC_SEQ_CST); diff --git a/projects/rccl/src/device/reduce_kernel.h b/projects/rccl/src/device/reduce_kernel.h index ca372a26ba..d3fd848371 100644 --- a/projects/rccl/src/device/reduce_kernel.h +++ b/projects/rccl/src/device/reduce_kernel.h @@ -207,22 +207,20 @@ struct Apply_Reduce, /*EltPerPack=*/4> { } }; -#ifdef NEED_CHECKING -template<> -struct Apply_Reduce, /*EltPerPack=*/4> { - __device__ static BytePack<4> reduce(FuncProd fn, BytePack<4> apack, BytePack<4> bpack) { - uint32_t a = apack.native; - uint32_t b = bpack.native; - uint32_t ab0 = (a*b) & 0xffu; - asm("mad.lo.u32 %0, %1, %2, %0;" : "+r"(ab0) : "r"(a&0xff00u), "r"(b&0xff00u)); - uint32_t ab1; - asm("mul.hi.u32 %0, %1, %2;" : "=r"(ab1) : "r"(a&0xff0000), "r"(b&0xff0000)); - asm("mad.hi.u32 %0, %1, %2, %0;" : "+r"(ab1) : "r"(a&0xff000000u), "r"(b&0xff000000u)); - apack.native = __byte_perm(ab0, ab1, 0x6420); - return apack; - } -}; -#endif +// template<> +// struct Apply_Reduce, /*EltPerPack=*/4> { +// __device__ static BytePack<4> reduce(FuncProd fn, BytePack<4> apack, BytePack<4> bpack) { +// uint32_t a = apack.native; +// uint32_t b = bpack.native; +// uint32_t ab0 = (a*b) & 0xffu; +// asm("mad.lo.u32 %0, %1, %2, %0;" : "+r"(ab0) : "r"(a&0xff00u), "r"(b&0xff00u)); +// uint32_t ab1; +// asm("mul.hi.u32 %0, %1, %2;" : "=r"(ab1) : "r"(a&0xff0000), "r"(b&0xff0000)); +// asm("mad.hi.u32 %0, %1, %2, %0;" : "+r"(ab1) : "r"(a&0xff000000u), "r"(b&0xff000000u)); +// apack.native = __byte_perm(ab0, ab1, 0x6420); +// return apack; +// } +// }; #define SPECIALIZE_REDUCE(Fn, T, EltPerPack, Vec, expr_of_fn_x_y) \ template<> \