AI/rocm-systems
2
0
포크 0
코드 이슈 풀 리퀘스트 액션 패키지 프로젝트 릴리즈 위키 활동

Disable unsupported ld/st instructions

소스 검색 
[ROCm/rccl commit: 6f3310605c]
이 커밋은 다음에 포함됨:
BertanDogancay
2024-02-15 13:58:16 -08:00
부모 7842411fb3
커밋 ef72944015
3개의 변경된 파일과 68개의 추가작업 그리고 92개의 파일을 삭제
패치 파일 다운로드 Diff 파일 다운로드 모든 파일 펼침 모든 파일 접기
projects/rccl/src/device/network/unpack/unpack.h
+2 -4
파일 보기
@@ -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
}
projects/rccl/src/device/op128.h
+52 -72
파일 보기
@@ -162,38 +162,22 @@ __device__ __forceinline__ T fromPack(typename BytePackOf<T>::Pack pack) {
// Load/store of BytePack<?> using integral addresses.
template<int Size> __device__ BytePack<Size> ld_global(uintptr_t addr);
// template<int Size> __device__ BytePack<Size> ld_shared(uint32_t addr);
template<int Size> __device__ BytePack<Size> ld_volatile_global(uintptr_t addr);
//template<int Size> __device__ BytePack<Size> ld_shared(uint32_t addr);
//template<int Size> __device__ BytePack<Size> ld_volatile_shared(uint32_t addr);
// template<int Size> __device__ BytePack<Size> ld_volatile_shared(uint32_t addr);
// template<int Size> __device__ BytePack<Size> ld_relaxed_gpu_global(uintptr_t addr);
template<int Size> __device__ void st_global(uintptr_t addr, BytePack<Size> value);
//template<int Size> __device__ void st_shared(uint32_t addr, BytePack<Size> value);
// template<int Size> __device__ void st_shared(uint32_t addr, BytePack<Size> value);
// template<int Size> __device__ void st_relaxed_gpu_global(uintptr_t addr, BytePack<Size> 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<int Size> __device__ BytePack<Size> ld_global(uintptr_t addr);
template<int Size> __device__ BytePack<Size> ld_shared(uint32_t addr);
template<int Size> __device__ BytePack<Size> ld_volatile_global(uintptr_t addr);
template<int Size> __device__ BytePack<Size> ld_volatile_shared(uint32_t addr);
template<int Size> __device__ BytePack<Size> ld_relaxed_gpu_global(uintptr_t addr);
template<int Size> __device__ void st_global(uintptr_t addr, BytePack<Size> value);
template<int Size> __device__ void st_shared(uint32_t addr, BytePack<Size> value);
template<int Size> __device__ void st_relaxed_gpu_global(uintptr_t addr, BytePack<Size> 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<bytes> ld_relaxed_gpu_global<bytes>(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<bytes> ans; \
ans.native = tmp; \
return ans; \
} \
template<> \
__device__ __forceinline__ void st_relaxed_gpu_global<bytes>(uintptr_t addr, BytePack<bytes> 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<bytes> ld_relaxed_gpu_global<bytes>(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<bytes> ans; \
// ans.native = tmp; \
// return ans; \
// } \
// template<> \
// __device__ __forceinline__ void st_relaxed_gpu_global<bytes>(uintptr_t addr, BytePack<bytes> 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);
projects/rccl/src/device/reduce_kernel.h
+14 -16
파일 보기
@@ -207,22 +207,20 @@ struct Apply_Reduce<FuncMinMax<uint8_t>, /*EltPerPack=*/4> {
}
};
#ifdef NEED_CHECKING
template<>
struct Apply_Reduce<FuncProd<uint8_t>, /*EltPerPack=*/4> {
__device__ static BytePack<4> reduce(FuncProd<uint8_t> 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<FuncProd<uint8_t>, /*EltPerPack=*/4> {
// __device__ static BytePack<4> reduce(FuncProd<uint8_t> 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<> \