AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

Update IPC object

Przeglądaj źródła 
* Update the IPC object in the context class with the instance created in the IPC backend


[ROCm/rocshmem commit: 45a8cb3354]
This commit is contained in:
avinashkethineedi
2024-08-28 08:14:38 -07:00
rodzic 1432a6be71
commit 1335d98f1f
2 zmienionych plików z 20 dodań i 22 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
projects/rocshmem/src/ipc/context_ipc_device.cpp
+20 -19
Wyświetl plik
@@ -39,7 +39,8 @@ namespace rocshmem {
__host__ IPCContext::IPCContext(Backend *b)
: Context(b, false) {
IPCBackend *backend{static_cast<IPCBackend *>(b)};
ipcImpl = &backend->ipcImpl;
ipcImpl_.ipc_bases = b->ipcImpl.ipc_bases;
ipcImpl_.shm_size = b->ipcImpl.shm_size;
auto *bp{backend->ipc_backend_proxy.get()};
@@ -59,21 +60,21 @@ __device__ void IPCContext::ctx_destroy(){
__device__ void IPCContext::putmem(void *dest, const void *source, size_t nelems,
int pe) {
// TODO (Avinash) check if PE is available for IPC using (isIpcAvailable)
int local_pe = pe % ipcImpl->shm_size;
int local_pe = pe % ipcImpl_.shm_size;
uint64_t L_offset =
reinterpret_cast<char *>(dest) - ipcImpl->ipc_bases[my_pe];
ipcImpl->ipcCopy(ipcImpl->ipc_bases[local_pe] + L_offset,
reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[my_pe];
ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[local_pe] + L_offset,
const_cast<void *>(source), nelems);
}
__device__ void IPCContext::getmem(void *dest, const void *source, size_t nelems,
int pe) {
// TODO (Avinash) check if PE is available for IPC using (isIpcAvailable)
int local_pe = pe % ipcImpl->shm_size;
int local_pe = pe % ipcImpl_.shm_size;
const char *src_typed = reinterpret_cast<const char *>(source);
uint64_t L_offset =
const_cast<char *>(src_typed) - ipcImpl->ipc_bases[my_pe];
ipcImpl->ipcCopy(dest, ipcImpl->ipc_bases[local_pe] + L_offset, nelems);
const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[my_pe];
ipcImpl_.ipcCopy(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems);
}
__device__ void IPCContext::putmem_nbi(void *dest, const void *source,
@@ -115,10 +116,10 @@ __device__ void IPCContext::sync(roc_shmem_team_t team) {
__device__ void IPCContext::putmem_wg(void *dest, const void *source,
size_t nelems, int pe) {
// TODO (Avinash) check if PE is available for IPC using (isIpcAvailable)
int local_pe = pe % ipcImpl->shm_size;
int local_pe = pe % ipcImpl_.shm_size;
uint64_t L_offset =
reinterpret_cast<char *>(dest) - ipcImpl->ipc_bases[my_pe];
ipcImpl->ipcCopy_wg(ipcImpl->ipc_bases[local_pe] + L_offset,
reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[my_pe];
ipcImpl_.ipcCopy_wg(ipcImpl_.ipc_bases[local_pe] + L_offset,
const_cast<void *>(source), nelems);
__syncthreads();
}
@@ -126,11 +127,11 @@ __device__ void IPCContext::putmem_wg(void *dest, const void *source,
__device__ void IPCContext::getmem_wg(void *dest, const void *source,
size_t nelems, int pe) {
// TODO (Avinash) check if PE is available for IPC using (isIpcAvailable)
int local_pe = pe % ipcImpl->shm_size;
int local_pe = pe % ipcImpl_.shm_size;
const char *src_typed = reinterpret_cast<const char *>(source);
uint64_t L_offset =
const_cast<char *>(src_typed) - ipcImpl->ipc_bases[my_pe];
ipcImpl->ipcCopy_wg(dest, ipcImpl->ipc_bases[local_pe] + L_offset, nelems);
const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[my_pe];
ipcImpl_.ipcCopy_wg(dest, ipcImpl_.ipc_bases[local_pe] + L_offset, nelems);
__syncthreads();
}
@@ -147,21 +148,21 @@ __device__ void IPCContext::getmem_nbi_wg(void *dest, const void *source,
__device__ void IPCContext::putmem_wave(void *dest, const void *source,
size_t nelems, int pe) {
// TODO (Avinash) check if PE is available for IPC using (isIpcAvailable)
int local_pe = pe % ipcImpl->shm_size;
int local_pe = pe % ipcImpl_.shm_size;
uint64_t L_offset =
reinterpret_cast<char *>(dest) - ipcImpl->ipc_bases[my_pe];
ipcImpl->ipcCopy_wave(ipcImpl->ipc_bases[local_pe] + L_offset,
reinterpret_cast<char *>(dest) - ipcImpl_.ipc_bases[my_pe];
ipcImpl_.ipcCopy_wave(ipcImpl_.ipc_bases[local_pe] + L_offset,
const_cast<void *>(source), nelems);
}
__device__ void IPCContext::getmem_wave(void *dest, const void *source,
size_t nelems, int pe) {
// TODO (Avinash) check if PE is available for IPC using (isIpcAvailable)
int local_pe = pe % ipcImpl->shm_size;
int local_pe = pe % ipcImpl_.shm_size;
const char *src_typed = reinterpret_cast<const char *>(source);
uint64_t L_offset =
const_cast<char *>(src_typed) - ipcImpl->ipc_bases[my_pe];
ipcImpl->ipcCopy_wave(dest, ipcImpl->ipc_bases[local_pe] + L_offset,
const_cast<char *>(src_typed) - ipcImpl_.ipc_bases[my_pe];
ipcImpl_.ipcCopy_wave(dest, ipcImpl_.ipc_bases[local_pe] + L_offset,
nelems);
}
projects/rocshmem/src/ipc/context_ipc_device.hpp
-3
Wyświetl plik
@@ -234,9 +234,6 @@ class IPCContext : public Context {
private:
//context class has IpcImpl object (ipcImpl_)
IpcImpl *ipcImpl{nullptr};
uint64_t* atomic_base_ptr{nullptr};
char* g_ret;