Skip to content

GitLab

Commit a117adf8 authored by lishen's avatar lishen
Browse files

Merge branch 'main' of http://112.11.119.99:10068//dcutoolkit/deeplearing/DeepEP

parents b33659dd b705eeca
Expand all Hide whitespace changes
Inline Side-by-side
Showing with 4 additions and 6888 deletions
.gitmodules 0 → 100644
View file @ a117adf8
[submodule "third-party/rocshmem"]
path = third-party/rocshmem
url = http://112.11.119.99:10068/dcutoolkit/deeplearing/rocshmem.git
rocshmem_dir/include/rocshmem/rocshmem.hpp deleted 100644 → 0
View file @ b33659dd
/******************************************************************************
* Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
*
* SPDX-License-Identifier: MIT
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to
* deal in the Software without restriction, including without limitation the
* rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
* sell copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice shall be included in
* all copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
* FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
* IN THE SOFTWARE.
*****************************************************************************/
#ifndef LIBRARY_INCLUDE_ROCSHMEM_HPP
#define LIBRARY_INCLUDE_ROCSHMEM_HPP
#include <hip/hip_runtime.h>
#include "rocshmem_config.h"
#include "rocshmem_common.hpp"
#include "rocshmem_RMA.hpp"
#include "rocshmem_AMO.hpp"
#include "rocshmem_SIG_OP.hpp"
#include "rocshmem_COLL.hpp"
#include "rocshmem_P2P_SYNC.hpp"
#include "rocshmem_RMA_X.hpp"
#if defined(HAVE_EXTERNAL_MPI)
#include <mpi.h>
#endif
/**
* @file rocshmem.hpp
* @brief Public header for rocSHMEM device and host libraries.
*
* This file contains all the callable functions and data structures for both
* the device-side runtime and host-side runtime.
*
* The comments on these functions are sparse, but the semantics are the same
* as those implemented in OpenSHMEM unless otherwise documented. Please see
* the OpenSHMEM 1.4 standards documentation for more details:
*
* http://openshmem.org/site/sites/default/site_files/OpenSHMEM-1.4.pdf
*/
namespace rocshmem {
constexpr char VERSION[] = "3.0.0";
/******************************************************************************
**************************** HOST INTERFACE **********************************
*****************************************************************************/
#if defined(HAVE_EXTERNAL_MPI)
/**
* @brief Initialize the rocSHMEM runtime and underlying transport layer.
*
* @param[in] comm MPI Communicator that rocSHMEM will be using
* If MPI_COMM_NULL, rocSHMEM will be using MPI_COMM_WORLD
*/
[[deprecated]] __host__ void rocshmem_init(MPI_Comm comm);
#endif
/**
* @brief Initialize the rocSHMEM runtime and underlying transport layer.
* This is equivalent to the previous function, using implicitely
* MPI_COMM_WORLD for initialization
*/
__host__ void rocshmem_init(void);
/**
* @brief Query rocSHMEM context from host API
*
* @param[out] ctx Returns ROCSHMEM_CTX_DEFAULT device pointer that users
* can query from one instance of rocshmem host library and
* use use later for dynamic module initialization in
* kernel bitcode device library in the same application
*/
__host__ void * rocshmem_get_device_ctx();
/**
* @brief Query rocSHMEM remote symmetric heap pointer
*
* @param[in] dest local symmetric heap allocation pointer for current pe/device
*
* @param[in] pe remote PE
*
* @param[out] ptr Returns remote symmetric heap device pointer from host-side API.
* This can be used to issue load/store from custom kernels
* instead of using rocshmem device side get/put APIs for RMA operations.
*/
__host__ void* rocshmem_ptr(const void *dest, int pe);
__device__ ATTR_NO_INLINE void* rocshmem_ptr(const void *dest, int pe);
#if defined(HAVE_EXTERNAL_MPI)
/**
* @brief Initialize the rocSHMEM runtime and underlying transport layer
* with an attempt to enable the requested thread support.
*
* @param[in] requested Requested thread mode (from rocshmem_thread_ops)
* for host-facing functions.
* @param[out] provided Thread mode selected by the runtime. May not be equal
* to requested thread mode.
* @param[in] comm (Optional) MPI Communicator that rocSHMEM will be using
* If MPI_COMM_NULL, rocSHMEM will be using MPI_COMM_WORLD
*
* @return int returns 0 upon success; otherwise, it returns a nonzero
* value
*/
[[deprecated]] __host__ int rocshmem_init_thread(int requested, int *provided,
MPI_Comm comm);
#endif
/**
* @brief Initialize the rocSHMEM runtime and underlying transport layer
* using the provided mode and attributes
*
* @param[in] flags initialization method to be used.
* Valid values are ROCSHMEM_INIT_WITH_UNIQUEID and
* ROCSHMEM_INIT_WITH_MPI_COMM
* @param[in] attr attribute structure specifying input characteristics
*
* @return int returns 0 upon success; otherwise, it returns a nonzero
* value
*/
__host__ int rocshmem_init_attr(unsigned int flags, rocshmem_init_attr_t *attr);
/**
* @brief Return a uniqueID
*
* @return int returns 0 upon success; otherwise, it returns a nonzero
* value
*/
__host__ int rocshmem_get_uniqueid(rocshmem_uniqueid_t *uid);
/**
* @brief Initalizes the rocshmem_init_attr_t struct
*
* @param[in] rank rank of the calling process
* @param[in] nranks number of pes
* @param[in] uid unique ID used to identify the group processes.
* All processes that
* @param[out] attr attribute structure to be passed to rocshmem_init_attr
*
* @return int returns 0 upon success; otherwise, it returns a nonzero
* value
*/
__host__ int rocshmem_set_attr_uniqueid_args(int rank, int nranks,
rocshmem_uniqueid_t *uid,
rocshmem_init_attr_t *attr);
/**
* @brief Query the thread mode used by the runtime.
*
* @param[out] provided Thread mode the runtime is operating in.
*
* @return void.
*/
__host__ void rocshmem_query_thread(int *provided);
/**
* @brief Function that dumps internal stats to stdout.
*/
__host__ void rocshmem_dump_stats();
/**
* @brief Reset all internal stats.
*/
__host__ void rocshmem_reset_stats();
/**
* @brief Finalize the rocSHMEM runtime.
*/
__host__ void rocshmem_finalize();
/**
* @brief Allocate memory of \p size bytes from the symmetric heap.
* This is a collective operation and must be called by all PEs.
*
* @param[in] size Memory allocation size in bytes.
*
* @return A pointer to the allocated memory on the symmetric heap.
*
* @todo Return error code instead of ptr.
*/
__host__ void *rocshmem_malloc(size_t size);
/**
* @brief Free a memory allocation from the symmetric heap.
* This is a collective operation and must be called by all PEs.
*
* @param[in] ptr Pointer to previously allocated memory on the symmetric heap.
*/
__host__ void rocshmem_free(void *ptr);
/**
* @brief Query for the number of PEs.
*
* @return Number of PEs.
*/
__host__ int rocshmem_n_pes();
/**
* @brief Query the PE ID of the caller.
*
* @return PE ID of the caller.
*/
__host__ int rocshmem_my_pe();
/**
* @brief Creates an OpenSHMEM context.
*
* @param[in] options Options for context creation. Ignored in current design.
* @param[out] ctx Context handle.
*
* @return Zero on success and nonzero otherwise.
*/
__host__ int rocshmem_ctx_create(int64_t options, rocshmem_ctx_t *ctx);
/**
* @brief Destroys an OpenSHMEM context.
*
* @param[out] ctx Context handle.
*
* @return void.
*/
__host__ void rocshmem_ctx_destroy(rocshmem_ctx_t ctx);
/**
* @brief Translate the PE in src_team to that in dest_team.
*
* @param[in] src_team Handle of the team from which to translate
* @param[in] src_pe PE-of-interest's index in src_team
* @param[in] dest_team Handle of the team to which to translate
*
* @return PE of src_pe in dest_team. If any input is invalid
* or if src_pe is not in both source and destination
* teams, a value of -1 is returned.
*/
__host__ int rocshmem_team_translate_pe(rocshmem_team_t src_team, int src_pe,
rocshmem_team_t dest_team);
/**
* @brief Query the number of PEs in a team.
*
* @param[in] team The team to query PE ID in.
*
* @return Number of PEs in the provided team.
*/
__host__ int rocshmem_team_n_pes(rocshmem_team_t team);
/**
* @brief Query the PE ID of the caller in a team.
*
* @param[in] team The team to query PE ID in.
*
* @return PE ID of the caller in the provided team.
*/
__host__ int rocshmem_team_my_pe(rocshmem_team_t team);
/**
* @brief Create a new a team of PEs. Must be called by all PEs
* in the parent team.
*
* @param[in] parent_team The team to split from.
* @param[in] start The lowest PE number of the subset of the PEs
* from the parent team that will form the new
* team.
* @param[in] stide The stride between team PE members in the
* parent team that comprise the subset of PEs
* that will form the new team.
* @param[in] size The number of PEs in the new team.
* @param[in] config Pointer to the config parameters for the new
* team.
* @param[in] config_mask Bitwise mask representing parameters to use
* from config
* @param[out] new_team Pointer to the newly created team. If an error
* occurs during team creation, or if the PE in
* the parent team is not in the new team, the
* value will be ROCSHMEM_TEAM_INVALID.
*
* @return Zero upon successful team creation; non-zero if erroneous.
*/
__host__ int rocshmem_team_split_strided(rocshmem_team_t parent_team,
int start, int stride, int size,
const rocshmem_team_config_t *config,
long config_mask,
rocshmem_team_t *new_team);
/**
* @brief Destroy a team. Must be called by all PEs in the team.
* The user must destroy all private contexts created in the
* team before destroying this team. Otherwise, the behavior
* is undefined. This call will destroy only the shareable contexts
* created from the referenced team.
*
* @param[in] team The team to destroy. The behavior is undefined if
* the input team is ROCSHMEM_TEAM_WORLD or any other
* invalid team. If the input is ROCSHMEM_TEAM_INVALID,
* this function will not perform any operation.
*
* @return None.
*/
__host__ void rocshmem_team_destroy(rocshmem_team_t team);
/**
* @brief Guarantees order between messages in this context in accordance with
* OpenSHMEM semantics.
*
* @param[in] ctx Context with which to perform this operation.
*
* @return void.
*/
__host__ void rocshmem_ctx_fence(rocshmem_ctx_t ctx);
__host__ void rocshmem_fence();
/**
* @brief Completes all previous operations posted on the host.
*
* @param[in] ctx Context with which to perform this operation.
*
* @return void.
*/
__host__ void rocshmem_ctx_quiet(rocshmem_ctx_t ctx);
__host__ void rocshmem_quiet();
/**
* @brief perform a collective barrier between all PEs in the system.
* The caller is blocked until the barrier is resolved.
*
* @return void
*/
__host__ void rocshmem_barrier_all();
/**
* @brief enqueues a collective barrier on given stream.
*
* @return void
*/
__host__ void rocshmem_barrier_all_on_stream(hipStream_t stream);
/**
* @brief registers the arrival of a PE at a barrier.
* The caller is blocked until the synchronization is resolved.
*
* In contrast with the shmem_barrier_all routine, shmem_sync_all only ensures
* completion and visibility of previously issued memory stores and does not
* ensure completion of remote memory updates issued via OpenSHMEM routines.
*
* @return void
*/
__host__ void rocshmem_sync_all();
/**
* @brief allows any PE to force the termination of an entire program.
*
* @param[in] status The exit status from the main program.
*
* @return void
*/
__host__ void rocshmem_global_exit(int status);
/******************************************************************************
**************************** DEVICE INTERFACE ********************************
*****************************************************************************/
/**
* @brief Initializes device-side rocSHMEM resources. Must be called before
* any threads in this work-group invoke other rocSHMEM functions.
*
* Must be called collectively by all threads in the work-group.
*
* @return void.
*/
[[deprecated]] __device__ void rocshmem_wg_init();
/**
* @brief Finalizes device-side rocSHMEM resources. Must be called before
* work-group completion if the work-group also called rocshmem_wg_init().
*
* Must be called collectively by all threads in the work-group.
*
* @return void.
*/
[[deprecated]] __device__ void rocshmem_wg_finalize();
/**
* @brief Initializes device-side rocSHMEM resources. Must be called before
* any threads in this work-group invoke other rocSHMEM functions. This is
* a variant of rocshmem_wg_init that allows the caller to request a
* threading mode.
*
* @param[in] requested Requested thread mode from rocshmem_thread_ops.
* @param[out] provided Thread mode selected by the runtime. May not be equal
* to requested thread mode.
*
* Must be called collectively by all threads in the work-group.
*
* @return void.
*/
[[deprecated]] __device__ void rocshmem_wg_init_thread(int requested, int *provided);
/**
* @brief Query the thread mode used by the runtime.
*
* @param[out] provided Thread mode the runtime is operating in.
*
* @return void.
*/
__device__ void rocshmem_query_thread(int *provided);
/**
* @brief Creates an OpenSHMEM context. By design, the context is private
* to the calling work-group.
*
* Must be called collectively by all threads in the work-group.
*
* @param[in] options Options for context creation. Ignored in current design.
* @param[out] ctx Context handle.
*
* @return All threads returns 0 if the context was created successfully. If any
* thread returns non-zero value, the operation failed and a higher number of
* `ROCSHMEM_MAX_NUM_CONTEXTS` is required.
*/
__device__ ATTR_NO_INLINE int rocshmem_wg_ctx_create(int64_t options,
rocshmem_ctx_t *ctx);
__device__ ATTR_NO_INLINE int rocshmem_wg_team_create_ctx(
rocshmem_team_t team, long options, rocshmem_ctx_t *ctx);
/**
* @brief Destroys an OpenSHMEM context.
*
* Must be called collectively by all threads in the work-group.
*
* @param[in] The context to destroy.
*
* @return void.
*/
__device__ ATTR_NO_INLINE void rocshmem_wg_ctx_destroy(rocshmem_ctx_t *ctx);
/**
* @brief Guarantees order between messages in this context in accordance with
* OpenSHMEM semantics.
*
* 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
* rocSHMEM function.
*
* @param[in] ctx Context with which to perform this operation.
*
* @return void.
*/
__device__ ATTR_NO_INLINE void rocshmem_ctx_fence(rocshmem_ctx_t ctx);
__device__ ATTR_NO_INLINE void rocshmem_fence();
/**
* @brief Guarantees order between messages in this context in accordance with
* OpenSHMEM semantics.
*
* This function is an extension as it is per PE. has same semantics as default
* API but it is per PE
*
* @param[in] ctx Context with which to perform this operation.
* @param[in] pe destination pe.
*
* @return void.
*/
__device__ ATTR_NO_INLINE void rocshmem_ctx_fence(rocshmem_ctx_t ctx, int pe);
__device__ ATTR_NO_INLINE void rocshmem_fence(int pe);
/**
* @brief Completes all previous operations posted to this context.
*
* 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
* rocSHMEM function.
*
* @param[in] ctx Context with which to perform this operation.
*
* @return void.
*/
__device__ ATTR_NO_INLINE void rocshmem_ctx_quiet(rocshmem_ctx_t ctx);
__device__ ATTR_NO_INLINE void rocshmem_quiet();
/**
* @brief Completes all previous operations posted to this context for PEs in the
* `target_pes` array.
*
* @param[in] ctx Context with which to perform this operation.
*
* @param[in] target_pes Address of target PE array where the operations need to be completed.
*
* @param[in] npes The number of PEs in the target PE array.
*
* @return void.
*/
__device__ ATTR_NO_INLINE void rocshmem_ctx_pe_quiet(rocshmem_ctx_t ctx, const int *target_pes, size_t npes);
__device__ ATTR_NO_INLINE void rocshmem_pe_quiet(const int *target_pes, size_t npes);
/**
* @brief Query the total number of PEs.
*
* Can be called per thread with no performance penalty.
*
* @param[in] ctx GPU side handle.
*
* @return Total number of PEs.
*/
__device__ int rocshmem_ctx_n_pes(rocshmem_ctx_t ctx);
__device__ int rocshmem_n_pes();
/**
* @brief Query the PE ID of the caller.
*
* Can be called per thread with no performance penalty.
*
* @param[in] ctx GPU side handle
*
* @return PE ID of the caller.
*/
__device__ int rocshmem_ctx_my_pe(rocshmem_ctx_t ctx);
__device__ int rocshmem_my_pe();
/**
* @brief Translate the PE in src_team to that in dest_team.
*
* @param[in] src_team Handle of the team from which to translate
* @param[in] src_pe PE-of-interest's index in src_team
* @param[in] dest_team Handle of the team to which to translate
*
* @return PE of src_pe in dest_team. If any input is invalid
* or if src_pe is not in both source and destination
* teams, a value of -1 is returned.
*/
__device__ int rocshmem_team_translate_pe(rocshmem_team_t src_team,
int src_pe,
rocshmem_team_t dest_team);
__device__ ATTR_NO_INLINE void rocshmem_ctx_threadfence_system(
rocshmem_ctx_t ctx);
__device__ ATTR_NO_INLINE void rocshmem_threadfence_system();
} // namespace rocshmem
#endif // LIBRARY_INCLUDE_ROCSHMEM_HPP
rocshmem_dir/include/rocshmem/rocshmem_AMO.hpp deleted 100644 → 0
View file @ b33659dd
This diff is collapsed.
rocshmem_dir/include/rocshmem/rocshmem_COLL.hpp deleted 100644 → 0
View file @ b33659dd
This diff is collapsed.
rocshmem_dir/include/rocshmem/rocshmem_P2P_SYNC.hpp deleted 100644 → 0
View file @ b33659dd
This diff is collapsed.
rocshmem_dir/include/rocshmem/rocshmem_RMA.hpp deleted 100644 → 0
View file @ b33659dd
This diff is collapsed.
rocshmem_dir/include/rocshmem/rocshmem_RMA_X.hpp deleted 100644 → 0
View file @ b33659dd
This diff is collapsed.
rocshmem_dir/include/rocshmem/rocshmem_SIG_OP.hpp deleted 100644 → 0
View file @ b33659dd
This diff is collapsed.
rocshmem_dir/include/rocshmem/rocshmem_common.hpp deleted 100644 → 0
View file @ b33659dd
/******************************************************************************
* Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
*
* SPDX-License-Identifier: MIT
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to
* deal in the Software without restriction, including without limitation the
* rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
* sell copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice shall be included in
* all copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
* FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
* IN THE SOFTWARE.
*****************************************************************************/
#ifndef LIBRARY_INCLUDE_ROCSHMEM_COMMON_HPP
#define LIBRARY_INCLUDE_ROCSHMEM_COMMON_HPP
namespace rocshmem {
#ifdef USE_FUNC_CALL
#define ATTR_NO_INLINE __attribute__((noinline))
#else
#define ATTR_NO_INLINE
#endif
enum ROCSHMEM_STATUS {
ROCSHMEM_SUCCESS = 0,
ROCSHMEM_ERROR = 1,
};
enum ROCSHMEM_OP {
ROCSHMEM_SUM,
ROCSHMEM_MAX,
ROCSHMEM_MIN,
ROCSHMEM_PROD,
ROCSHMEM_AND,
ROCSHMEM_OR,
ROCSHMEM_XOR,
ROCSHMEM_REPLACE
};
enum ROCSHMEM_SIGNAL_OPS {
ROCSHMEM_SIGNAL_SET,
ROCSHMEM_SIGNAL_ADD,
};
/**
* @brief Types defined for rocshmem_wait() operations.
*/
enum rocshmem_cmps {
ROCSHMEM_CMP_EQ,
ROCSHMEM_CMP_NE,
ROCSHMEM_CMP_GT,
ROCSHMEM_CMP_GE,
ROCSHMEM_CMP_LT,
ROCSHMEM_CMP_LE,
};
enum rocshmem_thread_ops {
ROCSHMEM_THREAD_SINGLE,
ROCSHMEM_THREAD_FUNNELED,
ROCSHMEM_THREAD_WG_FUNNELED,
ROCSHMEM_THREAD_SERIALIZED,
ROCSHMEM_THREAD_MULTIPLE
};
/**
* @brief Bitwise flags to mask configuration parameters.
*/
enum rocshmem_team_configs {
ROCSHMEM_TEAM_DEFAULT_CONFIGS,
ROCSHMEM_TEAM_NUM_CONTEXTS
};
typedef struct {
int num_contexts;
} rocshmem_team_config_t;
constexpr size_t ROCSHMEM_REDUCE_MIN_WRKDATA_SIZE = 1024;
constexpr size_t ROCSHMEM_BARRIER_SYNC_SIZE = 256;
constexpr size_t ROCSHMEM_REDUCE_SYNC_SIZE = 256;
// Internally calls sync function, which matches barrier implementation
constexpr size_t ROCSHMEM_BCAST_SYNC_SIZE = ROCSHMEM_BARRIER_SYNC_SIZE;
constexpr size_t ROCSHMEM_ALLTOALL_SYNC_SIZE = ROCSHMEM_BARRIER_SYNC_SIZE + 1;
constexpr size_t ROCSHMEM_FCOLLECT_SYNC_SIZE = ROCSHMEM_ALLTOALL_SYNC_SIZE;
constexpr size_t ROCSHMEM_SYNC_VALUE = 0;
const int ROCSHMEM_CTX_ZERO = 0;
const int ROCSHMEM_CTX_NOSTORE = 1;
const int ROCSHMEM_CTX_SERIALIZED = 2;
const int ROCSHMEM_CTX_WG_PRIVATE = 4;
const int ROCSHMEM_CTX_SHARED = 8;
/**
* @brief GPU side OpenSHMEM context created from each work-groups'
* rocshmem_wg_handle_t
*/
typedef struct rocshmem_ctx{
void *ctx_opaque;
void *team_opaque;
__host__ __device__ bool operator==(const struct rocshmem_ctx& other) const {
return (ctx_opaque == other.ctx_opaque &&
team_opaque == other.team_opaque);
}
__host__ __device__ bool operator!=(const struct rocshmem_ctx& other) const {
return !(*this == other);
}
} rocshmem_ctx_t;
/**
* Shmem default context.
*/
extern "C" __device__ rocshmem_ctx_t __attribute__((visibility("default"))) ROCSHMEM_CTX_DEFAULT;
/**
* A value corresponding to an invalid communication context. This value can be
* used to initialize or update context handles to indicate that they do not
* reference a valid context. When managed in this way, applications can use an
* equality comparison to test whether a given context handle references a
* valid context.
*/
extern __constant__ rocshmem_ctx_t ROCSHMEM_CTX_INVALID;
/**
* Used internally to set default context.
*/
void set_internal_ctx(rocshmem_ctx_t *ctx);
typedef uint64_t *rocshmem_team_t;
extern rocshmem_team_t ROCSHMEM_TEAM_WORLD;
const rocshmem_team_t ROCSHMEM_TEAM_INVALID = nullptr;
/**
* @brief Data structure defining the unqiueId
*/
/// Unique ID for a process. This is a ROCSHMEM_UNIQUE_ID_BYTES byte array that uniquely identifies a process.
#define ROCSHMEM_UNIQUE_ID_BYTES 128
using rocshmem_uniqueid_t = std::array<uint8_t, ROCSHMEM_UNIQUE_ID_BYTES>;
/**
* @brief Data structure used for attribute based
* initialization
*/
struct rocshmem_init_attr_t {
int32_t rank;
int32_t nranks;
rocshmem_uniqueid_t uid;
void* mpi_comm;
};
typedef struct rocshmem_init_attr_t rocshmem_init_attr_t;
constexpr unsigned int ROCSHMEM_INIT_WITH_MPI_COMM = 0;
constexpr unsigned int ROCSHMEM_INIT_WITH_UNIQUEID = 1;
} // namespace rocshmem
#endif // LIBRARY_INCLUDE_ROCSHMEM_COMMON_HPP
rocshmem_dir/include/rocshmem/rocshmem_config.h deleted 100644 → 0
View file @ b33659dd
This diff is collapsed.
rocshmem_dir/include/rocshmem/rocshmem_debug.hpp deleted 100644 → 0
View file @ b33659dd
This diff is collapsed.
rocshmem_dir/include/rocshmem/rocshmem_mpi.hpp deleted 100644 → 0
View file @ b33659dd
This diff is collapsed.
rocshmem_dir/share/doc/rocshmem/LICENSE.md deleted 100644 → 0
View file @ b33659dd
This diff is collapsed.
rocshmem @ f5a87af2
Subproject commit f5a87af2671b6daaea16ae766ca97db867ef996c
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment