AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge branch PR #55 into naming_scheme

Browse Source
This commit is contained in:
avinashkethineedi
2024-12-04 21:46:38 +00:00
parent 46dcfbbb9e 7e4aa5d601
commit d8ce066adc
180 changed files with 4436 additions and 4437 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/CMakeLists.txt
+2 -2
View File
@@ -31,8 +31,8 @@ target_sources(
context_host.cpp
context_device.cpp
mpi_init_singleton.cpp
roc_shmem_gpu.cpp
roc_shmem.cpp
rocshmem_gpu.cpp
rocshmem.cpp
team.cpp
team_tracker.cpp
util.cpp
src/backend_bc.cpp
+2 -2
View File
@@ -202,7 +202,7 @@ void Backend::reset_stats() {
reset_backend_stats();
}
__device__ bool Backend::create_ctx(int64_t option, roc_shmem_ctx_t* ctx) {
__device__ bool Backend::create_ctx(int64_t option, rocshmem_ctx_t* ctx) {
#ifdef USE_GPU_IB
return static_cast<GPUIBBackend*>(this)->create_ctx(option, ctx);
#elif defined(USE_RO)
@@ -212,7 +212,7 @@ __device__ bool Backend::create_ctx(int64_t option, roc_shmem_ctx_t* ctx) {
#endif
}
__device__ void Backend::destroy_ctx(roc_shmem_ctx_t* ctx) {
__device__ void Backend::destroy_ctx(rocshmem_ctx_t* ctx) {
#ifdef USE_GPU_IB
static_cast<GPUIBBackend*>(this)->destroy_ctx(ctx);
#elif defined(USE_RO)
src/backend_bc.hpp
+6 -6
View File
@@ -36,7 +36,7 @@
#include <vector>
#include "config.h" // NOLINT(build/include_subdir)
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#include "backend_type.hpp"
#include "ipc_policy.hpp"
#include "memory/symmetric_heap.hpp"
@@ -56,7 +56,7 @@ class TeamInfo;
* It uses this state to populate Context objects which the GPU may use to
* perform networking operations.
*
* The roc_shmem.cpp implementation file wraps many the Backend public
* The rocshmem.cpp implementation file wraps many the Backend public
* members to implement the library's public API.
*/
class Backend {
@@ -74,8 +74,8 @@ class Backend {
*/
virtual ~Backend();
__device__ bool create_ctx(int64_t option, roc_shmem_ctx_t* ctx);
__device__ void destroy_ctx(roc_shmem_ctx_t* ctx);
__device__ bool create_ctx(int64_t option, rocshmem_ctx_t* ctx);
__device__ void destroy_ctx(rocshmem_ctx_t* ctx);
/**
* @brief Create a new team object and initialize it.
@@ -93,14 +93,14 @@ class Backend {
TeamInfo* team_info_wrt_parent,
TeamInfo* team_info_wrt_world, int num_pes,
int my_pe_in_new_team, MPI_Comm team_comm,
roc_shmem_team_t* new_team) = 0;
rocshmem_team_t* new_team) = 0;
/**
* @brief Destruct a team
*
* @param[in] team Handle to the team to destroy.
*/
virtual void team_destroy(roc_shmem_team_t team) = 0;
virtual void team_destroy(rocshmem_team_t team) = 0;
/**
* @brief Reports processing element number id.
src/constants.hpp
+1 -1
View File
@@ -26,7 +26,7 @@
/**
* @file constants.hpp
*
* @brief Contains global constants for ROCSHMEM library
* @brief Contains global constants for rocSHMEM library
*/
namespace rocshmem {
src/containers/helper_macros.hpp
+4 -4
View File
@@ -30,11 +30,11 @@
#include <cstdio>
#include <iostream>
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#define BARRIER() rocshmem::roc_shmem_wg_barrier_all()
#define RANK rocshmem::roc_shmem_my_pe()
#define NPES rocshmem::roc_shmem_n_pes()
#define BARRIER() rocshmem::rocshmem_wg_barrier_all()
#define RANK rocshmem::rocshmem_my_pe()
#define NPES rocshmem::rocshmem_n_pes()
#define PE_BITS ((uint64_t)ceil(log(NPES) / log(2)))
#define PE_OF(X) ((X) >> (64 - PE_BITS))
src/context.hpp
+12 -12
View File
@@ -67,7 +67,7 @@ class Context {
* just removing the dispatch implementations.
*
* No comments for these guys since its basically the same as in the
* roc_shmem.hpp public header.
* rocshmem.hpp public header.
*/
/**************************************************************************
@@ -139,7 +139,7 @@ class Context {
__device__ void sync_all();
__device__ void sync(roc_shmem_team_t team);
__device__ void sync(rocshmem_team_t team);
template <typename T>
__device__ T amo_fetch(void* dst, T value, T cond, int pe, uint8_t atomic_op);
@@ -186,13 +186,13 @@ class Context {
template <typename T>
__device__ T g(T* source, int pe);
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__device__ void to_all(T* dest, const T* source, int nreduce, int PE_start,
int logPE_stride, int PE_size, T* pWrk,
long* pSync); // NOLINT(runtime/int)
template <typename T, ROC_SHMEM_OP Op>
__device__ int reduce(roc_shmem_team_t team, T* dest, const T* source, int nreduce);
template <typename T, ROCSHMEM_OP Op>
__device__ int reduce(rocshmem_team_t team, T* dest, const T* source, int nreduce);
template <typename T>
__device__ void put(T* dest, const T* source, size_t nelems, int pe);
@@ -207,15 +207,15 @@ class Context {
__device__ void get_nbi(T* dest, const T* source, size_t nelems, int pe);
template <typename T>
__device__ void alltoall(roc_shmem_team_t team, T* dest, const T* source,
__device__ void alltoall(rocshmem_team_t team, T* dest, const T* source,
int nelems);
template <typename T>
__device__ void fcollect(roc_shmem_team_t team, T* dest, const T* source,
__device__ void fcollect(rocshmem_team_t team, T* dest, const T* source,
int nelems);
template <typename T>
__device__ void broadcast(roc_shmem_team_t team, T* dest, const T* source,
__device__ void broadcast(rocshmem_team_t team, T* dest, const T* source,
int nelems, int pe_root);
template <typename T>
@@ -378,16 +378,16 @@ class Context {
long* p_sync); // NOLINT(runtime/int)
template <typename T>
__host__ void broadcast(roc_shmem_team_t team, T* dest, const T* source,
__host__ void broadcast(rocshmem_team_t team, T* dest, const T* source,
int nelems, int pe_root);
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__host__ void to_all(T* dest, const T* source, int nreduce, int PE_start,
int logPE_stride, int PE_size, T* pWrk,
long* pSync); // NOLINT(runtime/int)
template <typename T, ROC_SHMEM_OP Op>
__host__ int reduce(roc_shmem_team_t team, T* dest, const T* source, int nreduce);
template <typename T, ROCSHMEM_OP Op>
__host__ int reduce(rocshmem_team_t team, T* dest, const T* source, int nreduce);
template <typename T>
__host__ void wait_until(T *ivars, int cmp, T val);
src/context_device.cpp
+1 -1
View File
@@ -154,7 +154,7 @@ __device__ void Context::sync_all() {
DISPATCH(sync_all());
}
__device__ void Context::sync(roc_shmem_team_t team) {
__device__ void Context::sync(rocshmem_team_t team) {
ctxStats.incStat(NUM_SYNC_ALL);
DISPATCH(sync(team));
src/context_tmpl_device.hpp
+13 -13
View File
@@ -62,7 +62,7 @@ __device__ T Context::g(T *source, int pe) {
}
// The only way to get multi-arg templates to feed into a macro
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__device__ void Context::to_all(T *dest, const T *source, int nreduce,
int PE_start, int logPE_stride, int PE_size,
T *pWrk,
@@ -79,11 +79,11 @@ __device__ void Context::to_all(T *dest, const T *source, int nreduce,
PE_size, pWrk, pSync));
}
template <typename T, ROC_SHMEM_OP Op>
__device__ int Context::reduce(roc_shmem_team_t team, T *dest, const T *source,
template <typename T, ROCSHMEM_OP Op>
__device__ int Context::reduce(rocshmem_team_t team, T *dest, const T *source,
int nreduce) {
if (nreduce == 0) {
return ROC_SHMEM_SUCCESS;
return ROCSHMEM_SUCCESS;
}
if (is_thread_zero_in_block()) {
@@ -140,7 +140,7 @@ __device__ void Context::get_nbi(T *dest, const T *source, size_t nelems,
}
template <typename T>
__device__ void Context::alltoall(roc_shmem_team_t team, T *dest,
__device__ void Context::alltoall(rocshmem_team_t team, T *dest,
const T *source, int nelems) {
if (nelems == 0) {
return;
@@ -154,7 +154,7 @@ __device__ void Context::alltoall(roc_shmem_team_t team, T *dest,
}
template <typename T>
__device__ void Context::fcollect(roc_shmem_team_t team, T *dest,
__device__ void Context::fcollect(rocshmem_team_t team, T *dest,
const T *source, int nelems) {
if (nelems == 0) {
return;
@@ -168,7 +168,7 @@ __device__ void Context::fcollect(roc_shmem_team_t team, T *dest,
}
template <typename T>
__device__ void Context::broadcast(roc_shmem_team_t team, T *dest,
__device__ void Context::broadcast(rocshmem_team_t team, T *dest,
const T *source, int nelems, int pe_root) {
if (nelems == 0) {
return;
@@ -340,32 +340,32 @@ __device__ __forceinline__ int Context::test(T *ivars, int cmp,
int ret = 0;
volatile T *vol_ivars = reinterpret_cast<T *>(ivars);
switch (cmp) {
case ROC_SHMEM_CMP_EQ:
case ROCSHMEM_CMP_EQ:
if (uncached_load(vol_ivars) == val) {
ret = 1;
}
break;
case ROC_SHMEM_CMP_NE:
case ROCSHMEM_CMP_NE:
if (uncached_load(vol_ivars) != val) {
ret = 1;
}
break;
case ROC_SHMEM_CMP_GT:
case ROCSHMEM_CMP_GT:
if (uncached_load(vol_ivars) > val) {
ret = 1;
}
break;
case ROC_SHMEM_CMP_GE:
case ROCSHMEM_CMP_GE:
if (uncached_load(vol_ivars) >= val) {
ret = 1;
}
break;
case ROC_SHMEM_CMP_LT:
case ROCSHMEM_CMP_LT:
if (uncached_load(vol_ivars) < val) {
ret = 1;
}
break;
case ROC_SHMEM_CMP_LE:
case ROCSHMEM_CMP_LE:
if (uncached_load(vol_ivars) <= val) {
ret = 1;
}
src/context_tmpl_host.hpp
+5 -5
View File
@@ -194,7 +194,7 @@ __host__ void Context::broadcast(T *dest, const T *source, int nelems,
}
template <typename T>
__host__ void Context::broadcast(roc_shmem_team_t team, T *dest,
__host__ void Context::broadcast(rocshmem_team_t team, T *dest,
const T *source, int nelems,
int pe_root) { // NOLINT(runtime/int)
if (nelems == 0) {
@@ -206,7 +206,7 @@ __host__ void Context::broadcast(roc_shmem_team_t team, T *dest,
HOST_DISPATCH(broadcast<T>(team, dest, source, nelems, pe_root));
}
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__host__ void Context::to_all(T *dest, const T *source, int nreduce,
int PE_start, int logPE_stride, int PE_size,
T *pWrk,
@@ -221,11 +221,11 @@ __host__ void Context::to_all(T *dest, const T *source, int nreduce,
logPE_stride, PE_size, pWrk, pSync));
}
template <typename T, ROC_SHMEM_OP Op>
__host__ int Context::reduce(roc_shmem_team_t team, T *dest, const T *source,
template <typename T, ROCSHMEM_OP Op>
__host__ int Context::reduce(rocshmem_team_t team, T *dest, const T *source,
int nreduce) { // NOLINT(runtime/int)
if (nreduce == 0) {
return ROC_SHMEM_SUCCESS;
return ROCSHMEM_SUCCESS;
}
ctxHostStats.incStat(NUM_HOST_TO_ALL);
src/fence_policy.hpp
+2 -2
View File
@@ -23,7 +23,7 @@
#ifndef LIBRARY_SRC_FENCE_POLICY_HPP_
#define LIBRARY_SRC_FENCE_POLICY_HPP_
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
namespace rocshmem {
@@ -43,7 +43,7 @@ class Fence {
* @param[in] options interpreted as a bitfield using bitwise operations
*/
__host__ __device__ Fence(long option) {
if (option & ROC_SHMEM_CTX_NOSTORE) {
if (option & ROCSHMEM_CTX_NOSTORE) {
flush_ = false;
}
}
src/gpu_ib/backend_ib.cpp
+51 -51
View File
@@ -30,7 +30,7 @@
#include <cstdlib>
#include <mutex> // NOLINT(build/c++11)
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#include "../backend_type.hpp"
#include "../context_incl.hpp"
#include "gpu_ib_team.hpp"
@@ -47,10 +47,10 @@ namespace rocshmem {
} \
}
extern roc_shmem_ctx_t ROC_SHMEM_HOST_CTX_DEFAULT;
extern rocshmem_ctx_t ROCSHMEM_HOST_CTX_DEFAULT;
roc_shmem_team_t get_external_team(GPUIBTeam *team) {
return reinterpret_cast<roc_shmem_team_t>(team);
rocshmem_team_t get_external_team(GPUIBTeam *team) {
return reinterpret_cast<rocshmem_team_t>(team);
}
int get_ls_non_zero_bit(char *bitmask, int mask_length) {
@@ -68,7 +68,7 @@ int get_ls_non_zero_bit(char *bitmask, int mask_length) {
}
GPUIBBackend::GPUIBBackend(MPI_Comm comm) : Backend() {
if (auto maximum_num_contexts_str = getenv("ROC_SHMEM_MAX_NUM_CONTEXTS")) {
if (auto maximum_num_contexts_str = getenv("ROCSHMEM_MAX_NUM_CONTEXTS")) {
std::stringstream sstream(maximum_num_contexts_str);
sstream >> maximum_num_contexts_;
}
@@ -96,7 +96,7 @@ GPUIBBackend::GPUIBBackend(MPI_Comm comm) : Backend() {
setup_team_world();
roc_shmem_collective_init();
rocshmem_collective_init();
teams_init();
@@ -109,7 +109,7 @@ GPUIBBackend::GPUIBBackend(MPI_Comm comm) : Backend() {
#ifdef USE_HOST_SIDE_HDP_FLUSH
hdp_gpu_cpu_flush_flag_ =
static_cast<unsigned int *>(roc_shmem_malloc(sizeof(unsigned int)));
static_cast<unsigned int *>(rocshmem_malloc(sizeof(unsigned int)));
hdp_policy->set_flush_polling_ptr(hdp_gpu_cpu_flush_flag_);
hdp_flush_worker_thread = std::thread(&GPUIBBackend::hdp_flush_poll, this);
@@ -125,7 +125,7 @@ GPUIBBackend::GPUIBBackend(MPI_Comm comm) : Backend() {
}
__device__ bool GPUIBBackend::create_ctx(int64_t options,
roc_shmem_ctx_t *ctx) {
rocshmem_ctx_t *ctx) {
GPUIBContext *ctx_;
auto pop_result = ctx_free_list.get()->pop_front();
@@ -154,7 +154,7 @@ void GPUIBBackend::ctx_destroy(Context *ctx) {
delete gpu_ib_host_ctx;
}
__device__ void GPUIBBackend::destroy_ctx(roc_shmem_ctx_t *ctx) {
__device__ void GPUIBBackend::destroy_ctx(rocshmem_ctx_t *ctx) {
ctx_free_list.get()->push_back(static_cast<GPUIBContext *>(ctx->ctx_opaque));
}
@@ -167,7 +167,7 @@ GPUIBBackend::~GPUIBBackend() {
#ifdef USE_HOST_SIDE_HDP_FLUSH
hdp_flush_worker_thread.join();
hdp_policy->set_flush_polling_ptr(nullptr);
roc_shmem_free(hdp_gpu_cpu_flush_flag_);
rocshmem_free(hdp_gpu_cpu_flush_flag_);
#endif
/**
@@ -203,7 +203,7 @@ void GPUIBBackend::create_new_team([[maybe_unused]] Team *parent_team,
TeamInfo *team_info_wrt_parent,
TeamInfo *team_info_wrt_world, int num_pes,
int my_pe_in_new_team, MPI_Comm team_comm,
roc_shmem_team_t *new_team) {
rocshmem_team_t *new_team) {
/**
* Read the bit mask and find out a common index into
* the pool of available work arrays.
@@ -237,7 +237,7 @@ void GPUIBBackend::create_new_team([[maybe_unused]] Team *parent_team,
*new_team = get_external_team(new_team_obj);
}
void GPUIBBackend::team_destroy(roc_shmem_team_t team) {
void GPUIBBackend::team_destroy(rocshmem_team_t team) {
GPUIBTeam *team_obj = get_internal_gpu_ib_team(team);
/* Mark the pool as available */
@@ -263,7 +263,7 @@ void GPUIBBackend::initialize_network() { networkImpl.networkHostSetup(this); }
void GPUIBBackend::setup_default_host_ctx() {
default_host_ctx_ = new GPUIBHostContext(this, 0);
ROC_SHMEM_HOST_CTX_DEFAULT.ctx_opaque = default_host_ctx_;
ROCSHMEM_HOST_CTX_DEFAULT.ctx_opaque = default_host_ctx_;
}
void GPUIBBackend::setup_ctxs() {
@@ -288,19 +288,19 @@ void GPUIBBackend::setup_default_ctx() {
new (default_ctx_) GPUIBContext(this, true, 0);
/*
* Set the ROC_SHMEM_CTX_DEFAULT in constant memory.
* Set the ROCSHMEM_CTX_DEFAULT in constant memory.
*/
int *symbol_address;
CHECK_HIP(hipGetSymbolAddress(reinterpret_cast<void **>(&symbol_address),
HIP_SYMBOL(ROC_SHMEM_CTX_DEFAULT)));
HIP_SYMBOL(ROCSHMEM_CTX_DEFAULT)));
TeamInfo *tinfo = team_tracker.get_team_world()->tinfo_wrt_world;
roc_shmem_ctx_t ctx_default_host{default_ctx_, tinfo};
rocshmem_ctx_t ctx_default_host{default_ctx_, tinfo};
hipStream_t stream;
CHECK_HIP(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
CHECK_HIP(hipMemcpyAsync(symbol_address, &ctx_default_host,
sizeof(roc_shmem_ctx_t), hipMemcpyDefault, stream));
sizeof(rocshmem_ctx_t), hipMemcpyDefault, stream));
CHECK_HIP(hipStreamSynchronize(stream));
CHECK_HIP(hipStreamDestroy(stream));
}
@@ -328,9 +328,9 @@ void GPUIBBackend::setup_team_world() {
team_tracker.set_team_world(team_world);
/**
* Copy the address to ROC_SHMEM_TEAM_WORLD.
* Copy the address to ROCSHMEM_TEAM_WORLD.
*/
ROC_SHMEM_TEAM_WORLD = reinterpret_cast<roc_shmem_team_t>(team_world);
ROCSHMEM_TEAM_WORLD = reinterpret_cast<rocshmem_team_t>(team_world);
}
void GPUIBBackend::init_mpi_once(MPI_Comm comm) {
@@ -381,19 +381,19 @@ void GPUIBBackend::teams_init() {
* Allocate pools for the teams sync and work arrary from the SHEAP.
*/
auto max_num_teams{team_tracker.get_max_num_teams()};
barrier_pSync_pool = reinterpret_cast<long *>(roc_shmem_malloc(
sizeof(long) * ROC_SHMEM_BARRIER_SYNC_SIZE * max_num_teams));
reduce_pSync_pool = reinterpret_cast<long *>(roc_shmem_malloc(
sizeof(long) * ROC_SHMEM_REDUCE_SYNC_SIZE * max_num_teams));
bcast_pSync_pool = reinterpret_cast<long *>(roc_shmem_malloc(
sizeof(long) * ROC_SHMEM_BCAST_SYNC_SIZE * max_num_teams));
alltoall_pSync_pool = reinterpret_cast<long *>(roc_shmem_malloc(
sizeof(long) * ROC_SHMEM_ALLTOALL_SYNC_SIZE * max_num_teams));
barrier_pSync_pool = reinterpret_cast<long *>(rocshmem_malloc(
sizeof(long) * ROCSHMEM_BARRIER_SYNC_SIZE * max_num_teams));
reduce_pSync_pool = reinterpret_cast<long *>(rocshmem_malloc(
sizeof(long) * ROCSHMEM_REDUCE_SYNC_SIZE * max_num_teams));
bcast_pSync_pool = reinterpret_cast<long *>(rocshmem_malloc(
sizeof(long) * ROCSHMEM_BCAST_SYNC_SIZE * max_num_teams));
alltoall_pSync_pool = reinterpret_cast<long *>(rocshmem_malloc(
sizeof(long) * ROCSHMEM_ALLTOALL_SYNC_SIZE * max_num_teams));
/* Accommodating for largest possible data type for pWrk */
pWrk_pool = roc_shmem_malloc(
sizeof(double) * ROC_SHMEM_REDUCE_MIN_WRKDATA_SIZE * max_num_teams);
pAta_pool = roc_shmem_malloc(sizeof(double) * ROC_SHMEM_ATA_MAX_WRKDATA_SIZE *
pWrk_pool = rocshmem_malloc(
sizeof(double) * ROCSHMEM_REDUCE_MIN_WRKDATA_SIZE * max_num_teams);
pAta_pool = rocshmem_malloc(sizeof(double) * ROCSHMEM_ATA_MAX_WRKDATA_SIZE *
max_num_teams);
/**
@@ -402,25 +402,25 @@ void GPUIBBackend::teams_init() {
long *barrier_pSync, *reduce_pSync, *bcast_pSync, *alltoall_pSync;
for (int team_i = 0; team_i < max_num_teams; team_i++) {
barrier_pSync = reinterpret_cast<long *>(
&barrier_pSync_pool[team_i * ROC_SHMEM_BARRIER_SYNC_SIZE]);
&barrier_pSync_pool[team_i * ROCSHMEM_BARRIER_SYNC_SIZE]);
reduce_pSync = reinterpret_cast<long *>(
&reduce_pSync_pool[team_i * ROC_SHMEM_REDUCE_SYNC_SIZE]);
&reduce_pSync_pool[team_i * ROCSHMEM_REDUCE_SYNC_SIZE]);
bcast_pSync = reinterpret_cast<long *>(
&bcast_pSync_pool[team_i * ROC_SHMEM_BCAST_SYNC_SIZE]);
&bcast_pSync_pool[team_i * ROCSHMEM_BCAST_SYNC_SIZE]);
alltoall_pSync = reinterpret_cast<long *>(
&alltoall_pSync_pool[team_i * ROC_SHMEM_ALLTOALL_SYNC_SIZE]);
&alltoall_pSync_pool[team_i * ROCSHMEM_ALLTOALL_SYNC_SIZE]);
for (int i = 0; i < ROC_SHMEM_BARRIER_SYNC_SIZE; i++) {
barrier_pSync[i] = ROC_SHMEM_SYNC_VALUE;
for (int i = 0; i < ROCSHMEM_BARRIER_SYNC_SIZE; i++) {
barrier_pSync[i] = ROCSHMEM_SYNC_VALUE;
}
for (int i = 0; i < ROC_SHMEM_REDUCE_SYNC_SIZE; i++) {
reduce_pSync[i] = ROC_SHMEM_SYNC_VALUE;
for (int i = 0; i < ROCSHMEM_REDUCE_SYNC_SIZE; i++) {
reduce_pSync[i] = ROCSHMEM_SYNC_VALUE;
}
for (int i = 0; i < ROC_SHMEM_BCAST_SYNC_SIZE; i++) {
bcast_pSync[i] = ROC_SHMEM_SYNC_VALUE;
for (int i = 0; i < ROCSHMEM_BCAST_SYNC_SIZE; i++) {
bcast_pSync[i] = ROCSHMEM_SYNC_VALUE;
}
for (int i = 0; i < ROC_SHMEM_ALLTOALL_SYNC_SIZE; i++) {
alltoall_pSync[i] = ROC_SHMEM_SYNC_VALUE;
for (int i = 0; i < ROCSHMEM_ALLTOALL_SYNC_SIZE; i++) {
alltoall_pSync[i] = ROCSHMEM_SYNC_VALUE;
}
}
@@ -457,30 +457,30 @@ void GPUIBBackend::teams_init() {
}
void GPUIBBackend::teams_destroy() {
roc_shmem_free(barrier_pSync_pool);
roc_shmem_free(reduce_pSync_pool);
roc_shmem_free(bcast_pSync_pool);
roc_shmem_free(alltoall_pSync_pool);
roc_shmem_free(pWrk_pool);
roc_shmem_free(pAta_pool);
rocshmem_free(barrier_pSync_pool);
rocshmem_free(reduce_pSync_pool);
rocshmem_free(bcast_pSync_pool);
rocshmem_free(alltoall_pSync_pool);
rocshmem_free(pWrk_pool);
rocshmem_free(pAta_pool);
free(pool_bitmask_);
free(reduced_bitmask_);
}
void GPUIBBackend::roc_shmem_collective_init() {
void GPUIBBackend::rocshmem_collective_init() {
/*
* Allocate heap space for barrier_sync
*/
size_t one_sync_size_bytes{sizeof(*barrier_sync)};
size_t sync_size_bytes{one_sync_size_bytes * ROC_SHMEM_BARRIER_SYNC_SIZE};
size_t sync_size_bytes{one_sync_size_bytes * ROCSHMEM_BARRIER_SYNC_SIZE};
heap.malloc(reinterpret_cast<void **>(&barrier_sync), sync_size_bytes);
/*
* Initialize the barrier synchronization array with default values.
*/
for (int i = 0; i < num_pes; i++) {
barrier_sync[i] = ROC_SHMEM_SYNC_VALUE;
barrier_sync[i] = ROCSHMEM_SYNC_VALUE;
}
/*
src/gpu_ib/backend_ib.hpp
+11 -11
View File
@@ -72,19 +72,19 @@ class GPUIBBackend : public Backend {
void create_new_team(Team *parent_team, TeamInfo *team_info_wrt_parent,
TeamInfo *team_info_wrt_world, int num_pes,
int my_pe_in_new_team, MPI_Comm team_comm,
roc_shmem_team_t *new_team) override;
rocshmem_team_t *new_team) override;
/**
* @copydoc Backend::team_destroy(roc_shmem_team_t)
* @copydoc Backend::team_destroy(rocshmem_team_t)
*/
void team_destroy(roc_shmem_team_t team) override;
void team_destroy(rocshmem_team_t team) override;
/**
* @copydoc Backend::ctx_create
*/
void ctx_create(int64_t options, void **ctx) override;
__device__ bool create_ctx(int64_t options, roc_shmem_ctx_t *ctx);
__device__ bool create_ctx(int64_t options, rocshmem_ctx_t *ctx);
/**
* @copydoc Backend::ctx_destroy
@@ -94,7 +94,7 @@ class GPUIBBackend : public Backend {
/**
* @copydoc Backend::ctx_destroy
*/
__device__ void destroy_ctx(roc_shmem_ctx_t *ctx);
__device__ void destroy_ctx(rocshmem_ctx_t *ctx);
protected:
/**
@@ -151,10 +151,10 @@ class GPUIBBackend : public Backend {
void initialize_ipc();
/**
* @brief Allocate and initialize the ROC_SHMEM_CTX_DEFAULT variable.
* @brief Allocate and initialize the ROCSHMEM_CTX_DEFAULT variable.
*
* @todo The default_ctx member looks unused after it is copied into
* the ROC_SHMEM_CTX_DEFAULT variable.
* the ROCSHMEM_CTX_DEFAULT variable.
*/
void setup_default_ctx();
void setup_ctxs();
@@ -187,7 +187,7 @@ class GPUIBBackend : public Backend {
* When this method completes, the barrier_sync member will be available
* for use.
*/
void roc_shmem_collective_init();
void rocshmem_collective_init();
#ifdef USE_HOST_SIDE_HDP_FLUSH
/**
@@ -245,8 +245,8 @@ class GPUIBBackend : public Backend {
void *pAta_pool{nullptr};
/**
* @brief ROC_SHMEM's copy of MPI_COMM_WORLD (for interoperability
* with orthogonal MPI usage in an MPI+ROC_SHMEM program).
* @brief rocSHMEM's copy of MPI_COMM_WORLD (for interoperability
* with orthogonal MPI usage in an MPI+rocSHMEM program).
*/
MPI_Comm gpu_ib_comm_world{};
MPI_Comm backend_comm{};
@@ -334,7 +334,7 @@ class GPUIBBackend : public Backend {
* specification).
*
* @todo Remove this member from the backend class. There is another
* copy stored in ROC_SHMEM_CTX_DEFAULT.
* copy stored in ROCSHMEM_CTX_DEFAULT.
*/
GPUIBContext *default_ctx_{nullptr};
src/gpu_ib/connection.cpp
+4 -4
View File
@@ -39,19 +39,19 @@ int Connection::coherent_cq = 0;
Connection::Connection(GPUIBBackend* b, int k) : backend(b), key_offset(k) {
char* value = nullptr;
if ((value = getenv("ROC_SHMEM_USE_IB_HCA"))) {
if ((value = getenv("ROCSHMEM_USE_IB_HCA"))) {
requested_dev = value;
}
if ((value = getenv("ROC_SHMEM_SQ_SIZE"))) {
if ((value = getenv("ROCSHMEM_SQ_SIZE"))) {
sq_size = atoi(value);
}
if ((value = getenv("ROC_SHMEM_USE_CQ_GPU_MEM")) != nullptr) {
if ((value = getenv("ROCSHMEM_USE_CQ_GPU_MEM")) != nullptr) {
cq_use_gpu_mem = atoi(value);
}
if ((value = getenv("ROC_SHMEM_USE_SQ_GPU_MEM")) != nullptr) {
if ((value = getenv("ROCSHMEM_USE_SQ_GPU_MEM")) != nullptr) {
sq_use_gpu_mem = atoi(value);
}
}
src/gpu_ib/connection.hpp
+1 -1
View File
@@ -31,7 +31,7 @@ extern "C" {
#include <vector>
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#include "connection_policy.hpp"
namespace rocshmem {
src/gpu_ib/context_ib_device.cpp
+1 -1
View File
@@ -25,7 +25,7 @@
#include <hip/hip_runtime.h>
#include "config.h" // NOLINT(build/include_subdir)
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#include "../backend_type.hpp"
#include "../context_incl.hpp"
#include "backend_ib.hpp"
src/gpu_ib/context_ib_device.hpp
+17 -17
View File
@@ -69,7 +69,7 @@ class GPUIBContext : public Context {
__device__ void sync_all();
__device__ void sync(roc_shmem_team_t team);
__device__ void sync(rocshmem_team_t team);
template <typename T>
__device__ void amo_add(void *dst, T value, int pe);
@@ -113,13 +113,13 @@ class GPUIBContext : public Context {
template <typename T>
__device__ T g(const T *source, int pe);
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__device__ void to_all(T *dest, const T *source, int nreduce, int PE_start,
int logPE_stride, int PE_size, T *pWrk,
long *pSync); // NOLINT(runtime/int)
template <typename T, ROC_SHMEM_OP Op>
__device__ void to_all(roc_shmem_team_t team, T *dest, const T *source,
template <typename T, ROCSHMEM_OP Op>
__device__ void to_all(rocshmem_team_t team, T *dest, const T *source,
int nreduce);
template <typename T>
@@ -135,7 +135,7 @@ class GPUIBContext : public Context {
__device__ void get_nbi(T *dest, const T *source, size_t nelems, int pe);
template <typename T>
__device__ void broadcast(roc_shmem_team_t team, T *dest, const T *source,
__device__ void broadcast(rocshmem_team_t team, T *dest, const T *source,
int nelems, int pe_root);
template <typename T>
@@ -144,43 +144,43 @@ class GPUIBContext : public Context {
long *p_sync); // NOLINT(runtime/int)
template <typename T>
__device__ void alltoall(roc_shmem_team_t team, T *dest, const T *source,
__device__ void alltoall(rocshmem_team_t team, T *dest, const T *source,
int nelems);
template <typename T>
__device__ void alltoall_broadcast(roc_shmem_team_t team, T *dest,
__device__ void alltoall_broadcast(rocshmem_team_t team, T *dest,
const T *source, int nelems);
template <typename T>
__device__ void alltoall_brucks(roc_shmem_team_t team, T *dest,
__device__ void alltoall_brucks(rocshmem_team_t team, T *dest,
const T *source, int nelems);
template <typename T>
__device__ void alltoall_gcen(roc_shmem_team_t team, T *dest, const T *source,
__device__ void alltoall_gcen(rocshmem_team_t team, T *dest, const T *source,
int nelems);
template <typename T>
__device__ void alltoall_gcen2(roc_shmem_team_t team, T *dest,
__device__ void alltoall_gcen2(rocshmem_team_t team, T *dest,
const T *source, int nelems);
template <typename T>
__device__ void fcollect(roc_shmem_team_t team, T *dest, const T *source,
__device__ void fcollect(rocshmem_team_t team, T *dest, const T *source,
int nelems);
template <typename T>
__device__ void fcollect_broadcast(roc_shmem_team_t team, T *dest,
__device__ void fcollect_broadcast(rocshmem_team_t team, T *dest,
const T *source, int nelems);
template <typename T>
__device__ void fcollect_brucks(roc_shmem_team_t team, T *dest,
__device__ void fcollect_brucks(rocshmem_team_t team, T *dest,
const T *source, int nelems);
template <typename T>
__device__ void fcollect_gcen(roc_shmem_team_t team, T *dest, const T *source,
__device__ void fcollect_gcen(rocshmem_team_t team, T *dest, const T *source,
int nelems);
template <typename T>
__device__ void fcollect_gcen2(roc_shmem_team_t team, T *dest,
__device__ void fcollect_gcen2(rocshmem_team_t team, T *dest,
const T *source, int nelems);
__device__ void putmem_wg(void *dest, const void *source, size_t nelems,
@@ -232,13 +232,13 @@ class GPUIBContext : public Context {
__device__ void get_nbi_wave(T *dest, const T *source, size_t nelems, int pe);
private:
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__device__ void internal_direct_allreduce(
T *dst, const T *src, int nelems, int PE_start, int logPE_stride,
int PE_size, T *pWrk,
long *pSync); // NOLINT(runtime/int)
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__device__ void internal_ring_allreduce(T *dst, const T *src, int nelems,
int PE_start, int logPE_stride,
int PE_size, T *pWrk,
src/gpu_ib/context_ib_device_coll.cpp
+10 -10
View File
@@ -20,7 +20,7 @@
* IN THE SOFTWARE.
*****************************************************************************/
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#include "../context_incl.hpp"
#include "context_ib_tmpl_device.hpp"
#include "../util.hpp"
@@ -35,8 +35,8 @@ __device__ void GPUIBContext::internal_direct_barrier(int pe, int PE_start,
// Go through all PE offsets (except current offset = 0)
// and wait until they all reach
for (size_t i = 1; i < n_pes; i++) {
wait_until(&pSync[i], ROC_SHMEM_CMP_EQ, flag_val);
pSync[i] = ROC_SHMEM_SYNC_VALUE;
wait_until(&pSync[i], ROCSHMEM_CMP_EQ, flag_val);
pSync[i] = ROCSHMEM_SYNC_VALUE;
}
threadfence_system();
// Announce to other PEs that all have reached
@@ -48,8 +48,8 @@ __device__ void GPUIBContext::internal_direct_barrier(int pe, int PE_start,
// Mark current PE offset as reached
size_t pe_offset = (pe - PE_start) / stride;
put_nbi(&pSync[pe_offset], &flag_val, 1, PE_start);
wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, flag_val);
pSync[0] = ROC_SHMEM_SYNC_VALUE;
wait_until(&pSync[0], ROCSHMEM_CMP_EQ, flag_val);
pSync[0] = ROCSHMEM_SYNC_VALUE;
threadfence_system();
}
}
@@ -59,16 +59,16 @@ __device__ void GPUIBContext::internal_atomic_barrier(int pe, int PE_start,
int64_t *pSync) {
int64_t flag_val = 1;
if (pe == PE_start) {
wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, (int64_t)(n_pes - 1));
pSync[0] = ROC_SHMEM_SYNC_VALUE;
wait_until(&pSync[0], ROCSHMEM_CMP_EQ, (int64_t)(n_pes - 1));
pSync[0] = ROCSHMEM_SYNC_VALUE;
threadfence_system();
for (size_t i = 1, j = PE_start + stride; i < n_pes; ++i, j += stride) {
put_nbi(&pSync[0], &flag_val, 1, j);
}
} else {
amo_add<int64_t>(&pSync[0], flag_val, PE_start);
wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, flag_val);
pSync[0] = ROC_SHMEM_SYNC_VALUE;
wait_until(&pSync[0], ROCSHMEM_CMP_EQ, flag_val);
pSync[0] = ROCSHMEM_SYNC_VALUE;
threadfence_system();
}
}
@@ -88,7 +88,7 @@ __device__ void GPUIBContext::internal_sync(int pe, int PE_start, int stride,
__syncthreads();
}
__device__ void GPUIBContext::sync(roc_shmem_team_t team) {
__device__ void GPUIBContext::sync(rocshmem_team_t team) {
GPUIBTeam *team_obj = reinterpret_cast<GPUIBTeam *>(team);
double dbl_log_pe_stride = team_obj->tinfo_wrt_world->log_stride;
src/gpu_ib/context_ib_host.hpp
+4 -4
View File
@@ -86,16 +86,16 @@ class GPUIBHostContext : public Context {
long *p_sync); // NOLINT(runtime/int)
template <typename T>
__host__ void broadcast(roc_shmem_team_t team, T *dest, const T *source,
__host__ void broadcast(rocshmem_team_t team, T *dest, const T *source,
int nelems, int pe_root);
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__host__ void to_all(T *dest, const T *source, int nreduce, int pe_start,
int log_pe_stride, int pe_size, T *p_wrk,
long *p_sync); // NOLINT(runtime/int)
template <typename T, ROC_SHMEM_OP Op>
__host__ void to_all(roc_shmem_team_t team, T *dest, const T *source,
template <typename T, ROCSHMEM_OP Op>
__host__ void to_all(rocshmem_team_t team, T *dest, const T *source,
int nreduce);
template <typename T>
src/gpu_ib/context_ib_tmpl_device.hpp
+53 -53
View File
@@ -24,16 +24,16 @@
#define LIBRARY_SRC_GPU_IB_CONTEXT_IB_TMPL_DEVICE_HPP_
#include "config.h" // NOLINT(build/include_subdir)
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#include "context_ib_device.hpp"
#include "gpu_ib_team.hpp"
#include "queue_pair.hpp"
#include "../util.hpp"
#include "../roc_shmem_calc.hpp"
#include "../rocshmem_calc.hpp"
namespace rocshmem {
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__device__ void compute_reduce(T *src, T *dst, int size, int wg_id,
int wg_size) {
for (size_t i = wg_id; i < size; i += wg_size) {
@@ -47,7 +47,7 @@ __device__ void GPUIBContext::p(T *dest, T value, int pe) {
putmem_nbi(dest, &value, sizeof(T), pe);
}
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__device__ void GPUIBContext::internal_ring_allreduce(
T *dst, const T *src, int nelems, [[maybe_unused]] int PE_start,
[[maybe_unused]] int logPE_stride, [[maybe_unused]] int PE_size, T *pWrk,
@@ -81,7 +81,7 @@ __device__ void GPUIBContext::internal_ring_allreduce(
wait_val = seg + 100;
p(&pSync[round], wait_val, send_pe);
wait_until(&pSync[round], ROC_SHMEM_CMP_EQ, wait_val);
wait_until(&pSync[round], ROCSHMEM_CMP_EQ, wait_val);
__threadfence();
}
__syncthreads();
@@ -99,19 +99,19 @@ __device__ void GPUIBContext::internal_ring_allreduce(
fence();
wait_val = seg + 100;
p(&pSync[round], wait_val, send_pe);
wait_until(&pSync[round], ROC_SHMEM_CMP_EQ, wait_val);
wait_until(&pSync[round], ROCSHMEM_CMP_EQ, wait_val);
}
__syncthreads();
}
}
__syncthreads();
for (size_t i = wg_id; i < 2 * num_pes - 2; i += wg_size) {
pSync[i] = ROC_SHMEM_SYNC_VALUE;
pSync[i] = ROCSHMEM_SYNC_VALUE;
}
__syncthreads();
}
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__device__ void GPUIBContext::internal_direct_allreduce(
T *dst, const T *src, int nelems, int PE_start, int logPE_stride,
int PE_size, T *pWrk,
@@ -147,7 +147,7 @@ __device__ void GPUIBContext::internal_direct_allreduce(
if (i != pe) {
// Wait for leader thread to see that the buffer is ready.
if (is_thread_zero_in_block()) {
wait_until(&pSync[i], ROC_SHMEM_CMP_EQ, 1L);
wait_until(&pSync[i], ROCSHMEM_CMP_EQ, 1L);
}
__syncthreads();
@@ -159,14 +159,14 @@ __device__ void GPUIBContext::internal_direct_allreduce(
__syncthreads();
for (int i = wg_id; i < num_pes; i += wg_size) {
pSync[i] = ROC_SHMEM_SYNC_VALUE;
pSync[i] = ROCSHMEM_SYNC_VALUE;
}
__syncthreads();
}
template <typename T, ROC_SHMEM_OP Op>
__device__ void GPUIBContext::to_all(roc_shmem_team_t team, T *dest,
template <typename T, ROCSHMEM_OP Op>
__device__ void GPUIBContext::to_all(rocshmem_team_t team, T *dest,
const T *source, int nreduce) {
GPUIBTeam *team_obj = reinterpret_cast<GPUIBTeam *>(team);
@@ -189,7 +189,7 @@ __device__ void GPUIBContext::to_all(roc_shmem_team_t team, T *dest,
p_sync);
}
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__device__ void GPUIBContext::to_all(T *dest, const T *source, int nreduce,
int PE_start, int logPE_stride,
int PE_size, T *pWrk,
@@ -200,8 +200,8 @@ __device__ void GPUIBContext::to_all(T *dest, const T *source, int nreduce,
size_t ring_pSync = 2 * num_pes;
size_t provided_pWrk =
max(nreduce / 2 + 1, ROC_SHMEM_REDUCE_MIN_WRKDATA_SIZE);
size_t provided_pSync = ROC_SHMEM_REDUCE_SYNC_SIZE;
max(nreduce / 2 + 1, ROCSHMEM_REDUCE_MIN_WRKDATA_SIZE);
size_t provided_pSync = ROCSHMEM_REDUCE_SYNC_SIZE;
// TODO(bpotter):
// We basically do a direct reduce if pWrk is big enough, else we
@@ -212,12 +212,12 @@ __device__ void GPUIBContext::to_all(T *dest, const T *source, int nreduce,
internal_direct_allreduce<T, Op>(dest, source, nreduce, PE_start,
logPE_stride, PE_size, pWrk, pSync);
} else {
if (ring_pSync <= ROC_SHMEM_REDUCE_SYNC_SIZE) {
if (ring_pSync <= ROCSHMEM_REDUCE_SYNC_SIZE) {
int chunk_size = 1024;
size_t ring_pWrk = chunk_size * num_pes;
if (provided_pWrk < ring_pWrk) {
ring_pWrk = max(nreduce / 2, // NOLINT
ROC_SHMEM_REDUCE_MIN_WRKDATA_SIZE);
ROCSHMEM_REDUCE_MIN_WRKDATA_SIZE);
chunk_size = ring_pWrk / num_pes;
}
int seg_size = ring_pWrk;
@@ -434,7 +434,7 @@ __device__ void GPUIBContext::internal_get_broadcast(
}
template <typename T>
__device__ void GPUIBContext::broadcast(roc_shmem_team_t team, T *dst,
__device__ void GPUIBContext::broadcast(rocshmem_team_t team, T *dst,
const T *src, int nelems, int pe_root) {
GPUIBTeam *team_obj = reinterpret_cast<GPUIBTeam *>(team);
@@ -475,14 +475,14 @@ __device__ void GPUIBContext::broadcast(T *dst, const T *src, int nelems,
}
template <typename T>
__device__ void GPUIBContext::alltoall(roc_shmem_team_t team, T *dst,
__device__ void GPUIBContext::alltoall(rocshmem_team_t team, T *dst,
const T *src, int nelems) {
// Currently broadcast implementation performs the best
alltoall_broadcast(team, dst, src, nelems);
}
template <typename T>
__device__ void GPUIBContext::alltoall_broadcast(roc_shmem_team_t team, T *dst,
__device__ void GPUIBContext::alltoall_broadcast(rocshmem_team_t team, T *dst,
const T *src, int nelems) {
// Broadcast implementation of alltoall collective
GPUIBTeam *team_obj = reinterpret_cast<GPUIBTeam *>(team);
@@ -514,7 +514,7 @@ __device__ void GPUIBContext::alltoall_broadcast(roc_shmem_team_t team, T *dst,
}
template <typename T>
__device__ void GPUIBContext::alltoall_brucks(roc_shmem_team_t team, T *dst,
__device__ void GPUIBContext::alltoall_brucks(rocshmem_team_t team, T *dst,
const T *src, int nelems) {
// Brucks implementation of alltoall collective
GPUIBTeam *team_obj = reinterpret_cast<GPUIBTeam *>(team);
@@ -537,7 +537,7 @@ __device__ void GPUIBContext::alltoall_brucks(roc_shmem_team_t team, T *dst,
int blk_size = get_flat_block_size();
// Check if we have enough buffer space. If not, fail.
if (pe_size * nelems * 2 > ROC_SHMEM_ATA_MAX_WRKDATA_SIZE) {
if (pe_size * nelems * 2 > ROCSHMEM_ATA_MAX_WRKDATA_SIZE) {
GPU_DPRINTF("Unsupported alltoall size for gpu_ib.\n");
assert(false);
}
@@ -612,7 +612,7 @@ __device__ void GPUIBContext::alltoall_brucks(roc_shmem_team_t team, T *dst,
}
template <typename T>
__device__ void GPUIBContext::alltoall_gcen(roc_shmem_team_t team, T *dst,
__device__ void GPUIBContext::alltoall_gcen(rocshmem_team_t team, T *dst,
const T *src, int nelems) {
// GPU-centric implementation of alltoall collective
GPUIBTeam *team_obj = reinterpret_cast<GPUIBTeam *>(team);
@@ -629,12 +629,12 @@ __device__ void GPUIBContext::alltoall_gcen(roc_shmem_team_t team, T *dst,
int stride = 1 << log_pe_stride;
long *pSync = team_obj->alltoall_pSync;
int64_t *pSync2 = &team_obj->alltoall_pSync[ROC_SHMEM_BARRIER_SYNC_SIZE];
int64_t *pSync2 = &team_obj->alltoall_pSync[ROCSHMEM_BARRIER_SYNC_SIZE];
int my_pe_in_team = team_obj->my_pe;
// Check if we have enough buffer space. If not, fail.
T *pAta = reinterpret_cast<T *>(team_obj->pAta);
if (pe_size * nelems > ROC_SHMEM_ATA_MAX_WRKDATA_SIZE) {
if (pe_size * nelems > ROCSHMEM_ATA_MAX_WRKDATA_SIZE) {
GPU_DPRINTF("Unsupported alltoall size for gpu_ib.\n");
assert(false);
}
@@ -677,29 +677,29 @@ __device__ void GPUIBContext::alltoall_gcen(roc_shmem_team_t team, T *dst,
if (dest_pe2 != my_pe) amo_add<int64_t>(&pSync[0], flag_val, dest_pe2);
if (my_pe == dest_pe) {
wait_until(pSync2, ROC_SHMEM_CMP_EQ, flag_val * (clust_size - 1));
pSync2[0] = ROC_SHMEM_SYNC_VALUE;
wait_until(pSync2, ROCSHMEM_CMP_EQ, flag_val * (clust_size - 1));
pSync2[0] = ROCSHMEM_SYNC_VALUE;
__threadfence_system();
for (int i = 1; i < clust_size; ++i)
put_nbi(&pSync2[0], &flag_val, 1,
team_obj->get_pe_in_world(my_pe_in_team + i));
} else {
wait_until(pSync2, ROC_SHMEM_CMP_EQ, flag_val);
pSync2[0] = ROC_SHMEM_SYNC_VALUE;
wait_until(pSync2, ROCSHMEM_CMP_EQ, flag_val);
pSync2[0] = ROCSHMEM_SYNC_VALUE;
__threadfence_system();
}
if (my_pe == dest_pe2) {
wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, (int64_t)(num_clust - 1));
pSync[0] = ROC_SHMEM_SYNC_VALUE;
wait_until(&pSync[0], ROCSHMEM_CMP_EQ, (int64_t)(num_clust - 1));
pSync[0] = ROCSHMEM_SYNC_VALUE;
threadfence_system();
for (size_t i = 1, j = dest_pe2 + clust_size * stride; i < num_clust;
++i, j += clust_size * stride) {
put_nbi(&pSync[0], &flag_val, 1, j);
}
} else {
wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, flag_val);
pSync[0] = ROC_SHMEM_SYNC_VALUE;
wait_until(&pSync[0], ROCSHMEM_CMP_EQ, flag_val);
pSync[0] = ROCSHMEM_SYNC_VALUE;
threadfence_system();
}
}
@@ -707,7 +707,7 @@ __device__ void GPUIBContext::alltoall_gcen(roc_shmem_team_t team, T *dst,
}
template <typename T>
__device__ void GPUIBContext::alltoall_gcen2(roc_shmem_team_t team, T *dst,
__device__ void GPUIBContext::alltoall_gcen2(rocshmem_team_t team, T *dst,
const T *src, int nelems) {
// GPU-centric implementation of alltoall collective
// Uses in-place blocking sync
@@ -725,12 +725,12 @@ __device__ void GPUIBContext::alltoall_gcen2(roc_shmem_team_t team, T *dst,
int stride = 1 << log_pe_stride;
long *pSync = team_obj->alltoall_pSync;
int64_t *pSync2 = &team_obj->alltoall_pSync[ROC_SHMEM_BARRIER_SYNC_SIZE];
int64_t *pSync2 = &team_obj->alltoall_pSync[ROCSHMEM_BARRIER_SYNC_SIZE];
int my_pe_in_team = team_obj->my_pe;
// Check if we have enough buffer space. If not, fail.
T *pAta = reinterpret_cast<T *>(team_obj->pAta);
if (pe_size * nelems > ROC_SHMEM_ATA_MAX_WRKDATA_SIZE) {
if (pe_size * nelems > ROCSHMEM_ATA_MAX_WRKDATA_SIZE) {
GPU_DPRINTF("Unsupported alltoall size for gpu_ib.\n");
assert(false);
}
@@ -771,15 +771,15 @@ __device__ void GPUIBContext::alltoall_gcen2(roc_shmem_team_t team, T *dst,
if (is_thread_zero_in_block()) {
quiet();
if ((my_pe_in_team % clust_size) == 0) {
wait_until(pSync2, ROC_SHMEM_CMP_EQ, flag_val * (clust_size - 1));
pSync2[0] = ROC_SHMEM_SYNC_VALUE;
wait_until(pSync2, ROCSHMEM_CMP_EQ, flag_val * (clust_size - 1));
pSync2[0] = ROCSHMEM_SYNC_VALUE;
__threadfence_system();
for (int i = 1; i < clust_size; ++i)
put_nbi(&pSync2[0], &flag_val, 1,
team_obj->get_pe_in_world(my_pe_in_team + i));
} else {
wait_until(pSync2, ROC_SHMEM_CMP_EQ, flag_val);
pSync2[0] = ROC_SHMEM_SYNC_VALUE;
wait_until(pSync2, ROCSHMEM_CMP_EQ, flag_val);
pSync2[0] = ROCSHMEM_SYNC_VALUE;
__threadfence_system();
}
}
@@ -790,7 +790,7 @@ __device__ void GPUIBContext::alltoall_gcen2(roc_shmem_team_t team, T *dst,
}
template <typename T>
__device__ void GPUIBContext::fcollect(roc_shmem_team_t team, T *dst,
__device__ void GPUIBContext::fcollect(rocshmem_team_t team, T *dst,
const T *src, int nelems) {
// Main function for fcollect
// Broadcast version performs moderately well
@@ -799,7 +799,7 @@ __device__ void GPUIBContext::fcollect(roc_shmem_team_t team, T *dst,
}
template <typename T>
__device__ void GPUIBContext::fcollect_broadcast(roc_shmem_team_t team, T *dst,
__device__ void GPUIBContext::fcollect_broadcast(rocshmem_team_t team, T *dst,
const T *src, int nelems) {
// Broadcast implementation of fcollect collective
GPUIBTeam *team_obj = reinterpret_cast<GPUIBTeam *>(team);
@@ -832,7 +832,7 @@ __device__ void GPUIBContext::fcollect_broadcast(roc_shmem_team_t team, T *dst,
}
template <typename T>
__device__ void GPUIBContext::fcollect_brucks(roc_shmem_team_t team, T *dst,
__device__ void GPUIBContext::fcollect_brucks(rocshmem_team_t team, T *dst,
const T *src, int nelems) {
// Brucks implementation of fcollect collective
GPUIBTeam *team_obj = reinterpret_cast<GPUIBTeam *>(team);
@@ -855,7 +855,7 @@ __device__ void GPUIBContext::fcollect_brucks(roc_shmem_team_t team, T *dst,
int blk_size = get_flat_block_size();
// Check if we have enough buffer space. If not, fail.
if (pe_size * nelems > ROC_SHMEM_ATA_MAX_WRKDATA_SIZE) {
if (pe_size * nelems > ROCSHMEM_ATA_MAX_WRKDATA_SIZE) {
GPU_DPRINTF("Unsupported fcollect size for gpu_ib.\n");
assert(false);
}
@@ -895,7 +895,7 @@ __device__ void GPUIBContext::fcollect_brucks(roc_shmem_team_t team, T *dst,
}
template <typename T>
__device__ void GPUIBContext::fcollect_gcen(roc_shmem_team_t team, T *dst,
__device__ void GPUIBContext::fcollect_gcen(rocshmem_team_t team, T *dst,
const T *src, int nelems) {
// GPU-centric implementation of fcollect collective
GPUIBTeam *team_obj = reinterpret_cast<GPUIBTeam *>(team);
@@ -912,12 +912,12 @@ __device__ void GPUIBContext::fcollect_gcen(roc_shmem_team_t team, T *dst,
int stride = 1 << log_pe_stride;
long *pSync = team_obj->alltoall_pSync;
long *pSync2 = &team_obj->alltoall_pSync[ROC_SHMEM_BARRIER_SYNC_SIZE];
long *pSync2 = &team_obj->alltoall_pSync[ROCSHMEM_BARRIER_SYNC_SIZE];
int my_pe_in_team = team_obj->my_pe;
// Check if we have enough buffer space. If not, fail.
T *pAta = reinterpret_cast<T *>(team_obj->pAta);
if (pe_size * nelems > ROC_SHMEM_ATA_MAX_WRKDATA_SIZE) {
if (pe_size * nelems > ROCSHMEM_ATA_MAX_WRKDATA_SIZE) {
GPU_DPRINTF("Unsupported fcollect size for gpu_ib.\n");
assert(false);
}
@@ -957,15 +957,15 @@ __device__ void GPUIBContext::fcollect_gcen(roc_shmem_team_t team, T *dst,
if (is_thread_zero_in_block()) {
quiet();
if ((my_pe_in_team % clust_size) == 0) {
wait_until(pSync2, ROC_SHMEM_CMP_EQ, flag_val * (clust_size - 1));
pSync2[0] = ROC_SHMEM_SYNC_VALUE;
wait_until(pSync2, ROCSHMEM_CMP_EQ, flag_val * (clust_size - 1));
pSync2[0] = ROCSHMEM_SYNC_VALUE;
threadfence_system();
for (int i = 1; i < clust_size; ++i)
put_nbi(&pSync2[0], &flag_val, 1,
team_obj->get_pe_in_world(my_pe_in_team + i));
} else {
wait_until(pSync2, ROC_SHMEM_CMP_EQ, flag_val);
pSync2[0] = ROC_SHMEM_SYNC_VALUE;
wait_until(pSync2, ROCSHMEM_CMP_EQ, flag_val);
pSync2[0] = ROCSHMEM_SYNC_VALUE;
threadfence_system();
}
}
@@ -976,7 +976,7 @@ __device__ void GPUIBContext::fcollect_gcen(roc_shmem_team_t team, T *dst,
}
template <typename T>
__device__ void GPUIBContext::fcollect_gcen2(roc_shmem_team_t team, T *dst,
__device__ void GPUIBContext::fcollect_gcen2(rocshmem_team_t team, T *dst,
const T *src, int nelems) {
// GPU-centric implementation of fcollect collective
// Uses in-place blocking sync
@@ -998,7 +998,7 @@ __device__ void GPUIBContext::fcollect_gcen2(roc_shmem_team_t team, T *dst,
// Check if we have enough buffer space. If not, fail.
T *pAta = reinterpret_cast<T *>(team_obj->pAta);
if (pe_size * nelems > ROC_SHMEM_ATA_MAX_WRKDATA_SIZE) {
if (pe_size * nelems > ROCSHMEM_ATA_MAX_WRKDATA_SIZE) {
GPU_DPRINTF("Unsupported fcollect size for gpu_ib.\n");
assert(false);
}
src/gpu_ib/context_ib_tmpl_host.hpp
+4 -4
View File
@@ -93,13 +93,13 @@ __host__ void GPUIBHostContext::broadcast(
}
template <typename T>
__host__ void GPUIBHostContext::broadcast(roc_shmem_team_t team, T *dest,
__host__ void GPUIBHostContext::broadcast(rocshmem_team_t team, T *dest,
const T *source, int nelems,
int pe_root) {
host_interface->broadcast<T>(team, dest, source, nelems, pe_root);
}
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__host__ void GPUIBHostContext::to_all(T *dest, const T *source, int nreduce,
int pe_start, int log_pe_stride,
int pe_size, T *p_wrk,
@@ -108,8 +108,8 @@ __host__ void GPUIBHostContext::to_all(T *dest, const T *source, int nreduce,
pe_size, p_wrk, p_sync);
}
template <typename T, ROC_SHMEM_OP Op>
__host__ void GPUIBHostContext::to_all(roc_shmem_team_t team, T *dest,
template <typename T, ROCSHMEM_OP Op>
__host__ void GPUIBHostContext::to_all(rocshmem_team_t team, T *dest,
const T *source, int nreduce) {
host_interface->to_all<T, Op>(team, dest, source, nreduce);
}
src/gpu_ib/debug.cpp
+1 -1
View File
@@ -20,7 +20,7 @@
* IN THE SOFTWARE.
*****************************************************************************/
#include "roc_shmem/debug.hpp"
#include "rocshmem/debug.hpp"
#include "qe_dumper.hpp"
src/gpu_ib/dynamic_connection.cpp
+2 -2
View File
@@ -31,11 +31,11 @@ namespace rocshmem {
DynamicConnection::DynamicConnection(GPUIBBackend* b) : Connection(b, 4) {
char* value = nullptr;
if ((value = getenv("ROC_SHMEM_NUM_DCIs"))) {
if ((value = getenv("ROCSHMEM_NUM_DCIs"))) {
num_dcis = atoi(value);
}
if ((value = getenv("ROC_SHMEM_NUM_DCT"))) {
if ((value = getenv("ROCSHMEM_NUM_DCT"))) {
num_dct = atoi(value);
}
}
src/gpu_ib/gpu_ib_team.cpp
+6 -6
View File
@@ -38,17 +38,17 @@ GPUIBTeam::GPUIBTeam(Backend *backend, TeamInfo *team_info_parent,
pool_index_ = pool_index;
barrier_pSync =
&(b->barrier_pSync_pool[pool_index * ROC_SHMEM_BARRIER_SYNC_SIZE]);
&(b->barrier_pSync_pool[pool_index * ROCSHMEM_BARRIER_SYNC_SIZE]);
reduce_pSync =
&(b->reduce_pSync_pool[pool_index * ROC_SHMEM_REDUCE_SYNC_SIZE]);
bcast_pSync = &(b->bcast_pSync_pool[pool_index * ROC_SHMEM_BCAST_SYNC_SIZE]);
&(b->reduce_pSync_pool[pool_index * ROCSHMEM_REDUCE_SYNC_SIZE]);
bcast_pSync = &(b->bcast_pSync_pool[pool_index * ROCSHMEM_BCAST_SYNC_SIZE]);
alltoall_pSync =
&(b->alltoall_pSync_pool[pool_index * ROC_SHMEM_ALLTOALL_SYNC_SIZE]);
&(b->alltoall_pSync_pool[pool_index * ROCSHMEM_ALLTOALL_SYNC_SIZE]);
pWrk = reinterpret_cast<char *>(b->pWrk_pool) +
ROC_SHMEM_REDUCE_MIN_WRKDATA_SIZE * sizeof(double) * pool_index;
ROCSHMEM_REDUCE_MIN_WRKDATA_SIZE * sizeof(double) * pool_index;
pAta = reinterpret_cast<char *>(b->pAta_pool) +
ROC_SHMEM_ATA_MAX_WRKDATA_SIZE * sizeof(double) * pool_index;
ROCSHMEM_ATA_MAX_WRKDATA_SIZE * sizeof(double) * pool_index;
}
GPUIBTeam::~GPUIBTeam() {}
src/gpu_ib/network_policy.cpp
+2 -2
View File
@@ -297,7 +297,7 @@ void NetworkOnImpl::setup_gpu_qps(GPUIBBackend *B) {
}
}
void NetworkOnImpl::roc_shmem_g_init(SymmetricHeap *heap_handle,
void NetworkOnImpl::rocshmem_g_init(SymmetricHeap *heap_handle,
MPI_Comm thread_comm) {
init_g_ret(heap_handle, thread_comm, num_blocks, &g_ret);
}
@@ -327,7 +327,7 @@ __host__ void NetworkOnImpl::networkHostSetup(GPUIBBackend *B) {
connection->initialize_gpu_policy(&connection_policy, heap_rkey);
roc_shmem_g_init(&B->heap, B->thread_comm);
rocshmem_g_init(&B->heap, B->thread_comm);
connection->post_wqes();
src/gpu_ib/network_policy.hpp
+2 -2
View File
@@ -27,7 +27,7 @@
#include <mpi.h>
#include "config.h" // NOLINT(build/include_subdir)
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#include "connection_policy.hpp"
#include "queue_pair.hpp"
#include "../hdp_policy.hpp"
@@ -154,7 +154,7 @@ class NetworkOnImpl {
* @brief Allocate and initialize device-side memory that will be used for
* the return of g shmem ops (eg: shmem_int_g)
*/
void roc_shmem_g_init(SymmetricHeap *heap_handle, MPI_Comm thread_comm);
void rocshmem_g_init(SymmetricHeap *heap_handle, MPI_Comm thread_comm);
/**
* @brief The backend delegates some InfiniBand connection setup to
src/gpu_ib/queue_pair.hpp
+1 -1
View File
@@ -28,7 +28,7 @@
*
* @section DESCRIPTION
* An IB QueuePair (SQ and CQ) that the device can use to perform network
* operations. Most important ROC_SHMEM operations are performed by this
* operations. Most important rocSHMEM operations are performed by this
* class.
*/
src/gpu_ib/thread_policy.hpp
+2 -2
View File
@@ -32,7 +32,7 @@ class QueuePair;
/*
* GPU single-thread policy class. Only a single work-item per work-group
* is allowed to call into a ROC_SHMEM function (unless it is specifically
* is allowed to call into a rocSHMEM function (unless it is specifically
* called out as a collective API. This thread policy is the fastest but
* is not as flexible.
*/
@@ -59,7 +59,7 @@ class SingleThreadImpl {
/*
* GPU multi-thread policy class. Multiple work-items per work-group are
* allowed to call into a ROC_SHMEM function. A bit slower than its
* allowed to call into a rocSHMEM function. A bit slower than its
* single-thread counterpart but it enables a much more flexible user-facing
* API.
*/
src/host/host.cpp
+3 -3
View File
@@ -83,13 +83,13 @@ int HostInterface::find_win_info_in_pool(WindowInfo* window_info) {
}
__host__ HostInterface::HostInterface(HdpPolicy* hdp_policy,
MPI_Comm roc_shmem_comm,
MPI_Comm rocshmem_comm,
SymmetricHeap* heap) {
/*
* Duplicate a communicator from roc_shem's comm
* world for the host interface
*/
MPI_Comm_dup(roc_shmem_comm, &host_comm_world_);
MPI_Comm_dup(rocshmem_comm, &host_comm_world_);
MPI_Comm_rank(host_comm_world_, &my_pe_);
MPI_Comm_rank(host_comm_world_, &num_pes_);
@@ -103,7 +103,7 @@ __host__ HostInterface::HostInterface(HdpPolicy* hdp_policy,
* Allocate and initialize pool of windows for contexts
*/
char* value{nullptr};
if ((value = getenv("ROC_SHMEM_MAX_NUM_HOST_CONTEXTS"))) {
if ((value = getenv("ROCSHMEM_MAX_NUM_HOST_CONTEXTS"))) {
max_num_ctxs_ = atoi(value);
}
src/host/host.hpp
+8 -8
View File
@@ -36,7 +36,7 @@
#include <map>
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#include "../hdp_policy.hpp"
#include "../memory/symmetric_heap.hpp"
#include "../memory/window_info.hpp"
@@ -104,7 +104,7 @@ class HostInterface {
/**
* @brief Primary constructor
*/
__host__ HostInterface(HdpPolicy* hdp_policy, MPI_Comm roc_shmem_comm,
__host__ HostInterface(HdpPolicy* hdp_policy, MPI_Comm rocshmem_comm,
SymmetricHeap* heap);
/**
@@ -198,16 +198,16 @@ class HostInterface {
long* p_sync); // NOLINT(runtime/int)
template <typename T>
__host__ void broadcast(roc_shmem_team_t team, T* dest, const T* source,
__host__ void broadcast(rocshmem_team_t team, T* dest, const T* source,
int nelems, int pe_root);
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__host__ void to_all(T* dest, const T* source, int nreduce, int pe_start,
int log_pe_stride, int pe_size, T* p_wrk,
long* p_sync); // NOLINT(runtime/int)
template <typename T, ROC_SHMEM_OP Op>
__host__ int reduce(roc_shmem_team_t team, T* dest, const T* source, int nreduce);
template <typename T, ROCSHMEM_OP Op>
__host__ int reduce(rocshmem_team_t team, T* dest, const T* source, int nreduce);
template <typename T>
__host__ void wait_until(T *ivars, int cmp, T val,
@@ -288,7 +288,7 @@ class HostInterface {
__host__ MPI_Comm get_mpi_comm(int pe_start, int log_pe_stride, int pe_size);
__host__ MPI_Op get_mpi_op(ROC_SHMEM_OP Op);
__host__ MPI_Op get_mpi_op(ROCSHMEM_OP Op);
template <typename T>
__host__ MPI_Datatype get_mpi_type();
@@ -300,7 +300,7 @@ class HostInterface {
__host__ int test_and_compare(MPI_Aint offset, MPI_Datatype mpi_type,
int cmp, T val, MPI_Win win);
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__host__ void to_all_internal(MPI_Comm mpi_comm, T* dest, const T* source,
int nreduce);
src/host/host_templates.hpp
+22 -22
View File
@@ -200,7 +200,7 @@ __host__ void HostInterface::broadcast(T* dest, const T* source, int nelems,
}
template <typename T>
__host__ void HostInterface::broadcast(roc_shmem_team_t team, T* dest,
__host__ void HostInterface::broadcast(rocshmem_team_t team, T* dest,
const T* source, int nelems,
int pe_root) {
DPRINTF("Function: Team-based host_broadcast\n");
@@ -216,24 +216,24 @@ __host__ void HostInterface::broadcast(roc_shmem_team_t team, T* dest,
return;
}
__host__ inline MPI_Op HostInterface::get_mpi_op(ROC_SHMEM_OP Op) {
__host__ inline MPI_Op HostInterface::get_mpi_op(ROCSHMEM_OP Op) {
switch (Op) {
case ROC_SHMEM_SUM:
case ROCSHMEM_SUM:
return MPI_SUM;
case ROC_SHMEM_MAX:
case ROCSHMEM_MAX:
return MPI_MAX;
case ROC_SHMEM_MIN:
case ROCSHMEM_MIN:
return MPI_MIN;
case ROC_SHMEM_PROD:
case ROCSHMEM_PROD:
return MPI_PROD;
case ROC_SHMEM_AND:
case ROCSHMEM_AND:
return MPI_BAND;
case ROC_SHMEM_OR:
case ROCSHMEM_OR:
return MPI_BOR;
case ROC_SHMEM_XOR:
case ROCSHMEM_XOR:
return MPI_BXOR;
default:
fprintf(stderr, "Unknown ROC_SHMEM op MPI conversion %d\n", Op);
fprintf(stderr, "Unknown rocSHMEM op MPI conversion %d\n", Op);
abort();
return 0;
}
@@ -330,7 +330,7 @@ __host__ T HostInterface::amo_fetch_cas(void* dst, T value, T cond, int pe,
return ret;
}
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__host__ void HostInterface::to_all_internal(MPI_Comm mpi_comm, T* dest,
const T* source, int nreduce) {
DPRINTF("Function: host_to_all_internal\n");
@@ -356,7 +356,7 @@ __host__ void HostInterface::to_all_internal(MPI_Comm mpi_comm, T* dest,
return;
}
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__host__ void HostInterface::to_all(T* dest, const T* source, int nreduce,
int pe_start, int log_pe_stride,
int pe_size, [[maybe_unused]] T* p_wrk,
@@ -375,8 +375,8 @@ __host__ void HostInterface::to_all(T* dest, const T* source, int nreduce,
return;
}
template <typename T, ROC_SHMEM_OP Op>
__host__ int HostInterface::reduce(roc_shmem_team_t team, T* dest,
template <typename T, ROCSHMEM_OP Op>
__host__ int HostInterface::reduce(rocshmem_team_t team, T* dest,
const T* source, int nreduce) {
DPRINTF("Function: Team-based host_reduce\n");
@@ -388,7 +388,7 @@ __host__ int HostInterface::reduce(roc_shmem_team_t team, T* dest,
to_all_internal<T, Op>(mpi_comm, dest, source, nreduce);
return ROC_SHMEM_SUCCESS;
return ROCSHMEM_SUCCESS;
}
template <typename T>
@@ -397,26 +397,26 @@ __host__ inline int HostInterface::compare(int cmp, T input_val,
int cond_satisfied{0};
switch (cmp) {
case ROC_SHMEM_CMP_EQ:
case ROCSHMEM_CMP_EQ:
cond_satisfied = (input_val == target_val) ? 1 : 0;
break;
case ROC_SHMEM_CMP_NE:
case ROCSHMEM_CMP_NE:
cond_satisfied = (input_val != target_val) ? 1 : 0;
break;
case ROC_SHMEM_CMP_GT:
case ROCSHMEM_CMP_GT:
cond_satisfied = (input_val > target_val) ? 1 : 0;
break;
case ROC_SHMEM_CMP_GE:
case ROCSHMEM_CMP_GE:
cond_satisfied = (input_val >= target_val) ? 1 : 0;
break;
case ROC_SHMEM_CMP_LT:
case ROCSHMEM_CMP_LT:
cond_satisfied = (input_val < target_val) ? 1 : 0;
break;
case ROC_SHMEM_CMP_LE:
case ROCSHMEM_CMP_LE:
cond_satisfied = (input_val <= target_val) ? 1 : 0;
break;
default:
assert(cmp >= ROC_SHMEM_CMP_EQ && cmp <= ROC_SHMEM_CMP_LE);
assert(cmp >= ROCSHMEM_CMP_EQ && cmp <= ROCSHMEM_CMP_LE);
break;
}
src/ipc/backend_ipc.cpp
+40 -40
View File
@@ -33,10 +33,10 @@ namespace rocshmem {
} \
}
extern roc_shmem_ctx_t ROC_SHMEM_HOST_CTX_DEFAULT;
extern rocshmem_ctx_t ROCSHMEM_HOST_CTX_DEFAULT;
roc_shmem_team_t get_external_team(GPUIBTeam *team) {
return reinterpret_cast<roc_shmem_team_t>(team);
rocshmem_team_t get_external_team(GPUIBTeam *team) {
return reinterpret_cast<rocshmem_team_t>(team);
}
int get_ls_non_zero_bit(char *bitmask, int mask_length) {
@@ -57,7 +57,7 @@ IPCBackend::IPCBackend(MPI_Comm comm)
: Backend() {
type = BackendType::IPC_BACKEND;
if (auto maximum_num_contexts_str = getenv("ROC_SHMEM_MAX_NUM_CONTEXTS")) {
if (auto maximum_num_contexts_str = getenv("ROCSHMEM_MAX_NUM_CONTEXTS")) {
std::stringstream sstream(maximum_num_contexts_str);
sstream >> maximum_num_contexts_;
}
@@ -82,7 +82,7 @@ IPCBackend::IPCBackend(MPI_Comm comm)
default_host_ctx = std::make_unique<IPCHostContext>(this, 0);
ROC_SHMEM_HOST_CTX_DEFAULT.ctx_opaque = default_host_ctx.get();
ROCSHMEM_HOST_CTX_DEFAULT.ctx_opaque = default_host_ctx.get();
init_g_ret(&heap, thread_comm, MAX_NUM_BLOCKS, &bp->g_ret);
@@ -92,7 +92,7 @@ IPCBackend::IPCBackend(MPI_Comm comm)
init_wrk_sync_buffer();
roc_shmem_collective_init();
rocshmem_collective_init();
setup_fence_buffer();
@@ -143,7 +143,7 @@ void IPCBackend::setup_ctxs() {
}
}
__device__ bool IPCBackend::create_ctx(int64_t options, roc_shmem_ctx_t *ctx) {
__device__ bool IPCBackend::create_ctx(int64_t options, rocshmem_ctx_t *ctx) {
IPCContext *ctx_{nullptr};
auto pop_result = ctx_free_list.get()->pop_front();
@@ -158,7 +158,7 @@ __device__ bool IPCBackend::create_ctx(int64_t options, roc_shmem_ctx_t *ctx) {
return true;
}
__device__ void IPCBackend::destroy_ctx(roc_shmem_ctx_t *ctx) {
__device__ void IPCBackend::destroy_ctx(rocshmem_ctx_t *ctx) {
ctx_free_list.get()->push_back(static_cast<IPCContext *>(ctx->ctx_opaque));
}
@@ -182,9 +182,9 @@ void IPCBackend::setup_team_world() {
team_tracker.set_team_world(team_world);
/**
* Copy the address to ROC_SHMEM_TEAM_WORLD.
* Copy the address to ROCSHMEM_TEAM_WORLD.
*/
ROC_SHMEM_TEAM_WORLD = reinterpret_cast<roc_shmem_team_t>(team_world);
ROCSHMEM_TEAM_WORLD = reinterpret_cast<rocshmem_team_t>(team_world);
}
void IPCBackend::init_mpi_once(MPI_Comm comm) {
@@ -205,7 +205,7 @@ void IPCBackend::init_mpi_once(MPI_Comm comm) {
NET_CHECK(MPI_Comm_rank(thread_comm, &my_pe));
}
void IPCBackend::team_destroy(roc_shmem_team_t team) {
void IPCBackend::team_destroy(rocshmem_team_t team) {
IPCTeam *team_obj = get_internal_ipc_team(team);
/* Mark the pool as available */
@@ -221,7 +221,7 @@ void IPCBackend::create_new_team([[maybe_unused]] Team *parent_team,
TeamInfo *team_info_wrt_parent,
TeamInfo *team_info_wrt_world, int num_pes,
int my_pe_in_new_team, MPI_Comm team_comm,
roc_shmem_team_t *new_team) {
rocshmem_team_t *new_team) {
/**
* Read the bit mask and find out a common index into
* the pool of available work arrays.
@@ -303,24 +303,24 @@ void IPCBackend::init_wrk_sync_buffer() {
/**
* size of barrier sync
*/
Wrk_Sync_buffer_size_ += sizeof(*barrier_sync) * ROC_SHMEM_BARRIER_SYNC_SIZE;
Wrk_Sync_buffer_size_ += sizeof(*barrier_sync) * ROCSHMEM_BARRIER_SYNC_SIZE;
/**
* Size of sync arrays for the teams
*/
Wrk_Sync_buffer_size_ += sizeof(long) * max_num_teams *
(ROC_SHMEM_BARRIER_SYNC_SIZE +
ROC_SHMEM_REDUCE_SYNC_SIZE +
ROC_SHMEM_BCAST_SYNC_SIZE +
ROC_SHMEM_ALLTOALL_SYNC_SIZE);
(ROCSHMEM_BARRIER_SYNC_SIZE +
ROCSHMEM_REDUCE_SYNC_SIZE +
ROCSHMEM_BCAST_SYNC_SIZE +
ROCSHMEM_ALLTOALL_SYNC_SIZE);
/**
* Size of work arrays for the teams
* Accommodate largest possible data type for pWrk
*/
Wrk_Sync_buffer_size_ += sizeof(double) * max_num_teams *
(ROC_SHMEM_REDUCE_MIN_WRKDATA_SIZE +
ROC_SHMEM_ATA_MAX_WRKDATA_SIZE);
(ROCSHMEM_REDUCE_MIN_WRKDATA_SIZE +
ROCSHMEM_ATA_MAX_WRKDATA_SIZE);
/**
* Size of fence array
@@ -397,12 +397,12 @@ void IPCBackend::setup_fence_buffer() {
temp_Wrk_Sync_buff_ptr_ += sizeof(int) * num_pes;
}
void IPCBackend::roc_shmem_collective_init() {
void IPCBackend::rocshmem_collective_init() {
/*
* Allocate heap space for barrier_sync
*/
size_t one_sync_size_bytes{sizeof(*barrier_sync)};
size_t sync_size_bytes{one_sync_size_bytes * ROC_SHMEM_BARRIER_SYNC_SIZE};
size_t sync_size_bytes{one_sync_size_bytes * ROCSHMEM_BARRIER_SYNC_SIZE};
barrier_sync = reinterpret_cast<int64_t*>(temp_Wrk_Sync_buff_ptr_);
temp_Wrk_Sync_buff_ptr_ += sync_size_bytes;
@@ -410,7 +410,7 @@ void IPCBackend::roc_shmem_collective_init() {
* Initialize the barrier synchronization array with default values.
*/
for (int i = 0; i < num_pes; i++) {
barrier_sync[i] = ROC_SHMEM_SYNC_VALUE;
barrier_sync[i] = ROCSHMEM_SYNC_VALUE;
}
/*
@@ -427,29 +427,29 @@ void IPCBackend::teams_init() {
auto max_num_teams{team_tracker.get_max_num_teams()};
barrier_pSync_pool = reinterpret_cast<long *>(temp_Wrk_Sync_buff_ptr_);
temp_Wrk_Sync_buff_ptr_ += sizeof(long) * ROC_SHMEM_BARRIER_SYNC_SIZE
temp_Wrk_Sync_buff_ptr_ += sizeof(long) * ROCSHMEM_BARRIER_SYNC_SIZE
* max_num_teams;
reduce_pSync_pool = reinterpret_cast<long *>(temp_Wrk_Sync_buff_ptr_);
temp_Wrk_Sync_buff_ptr_ += sizeof(long) * ROC_SHMEM_REDUCE_SYNC_SIZE
temp_Wrk_Sync_buff_ptr_ += sizeof(long) * ROCSHMEM_REDUCE_SYNC_SIZE
* max_num_teams;
bcast_pSync_pool = reinterpret_cast<long *>(temp_Wrk_Sync_buff_ptr_);
temp_Wrk_Sync_buff_ptr_ += sizeof(long) * ROC_SHMEM_BCAST_SYNC_SIZE
temp_Wrk_Sync_buff_ptr_ += sizeof(long) * ROCSHMEM_BCAST_SYNC_SIZE
* max_num_teams;
alltoall_pSync_pool = reinterpret_cast<long *>(temp_Wrk_Sync_buff_ptr_);
temp_Wrk_Sync_buff_ptr_ += sizeof(long) * ROC_SHMEM_BCAST_SYNC_SIZE
temp_Wrk_Sync_buff_ptr_ += sizeof(long) * ROCSHMEM_BCAST_SYNC_SIZE
* max_num_teams;
/* Accommodating for largest possible data type for pWrk */
pWrk_pool = reinterpret_cast<void *>(temp_Wrk_Sync_buff_ptr_);
temp_Wrk_Sync_buff_ptr_ += sizeof(double) * ROC_SHMEM_REDUCE_MIN_WRKDATA_SIZE
temp_Wrk_Sync_buff_ptr_ += sizeof(double) * ROCSHMEM_REDUCE_MIN_WRKDATA_SIZE
* max_num_teams;
pAta_pool = reinterpret_cast<void *>(temp_Wrk_Sync_buff_ptr_);
temp_Wrk_Sync_buff_ptr_ += sizeof(double) * ROC_SHMEM_ATA_MAX_WRKDATA_SIZE
temp_Wrk_Sync_buff_ptr_ += sizeof(double) * ROCSHMEM_ATA_MAX_WRKDATA_SIZE
* max_num_teams;
/**
@@ -458,25 +458,25 @@ void IPCBackend::teams_init() {
long *barrier_pSync, *reduce_pSync, *bcast_pSync, *alltoall_pSync;
for (int team_i = 0; team_i < max_num_teams; team_i++) {
barrier_pSync = reinterpret_cast<long *>(
&barrier_pSync_pool[team_i * ROC_SHMEM_BARRIER_SYNC_SIZE]);
&barrier_pSync_pool[team_i * ROCSHMEM_BARRIER_SYNC_SIZE]);
reduce_pSync = reinterpret_cast<long *>(
&reduce_pSync_pool[team_i * ROC_SHMEM_REDUCE_SYNC_SIZE]);
&reduce_pSync_pool[team_i * ROCSHMEM_REDUCE_SYNC_SIZE]);
bcast_pSync = reinterpret_cast<long *>(
&bcast_pSync_pool[team_i * ROC_SHMEM_BCAST_SYNC_SIZE]);
&bcast_pSync_pool[team_i * ROCSHMEM_BCAST_SYNC_SIZE]);
alltoall_pSync = reinterpret_cast<long *>(
&alltoall_pSync_pool[team_i * ROC_SHMEM_ALLTOALL_SYNC_SIZE]);
&alltoall_pSync_pool[team_i * ROCSHMEM_ALLTOALL_SYNC_SIZE]);
for (int i = 0; i < ROC_SHMEM_BARRIER_SYNC_SIZE; i++) {
barrier_pSync[i] = ROC_SHMEM_SYNC_VALUE;
for (int i = 0; i < ROCSHMEM_BARRIER_SYNC_SIZE; i++) {
barrier_pSync[i] = ROCSHMEM_SYNC_VALUE;
}
for (int i = 0; i < ROC_SHMEM_REDUCE_SYNC_SIZE; i++) {
reduce_pSync[i] = ROC_SHMEM_SYNC_VALUE;
for (int i = 0; i < ROCSHMEM_REDUCE_SYNC_SIZE; i++) {
reduce_pSync[i] = ROCSHMEM_SYNC_VALUE;
}
for (int i = 0; i < ROC_SHMEM_BCAST_SYNC_SIZE; i++) {
bcast_pSync[i] = ROC_SHMEM_SYNC_VALUE;
for (int i = 0; i < ROCSHMEM_BCAST_SYNC_SIZE; i++) {
bcast_pSync[i] = ROCSHMEM_SYNC_VALUE;
}
for (int i = 0; i < ROC_SHMEM_ALLTOALL_SYNC_SIZE; i++) {
alltoall_pSync[i] = ROC_SHMEM_SYNC_VALUE;
for (int i = 0; i < ROCSHMEM_ALLTOALL_SYNC_SIZE; i++) {
alltoall_pSync[i] = ROCSHMEM_SYNC_VALUE;
}
}
src/ipc/backend_ipc.hpp
+7 -7
View File
@@ -48,13 +48,13 @@ class IPCBackend : public Backend {
*/
virtual ~IPCBackend();
__device__ bool create_ctx(int64_t options, roc_shmem_ctx_t *ctx);
__device__ bool create_ctx(int64_t options, rocshmem_ctx_t *ctx);
/**
* @brief Destroy a `roc_shmem_ctx_t` context and returns it back to the
* @brief Destroy a `rocshmem_ctx_t` context and returns it back to the
* context free list.
*/
__device__ void destroy_ctx(roc_shmem_ctx_t *ctx);
__device__ void destroy_ctx(rocshmem_ctx_t *ctx);
/**
* @copydoc Backend::ctx_create
@@ -103,12 +103,12 @@ class IPCBackend : public Backend {
void create_new_team(Team *parent_team, TeamInfo *team_info_wrt_parent,
TeamInfo *team_info_wrt_world, int num_pes,
int my_pe_in_new_team, MPI_Comm team_comm,
roc_shmem_team_t *new_team) override;
rocshmem_team_t *new_team) override;
/**
* @copydoc Backend::team_destroy(roc_shmem_team_t)
* @copydoc Backend::team_destroy(rocshmem_team_t)
*/
void team_destroy(roc_shmem_team_t team) override;
void team_destroy(rocshmem_team_t team) override;
/**
* @brief Accessor for work/sync bases
@@ -213,7 +213,7 @@ class IPCBackend : public Backend {
* When this method completes, the barrier_sync member will be available
* for use.
*/
void roc_shmem_collective_init();
void rocshmem_collective_init();
/**
* @brief Allocate buffer for fence/quiet operation
src/ipc/context_ipc_device.cpp
+7 -7
View File
@@ -31,7 +31,7 @@
#include <cstdlib>
#include "config.h" // NOLINT(build/include_subdir)
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#include "backend_ipc.hpp"
namespace rocshmem {
@@ -229,10 +229,10 @@ __device__ void IPCContext::putmem_signal(void *dest, const void *source, size_t
fence();
switch (sig_op) {
case ROC_SHMEM_SIGNAL_SET:
case ROCSHMEM_SIGNAL_SET:
amo_set<uint64_t>(static_cast<void*>(sig_addr), signal, pe);
break;
case ROC_SHMEM_SIGNAL_ADD:
case ROCSHMEM_SIGNAL_ADD:
amo_add<uint64_t>(static_cast<void*>(dest), signal, pe);
break;
default:
@@ -249,10 +249,10 @@ __device__ void IPCContext::putmem_signal_wg(void *dest, const void *source, siz
if (is_thread_zero_in_block()) {
switch (sig_op) {
case ROC_SHMEM_SIGNAL_SET:
case ROCSHMEM_SIGNAL_SET:
amo_set<uint64_t>(static_cast<void*>(sig_addr), signal, pe);
break;
case ROC_SHMEM_SIGNAL_ADD:
case ROCSHMEM_SIGNAL_ADD:
amo_add<uint64_t>(static_cast<void*>(dest), signal, pe);
break;
default:
@@ -270,10 +270,10 @@ __device__ void IPCContext::putmem_signal_wave(void *dest, const void *source, s
if (is_thread_zero_in_wave()) {
switch (sig_op) {
case ROC_SHMEM_SIGNAL_SET:
case ROCSHMEM_SIGNAL_SET:
amo_set<uint64_t>(static_cast<void*>(sig_addr), signal, pe);
break;
case ROC_SHMEM_SIGNAL_ADD:
case ROCSHMEM_SIGNAL_ADD:
amo_add<uint64_t>(static_cast<void*>(dest), signal, pe);
break;
default:
src/ipc/context_ipc_device.hpp
+10 -10
View File
@@ -63,7 +63,7 @@ class IPCContext : public Context {
__device__ void sync_all();
__device__ void sync(roc_shmem_team_t team);
__device__ void sync(rocshmem_team_t team);
template <typename T>
__device__ void p(T *dest, T value, int pe);
@@ -121,18 +121,18 @@ class IPCContext : public Context {
__device__ T amo_fetch_cas(void *dst, T value, T cond, int pe);
// Collectives
template <typename T, ROC_SHMEM_OP Op>
__device__ int reduce(roc_shmem_team_t team, T *dest, const T *source, int nreduce);
template <typename T, ROCSHMEM_OP Op>
__device__ int reduce(rocshmem_team_t team, T *dest, const T *source, int nreduce);
template <typename T>
__device__ void broadcast(roc_shmem_team_t team, T *dest, const T *source,
__device__ void broadcast(rocshmem_team_t team, T *dest, const T *source,
int nelems, int pe_root);
template <typename T>
__device__ void alltoall(roc_shmem_team_t team, T *dest, const T *source,
__device__ void alltoall(rocshmem_team_t team, T *dest, const T *source,
int nelems);
template <typename T>
__device__ void fcollect(roc_shmem_team_t team, T *dest, const T *source,
__device__ void fcollect(rocshmem_team_t team, T *dest, const T *source,
int nelems);
@@ -232,11 +232,11 @@ class IPCContext : public Context {
int pe_root); // NOLINT(runtime/int)
template <typename T>
__device__ void fcollect_linear(roc_shmem_team_t team, T *dest,
__device__ void fcollect_linear(rocshmem_team_t team, T *dest,
const T *source, int nelems);
template <typename T>
__device__ void alltoall_linear(roc_shmem_team_t team, T *dest,
__device__ void alltoall_linear(rocshmem_team_t team, T *dest,
const T *source, int nelems);
__device__ void internal_sync(int pe, int PE_start, int stride, int PE_size,
@@ -248,10 +248,10 @@ class IPCContext : public Context {
__device__ void internal_atomic_barrier(int pe, int PE_start, int stride,
int n_pes, int64_t *pSync);
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__device__ void internal_direct_allreduce(T *dst, const T *src,
int nelems, IPCTeam *team_obj);
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__device__ void internal_ring_allreduce(T *dst, const T *src,
int nelems, IPCTeam *team_obj,
int n_seg, int seg_size, int chunk_size);
src/ipc/context_ipc_device_coll.cpp
+10 -10
View File
@@ -20,7 +20,7 @@
* IN THE SOFTWARE.
*****************************************************************************/
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#include "../context_incl.hpp"
#include "context_ipc_tmpl_device.hpp"
#include "../util.hpp"
@@ -39,8 +39,8 @@ __device__ void IPCContext::internal_direct_barrier(int pe, int PE_start,
__threadfence_system();
#endif /* __gfx90a__ */
for (size_t i = 1; i < n_pes; i++) {
wait_until(&pSync[i], ROC_SHMEM_CMP_EQ, flag_val);
pSync[i] = ROC_SHMEM_SYNC_VALUE;
wait_until(&pSync[i], ROCSHMEM_CMP_EQ, flag_val);
pSync[i] = ROCSHMEM_SYNC_VALUE;
}
threadfence_system();
@@ -58,8 +58,8 @@ __device__ void IPCContext::internal_direct_barrier(int pe, int PE_start,
#if defined(__gfx90a__)
__threadfence_system();
#endif /* __gfx90a__ */
wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, flag_val);
pSync[0] = ROC_SHMEM_SYNC_VALUE;
wait_until(&pSync[0], ROCSHMEM_CMP_EQ, flag_val);
pSync[0] = ROCSHMEM_SYNC_VALUE;
threadfence_system();
}
}
@@ -69,8 +69,8 @@ __device__ void IPCContext::internal_atomic_barrier(int pe, int PE_start,
int64_t *pSync) {
int64_t flag_val = 1;
if (pe == PE_start) {
wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, (int64_t)(n_pes - 1));
pSync[0] = ROC_SHMEM_SYNC_VALUE;
wait_until(&pSync[0], ROCSHMEM_CMP_EQ, (int64_t)(n_pes - 1));
pSync[0] = ROCSHMEM_SYNC_VALUE;
threadfence_system();
for (size_t i = 1, j = PE_start + stride; i < n_pes; ++i, j += stride) {
@@ -78,8 +78,8 @@ __device__ void IPCContext::internal_atomic_barrier(int pe, int PE_start,
}
} else {
amo_add<int64_t>(&pSync[0], flag_val, PE_start);
wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, flag_val);
pSync[0] = ROC_SHMEM_SYNC_VALUE;
wait_until(&pSync[0], ROCSHMEM_CMP_EQ, flag_val);
pSync[0] = ROCSHMEM_SYNC_VALUE;
threadfence_system();
}
}
@@ -98,7 +98,7 @@ __device__ void IPCContext::internal_sync(int pe, int PE_start, int stride,
__syncthreads();
}
__device__ void IPCContext::sync(roc_shmem_team_t team) {
__device__ void IPCContext::sync(rocshmem_team_t team) {
IPCTeam *team_obj = reinterpret_cast<IPCTeam *>(team);
int pe = team_obj->my_pe_in_world;
src/ipc/context_ipc_host.hpp
+4 -4
View File
@@ -86,16 +86,16 @@ class IPCHostContext : public Context {
long *p_sync);
template <typename T>
__host__ void broadcast(roc_shmem_team_t team, T *dest, const T *source,
__host__ void broadcast(rocshmem_team_t team, T *dest, const T *source,
int nelems, int pe_root);
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__host__ void to_all(T *dest, const T *source, int nreduce, int pe_start,
int log_pe_stride, int pe_size, T *p_wrk,
long *p_sync);
template <typename T, ROC_SHMEM_OP Op>
__host__ int reduce(roc_shmem_team_t team, T *dest, const T *source, int nreduce);
template <typename T, ROCSHMEM_OP Op>
__host__ int reduce(rocshmem_team_t team, T *dest, const T *source, int nreduce);
template <typename T>
__host__ void wait_until(T *ivars, int cmp, T val);
src/ipc/context_ipc_tmpl_device.hpp
+23 -23
View File
@@ -24,11 +24,11 @@
#define LIBRARY_SRC_IPC_CONTEXT_TMPL_DEVICE_HPP_
#include "config.h" // NOLINT(build/include_subdir)
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#include "context_ipc_device.hpp"
#include "../util.hpp"
#include "ipc_team.hpp"
#include "../roc_shmem_calc.hpp"
#include "../rocshmem_calc.hpp"
namespace rocshmem {
@@ -153,7 +153,7 @@ __device__ T IPCContext::amo_fetch_cas(void *dest, T value, T cond, int pe) {
}
// Collectives
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__device__ void compute_reduce(T *src, T *dst, int size, int wg_id,
int wg_size) {
for (size_t i = wg_id; i < size; i += wg_size) {
@@ -162,7 +162,7 @@ __device__ void compute_reduce(T *src, T *dst, int size, int wg_id,
__syncthreads();
}
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__device__ void IPCContext::internal_direct_allreduce(
T *dst, const T *src, int nelems, IPCTeam *team_obj) { // NOLINT(runtime/int)
@@ -203,7 +203,7 @@ __device__ void IPCContext::internal_direct_allreduce(
if (i != pe) {
// Wait for leader thread to see that the buffer is ready.
if (is_thread_zero_in_block()) {
wait_until(&pSync[i], ROC_SHMEM_CMP_EQ, flag_val);
wait_until(&pSync[i], ROCSHMEM_CMP_EQ, flag_val);
}
__syncthreads();
@@ -216,7 +216,7 @@ __device__ void IPCContext::internal_direct_allreduce(
__syncthreads();
for (int i = wg_id; i < num_pes; i += wg_size) {
pSync[i] = ROC_SHMEM_SYNC_VALUE;
pSync[i] = ROCSHMEM_SYNC_VALUE;
}
threadfence_system();
__syncthreads();
@@ -278,7 +278,7 @@ __device__ void IPCContext::internal_direct_allreduce(
* [02+12+22+32] [02+12+22+32] [02+12+22+32] [02+12+22+32]
* [03+13+23+33] [03+13+23+33] [03+13+23+33] [03+13+23+33]
*/
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__device__ void IPCContext::internal_ring_allreduce(
T *dst, const T *src, int nelems, IPCTeam *team_obj, // NOLINT(runtime/int)
int n_seg, int seg_size, int chunk_size) {
@@ -323,7 +323,7 @@ __device__ void IPCContext::internal_ring_allreduce(
#if defined(__gfx90a__)
__threadfence_system();
#endif /* __gfx90a__ */
wait_until(&pSync[iter], ROC_SHMEM_CMP_EQ, wait_val);
wait_until(&pSync[iter], ROCSHMEM_CMP_EQ, wait_val);
}
__syncthreads();
compute_reduce<T, Op>(&pWrk[off_recv], &dst[off_seg + off_recv],
@@ -344,7 +344,7 @@ __device__ void IPCContext::internal_ring_allreduce(
#if defined(__gfx90a__)
__threadfence_system();
#endif /* __gfx90a__ */
wait_until(&pSync[iter], ROC_SHMEM_CMP_EQ, wait_val);
wait_until(&pSync[iter], ROCSHMEM_CMP_EQ, wait_val);
}
__syncthreads();
}
@@ -352,13 +352,13 @@ __device__ void IPCContext::internal_ring_allreduce(
__syncthreads();
for (size_t i = wg_id; i < 2 * num_pes - 2; i += wg_size) {
pSync[i] = ROC_SHMEM_SYNC_VALUE;
pSync[i] = ROCSHMEM_SYNC_VALUE;
}
__syncthreads();
}
template <typename T, ROC_SHMEM_OP Op>
__device__ int IPCContext::reduce(roc_shmem_team_t team, T *dest,
template <typename T, ROCSHMEM_OP Op>
__device__ int IPCContext::reduce(rocshmem_team_t team, T *dest,
const T *source, int nreduce) {
IPCTeam *team_obj = reinterpret_cast<IPCTeam *>(team);
@@ -367,14 +367,14 @@ __device__ int IPCContext::reduce(roc_shmem_team_t team, T *dest,
size_t direct_pWrk = PE_size * nreduce;
size_t direct_pSync = PE_size;
size_t ring_pSync = 2 * PE_size;
size_t provided_pWrk = max(nreduce / 2 + 1, ROC_SHMEM_REDUCE_MIN_WRKDATA_SIZE);
size_t provided_pSync = ROC_SHMEM_REDUCE_SYNC_SIZE;
size_t provided_pWrk = max(nreduce / 2 + 1, ROCSHMEM_REDUCE_MIN_WRKDATA_SIZE);
size_t provided_pSync = ROCSHMEM_REDUCE_SYNC_SIZE;
if (provided_pWrk >= direct_pWrk && provided_pSync >= direct_pSync) {
internal_direct_allreduce<T, Op>(dest, source, nreduce, team_obj);
} else {
if (ring_pSync <= ROC_SHMEM_REDUCE_SYNC_SIZE) {
size_t ring_pWrk = ROC_SHMEM_REDUCE_MIN_WRKDATA_SIZE;
if (ring_pSync <= ROCSHMEM_REDUCE_SYNC_SIZE) {
size_t ring_pWrk = ROCSHMEM_REDUCE_MIN_WRKDATA_SIZE;
// integer division truncating value
int chunk_size = ring_pWrk / PE_size;
int seg_size = chunk_size * PE_size;
@@ -410,10 +410,10 @@ __device__ int IPCContext::reduce(roc_shmem_team_t team, T *dest,
}
} else {
GPU_DPRINTF("Unsupported reduction size for IPC conduit.\n");
return ROC_SHMEM_ERROR;
return ROCSHMEM_ERROR;
}
}
return ROC_SHMEM_SUCCESS;
return ROCSHMEM_SUCCESS;
}
template <typename T>
@@ -439,7 +439,7 @@ __device__ void IPCContext::internal_get_broadcast(
}
template <typename T>
__device__ void IPCContext::broadcast(roc_shmem_team_t team, T *dst,
__device__ void IPCContext::broadcast(rocshmem_team_t team, T *dst,
const T *src, int nelems, int pe_root) {
IPCTeam *team_obj = reinterpret_cast<IPCTeam *>(team);
@@ -471,13 +471,13 @@ __device__ void IPCContext::internal_broadcast(T *dst, const T *src, int nelems,
}
template <typename T>
__device__ void IPCContext::alltoall(roc_shmem_team_t team, T *dst,
__device__ void IPCContext::alltoall(rocshmem_team_t team, T *dst,
const T *src, int nelems) {
alltoall_linear(team, dst, src, nelems);
}
template <typename T>
__device__ void IPCContext::alltoall_linear(roc_shmem_team_t team, T *dst,
__device__ void IPCContext::alltoall_linear(rocshmem_team_t team, T *dst,
const T *src, int nelems) {
IPCTeam *team_obj = reinterpret_cast<IPCTeam *>(team);
@@ -500,13 +500,13 @@ __device__ void IPCContext::alltoall_linear(roc_shmem_team_t team, T *dst,
}
template <typename T>
__device__ void IPCContext::fcollect(roc_shmem_team_t team, T *dst,
__device__ void IPCContext::fcollect(rocshmem_team_t team, T *dst,
const T *src, int nelems) {
fcollect_linear(team, dst, src, nelems);
}
template <typename T>
__device__ void IPCContext::fcollect_linear(roc_shmem_team_t team, T *dst,
__device__ void IPCContext::fcollect_linear(rocshmem_team_t team, T *dst,
const T *src, int nelems) {
IPCTeam *team_obj = reinterpret_cast<IPCTeam *>(team);
src/ipc/context_ipc_tmpl_host.hpp
+4 -4
View File
@@ -93,13 +93,13 @@ __host__ void IPCHostContext::broadcast(
}
template <typename T>
__host__ void IPCHostContext::broadcast(roc_shmem_team_t team, T *dest,
__host__ void IPCHostContext::broadcast(rocshmem_team_t team, T *dest,
const T *source, int nelems,
int pe_root) {
host_interface->broadcast<T>(team, dest, source, nelems, pe_root);
}
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__host__ void IPCHostContext::to_all(T *dest, const T *source, int nreduce,
int pe_start, int log_pe_stride,
int pe_size, T *p_wrk,
@@ -108,8 +108,8 @@ __host__ void IPCHostContext::to_all(T *dest, const T *source, int nreduce,
pe_size, p_wrk, p_sync);
}
template <typename T, ROC_SHMEM_OP Op>
__host__ int IPCHostContext::reduce(roc_shmem_team_t team, T *dest,
template <typename T, ROCSHMEM_OP Op>
__host__ int IPCHostContext::reduce(rocshmem_team_t team, T *dest,
const T *source, int nreduce) {
return host_interface->reduce<T, Op>(team, dest, source, nreduce);
}
src/ipc/ipc_context_proxy.hpp
+2 -2
View File
@@ -46,7 +46,7 @@ class IPCDefaultContextProxy {
auto ctx{proxy_.get()};
new (ctx) IPCContext(reinterpret_cast<Backend*>(backend));
ctx->tinfo = tinfo;
roc_shmem_ctx_t local{ctx, tinfo};
rocshmem_ctx_t local{ctx, tinfo};
set_internal_ctx(&local);
}
@@ -89,4 +89,4 @@ using IPCDefaultContextProxyT = IPCDefaultContextProxy<HIPAllocator>;
} // namespace rocshmem
#endif // LIBRARY_SRC_IPC_CONTEXT_PROXY_HPP_
#endif // LIBRARY_SRC_IPC_CONTEXT_PROXY_HPP_
src/ipc/ipc_team.cpp
+6 -6
View File
@@ -38,17 +38,17 @@ IPCTeam::IPCTeam(Backend *backend, TeamInfo *team_info_parent,
pool_index_ = pool_index;
barrier_pSync =
&(b->barrier_pSync_pool[pool_index * ROC_SHMEM_BARRIER_SYNC_SIZE]);
&(b->barrier_pSync_pool[pool_index * ROCSHMEM_BARRIER_SYNC_SIZE]);
reduce_pSync =
&(b->reduce_pSync_pool[pool_index * ROC_SHMEM_REDUCE_SYNC_SIZE]);
bcast_pSync = &(b->bcast_pSync_pool[pool_index * ROC_SHMEM_BCAST_SYNC_SIZE]);
&(b->reduce_pSync_pool[pool_index * ROCSHMEM_REDUCE_SYNC_SIZE]);
bcast_pSync = &(b->bcast_pSync_pool[pool_index * ROCSHMEM_BCAST_SYNC_SIZE]);
alltoall_pSync =
&(b->alltoall_pSync_pool[pool_index * ROC_SHMEM_ALLTOALL_SYNC_SIZE]);
&(b->alltoall_pSync_pool[pool_index * ROCSHMEM_ALLTOALL_SYNC_SIZE]);
pWrk = reinterpret_cast<char *>(b->pWrk_pool) +
ROC_SHMEM_REDUCE_MIN_WRKDATA_SIZE * sizeof(double) * pool_index;
ROCSHMEM_REDUCE_MIN_WRKDATA_SIZE * sizeof(double) * pool_index;
pAta = reinterpret_cast<char *>(b->pAta_pool) +
ROC_SHMEM_ATA_MAX_WRKDATA_SIZE * sizeof(double) * pool_index;
ROCSHMEM_ATA_MAX_WRKDATA_SIZE * sizeof(double) * pool_index;
}
IPCTeam::~IPCTeam() {}
src/ipc_policy.hpp
+1 -1
View File
@@ -123,7 +123,7 @@ class IpcOffImpl {
__device__ bool isIpcAvailable(int my_pe, int target_pe) { return false; }
__device__ void ipcGpuInit(Backend *roc_shmem_handle, Context *ctx,
__device__ void ipcGpuInit(Backend *rocshmem_handle, Context *ctx,
int thread_id) {}
__device__ void ipcCopy(void *dst, void *src, size_t size) {}
src/memory/single_heap.cpp
+1 -1
View File
@@ -27,7 +27,7 @@
namespace rocshmem {
SingleHeap::SingleHeap() {
if (auto heap_size_cstr = getenv("ROC_SHMEM_HEAP_SIZE")) {
if (auto heap_size_cstr = getenv("ROCSHMEM_HEAP_SIZE")) {
std::stringstream sstream(heap_size_cstr);
size_t heap_size;
sstream >> heap_size;
src/memory/slab_heap.cpp
+1 -1
View File
@@ -29,7 +29,7 @@
namespace rocshmem {
SlabHeap::SlabHeap() {
if (auto slab_size_cstr = getenv("ROC_SHMEM_SLAB_SIZE")) {
if (auto slab_size_cstr = getenv("ROCSHMEM_SLAB_SIZE")) {
std::stringstream sstream(slab_size_cstr);
size_t slab_size;
sstream >> slab_size;
src/mpi_init_singleton.hpp
+1 -1
View File
@@ -81,7 +81,7 @@ class MPIInitSingleton {
int nprocs_{-1};
/**
* @brief Was MPI initialized before ROCSHMEM_init call
* @brief Was MPI initialized before rocshmem_init call
*/
int pre_init_done{0};
src/reverse_offload/backend_ro.cpp
+10 -10
View File
@@ -32,7 +32,7 @@
#include <memory>
#include <thread> // NOLINT
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#include "../atomic_return.hpp"
#include "../backend_type.hpp"
#include "../context_incl.hpp"
@@ -42,13 +42,13 @@
namespace rocshmem {
extern roc_shmem_ctx_t ROC_SHMEM_HOST_CTX_DEFAULT;
extern rocshmem_ctx_t ROCSHMEM_HOST_CTX_DEFAULT;
ROBackend::ROBackend(MPI_Comm comm)
: profiler_proxy_(MAX_NUM_BLOCKS), Backend() {
type = BackendType::RO_BACKEND;
if (auto maximum_num_contexts_str = getenv("ROC_SHMEM_MAX_NUM_CONTEXTS")) {
if (auto maximum_num_contexts_str = getenv("ROCSHMEM_MAX_NUM_CONTEXTS")) {
std::stringstream sstream(maximum_num_contexts_str);
sstream >> maximum_num_contexts_;
}
@@ -83,14 +83,14 @@ ROBackend::ROBackend(MPI_Comm comm)
default_host_ctx = std::make_unique<ROHostContext>(this, 0);
ROC_SHMEM_HOST_CTX_DEFAULT.ctx_opaque = default_host_ctx.get();
ROCSHMEM_HOST_CTX_DEFAULT.ctx_opaque = default_host_ctx.get();
team_world_proxy_ = new ROTeamProxy<HIPAllocator>(
this, transport_->get_world_comm(), my_pe, num_pes);
team_tracker.set_team_world(team_world_proxy_->get());
ROC_SHMEM_TEAM_WORLD =
reinterpret_cast<roc_shmem_team_t>(team_world_proxy_->get());
ROCSHMEM_TEAM_WORLD =
reinterpret_cast<rocshmem_team_t>(team_world_proxy_->get());
default_block_handle_proxy_ = DefaultBlockHandleProxyT(
bp->g_ret, bp->atomic_ret, &queue_, &ipcImpl, hdp_proxy_.get());
@@ -120,7 +120,7 @@ ROBackend::~ROBackend() {
CHECK_HIP(hipFree(ctx_array));
}
__device__ bool ROBackend::create_ctx(int64_t options, roc_shmem_ctx_t *ctx) {
__device__ bool ROBackend::create_ctx(int64_t options, rocshmem_ctx_t *ctx) {
ROContext *ctx_;
auto pop_result = ctx_free_list.get()->pop_front();
@@ -133,11 +133,11 @@ __device__ bool ROBackend::create_ctx(int64_t options, roc_shmem_ctx_t *ctx) {
return true;
}
__device__ void ROBackend::destroy_ctx(roc_shmem_ctx_t *ctx) {
__device__ void ROBackend::destroy_ctx(rocshmem_ctx_t *ctx) {
ctx_free_list.get()->push_back(static_cast<ROContext *>(ctx->ctx_opaque));
}
void ROBackend::team_destroy(roc_shmem_team_t team) {
void ROBackend::team_destroy(rocshmem_team_t team) {
ROTeam *team_obj{get_internal_ro_team(team)};
team_obj->~ROTeam();
@@ -148,7 +148,7 @@ void ROBackend::create_new_team(Team *parent_team,
TeamInfo *team_info_wrt_parent,
TeamInfo *team_info_wrt_world, int num_pes,
int my_pe_in_new_team, MPI_Comm team_comm,
roc_shmem_team_t *new_team) {
rocshmem_team_t *new_team) {
transport_->createNewTeam(this, parent_team, team_info_wrt_parent,
team_info_wrt_world, num_pes, my_pe_in_new_team,
team_comm, new_team);
src/reverse_offload/backend_ro.hpp
+6 -6
View File
@@ -85,20 +85,20 @@ class ROBackend : public Backend {
void create_new_team(Team *parent_team, TeamInfo *team_info_wrt_parent,
TeamInfo *team_info_wrt_world, int num_pes,
int my_pe_in_new_team, MPI_Comm team_comm,
roc_shmem_team_t *new_team) override;
rocshmem_team_t *new_team) override;
/**
* @copydoc Backend::team_destroy(roc_shmem_team_t)
* @copydoc Backend::team_destroy(rocshmem_team_t)
*/
void team_destroy(roc_shmem_team_t team) override;
void team_destroy(rocshmem_team_t team) override;
__device__ bool create_ctx(int64_t options, roc_shmem_ctx_t *ctx);
__device__ bool create_ctx(int64_t options, rocshmem_ctx_t *ctx);
/**
* @brief Destroy a `roc_shmem_ctx_t` context and returns it back to the
* @brief Destroy a `rocshmem_ctx_t` context and returns it back to the
* context free list.
*/
__device__ void destroy_ctx(roc_shmem_ctx_t *ctx);
__device__ void destroy_ctx(rocshmem_ctx_t *ctx);
/**
* @copydoc Backend::ctx_create
src/reverse_offload/context_proxy.hpp
+2 -2
View File
@@ -23,7 +23,7 @@
#ifndef LIBRARY_SRC_REVERSE_OFFLOAD_CONTEXT_PROXY_HPP_
#define LIBRARY_SRC_REVERSE_OFFLOAD_CONTEXT_PROXY_HPP_
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#include "../device_proxy.hpp"
#include "../memory/hip_allocator.hpp"
#include "context_ro_device.hpp"
@@ -46,7 +46,7 @@ class DefaultContextProxy {
: constructed_{true} {
auto ctx{proxy_.get()};
new (ctx) ROContext(reinterpret_cast<Backend*>(backend), -1);
roc_shmem_ctx_t local{ctx, tinfo};
rocshmem_ctx_t local{ctx, tinfo};
set_internal_ctx(&local);
}
src/reverse_offload/context_ro_device.cpp
+3 -3
View File
@@ -30,7 +30,7 @@
#include <cstdlib>
#include "config.h" // NOLINT(build/include_subdir)
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#include "../backend_type.hpp"
#include "../hdp_policy.hpp"
#include "backend_proxy.hpp"
@@ -176,7 +176,7 @@ __device__ void ROContext::sync_all() {
__syncthreads();
}
__device__ void ROContext::sync(roc_shmem_team_t team) {
__device__ void ROContext::sync(rocshmem_team_t team) {
ROTeam *team_obj = reinterpret_cast<ROTeam *>(team);
if (is_thread_zero_in_block()) {
build_queue_element(RO_NET_SYNC, nullptr, nullptr, 0, 0, 0, 0, 0, nullptr,
@@ -472,7 +472,7 @@ __device__ void build_queue_element(
ro_net_cmds type, void *dst, void *src, size_t size, int pe,
int logPE_stride, int PE_size, int PE_root, void *pWrk, long *pSync,
MPI_Comm team_comm, int ro_net_win_id, BlockHandle *handle,
bool blocking, ROC_SHMEM_OP op, ro_net_types datatype) {
bool blocking, ROCSHMEM_OP op, ro_net_types datatype) {
auto write_slot{next_write_slot(handle)};
auto queue_element = &handle->queue[write_slot];
src/reverse_offload/context_ro_device.hpp
+16 -16
View File
@@ -34,7 +34,7 @@ __device__ void build_queue_element(
ro_net_cmds type, void *dst, void *src, size_t size, int pe,
int logPE_stride, int PE_size, int PE_root, void *pWrk, long *pSync,
MPI_Comm team_comm, int ro_net_win_id, BlockHandle *handle,
bool blocking, ROC_SHMEM_OP op = ROC_SHMEM_SUM,
bool blocking, ROCSHMEM_OP op = ROCSHMEM_SUM,
ro_net_types datatype = RO_NET_INT);
class ROContext : public Context {
@@ -67,7 +67,7 @@ class ROContext : public Context {
__device__ void sync_all();
__device__ void sync(roc_shmem_team_t team);
__device__ void sync(rocshmem_team_t team);
template <typename T>
__device__ void p(T *dest, T value, int pe);
@@ -75,13 +75,13 @@ class ROContext : public Context {
template <typename T>
__device__ T g(const T *source, int pe);
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__device__ void to_all(T *dest, const T *source, int nreduce, int PE_start,
int logPE_stride, int PE_size, T *pWrk,
long *pSync); // NOLINT(runtime/int)
template <typename T, ROC_SHMEM_OP Op>
__device__ void to_all(roc_shmem_team_t team, T *dest, const T *source,
template <typename T, ROCSHMEM_OP Op>
__device__ void to_all(rocshmem_team_t team, T *dest, const T *source,
int nreduce);
template <typename T>
@@ -133,7 +133,7 @@ class ROContext : public Context {
__device__ void amo_xor(void *dst, T value, int pe);
template <typename T>
__device__ void broadcast(roc_shmem_team_t team, T *dest, const T *source,
__device__ void broadcast(rocshmem_team_t team, T *dest, const T *source,
int nelems, int pe_root);
template <typename T>
@@ -142,43 +142,43 @@ class ROContext : public Context {
long *p_sync); // NOLINT(runtime/int)
template <typename T>
__device__ void alltoall(roc_shmem_team_t team, T *dest, const T *source,
__device__ void alltoall(rocshmem_team_t team, T *dest, const T *source,
int nelems);
template <typename T>
__device__ void alltoall_broadcast(roc_shmem_team_t team, T *dest,
__device__ void alltoall_broadcast(rocshmem_team_t team, T *dest,
const T *source, int nelems);
template <typename T>
__device__ void alltoall_mpi(roc_shmem_team_t team, T *dest, const T *source,
__device__ void alltoall_mpi(rocshmem_team_t team, T *dest, const T *source,
int nelems);
template <typename T>
__device__ void alltoall_gcen(roc_shmem_team_t team, T *dest, const T *source,
__device__ void alltoall_gcen(rocshmem_team_t team, T *dest, const T *source,
int nelems);
template <typename T>
__device__ void alltoall_gcen2(roc_shmem_team_t team, T *dest,
__device__ void alltoall_gcen2(rocshmem_team_t team, T *dest,
const T *source, int nelems);
template <typename T>
__device__ void fcollect(roc_shmem_team_t team, T *dest, const T *source,
__device__ void fcollect(rocshmem_team_t team, T *dest, const T *source,
int nelems);
template <typename T>
__device__ void fcollect_broadcast(roc_shmem_team_t team, T *dest,
__device__ void fcollect_broadcast(rocshmem_team_t team, T *dest,
const T *source, int nelems);
template <typename T>
__device__ void fcollect_mpi(roc_shmem_team_t team, T *dest, const T *source,
__device__ void fcollect_mpi(rocshmem_team_t team, T *dest, const T *source,
int nelems);
template <typename T>
__device__ void fcollect_gcen(roc_shmem_team_t team, T *dest, const T *source,
__device__ void fcollect_gcen(rocshmem_team_t team, T *dest, const T *source,
int nelems);
template <typename T>
__device__ void fcollect_gcen2(roc_shmem_team_t team, T *dest,
__device__ void fcollect_gcen2(rocshmem_team_t team, T *dest,
const T *source, int nelems);
__device__ void putmem_wg(void *dest, const void *source, size_t nelems,
src/reverse_offload/context_ro_host.hpp
+4 -4
View File
@@ -135,16 +135,16 @@ class ROHostContext : public Context {
long *p_sync); // NOLINT(runtime/int)
template <typename T>
__host__ void broadcast(roc_shmem_team_t team, T *dest, const T *source,
__host__ void broadcast(rocshmem_team_t team, T *dest, const T *source,
int nelems, int pe_root);
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__host__ void to_all(T *dest, const T *source, int nreduce, int pe_start,
int log_pe_stride, int pe_size, T *p_wrk,
long *p_sync); // NOLINT(runtime/int)
template <typename T, ROC_SHMEM_OP Op>
__host__ void to_all(roc_shmem_team_t team, T *dest, const T *source,
template <typename T, ROCSHMEM_OP Op>
__host__ void to_all(rocshmem_team_t team, T *dest, const T *source,
int nreduce);
template <typename T>
src/reverse_offload/context_ro_tmpl_device.hpp
+16 -16
View File
@@ -108,8 +108,8 @@ struct GetROType<long double> {
********************************* DEVICE API *********************************
*****************************************************************************/
template <typename T, ROC_SHMEM_OP Op>
__device__ void ROContext::to_all(roc_shmem_team_t team, T *dest,
template <typename T, ROCSHMEM_OP Op>
__device__ void ROContext::to_all(rocshmem_team_t team, T *dest,
const T *source, int nreduce) {
if (!is_thread_zero_in_block()) {
__syncthreads();
@@ -125,7 +125,7 @@ __device__ void ROContext::to_all(roc_shmem_team_t team, T *dest,
__syncthreads();
}
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__device__ void ROContext::to_all(T *dest, const T *source, int nreduce,
int PE_start, int logPE_stride, int PE_size,
T *pWrk, long *pSync) {
@@ -210,7 +210,7 @@ __device__ T ROContext::amo_fetch_cas(void *dst, T value, T cond, int pe) {
value, pe, 0, 0, 0,
reinterpret_cast<void *>(static_cast<long long>(cond)),
nullptr, (MPI_Comm)NULL, ro_net_win_id, block_handle, true,
ROC_SHMEM_SUM, GetROType<T>::Type);
ROCSHMEM_SUM, GetROType<T>::Type);
__threadfence();
return *source;
}
@@ -225,7 +225,7 @@ __device__ T ROContext::amo_fetch_add(void *dst, T value, int pe) {
auto source{get_unused_atomic()};
build_queue_element(RO_NET_AMO_FOP, dst, reinterpret_cast<T *>(source), value,
pe, 0, 0, 0, nullptr, nullptr, (MPI_Comm)NULL,
ro_net_win_id, block_handle, true, ROC_SHMEM_SUM,
ro_net_win_id, block_handle, true, ROCSHMEM_SUM,
GetROType<T>::Type);
__threadfence();
return *source;
@@ -241,7 +241,7 @@ __device__ T ROContext::amo_swap(void *dst, T value, int pe) {
auto source{get_unused_atomic()};
build_queue_element(RO_NET_AMO_FOP, dst, reinterpret_cast<void *>(source),
value, pe, 0, 0, 0, nullptr, nullptr, (MPI_Comm)NULL,
ro_net_win_id, block_handle, true, ROC_SHMEM_REPLACE,
ro_net_win_id, block_handle, true, ROCSHMEM_REPLACE,
GetROType<T>::Type);
__threadfence();
return *source;
@@ -257,7 +257,7 @@ __device__ T ROContext::amo_fetch_and(void *dst, T value, int pe) {
auto source{get_unused_atomic()};
build_queue_element(RO_NET_AMO_FOP, dst, reinterpret_cast<void *>(source),
value, pe, 0, 0, 0, nullptr, nullptr, (MPI_Comm)NULL,
ro_net_win_id, block_handle, true, ROC_SHMEM_AND,
ro_net_win_id, block_handle, true, ROCSHMEM_AND,
GetROType<T>::Type);
__threadfence();
return *source;
@@ -273,7 +273,7 @@ __device__ T ROContext::amo_fetch_or(void *dst, T value, int pe) {
auto source{get_unused_atomic()};
build_queue_element(RO_NET_AMO_FOP, dst, reinterpret_cast<void *>(source),
value, pe, 0, 0, 0, nullptr, nullptr, (MPI_Comm)NULL,
ro_net_win_id, block_handle, true, ROC_SHMEM_OR,
ro_net_win_id, block_handle, true, ROCSHMEM_OR,
GetROType<T>::Type);
__threadfence();
return *source;
@@ -289,7 +289,7 @@ __device__ T ROContext::amo_fetch_xor(void *dst, T value, int pe) {
auto source{get_unused_atomic()};
build_queue_element(RO_NET_AMO_FOP, dst, reinterpret_cast<void *>(source),
value, pe, 0, 0, 0, nullptr, nullptr, (MPI_Comm)NULL,
ro_net_win_id, block_handle, true, ROC_SHMEM_XOR,
ro_net_win_id, block_handle, true, ROCSHMEM_XOR,
GetROType<T>::Type);
__threadfence();
return *source;
@@ -301,7 +301,7 @@ __device__ void ROContext::amo_xor(void *dst, T value, int pe) {
}
template <typename T>
__device__ void ROContext::broadcast(roc_shmem_team_t team, T *dest,
__device__ void ROContext::broadcast(rocshmem_team_t team, T *dest,
const T *source, int nelems, int pe_root) {
if (!is_thread_zero_in_block()) {
__syncthreads();
@@ -313,7 +313,7 @@ __device__ void ROContext::broadcast(roc_shmem_team_t team, T *dest,
build_queue_element(RO_NET_TEAM_BROADCAST, dest, const_cast<T *>(source),
nelems, 0, 0, 0, pe_root, nullptr, nullptr,
team_obj->mpi_comm, ro_net_win_id, block_handle, true,
ROC_SHMEM_SUM, GetROType<T>::Type);
ROCSHMEM_SUM, GetROType<T>::Type);
__syncthreads();
}
@@ -331,13 +331,13 @@ __device__ void ROContext::broadcast(T *dest, const T *source, int nelems,
build_queue_element(RO_NET_BROADCAST, dest, const_cast<T *>(source), nelems,
pe_start, log_pe_stride, pe_size, pe_root, nullptr,
p_sync, (MPI_Comm)NULL, ro_net_win_id, block_handle, true,
ROC_SHMEM_SUM, GetROType<T>::Type);
ROCSHMEM_SUM, GetROType<T>::Type);
__syncthreads();
}
template <typename T>
__device__ void ROContext::alltoall(roc_shmem_team_t team, T *dest,
__device__ void ROContext::alltoall(rocshmem_team_t team, T *dest,
const T *source, int nelems) {
if (!is_thread_zero_in_block()) {
__syncthreads();
@@ -349,13 +349,13 @@ __device__ void ROContext::alltoall(roc_shmem_team_t team, T *dest,
build_queue_element(RO_NET_ALLTOALL, dest, const_cast<T *>(source), nelems, 0,
0, 0, 0, team_obj->ata_buffer, nullptr,
team_obj->mpi_comm, ro_net_win_id, block_handle, true,
ROC_SHMEM_SUM, GetROType<T>::Type);
ROCSHMEM_SUM, GetROType<T>::Type);
__syncthreads();
}
template <typename T>
__device__ void ROContext::fcollect(roc_shmem_team_t team, T *dest,
__device__ void ROContext::fcollect(rocshmem_team_t team, T *dest,
const T *source, int nelems) {
if (!is_thread_zero_in_block()) {
__syncthreads();
@@ -367,7 +367,7 @@ __device__ void ROContext::fcollect(roc_shmem_team_t team, T *dest,
build_queue_element(RO_NET_FCOLLECT, dest, const_cast<T *>(source), nelems, 0,
0, 0, 0, team_obj->ata_buffer, nullptr,
team_obj->mpi_comm, ro_net_win_id, block_handle, true,
ROC_SHMEM_SUM, GetROType<T>::Type);
ROCSHMEM_SUM, GetROType<T>::Type);
__syncthreads();
}
src/reverse_offload/context_ro_tmpl_host.hpp
+4 -4
View File
@@ -114,7 +114,7 @@ __host__ void ROHostContext::broadcast(T *dest, const T *source, int nelems,
}
template <typename T>
__host__ void ROHostContext::broadcast(roc_shmem_team_t team, T *dest,
__host__ void ROHostContext::broadcast(rocshmem_team_t team, T *dest,
const T *source, int nelems,
int pe_root) {
DPRINTF("Function: Team-based ro_net_host_broadcast\n");
@@ -122,7 +122,7 @@ __host__ void ROHostContext::broadcast(roc_shmem_team_t team, T *dest,
host_interface->broadcast<T>(team, dest, source, nelems, pe_root);
}
template <typename T, ROC_SHMEM_OP Op>
template <typename T, ROCSHMEM_OP Op>
__host__ void ROHostContext::to_all(T *dest, const T *source, int nreduce,
int pe_start, int log_pe_stride,
int pe_size, T *p_wrk, long *p_sync) {
@@ -132,8 +132,8 @@ __host__ void ROHostContext::to_all(T *dest, const T *source, int nreduce,
pe_size, p_wrk, p_sync);
}
template <typename T, ROC_SHMEM_OP Op>
__host__ void ROHostContext::to_all(roc_shmem_team_t team, T *dest,
template <typename T, ROCSHMEM_OP Op>
__host__ void ROHostContext::to_all(rocshmem_team_t team, T *dest,
const T *source, int nreduce) {
DPRINTF("Function: Team-based ro_net_host_to_all\n");
src/reverse_offload/mpi_transport.cpp
+20 -20
View File
@@ -142,7 +142,7 @@ void MPITransport::submitRequestsToMPI() {
const_cast<unsigned long long *>(&next_element.ol1.atomic_value),
next_element.PE, next_element.ro_net_win_id, queue_idx,
next_element.threadId, true,
static_cast<ROC_SHMEM_OP>(next_element.op),
static_cast<ROCSHMEM_OP>(next_element.op),
static_cast<ro_net_types>(next_element.datatype));
DPRINTF("Received AMO dst %p src %p Val %llu pe %d\n", next_element.dst,
next_element.src, next_element.ol1.atomic_value, next_element.PE);
@@ -163,7 +163,7 @@ void MPITransport::submitRequestsToMPI() {
team_reduction(next_element.dst, next_element.src, next_element.ol1.size,
next_element.ro_net_win_id, queue_idx,
next_element.team_comm,
static_cast<ROC_SHMEM_OP>(next_element.op),
static_cast<ROCSHMEM_OP>(next_element.op),
static_cast<ro_net_types>(next_element.datatype),
next_element.threadId, true);
DPRINTF("Received FLOAT_SUM_TEAM_TO_ALL dst %p src %p size %lu team %d\n",
@@ -175,7 +175,7 @@ void MPITransport::submitRequestsToMPI() {
next_element.PE, next_element.ro_net_win_id, queue_idx,
next_element.PE, next_element.logPE_stride,
next_element.PE_size, next_element.ol2.pWrk, next_element.pSync,
static_cast<ROC_SHMEM_OP>(next_element.op),
static_cast<ROCSHMEM_OP>(next_element.op),
static_cast<ro_net_types>(next_element.datatype),
next_element.threadId, true);
DPRINTF(
@@ -275,15 +275,15 @@ void MPITransport::finalizeTransport() {
delete host_interface;
}
roc_shmem_team_t get_external_team(ROTeam *team) {
return reinterpret_cast<roc_shmem_team_t>(team);
rocshmem_team_t get_external_team(ROTeam *team) {
return reinterpret_cast<rocshmem_team_t>(team);
}
void MPITransport::createNewTeam(ROBackend *backend, Team *parent_team,
TeamInfo *team_info_wrt_parent,
TeamInfo *team_info_wrt_world, int num_pes,
int my_pe_in_new_team, MPI_Comm team_comm,
roc_shmem_team_t *new_team) {
rocshmem_team_t *new_team) {
ROTeam *new_team_obj{nullptr};
CHECK_HIP(hipMalloc(&new_team_obj, sizeof(ROTeam)));
@@ -342,26 +342,26 @@ void MPITransport::barrier(int blockId, int threadId, bool blocking,
outstanding[blockId]++;
}
MPI_Op MPITransport::get_mpi_op(ROC_SHMEM_OP op) {
MPI_Op MPITransport::get_mpi_op(ROCSHMEM_OP op) {
switch (op) {
case ROC_SHMEM_SUM:
case ROCSHMEM_SUM:
return MPI_SUM;
case ROC_SHMEM_MAX:
case ROCSHMEM_MAX:
return MPI_MAX;
case ROC_SHMEM_MIN:
case ROCSHMEM_MIN:
return MPI_MIN;
case ROC_SHMEM_PROD:
case ROCSHMEM_PROD:
return MPI_PROD;
case ROC_SHMEM_AND:
case ROCSHMEM_AND:
return MPI_BAND;
case ROC_SHMEM_OR:
case ROCSHMEM_OR:
return MPI_BOR;
case ROC_SHMEM_XOR:
case ROCSHMEM_XOR:
return MPI_BXOR;
case ROC_SHMEM_REPLACE:
case ROCSHMEM_REPLACE:
return MPI_REPLACE;
default:
fprintf(stderr, "Unknown ROC_SHMEM op MPI conversion %d\n", op);
fprintf(stderr, "Unknown rocSHMEM op MPI conversion %d\n", op);
abort();
}
}
@@ -383,7 +383,7 @@ static MPI_Datatype convertType(ro_net_types type) {
case RO_NET_LONG_DOUBLE:
return MPI_LONG_DOUBLE;
default:
fprintf(stderr, "Unknown ROC_SHMEM type MPI conversion %d\n", type);
fprintf(stderr, "Unknown rocSHMEM type MPI conversion %d\n", type);
abort();
}
}
@@ -391,7 +391,7 @@ static MPI_Datatype convertType(ro_net_types type) {
void MPITransport::reduction(void *dst, void *src, int size, int pe,
int win_id, int blockId, int start, int logPstride,
int sizePE, void *pWrk, long *pSync,
ROC_SHMEM_OP op, ro_net_types type, int threadId,
ROCSHMEM_OP op, ro_net_types type, int threadId,
bool blocking) {
MPI_Request request{};
MPI_Op mpi_op{get_mpi_op(op)};
@@ -435,7 +435,7 @@ void MPITransport::broadcast(void *dst, void *src, int size, int pe,
}
void MPITransport::team_reduction(void *dst, void *src, int size, int win_id,
int blockId, MPI_Comm team, ROC_SHMEM_OP op,
int blockId, MPI_Comm team, ROCSHMEM_OP op,
ro_net_types type, int threadId,
bool blocking) {
MPI_Request request{};
@@ -1046,7 +1046,7 @@ void MPITransport::putMem(void *dst, void *src, int size, int pe, int win_id,
void MPITransport::amoFOP(void *dst, void *src, void *val, int pe, int win_id,
int blockId, int threadId, bool blocking,
ROC_SHMEM_OP op, ro_net_types type) {
ROCSHMEM_OP op, ro_net_types type) {
queue->flush_hdp();
auto *bp{backend_proxy->get()};
src/reverse_offload/mpi_transport.hpp
+5 -5
View File
@@ -49,18 +49,18 @@ class MPITransport : public Transport {
TeamInfo *team_info_wrt_parent,
TeamInfo *team_info_wrt_world, int num_pes,
int my_pe_in_new_team, MPI_Comm team_comm,
roc_shmem_team_t *new_team) override;
rocshmem_team_t *new_team) override;
void barrier(int blockId, int threadId, bool blocking,
MPI_Comm team) override;
void reduction(void *dst, void *src, int size, int pe, int win_id,
int blockId, int start, int logPstride, int sizePE, void *pWrk,
long *pSync, ROC_SHMEM_OP op, ro_net_types type,
long *pSync, ROCSHMEM_OP op, ro_net_types type,
int threadId, bool blocking) override;
void team_reduction(void *dst, void *src, int size, int win_id, int blockId,
MPI_Comm team, ROC_SHMEM_OP op, ro_net_types type,
MPI_Comm team, ROCSHMEM_OP op, ro_net_types type,
int threadId, bool blocking) override;
void broadcast(void *dst, void *src, int size, int pe, int win_id,
@@ -116,7 +116,7 @@ class MPITransport : public Transport {
int threadId, bool blocking, bool inline_data = false) override;
void amoFOP(void *dst, void *src, void *val, int pe, int win_id, int blockId,
int threadId, bool blocking, ROC_SHMEM_OP op,
int threadId, bool blocking, ROCSHMEM_OP op,
ro_net_types type) override;
void amoFCAS(void *dst, void *src, void *val, int pe, int win_id, int blockId,
@@ -193,7 +193,7 @@ class MPITransport : public Transport {
void submitRequestsToMPI();
MPI_Op get_mpi_op(ROC_SHMEM_OP op);
MPI_Op get_mpi_op(ROCSHMEM_OP op);
Queue *queue{nullptr};
src/reverse_offload/transport.hpp
+5 -5
View File
@@ -27,7 +27,7 @@
#include <cassert>
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#include "backend_proxy.hpp"
#include "ro_net_team.hpp"
@@ -48,18 +48,18 @@ class Transport {
TeamInfo *team_info_wrt_parent,
TeamInfo *team_info_wrt_world, int num_pes,
int my_pe_in_new_team, MPI_Comm team_comm,
roc_shmem_team_t *new_team) = 0;
rocshmem_team_t *new_team) = 0;
virtual void barrier(int wg_id, int threadId, bool blocking,
MPI_Comm team) = 0;
virtual void reduction(void *dst, void *src, int size, int pe, int win_id,
int wg_id, int start, int logPstride, int sizePE,
void *pWrk, long *pSync, ROC_SHMEM_OP op,
void *pWrk, long *pSync, ROCSHMEM_OP op,
ro_net_types type, int threadId, bool blocking) = 0;
virtual void team_reduction(void *dst, void *src, int size, int win_id,
int wg_id, MPI_Comm team, ROC_SHMEM_OP op,
int wg_id, MPI_Comm team, ROCSHMEM_OP op,
ro_net_types type, int threadId,
bool blocking) = 0;
@@ -89,7 +89,7 @@ class Transport {
int wg_id, int threadId, bool blocking) = 0;
virtual void amoFOP(void *dst, void *src, void *val, int pe, int win_id,
int wg_id, int threadId, bool blocking, ROC_SHMEM_OP op,
int wg_id, int threadId, bool blocking, ROCSHMEM_OP op,
ro_net_types type) = 0;
virtual void amoFCAS(void *dst, void *src, void *val, int pe, int win_id,
src/roc_shmem.cpp → src/rocshmem.cpp
+508 -508
View File
File diff suppressed because it is too large Load Diff
src/roc_shmem_calc.hpp → src/rocshmem_calc.hpp
+11 -11
View File
@@ -20,14 +20,14 @@
* IN THE SOFTWARE.
*****************************************************************************/
#ifndef LIBRARY_SRC_ROC_SHMEM_CALC_HPP_
#define LIBRARY_SRC_ROC_SHMEM_CALC_HPP_
#ifndef LIBRARY_SRC_ROCSHMEM_CALC_HPP_
#define LIBRARY_SRC_ROCSHMEM_CALC_HPP_
namespace rocshmem {
// clang-format off
NOWARN(-Wunused-parameter,
template <ROC_SHMEM_OP Op>
template <ROCSHMEM_OP Op>
struct OpWrap {
template <typename T>
__device__ static void Calc(T *src, T *dst, int i) {
@@ -41,7 +41,7 @@ struct OpWrap {
************************** TEMPLATE SPECIALIZATIONS **************************
*****************************************************************************/
template <>
struct OpWrap<ROC_SHMEM_SUM> {
struct OpWrap<ROCSHMEM_SUM> {
template <typename T>
__device__ static void Calc(T *src, T *dst, int i) {
dst[i] += src[i];
@@ -49,7 +49,7 @@ struct OpWrap<ROC_SHMEM_SUM> {
};
template <>
struct OpWrap<ROC_SHMEM_MAX> {
struct OpWrap<ROCSHMEM_MAX> {
template <typename T>
__device__ static void Calc(T *src, T *dst, int i) {
dst[i] = max(dst[i], src[i]);
@@ -57,7 +57,7 @@ struct OpWrap<ROC_SHMEM_MAX> {
};
template <>
struct OpWrap<ROC_SHMEM_MIN> {
struct OpWrap<ROCSHMEM_MIN> {
template <typename T>
__device__ static void Calc(T *src, T *dst, int i) {
dst[i] = min(dst[i], src[i]);
@@ -65,7 +65,7 @@ struct OpWrap<ROC_SHMEM_MIN> {
};
template <>
struct OpWrap<ROC_SHMEM_PROD> {
struct OpWrap<ROCSHMEM_PROD> {
template <typename T>
__device__ static void Calc(T *src, T *dst, int i) {
dst[i] *= src[i];
@@ -73,7 +73,7 @@ struct OpWrap<ROC_SHMEM_PROD> {
};
template <>
struct OpWrap<ROC_SHMEM_AND> {
struct OpWrap<ROCSHMEM_AND> {
template <typename T>
__device__ static void Calc(T *src, T *dst, int i) {
dst[i] &= src[i];
@@ -81,7 +81,7 @@ struct OpWrap<ROC_SHMEM_AND> {
};
template <>
struct OpWrap<ROC_SHMEM_OR> {
struct OpWrap<ROCSHMEM_OR> {
template <typename T>
__device__ static void Calc(T *src, T *dst, int i) {
dst[i] |= src[i];
@@ -89,7 +89,7 @@ struct OpWrap<ROC_SHMEM_OR> {
};
template <>
struct OpWrap<ROC_SHMEM_XOR> {
struct OpWrap<ROCSHMEM_XOR> {
template <typename T>
__device__ static void Calc(T *src, T *dst, int i) {
dst[i] ^= src[i];
@@ -97,4 +97,4 @@ struct OpWrap<ROC_SHMEM_XOR> {
};
}
#endif // LIBRARY_SRC_ROC_SHMEM_CALC_HPP_
#endif // LIBRARY_SRC_ROCSHMEM_CALC_HPP_
src/roc_shmem_gpu.cpp → src/rocshmem_gpu.cpp
+658 -658
View File
File diff suppressed because it is too large Load Diff
src/stats.hpp
+5 -5
View File
@@ -27,12 +27,12 @@
#include <atomic>
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#include "util.hpp"
namespace rocshmem {
enum roc_shmem_stats {
enum rocshmem_stats {
NUM_PUT = 0,
NUM_PUT_NBI,
NUM_P,
@@ -91,7 +91,7 @@ enum roc_shmem_stats {
NUM_STATS
};
enum roc_shmem_host_stats {
enum rocshmem_host_stats {
NUM_HOST_PUT = 0,
NUM_HOST_PUT_NBI,
NUM_HOST_P,
@@ -141,10 +141,10 @@ class Stats {
StatType stats[I] = {0};
public:
__device__ uint64_t startTimer() const { return roc_shmem_timer(); }
__device__ uint64_t startTimer() const { return rocshmem_timer(); }
__device__ void endTimer(uint64_t start, int index) {
incStat(index, roc_shmem_timer() - start);
incStat(index, rocshmem_timer() - start);
}
__device__ void incStat(int index, int value = 1) {
src/team.cpp
+10 -10
View File
@@ -24,34 +24,34 @@
#include <cmath>
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#include "backend_bc.hpp"
#include "util.hpp"
namespace rocshmem {
roc_shmem_team_t ROC_SHMEM_TEAM_WORLD;
rocshmem_team_t ROCSHMEM_TEAM_WORLD;
__host__ __device__ Team* get_internal_team(roc_shmem_team_t team) {
__host__ __device__ Team* get_internal_team(rocshmem_team_t team) {
return reinterpret_cast<Team*>(team);
}
GPUIBTeam* get_internal_gpu_ib_team(roc_shmem_team_t team) {
GPUIBTeam* get_internal_gpu_ib_team(rocshmem_team_t team) {
return reinterpret_cast<GPUIBTeam*>(team);
}
ROTeam* get_internal_ro_team(roc_shmem_team_t team) {
ROTeam* get_internal_ro_team(rocshmem_team_t team) {
return reinterpret_cast<ROTeam*>(team);
}
IPCTeam* get_internal_ipc_team(roc_shmem_team_t team) {
IPCTeam* get_internal_ipc_team(rocshmem_team_t team) {
return reinterpret_cast<IPCTeam*>(team);
}
__host__ __device__ int team_translate_pe(roc_shmem_team_t src_team, int src_pe,
roc_shmem_team_t dst_team) {
if (src_team == ROC_SHMEM_TEAM_INVALID ||
dst_team == ROC_SHMEM_TEAM_INVALID) {
__host__ __device__ int team_translate_pe(rocshmem_team_t src_team, int src_pe,
rocshmem_team_t dst_team) {
if (src_team == ROCSHMEM_TEAM_INVALID ||
dst_team == ROCSHMEM_TEAM_INVALID) {
return -1;
}
src/team.hpp
+7 -7
View File
@@ -25,7 +25,7 @@
#include <mpi.h>
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
#include "backend_type.hpp"
namespace rocshmem {
@@ -157,16 +157,16 @@ class Team {
BackendType type{BackendType::GPU_IB_BACKEND};
};
__host__ __device__ Team* get_internal_team(roc_shmem_team_t team);
__host__ __device__ Team* get_internal_team(rocshmem_team_t team);
GPUIBTeam* get_internal_gpu_ib_team(roc_shmem_team_t team);
GPUIBTeam* get_internal_gpu_ib_team(rocshmem_team_t team);
ROTeam* get_internal_ro_team(roc_shmem_team_t team);
ROTeam* get_internal_ro_team(rocshmem_team_t team);
IPCTeam* get_internal_ipc_team(roc_shmem_team_t team);
IPCTeam* get_internal_ipc_team(rocshmem_team_t team);
__host__ __device__ int team_translate_pe(roc_shmem_team_t src_team, int src_pe,
roc_shmem_team_t dst_team);
__host__ __device__ int team_translate_pe(rocshmem_team_t src_team, int src_pe,
rocshmem_team_t dst_team);
} // namespace rocshmem
src/team_tracker.cpp
+4 -4
View File
@@ -28,19 +28,19 @@ namespace rocshmem {
TeamTracker::TeamTracker() {
char* value{nullptr};
if ((value = getenv("ROC_SHMEM_MAX_NUM_TEAMS"))) {
if ((value = getenv("ROCSHMEM_MAX_NUM_TEAMS"))) {
max_num_teams_ = atoi(value);
}
}
void TeamTracker::track(roc_shmem_team_t team) {
if (team == ROC_SHMEM_TEAM_INVALID) {
void TeamTracker::track(rocshmem_team_t team) {
if (team == ROCSHMEM_TEAM_INVALID) {
return;
}
teams_.push_back(team);
}
void TeamTracker::untrack(roc_shmem_team_t team) {
void TeamTracker::untrack(rocshmem_team_t team) {
auto it{std::find(teams_.begin(), teams_.end(), team)};
assert(it != teams_.end());
teams_.erase(it);
src/team_tracker.hpp
+5 -5
View File
@@ -30,7 +30,7 @@
#include <vector>
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
namespace rocshmem {
@@ -55,7 +55,7 @@ class TeamTracker {
*
* @param void
*/
void track(roc_shmem_team_t team);
void track(rocshmem_team_t team);
/**
* @brief Remove team from the list of user-created teams
@@ -64,7 +64,7 @@ class TeamTracker {
*
* @return void
*/
void untrack(roc_shmem_team_t team);
void untrack(rocshmem_team_t team);
/**
* @brief Remove all teams from the list of user-created teams
@@ -113,7 +113,7 @@ class TeamTracker {
/**
* @brief List of teams created by the user.
*/
std::vector<roc_shmem_team_t> teams_{};
std::vector<rocshmem_team_t> teams_{};
/**
* @brief The maximum number of teams the user can create.
@@ -125,7 +125,7 @@ class TeamTracker {
int max_num_teams_{40};
/**
* @brief Pointer to implementation of ROC_SHMEM_TEAM_WORLD
* @brief Pointer to implementation of ROCSHMEM_TEAM_WORLD
*/
Team* team_world_{nullptr};
};
src/templates.hpp
+106 -107
View File
@@ -23,11 +23,11 @@
#ifndef LIBRARY_SRC_TEMPLATES_HPP_
#define LIBRARY_SRC_TEMPLATES_HPP_
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
/**
* @file templates.hpp
* @brief Internal header that declares templates for ROC_SHMEM's implentation
* @brief Internal header that declares templates for rocSHMEM's implentation
* of the user-facing device APIs.
*
* This file contains templates for the OpenSHMEM APIs that take have
@@ -35,8 +35,7 @@
*/
/******************************************************************************
**************************** DEVICE FUNCTIONS
***********************************
**************************** DEVICE FUNCTIONS ********************************
*****************************************************************************/
namespace rocshmem {
@@ -45,12 +44,12 @@ namespace rocshmem {
* @brief Writes contiguous data of \p nelems elements from \p source on the
* calling PE to \p dest at \p pe. The caller will block until the operation
* completes locally (it is safe to reuse \p source). The caller must
* call into roc_shmem_quiet() if remote completion is required.
* call into rocshmem_quiet() if remote completion is required.
*
* This function can be called from divergent control paths at per-thread
* granularity. However, performance may be improved if the caller can
* coalesce contiguous messages and elect a leader thread to call into the
* ROC_SHMEM function.
* rocSHMEM function.
*
* @param[in] ctx Context with which to perform this operation.
* @param[in] dest Destination address. Must be an address on the symmetric
@@ -63,21 +62,21 @@ namespace rocshmem {
*
*/
template <typename T>
__device__ void roc_shmem_put(roc_shmem_ctx_t ctx, T *dest, const T *source,
__device__ void rocshmem_put(rocshmem_ctx_t ctx, T *dest, const T *source,
size_t nelems, int pe);
template <typename T>
__device__ void roc_shmem_put(T *dest, const T *source, size_t nelems, int pe);
__device__ void rocshmem_put(T *dest, const T *source, size_t nelems, int pe);
/**
* @brief Writes a single value to \p dest at \p pe PE to \p dst at \p pe.
* The caller must call into roc_shmem_quiet() if remote completion is
* The caller must call into rocshmem_quiet() if remote completion is
* required.
*
* This function can be called from divergent control paths at per-thread
* granularity. However, performance may be improved if the caller can
* coalesce contiguous messages and elect a leader thread to call into the
* ROC_SHMEM function.
* rocSHMEM function.
*
* @param[in] ctx Context with which to perform this operation.
* @param[in] dest Destination address. Must be an address on the symmetric
@@ -89,10 +88,10 @@ __device__ void roc_shmem_put(T *dest, const T *source, size_t nelems, int pe);
*
*/
template <typename T>
__device__ void roc_shmem_p(roc_shmem_ctx_t ctx, T *dest, T value, int pe);
__device__ void rocshmem_p(rocshmem_ctx_t ctx, T *dest, T value, int pe);
template <typename T>
__device__ void roc_shmem_p(T *dest, T value, int pe);
__device__ void rocshmem_p(T *dest, T value, int pe);
/**
* @brief Reads contiguous data of \p nelems elements from \p source on \p pe
@@ -102,7 +101,7 @@ __device__ void roc_shmem_p(T *dest, T value, int pe);
* This function can be called from divergent control paths at per-thread
* granularity. However, performance may be improved if the caller can
* coalesce contiguous messages and elect a leader thread to call into the
* ROC_SHMEM function.
* rocSHMEM function.
*
* @param[in] ctx Context with which to perform this operation.
* @param[in] dest Destination address. Must be an address on the symmetric
@@ -115,11 +114,11 @@ __device__ void roc_shmem_p(T *dest, T value, int pe);
*
*/
template <typename T>
__device__ void roc_shmem_get(roc_shmem_ctx_t ctx, T *dest, const T *source,
__device__ void rocshmem_get(rocshmem_ctx_t ctx, T *dest, const T *source,
size_t nelems, int pe);
template <typename T>
__device__ void roc_shmem_get(T *dest, const T *source, size_t nelems, int pe);
__device__ void rocshmem_get(T *dest, const T *source, size_t nelems, int pe);
/**
* @brief reads and returns single value from \p source at \p pe.
@@ -128,7 +127,7 @@ __device__ void roc_shmem_get(T *dest, const T *source, size_t nelems, int pe);
* This function can be called from divergent control paths at per-thread
* granularity. However, performance may be improved if the caller can
* coalesce contiguous messages and elect a leader thread to call into the
* ROC_SHMEM function.
* rocSHMEM function.
*
* @param[in] ctx Context with which to perform this operation.
* @param[in] source sourcen address. Must be an address on the symmetric
@@ -139,22 +138,22 @@ __device__ void roc_shmem_get(T *dest, const T *source, size_t nelems, int pe);
*
*/
template <typename T>
__device__ T roc_shmem_g(roc_shmem_ctx_t ctx, const T *source, int pe);
__device__ T rocshmem_g(rocshmem_ctx_t ctx, const T *source, int pe);
template <typename T>
__device__ T roc_shmem_g(const T *source, int pe);
__device__ T rocshmem_g(const T *source, int pe);
/**
* @brief Writes contiguous data of \p nelems elements from \p source on the
* calling PE to \p dest on \p pe. The operation is not blocking. The caller
* will return as soon as the request is posted. The caller must call
* roc_shmem_quiet() on the same context if completion notification is
* rocshmem_quiet() on the same context if completion notification is
* required.
*
* This function can be called from divergent control paths at per-thread
* granularity. However, performance may be improved if the caller can
* coalesce contiguous messages and elect a leader thread to call into the
* ROC_SHMEM function.
* rocSHMEM function.
*
* @param[in] ctx Context with which to perform this operation.
* @param[in] dest Destination address. Must be an address on the symmetric
@@ -167,23 +166,23 @@ __device__ T roc_shmem_g(const T *source, int pe);
*
*/
template <typename T>
__device__ void roc_shmem_put_nbi(roc_shmem_ctx_t ctx, T *dest, const T *src,
__device__ void rocshmem_put_nbi(rocshmem_ctx_t ctx, T *dest, const T *src,
size_t nelems, int pe);
template <typename T>
__device__ void roc_shmem_put_nbi(T *dest, const T *src, size_t nelems, int pe);
__device__ void rocshmem_put_nbi(T *dest, const T *src, size_t nelems, int pe);
/**
* @brief Reads contiguous data of \p nelems elements from \p source on \p pe
* to \p dest on the calling PE. The operation is not blocking. The caller will
* return as soon as the request is posted. The caller must call
* roc_shmem_quiet() on the same context if completion notification is
* rocshmem_quiet() on the same context if completion notification is
* required.
*
* This function can be called from divergent control paths at per-thread
* granularity. However, performance may be improved if the caller can
* coalesce contiguous messages and elect a leader thread to call into the
* ROC_SHMEM function.
* rocSHMEM function.
*
* @param[in] ctx Context with which to perform this operation.
* @param[in] dest Destination address. Must be an address on the symmetric
@@ -196,11 +195,11 @@ __device__ void roc_shmem_put_nbi(T *dest, const T *src, size_t nelems, int pe);
*
*/
template <typename T>
__device__ void roc_shmem_get_nbi(roc_shmem_ctx_t ctx, T *dest, const T *source,
__device__ void rocshmem_get_nbi(rocshmem_ctx_t ctx, T *dest, const T *source,
size_t nelems, int pe);
template <typename T>
__device__ void roc_shmem_get_nbi(T *dest, const T *source, size_t nelems,
__device__ void rocshmem_get_nbi(T *dest, const T *source, size_t nelems,
int pe);
/**
@@ -222,11 +221,11 @@ __device__ void roc_shmem_get_nbi(T *dest, const T *source, size_t nelems,
*
*/
template <typename T>
__device__ T roc_shmem_atomic_fetch_add(roc_shmem_ctx_t ctx, T *dest, T val,
__device__ T rocshmem_atomic_fetch_add(rocshmem_ctx_t ctx, T *dest, T val,
int pe);
template <typename T>
__device__ T roc_shmem_atomic_fetch_add(T *dest, T val, int pe);
__device__ T rocshmem_atomic_fetch_add(T *dest, T val, int pe);
/**
* @brief Atomically compares if the value in \p dest with \p cond is equal
@@ -249,11 +248,11 @@ __device__ T roc_shmem_atomic_fetch_add(T *dest, T val, int pe);
*
*/
template <typename T>
__device__ T roc_shmem_atomic_compare_swap(roc_shmem_ctx_t ctx, T *dest, T cond,
__device__ T rocshmem_atomic_compare_swap(rocshmem_ctx_t ctx, T *dest, T cond,
T val, int pe);
template <typename T>
__device__ T roc_shmem_atomic_compare_swap(T *dest, T cond, T val, int pe);
__device__ T rocshmem_atomic_compare_swap(T *dest, T cond, T val, int pe);
/**
* @brief Atomically add 1 to \p dest on \p pe. The operation
@@ -273,10 +272,10 @@ __device__ T roc_shmem_atomic_compare_swap(T *dest, T cond, T val, int pe);
*
*/
template <typename T>
__device__ T roc_shmem_atomic_fetch_inc(roc_shmem_ctx_t ctx, T *dest, int pe);
__device__ T rocshmem_atomic_fetch_inc(rocshmem_ctx_t ctx, T *dest, int pe);
template <typename T>
__device__ T roc_shmem_atomic_fetch_inc(T *dest, int pe);
__device__ T rocshmem_atomic_fetch_inc(T *dest, int pe);
/**
* @brief Atomically return the value of \p dest to the calling PE.
@@ -296,10 +295,10 @@ __device__ T roc_shmem_atomic_fetch_inc(T *dest, int pe);
*
*/
template <typename T>
__device__ T roc_shmem_atomic_fetch(roc_shmem_ctx_t ctx, T *source, int pe);
__device__ T rocshmem_atomic_fetch(rocshmem_ctx_t ctx, T *source, int pe);
template <typename T>
__device__ T roc_shmem_atomic_fetch(T *source, int pe);
__device__ T rocshmem_atomic_fetch(T *source, int pe);
/**
* @brief Atomically add the value \p val to \p dest on \p pe.
@@ -319,11 +318,11 @@ __device__ T roc_shmem_atomic_fetch(T *source, int pe);
*
*/
template <typename T>
__device__ void roc_shmem_atomic_add(roc_shmem_ctx_t ctx, T *dest, T val,
__device__ void rocshmem_atomic_add(rocshmem_ctx_t ctx, T *dest, T val,
int pe);
template <typename T>
__device__ void roc_shmem_atomic_add(T *dest, T val, int pe);
__device__ void rocshmem_atomic_add(T *dest, T val, int pe);
/**
* @brief Atomically add 1 to \p dest on \p pe.
@@ -342,10 +341,10 @@ __device__ void roc_shmem_atomic_add(T *dest, T val, int pe);
*
*/
template <typename T>
__device__ void roc_shmem_atomic_inc(roc_shmem_ctx_t ctx, T *dest, int pe);
__device__ void rocshmem_atomic_inc(rocshmem_ctx_t ctx, T *dest, int pe);
template <typename T>
__device__ void roc_shmem_atomic_inc(T *dest, int pe);
__device__ void rocshmem_atomic_inc(T *dest, int pe);
/**
* @brief Atomically set value for \p dest on \p pe.
@@ -365,11 +364,11 @@ __device__ void roc_shmem_atomic_inc(T *dest, int pe);
*
*/
template <typename T>
__device__ void roc_shmem_atomic_set(roc_shmem_ctx_t ctx, T *dest, T value,
__device__ void rocshmem_atomic_set(rocshmem_ctx_t ctx, T *dest, T value,
int pe);
template <typename T>
__device__ void roc_shmem_atomic_set(T *dest, T value, int pe);
__device__ void rocshmem_atomic_set(T *dest, T value, int pe);
/**
* @brief Block the caller until the condition (* \p ptr \p cmps \p val) is
@@ -378,7 +377,7 @@ __device__ void roc_shmem_atomic_set(T *dest, T value, int pe);
* This function can be called from divergent control paths at per-thread
* granularity. However, performance may be improved if the caller can
* coalesce contiguous messages and elect a leader thread to call into the
* ROC_SHMEM function.
* rocSHMEM function.
*
* @param[in] ivars Pointer to memory on the symmetric heap to wait for.
* @param[in] cmp Operation for the comparison.
@@ -388,7 +387,7 @@ __device__ void roc_shmem_atomic_set(T *dest, T value, int pe);
*
*/
template <typename T>
__device__ void roc_shmem_wait_until(T *ivars, int cmp, T val);
__device__ void rocshmem_wait_until(T *ivars, int cmp, T val);
/**
* @brief test if the condition (* \p ptr \p cmps \p val) is
@@ -397,7 +396,7 @@ __device__ void roc_shmem_wait_until(T *ivars, int cmp, T val);
* This function can be called from divergent control paths at per-thread
* granularity. However, performance may be improved if the caller can
* coalesce contiguous messages and elect a leader thread to call into the
* ROC_SHMEM function.
* rocSHMEM function.
*
* @param[in] ivars Pointer to memory on the symmetric heap to wait for.
* @param[in] cmp Operation for the comparison.
@@ -407,7 +406,7 @@ __device__ void roc_shmem_wait_until(T *ivars, int cmp, T val);
*
*/
template <typename T>
__device__ int roc_shmem_test(T *ivars, int cmp, T val);
__device__ int rocshmem_test(T *ivars, int cmp, T val);
/**
* @brief Perform a broadcast between PEs in the active set. The caller
@@ -425,14 +424,14 @@ __device__ int roc_shmem_test(T *ivars, int cmp, T val);
* @param[in] PE_start PE to start the reduction.
* @param[in] logPE_stride Stride of PEs participating in the reduction.
* @param[in] PE_size Number PEs participating in the reduction.
* @param[in] pSync Temporary sync buffer provided to ROC_SHMEM. Must
be of size at least ROC_SHMEM_REDUCE_SYNC_SIZE.
* @param[in] pSync Temporary sync buffer provided to rocSHMEM. Must
be of size at least ROCSHMEM_REDUCE_SYNC_SIZE.
*
* @return void
*
*/
template <typename T>
__device__ void roc_shmem_wg_broadcast(roc_shmem_ctx_t ctx, T *dest,
__device__ void rocshmem_wg_broadcast(rocshmem_ctx_t ctx, T *dest,
const T *source, int nelement,
int PE_root, int PE_start,
int logPE_stride, int PE_size,
@@ -452,18 +451,18 @@ __device__ void roc_shmem_wg_broadcast(roc_shmem_ctx_t ctx, T *dest,
* @param[in] PE_start PE to start the reduction.
* @param[in] logPE_stride Stride of PEs participating in the reduction.
* @param[in] PE_size Number PEs participating in the reduction.
* @param[in] pWrk Temporary work buffer provided to ROC_SHMEM. Must
* @param[in] pWrk Temporary work buffer provided to rocSHMEM. Must
* be of size at least max(size/2 + 1,
ROC_SHMEM_REDUCE_MIN_WRKDATA_SIZE).
* @param[in] pSync Temporary sync buffer provided to ROC_SHMEM. Must
be of size at least ROC_SHMEM_REDUCE_SYNC_SIZE.
ROCSHMEM_REDUCE_MIN_WRKDATA_SIZE).
* @param[in] pSync Temporary sync buffer provided to rocSHMEM. Must
be of size at least ROCSHMEM_REDUCE_SYNC_SIZE.
* @param[in] handle GPU side handle.
*
* @return void
*
*/
template <typename T, ROC_SHMEM_OP Op>
__device__ void roc_shmem_wg_to_all(roc_shmem_ctx_t ctx, T *dest,
template <typename T, ROCSHMEM_OP Op>
__device__ void rocshmem_wg_to_all(rocshmem_ctx_t ctx, T *dest,
const T *source, int nreduce, int PE_start,
int logPE_stride, int PE_size, T *pWrk,
long *pSync);
@@ -472,7 +471,7 @@ __device__ void roc_shmem_wg_to_all(roc_shmem_ctx_t ctx, T *dest,
* @brief Writes contiguous data of \p nelems elements from \p source on the
* calling PE to \p dest at \p pe. The caller will block until the operation
* completes locally (it is safe to reuse \p source). The caller must
* call into roc_shmem_quiet() if remote completion is required.
* call into rocshmem_quiet() if remote completion is required.
*
* This function can be called from divergent control paths at per-wave
* granularity. However, all threads in a wave must collectivlily participate in
@@ -489,18 +488,18 @@ __device__ void roc_shmem_wg_to_all(roc_shmem_ctx_t ctx, T *dest,
*
*/
template <typename T>
__device__ void roc_shmem_put_wave(roc_shmem_ctx_t ctx, T *dest,
__device__ void rocshmem_put_wave(rocshmem_ctx_t ctx, T *dest,
const T *source, size_t nelems, int pe);
template <typename T>
__device__ void roc_shmem_put_wave(T *dest, const T *source, size_t nelems,
__device__ void rocshmem_put_wave(T *dest, const T *source, size_t nelems,
int pe);
/**
* @brief Writes contiguous data of \p nelems elements from \p source on the
* calling PE to \p dest at \p pe. The caller will block until the operation
* completes locally (it is safe to reuse \p source). The caller must
* call into roc_shmem_quiet() if remote completion is required.
* call into rocshmem_quiet() if remote completion is required.
*
* This function can be called from divergent control paths at per-workgroub
* (WG) granularity. However, All threads in a WG must collectivelly participate
@@ -517,11 +516,11 @@ __device__ void roc_shmem_put_wave(T *dest, const T *source, size_t nelems,
*
*/
template <typename T>
__device__ void roc_shmem_put_wg(roc_shmem_ctx_t ctx, T *dest, const T *source,
__device__ void rocshmem_put_wg(rocshmem_ctx_t ctx, T *dest, const T *source,
size_t nelems, int pe);
template <typename T>
__device__ void roc_shmem_put_wg(T *dest, const T *source, size_t nelems,
__device__ void rocshmem_put_wg(T *dest, const T *source, size_t nelems,
int pe);
/**
@@ -544,11 +543,11 @@ __device__ void roc_shmem_put_wg(T *dest, const T *source, size_t nelems,
*
*/
template <typename T>
__device__ void roc_shmem_get_wave(roc_shmem_ctx_t ctx, T *dest,
__device__ void rocshmem_get_wave(rocshmem_ctx_t ctx, T *dest,
const T *source, size_t nelems, int pe);
template <typename T>
__device__ void roc_shmem_get_wave(T *dest, const T *source, size_t nelems,
__device__ void rocshmem_get_wave(T *dest, const T *source, size_t nelems,
int pe);
/**
@@ -571,18 +570,18 @@ __device__ void roc_shmem_get_wave(T *dest, const T *source, size_t nelems,
*
*/
template <typename T>
__device__ void roc_shmem_get_wg(roc_shmem_ctx_t ctx, T *dest, const T *source,
__device__ void rocshmem_get_wg(rocshmem_ctx_t ctx, T *dest, const T *source,
size_t nelems, int pe);
template <typename T>
__device__ void roc_shmem_get_wg(T *dest, const T *source, size_t nelems,
__device__ void rocshmem_get_wg(T *dest, const T *source, size_t nelems,
int pe);
/**
* @brief Writes contiguous data of \p nelems elements from \p source on the
* calling PE to \p dest on \p pe. The operation is not blocking. The caller
* will return as soon as the request is posted. The caller must call
* roc_shmem_quiet() on the same context if completion notification is
* rocshmem_quiet() on the same context if completion notification is
* required.
*
* This function can be called from divergent control paths at per-wave
@@ -599,18 +598,18 @@ __device__ void roc_shmem_get_wg(T *dest, const T *source, size_t nelems,
*
*/
template <typename T>
__device__ void roc_shmem_put_nbi_wave(roc_shmem_ctx_t ctx, T *dest,
__device__ void rocshmem_put_nbi_wave(rocshmem_ctx_t ctx, T *dest,
const T *src, size_t nelems, int pe);
template <typename T>
__device__ void roc_shmem_put_nbi_wave(T *dest, const T *src, size_t nelems,
__device__ void rocshmem_put_nbi_wave(T *dest, const T *src, size_t nelems,
int pe);
/**
* @brief Writes contiguous data of \p nelems elements from \p source on the
* calling PE to \p dest on \p pe. The operation is not blocking. The caller
* will return as soon as the request is posted. The caller must call
* roc_shmem_quiet() on the same context if completion notification is
* rocshmem_quiet() on the same context if completion notification is
* required.
*
* This function can be called from divergent control paths at per-workgroup
@@ -627,18 +626,18 @@ __device__ void roc_shmem_put_nbi_wave(T *dest, const T *src, size_t nelems,
*
*/
template <typename T>
__device__ void roc_shmem_put_nbi_wg(roc_shmem_ctx_t ctx, T *dest,
__device__ void rocshmem_put_nbi_wg(rocshmem_ctx_t ctx, T *dest,
const T *src, size_t nelems, int pe);
template <typename T>
__device__ void roc_shmem_put_nbi_wg(T *dest, const T *src, size_t nelems,
__device__ void rocshmem_put_nbi_wg(T *dest, const T *src, size_t nelems,
int pe);
/**
* @brief Reads contiguous data of \p nelems elements from \p source on \p pe
* to \p dest on the calling PE. The operation is not blocking. The caller will
* return as soon as the request is posted. The caller must call
* roc_shmem_quiet() on the same context if completion notification is
* rocshmem_quiet() on the same context if completion notification is
* required.
*
* This function can be called from divergent control paths at per-wave
@@ -655,18 +654,18 @@ __device__ void roc_shmem_put_nbi_wg(T *dest, const T *src, size_t nelems,
*
*/
template <typename T>
__device__ void roc_shmem_get_nbi_wave(roc_shmem_ctx_t ctx, T *dest,
__device__ void rocshmem_get_nbi_wave(rocshmem_ctx_t ctx, T *dest,
const T *source, size_t nelems, int pe);
template <typename T>
__device__ void roc_shmem_get_nbi_wave(T *dest, const T *source, size_t nelems,
__device__ void rocshmem_get_nbi_wave(T *dest, const T *source, size_t nelems,
int pe);
/**
* @brief Reads contiguous data of \p nelems elements from \p source on \p pe
* to \p dest on the calling PE. The operation is not blocking. The caller will
* return as soon as the request is posted. The caller must call
* roc_shmem_quiet() on the same context if completion notification is
* rocshmem_quiet() on the same context if completion notification is
* required.
*
* This function can be called from divergent control paths at per-workgroup
@@ -683,101 +682,101 @@ __device__ void roc_shmem_get_nbi_wave(T *dest, const T *source, size_t nelems,
*
*/
template <typename T>
__device__ void roc_shmem_get_nbi_wg(roc_shmem_ctx_t ctx, T *dest,
__device__ void rocshmem_get_nbi_wg(rocshmem_ctx_t ctx, T *dest,
const T *source, size_t nelems, int pe);
template <typename T>
__device__ void roc_shmem_get_nbi_wg(T *dest, const T *source, size_t nelems,
__device__ void rocshmem_get_nbi_wg(T *dest, const T *source, size_t nelems,
int pe);
__device__ void roc_shmem_putmem_wave(void *dest, const void *source,
__device__ void rocshmem_putmem_wave(void *dest, const void *source,
size_t nelems, int pe) {
roc_shmem_ctx_putmem_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe);
rocshmem_ctx_putmem_wave(ROCSHMEM_CTX_DEFAULT, dest, source, nelems, pe);
}
__device__ void roc_shmem_putmem_wg(void *dest, const void *source,
__device__ void rocshmem_putmem_wg(void *dest, const void *source,
size_t nelems, int pe) {
roc_shmem_ctx_putmem_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe);
rocshmem_ctx_putmem_wg(ROCSHMEM_CTX_DEFAULT, dest, source, nelems, pe);
}
template <typename T>
__device__ void roc_shmem_put_wave(T *dest, const T *source, size_t nelems,
__device__ void rocshmem_put_wave(T *dest, const T *source, size_t nelems,
int pe) {
roc_shmem_put_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe);
rocshmem_put_wave(ROCSHMEM_CTX_DEFAULT, dest, source, nelems, pe);
}
template <typename T>
__device__ void roc_shmem_put_wg(T *dest, const T *source, size_t nelems,
__device__ void rocshmem_put_wg(T *dest, const T *source, size_t nelems,
int pe) {
roc_shmem_put_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe);
rocshmem_put_wg(ROCSHMEM_CTX_DEFAULT, dest, source, nelems, pe);
}
__device__ void roc_shmem_getmem_wg(void *dest, const void *source,
__device__ void rocshmem_getmem_wg(void *dest, const void *source,
size_t nelems, int pe) {
roc_shmem_ctx_getmem_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe);
rocshmem_ctx_getmem_wg(ROCSHMEM_CTX_DEFAULT, dest, source, nelems, pe);
}
template <typename T>
__device__ void roc_shmem_get_wg(T *dest, const T *source, size_t nelems,
__device__ void rocshmem_get_wg(T *dest, const T *source, size_t nelems,
int pe) {
roc_shmem_get_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe);
rocshmem_get_wg(ROCSHMEM_CTX_DEFAULT, dest, source, nelems, pe);
}
__device__ void roc_shmem_getmem_wave(void *dest, const void *source,
__device__ void rocshmem_getmem_wave(void *dest, const void *source,
size_t nelems, int pe) {
roc_shmem_ctx_getmem_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe);
rocshmem_ctx_getmem_wave(ROCSHMEM_CTX_DEFAULT, dest, source, nelems, pe);
}
template <typename T>
__device__ void roc_shmem_get_wave(T *dest, const T *source, size_t nelems,
__device__ void rocshmem_get_wave(T *dest, const T *source, size_t nelems,
int pe) {
roc_shmem_get_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe);
rocshmem_get_wave(ROCSHMEM_CTX_DEFAULT, dest, source, nelems, pe);
}
__device__ void roc_shmem_putmem_nbi_wg(void *dest, const void *source,
__device__ void rocshmem_putmem_nbi_wg(void *dest, const void *source,
size_t nelems, int pe) {
roc_shmem_ctx_putmem_nbi_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe);
rocshmem_ctx_putmem_nbi_wg(ROCSHMEM_CTX_DEFAULT, dest, source, nelems, pe);
}
template <typename T>
__device__ void roc_shmem_put_nbi_wg(T *dest, const T *source, size_t nelems,
__device__ void rocshmem_put_nbi_wg(T *dest, const T *source, size_t nelems,
int pe) {
roc_shmem_put_nbi_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe);
rocshmem_put_nbi_wg(ROCSHMEM_CTX_DEFAULT, dest, source, nelems, pe);
}
__device__ void roc_shmem_putmem_nbi_wave(void *dest, const void *source,
__device__ void rocshmem_putmem_nbi_wave(void *dest, const void *source,
size_t nelems, int pe) {
roc_shmem_ctx_putmem_nbi_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems,
rocshmem_ctx_putmem_nbi_wave(ROCSHMEM_CTX_DEFAULT, dest, source, nelems,
pe);
}
template <typename T>
__device__ void roc_shmem_put_nbi_wave(T *dest, const T *source, size_t nelems,
__device__ void rocshmem_put_nbi_wave(T *dest, const T *source, size_t nelems,
int pe) {
roc_shmem_put_nbi_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe);
rocshmem_put_nbi_wave(ROCSHMEM_CTX_DEFAULT, dest, source, nelems, pe);
}
__device__ void roc_shmem_getmem_nbi_wg(void *dest, const void *source,
__device__ void rocshmem_getmem_nbi_wg(void *dest, const void *source,
size_t nelems, int pe) {
roc_shmem_ctx_getmem_nbi_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe);
rocshmem_ctx_getmem_nbi_wg(ROCSHMEM_CTX_DEFAULT, dest, source, nelems, pe);
}
template <typename T>
__device__ void roc_shmem_get_nbi_wg(T *dest, const T *source, size_t nelems,
__device__ void rocshmem_get_nbi_wg(T *dest, const T *source, size_t nelems,
int pe) {
roc_shmem_get_nbi_wg(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe);
rocshmem_get_nbi_wg(ROCSHMEM_CTX_DEFAULT, dest, source, nelems, pe);
}
__device__ void roc_shmem_getmem_nbi_wave(void *dest, const void *source,
__device__ void rocshmem_getmem_nbi_wave(void *dest, const void *source,
size_t nelems, int pe) {
roc_shmem_ctx_getmem_nbi_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems,
rocshmem_ctx_getmem_nbi_wave(ROCSHMEM_CTX_DEFAULT, dest, source, nelems,
pe);
}
template <typename T>
__device__ void roc_shmem_get_nbi_wave(T *dest, const T *source, size_t nelems,
__device__ void rocshmem_get_nbi_wave(T *dest, const T *source, size_t nelems,
int pe) {
roc_shmem_get_nbi_wave(ROC_SHMEM_CTX_DEFAULT, dest, source, nelems, pe);
rocshmem_get_nbi_wave(ROCSHMEM_CTX_DEFAULT, dest, source, nelems, pe);
}
} // namespace rocshmem
src/templates_host.hpp
+33 -33
View File
@@ -23,11 +23,11 @@
#ifndef LIBRARY_SRC_TEMPLATES_HOST_HPP_
#define LIBRARY_SRC_TEMPLATES_HOST_HPP_
#include "roc_shmem/roc_shmem.hpp"
#include "rocshmem/rocshmem.hpp"
/**
* @file templates_host.hpp
* @brief Internal header that declares templates for ROC_SHMEM's implentation
* @brief Internal header that declares templates for rocSHMEM's implementation
* of the user-facing host APIs.
*
* This file contains templates for the OpenSHMEM APIs that take have
@@ -41,102 +41,102 @@
namespace rocshmem {
template <typename T>
__host__ void roc_shmem_put(roc_shmem_ctx_t ctx, T *dest, const T *source,
__host__ void rocshmem_put(rocshmem_ctx_t ctx, T *dest, const T *source,
size_t nelems, int pe);
template <typename T>
__host__ void roc_shmem_put(T *dest, const T *source, size_t nelems, int pe);
__host__ void rocshmem_put(T *dest, const T *source, size_t nelems, int pe);
template <typename T>
__host__ void roc_shmem_p(roc_shmem_ctx_t ctx, T *dest, T value, int pe);
__host__ void rocshmem_p(rocshmem_ctx_t ctx, T *dest, T value, int pe);
template <typename T>
__host__ void roc_shmem_p(T *dest, T value, int pe);
__host__ void rocshmem_p(T *dest, T value, int pe);
template <typename T>
__host__ void roc_shmem_get(roc_shmem_ctx_t ctx, T *dest, const T *source,
__host__ void rocshmem_get(rocshmem_ctx_t ctx, T *dest, const T *source,
size_t nelems, int pe);
template <typename T>
__host__ void roc_shmem_get(T *dest, const T *source, size_t nelems, int pe);
__host__ void rocshmem_get(T *dest, const T *source, size_t nelems, int pe);
template <typename T>
__host__ T roc_shmem_g(roc_shmem_ctx_t ctx, const T *source, int pe);
__host__ T rocshmem_g(rocshmem_ctx_t ctx, const T *source, int pe);
template <typename T>
__host__ T roc_shmem_g(const T *source, int pe);
__host__ T rocshmem_g(const T *source, int pe);
template <typename T>
__host__ void roc_shmem_put_nbi(roc_shmem_ctx_t ctx, T *dest, const T *src,
__host__ void rocshmem_put_nbi(rocshmem_ctx_t ctx, T *dest, const T *src,
size_t nelems, int pe);
template <typename T>
__host__ void roc_shmem_put_nbi(T *dest, const T *src, size_t nelems, int pe);
__host__ void rocshmem_put_nbi(T *dest, const T *src, size_t nelems, int pe);
template <typename T>
__host__ void roc_shmem_get_nbi(roc_shmem_ctx_t ctx, T *dest, const T *source,
__host__ void rocshmem_get_nbi(rocshmem_ctx_t ctx, T *dest, const T *source,
size_t nelems, int pe);
template <typename T>
__host__ void roc_shmem_get_nbi(T *dest, const T *source, size_t nelems,
__host__ void rocshmem_get_nbi(T *dest, const T *source, size_t nelems,
int pe);
template <typename T>
__host__ T roc_shmem_atomic_fetch_add(roc_shmem_ctx_t ctx, T *dest, T val,
__host__ T rocshmem_atomic_fetch_add(rocshmem_ctx_t ctx, T *dest, T val,
int pe);
template <typename T>
__host__ T roc_shmem_atomic_fetch_add(T *dest, T val, int pe);
__host__ T rocshmem_atomic_fetch_add(T *dest, T val, int pe);
template <typename T>
__host__ T roc_shmem_atomic_compare_swap(roc_shmem_ctx_t ctx, T *dest, T cond,
__host__ T rocshmem_atomic_compare_swap(rocshmem_ctx_t ctx, T *dest, T cond,
T val, int pe);
template <typename T>
__host__ T roc_shmem_atomic_compare_swap(T *dest, T cond, T val, int pe);
__host__ T rocshmem_atomic_compare_swap(T *dest, T cond, T val, int pe);
template <typename T>
__host__ T roc_shmem_atomic_fetch_inc(roc_shmem_ctx_t ctx, T *dest, int pe);
__host__ T rocshmem_atomic_fetch_inc(rocshmem_ctx_t ctx, T *dest, int pe);
template <typename T>
__host__ T roc_shmem_atomic_fetch_inc(T *dest, int pe);
__host__ T rocshmem_atomic_fetch_inc(T *dest, int pe);
template <typename T>
__host__ T roc_shmem_atomic_fetch(roc_shmem_ctx_t ctx, T *source, int pe);
__host__ T rocshmem_atomic_fetch(rocshmem_ctx_t ctx, T *source, int pe);
template <typename T>
__host__ T roc_shmem_atomic_fetch(T *source, int pe);
__host__ T rocshmem_atomic_fetch(T *source, int pe);
template <typename T>
__host__ void roc_shmem_atomic_add(roc_shmem_ctx_t ctx, T *dest, T val, int pe);
__host__ void rocshmem_atomic_add(rocshmem_ctx_t ctx, T *dest, T val, int pe);
template <typename T>
__host__ void roc_shmem_atomic_add(T *dest, T val, int pe);
__host__ void rocshmem_atomic_add(T *dest, T val, int pe);
template <typename T>
__host__ void roc_shmem_atomic_inc(roc_shmem_ctx_t ctx, T *dest, int pe);
__host__ void rocshmem_atomic_inc(rocshmem_ctx_t ctx, T *dest, int pe);
template <typename T>
__host__ void roc_shmem_atomic_inc(T *dest, int pe);
__host__ void rocshmem_atomic_inc(T *dest, int pe);
template <typename T>
__host__ void roc_shmem_atomic_set(T *dest, T val, int pe);
__host__ void rocshmem_atomic_set(T *dest, T val, int pe);
template <typename T>
__host__ void roc_shmem_atomic_set(roc_shmem_ctx_t ctx, T *dest, T val, int pe);
__host__ void rocshmem_atomic_set(rocshmem_ctx_t ctx, T *dest, T val, int pe);
template <typename T>
__host__ void roc_shmem_broadcast(roc_shmem_ctx_t ctx, T *dest, const T *source,
__host__ void rocshmem_broadcast(rocshmem_ctx_t ctx, T *dest, const T *source,
int nelement, int PE_root, int PE_start,
int logPE_stride, int PE_size, long *pSync);
template <typename T, ROC_SHMEM_OP Op>
__host__ void roc_shmem_to_all(roc_shmem_ctx_t ctx, T *dest, const T *source,
template <typename T, ROCSHMEM_OP Op>
__host__ void rocshmem_to_all(rocshmem_ctx_t ctx, T *dest, const T *source,
int nreduce, int PE_start, int logPE_stride,
int PE_size, T *pWrk, long *pSync);
template <typename T>
__host__ void roc_shmem_wait_until(T *ivars, int cmp, T val);
__host__ void rocshmem_wait_until(T *ivars, int cmp, T val);
template <typename T>
__host__ void wait_until_all(T* ivars, size_t nelems, const int *status,
@@ -164,7 +164,7 @@ __host__ size_t wait_until_some_vector(T* ivars, size_t nelems,
int cmp, T* vals);
template <typename T>
__host__ int roc_shmem_test(T *ivars, int cmp, T val);
__host__ int rocshmem_test(T *ivars, int cmp, T val);
} // namespace rocshmem