AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

use correct id when accessing ipc-bases (#88)

Просмотр исходного кода 
we need to use the position of that processes in the local ipc-bases
array, not the global rank.
Этот коммит содержится в:
Edgar Gabriel
2025-04-17 10:11:32 -05:00
коммит произвёл GitHub
родитель 9befbe8293
Коммит e3b0353fa9
3 изменённых файлов: 57 добавлений и 48 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
src/ipc_policy.hpp
+8 -3
Просмотреть файл
@@ -57,11 +57,14 @@ class IpcOnImpl {
__host__ void ipcHostStop();
__device__ bool isIpcAvailable(int my_pe, int target_pe) {
__device__ bool isIpcAvailable(int my_pe, int target_pe, int *local_target_pe) {
if (nullptr == pes_with_ipc_avail) { return false; }
for (int i=0; i<shm_size; i++) {
if (pes_with_ipc_avail[i] == target_pe) { return true; }
if (pes_with_ipc_avail[i] == target_pe) {
*local_target_pe = i;
return true;
}
}
return false;
@@ -123,6 +126,8 @@ class IpcOffImpl {
using HEAP_BASES_T = std::vector<char *, StdAllocatorHIP<char *>>;
public:
int shm_rank{0};
uint32_t shm_size{0};
char **ipc_bases{nullptr};
@@ -134,7 +139,7 @@ class IpcOffImpl {
__host__ void ipcHostStop() {}
__device__ bool isIpcAvailable(int my_pe, int target_pe) { return false; }
__device__ bool isIpcAvailable(int my_pe, int target_pe, int *local_target_pe) { return false; }
__device__ void ipcGpuInit(Backend *rocshmem_handle, Context *ctx,
int thread_id) {}
src/reverse_offload/context_ro_device.cpp
+41 -39
Просмотреть файл
@@ -60,15 +60,16 @@ __host__ ROContext::ROContext(Backend *b,
ipcImpl_.ipc_bases = b->ipcImpl.ipc_bases;
ipcImpl_.shm_size = b->ipcImpl.shm_size;
ipcImpl_.shm_rank = b->ipcImpl.shm_rank;
ipcImpl_.pes_with_ipc_avail = b->ipcImpl.pes_with_ipc_avail;
}
__device__ void ROContext::putmem(void *dest, const void *source, size_t nelems,
int pe) {
if (ipcImpl_.isIpcAvailable(my_pe, pe)) {
int local_pe = pe % ipcImpl_.shm_size;
int local_pe{-1};
if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) {
uint64_t L_offset =
reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[my_pe];
reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank];
ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[local_pe] + L_offset,
const_cast<void *>(source), nelems);
} else {
@@ -85,11 +86,11 @@ __device__ void ROContext::putmem(void *dest, const void *source, size_t nelems,
__device__ void ROContext::getmem(void *dest, const void *source, size_t nelems,
int pe) {
if (ipcImpl_.isIpcAvailable(my_pe, pe)) {
int local_pe = pe % ipcImpl_.shm_size;
int local_pe{-1};
if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) {
const char *src_typed = reinterpret_cast<const char *>(source);
uint64_t L_offset =
const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[my_pe];
const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank];
ipcImpl_.ipcCopy(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems);
} else {
bool must_send_message = wf_coal_.coalesce(pe, source, dest, &nelems);
@@ -105,10 +106,10 @@ __device__ void ROContext::getmem(void *dest, const void *source, size_t nelems,
__device__ void ROContext::putmem_nbi(void *dest, const void *source,
size_t nelems, int pe) {
if (ipcImpl_.isIpcAvailable(my_pe, pe)) {
int local_pe = pe % ipcImpl_.shm_size;
int local_pe{-1};
if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) {
uint64_t L_offset =
reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[my_pe];
reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank];
ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[local_pe] + L_offset,
const_cast<void *>(source), nelems);
} else {
@@ -124,11 +125,11 @@ __device__ void ROContext::putmem_nbi(void *dest, const void *source,
__device__ void ROContext::getmem_nbi(void *dest, const void *source,
size_t nelems, int pe) {
if (ipcImpl_.isIpcAvailable(my_pe, pe)) {
int local_pe = pe % ipcImpl_.shm_size;
int local_pe{-1};
if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) {
const char *src_typed = reinterpret_cast<const char *>(source);
uint64_t L_offset =
const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[my_pe];
const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank];
ipcImpl_.ipcCopy(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems);
} else {
bool must_send_message = wf_coal_.coalesce(pe, source, dest, &nelems);
@@ -162,11 +163,12 @@ __device__ void ROContext::quiet() {
__device__ void *ROContext::shmem_ptr(const void *dest, int pe) {
void *ret = nullptr;
if (ipcImpl_.isIpcAvailable(my_pe, pe)) {
int local_pe{-1};
if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) {
void *dst = const_cast<void *>(dest);
uint64_t L_offset =
reinterpret_cast<char *>(dst) - ipcImpl_.ipc_bases[my_pe];
ret = ipcImpl_.ipc_bases[pe] + L_offset;
reinterpret_cast<char *>(dst) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank];
ret = ipcImpl_.ipc_bases[local_pe] + L_offset;
}
return ret;
}
@@ -291,10 +293,10 @@ __device__ void ROContext::ctx_destroy() {
__device__ void ROContext::putmem_wg(void *dest, const void *source,
size_t nelems, int pe) {
if (ipcImpl_.isIpcAvailable(my_pe, pe)) {
int local_pe = pe % ipcImpl_.shm_size;
int local_pe{-1};
if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) {
uint64_t L_offset =
reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[my_pe];
reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank];
ipcImpl_.ipcCopy_wg(ipcImpl_.ipc_bases[local_pe] + L_offset,
const_cast<void *>(source), nelems);
} else {
@@ -310,11 +312,11 @@ __device__ void ROContext::putmem_wg(void *dest, const void *source,
__device__ void ROContext::getmem_wg(void *dest, const void *source,
size_t nelems, int pe) {
if (ipcImpl_.isIpcAvailable(my_pe, pe)) {
int local_pe = pe % ipcImpl_.shm_size;
int local_pe{-1};
if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) {
const char *src_typed = reinterpret_cast<const char *>(source);
uint64_t L_offset =
const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[my_pe];
const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank];
ipcImpl_.ipcCopy_wg(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems);
} else {
if (is_thread_zero_in_block()) {
@@ -329,10 +331,10 @@ __device__ void ROContext::getmem_wg(void *dest, const void *source,
__device__ void ROContext::putmem_nbi_wg(void *dest, const void *source,
size_t nelems, int pe) {
if (ipcImpl_.isIpcAvailable(my_pe, pe)) {
int local_pe = pe % ipcImpl_.shm_size;
int local_pe{-1};
if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) {
uint64_t L_offset =
reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[my_pe];
reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank];
ipcImpl_.ipcCopy_wg(ipcImpl_.ipc_bases[local_pe] + L_offset,
const_cast<void *>(source), nelems);
} else {
@@ -347,11 +349,11 @@ __device__ void ROContext::putmem_nbi_wg(void *dest, const void *source,
__device__ void ROContext::getmem_nbi_wg(void *dest, const void *source,
size_t nelems, int pe) {
if (ipcImpl_.isIpcAvailable(my_pe, pe)) {
int local_pe = pe % ipcImpl_.shm_size;
int local_pe{-1};
if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) {
const char *src_typed = reinterpret_cast<const char *>(source);
uint64_t L_offset =
const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[my_pe];
const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank];
ipcImpl_.ipcCopy_wg(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems);
} else {
if (is_thread_zero_in_block()) {
@@ -365,10 +367,10 @@ __device__ void ROContext::getmem_nbi_wg(void *dest, const void *source,
__device__ void ROContext::putmem_wave(void *dest, const void *source,
size_t nelems, int pe) {
if (ipcImpl_.isIpcAvailable(my_pe, pe)) {
int local_pe = pe % ipcImpl_.shm_size;
int local_pe{-1};
if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) {
uint64_t L_offset =
reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[my_pe];
reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank];
ipcImpl_.ipcCopy_wave(ipcImpl_.ipc_bases[local_pe] + L_offset,
const_cast<void *>(source), nelems);
} else {
@@ -383,11 +385,11 @@ __device__ void ROContext::putmem_wave(void *dest, const void *source,
__device__ void ROContext::getmem_wave(void *dest, const void *source,
size_t nelems, int pe) {
if (ipcImpl_.isIpcAvailable(my_pe, pe)) {
int local_pe = pe % ipcImpl_.shm_size;
int local_pe{-1};
if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) {
const char *src_typed = reinterpret_cast<const char *>(source);
uint64_t L_offset =
const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[my_pe];
const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank];
ipcImpl_.ipcCopy_wave(dest, ipcImpl_.ipc_bases[local_pe] + L_offset,
nelems);
} else {
@@ -402,10 +404,10 @@ __device__ void ROContext::getmem_wave(void *dest, const void *source,
__device__ void ROContext::putmem_nbi_wave(void *dest, const void *source,
size_t nelems, int pe) {
if (ipcImpl_.isIpcAvailable(my_pe, pe)) {
int local_pe = pe % ipcImpl_.shm_size;
int local_pe{-1};
if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) {
uint64_t L_offset =
reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[my_pe];
reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank];
ipcImpl_.ipcCopy_wave(ipcImpl_.ipc_bases[local_pe] + L_offset,
const_cast<void *>(source), nelems);
} else {
@@ -419,11 +421,11 @@ __device__ void ROContext::putmem_nbi_wave(void *dest, const void *source,
__device__ void ROContext::getmem_nbi_wave(void *dest, const void *source,
size_t nelems, int pe) {
if (ipcImpl_.isIpcAvailable(my_pe, pe)) {
int local_pe = pe % ipcImpl_.shm_size;
int local_pe{-1};
if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) {
const char *src_typed = reinterpret_cast<const char *>(source);
uint64_t L_offset =
const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[my_pe];
const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank];
ipcImpl_.ipcCopy_wave(dest, ipcImpl_.ipc_bases[local_pe] + L_offset,
nelems);
} else {
src/reverse_offload/context_ro_tmpl_device.hpp
+8 -6
Просмотреть файл
@@ -145,9 +145,10 @@ __device__ void ROContext::put_nbi(T *dest, const T *source, size_t nelems,
template <typename T>
__device__ void ROContext::p(T *dest, T value, int pe) {
if (ipcImpl_.isIpcAvailable(my_pe, pe)) {
long L_offset{reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[my_pe]};
ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[pe] + L_offset,
int local_pe{-1};
if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) {
long L_offset{reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]};
ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[local_pe] + L_offset,
reinterpret_cast<void *>(&value), sizeof(T));
} else {
build_queue_element(RO_NET_P, dest, &value, sizeof(T), pe, 0, 0, 0, nullptr,
@@ -158,11 +159,12 @@ __device__ void ROContext::p(T *dest, T value, int pe) {
template <typename T>
__device__ T ROContext::g(const T *source, int pe) {
if (ipcImpl_.isIpcAvailable(my_pe, pe)) {
int local_pe{-1};
if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) {
const char *src_typed{reinterpret_cast<const char *>(source)};
long L_offset{const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[my_pe]};
long L_offset{const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank]};
T dest;
ipcImpl_.ipcCopy(&dest, ipcImpl_.ipc_bases[pe] + L_offset, sizeof(T));
ipcImpl_.ipcCopy(&dest, ipcImpl_.ipc_bases[local_pe] + L_offset, sizeof(T));
return dest;
} else {
auto dest{get_g_ret_buf()};