AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

Limit max channels for ring graph on single node Rome (#347)

Telusuri Sumber 
* Limit max channels for ring graph on single node Rome
* Partially revert "Use non-temporal access for streaming data (#341)"
This commit is contained in:
Wenkai Du
2021-04-14 10:14:54 -07:00
committed by GitHub
orang tua 1fe031402a
melakukan a79f74082e
4 mengubah file dengan 13 tambahan dan 10 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
src/collectives/device/common_kernel.h
+4 -5
Melihat File
@@ -267,17 +267,16 @@ struct MULTI128 {
inline __device__ void Fetch128(Pack128& v, const Pack128* p) {
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
v.x = __builtin_nontemporal_load(&p->x);
v.y = __builtin_nontemporal_load(&p->y);
v.x = p->x;
v.y = p->y;
#else
asm volatile("ld.volatile.global.v2.u64 {%0,%1}, [%2];" : "=l"(v.x), "=l"(v.y) : "l"(p) : "memory");
#endif
}
inline __device__ void Store128(Pack128* p, Pack128& v) {
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
__builtin_nontemporal_store(v.x, &p->x);
__builtin_nontemporal_store(v.y, &p->y);
p->x = v.x;
p->y = v.y;
#else
asm volatile("st.volatile.global.v2.u64 [%0], {%1,%2};" :: "l"(p), "l"(v.x), "l"(v.y) : "memory");
#endif