diff --git a/projects/rocshmem/src/context.hpp b/projects/rocshmem/src/context.hpp index 619b8a343c..5eb82ef007 100644 --- a/projects/rocshmem/src/context.hpp +++ b/projects/rocshmem/src/context.hpp @@ -28,7 +28,6 @@ #include #include "backend_type.hpp" -#include "fence_policy.hpp" #include "host/host.hpp" #include "ipc_policy.hpp" #include "stats.hpp" @@ -55,9 +54,9 @@ class Backend; */ class Context { public: - __host__ Context(Backend* handle, bool shareable); + __host__ Context(Backend* handle); - __device__ Context(Backend* handle, bool shareable); + __device__ Context(Backend* handle); __host__ virtual ~Context(); @@ -459,13 +458,6 @@ class Context { __host__ int test(T *ivars, int cmp, T val); public: - /** - * @brief Set the fence policy using a runtime option - * - * @param[in] options interpreted as a bitfield using bitwise operations - */ - __device__ void setFence(long options) { fence_ = Fence(options); } - /************************************************************************** ***************************** PUBLIC MEMBERS ***************************** *************************************************************************/ @@ -504,11 +496,6 @@ class Context { */ WavefrontCoalescer wf_coal_{}; - /** - * @brief Controls fence behavior in device code - */ - Fence fence_{}; - public: /** * @brief Inter-Process Communication (IPC) interface for context class diff --git a/projects/rocshmem/src/context_device.cpp b/projects/rocshmem/src/context_device.cpp index 5836da5471..12016e15b3 100644 --- a/projects/rocshmem/src/context_device.cpp +++ b/projects/rocshmem/src/context_device.cpp @@ -29,10 +29,9 @@ namespace rocshmem { -__device__ Context::Context(Backend* handle, bool shareable) +__device__ Context::Context(Backend* handle) : num_pes(handle->getNumPEs()), my_pe(handle->getMyPE()), - fence_(shareable), btype(handle->type) { /* * Device-side context constructor is a work-group collective, so make diff --git a/projects/rocshmem/src/context_host.cpp b/projects/rocshmem/src/context_host.cpp index ead3485b45..653f6ddaf0 100644 --- a/projects/rocshmem/src/context_host.cpp +++ b/projects/rocshmem/src/context_host.cpp @@ -28,10 +28,9 @@ namespace rocshmem { -__host__ Context::Context(Backend* handle, bool shareable) +__host__ Context::Context(Backend* handle) : num_pes(handle->getNumPEs()), my_pe(handle->getMyPE()), - fence_(shareable), btype(handle->type) { } diff --git a/projects/rocshmem/src/fence_policy.hpp b/projects/rocshmem/src/fence_policy.hpp deleted file mode 100644 index 1e630c4716..0000000000 --- a/projects/rocshmem/src/fence_policy.hpp +++ /dev/null @@ -1,78 +0,0 @@ -/****************************************************************************** - * 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_SRC_FENCE_POLICY_HPP_ -#define LIBRARY_SRC_FENCE_POLICY_HPP_ - -#include "rocshmem/rocshmem.hpp" - -namespace rocshmem { - -/** - * @brief Controls the behavior of device code which may need to stall - */ -class Fence { - public: - /** - * Secondary constructor - */ - __host__ __device__ Fence() = default; - - /** - * Primary constructor - * - * @param[in] options interpreted as a bitfield using bitwise operations - */ - __host__ __device__ Fence(long option) { - if (option & ROCSHMEM_CTX_NOSTORE) { - flush_ = false; - } - } - - /** - * @brief Wait for outstanding memory operations to complete - * - * This can be useful when code needs guarantees about visibility - * before moving past the flush. - * - * @return void - */ - __device__ void flush() { - if (flush_) { - __threadfence(); - } - } - - private: - /** - * @brief Used to toggle flushes behavior on and off - * - * @note By default, flushing is enabled. - */ - bool flush_{true}; -}; - -} // namespace rocshmem - -#endif // LIBRARY_SRC_FENCE_POLICY_HPP_ diff --git a/projects/rocshmem/src/gda/context_gda_device.cpp b/projects/rocshmem/src/gda/context_gda_device.cpp index e179ba29cd..20876ebbba 100644 --- a/projects/rocshmem/src/gda/context_gda_device.cpp +++ b/projects/rocshmem/src/gda/context_gda_device.cpp @@ -35,7 +35,7 @@ namespace rocshmem { __host__ GDAContext::GDAContext(Backend *b, unsigned int ctx_id, int gda_provider) - : Context(b, false) { + : Context(b) { GDABackend *backend{static_cast(b)}; base_heap = backend->heap.get_heap_bases().data(); diff --git a/projects/rocshmem/src/gda/context_gda_host.cpp b/projects/rocshmem/src/gda/context_gda_host.cpp index 2345d0f9ad..61c1e59cec 100644 --- a/projects/rocshmem/src/gda/context_gda_host.cpp +++ b/projects/rocshmem/src/gda/context_gda_host.cpp @@ -34,7 +34,7 @@ namespace rocshmem { __host__ GDAHostContext::GDAHostContext(Backend *backend, [[maybe_unused]] int64_t options) - : Context(backend, true) { + : Context(backend) { GDABackend *b{static_cast(backend)}; host_interface = b->host_interface; diff --git a/projects/rocshmem/src/ipc/context_ipc_device.cpp b/projects/rocshmem/src/ipc/context_ipc_device.cpp index 6cda9876d2..bc0ef91d09 100644 --- a/projects/rocshmem/src/ipc/context_ipc_device.cpp +++ b/projects/rocshmem/src/ipc/context_ipc_device.cpp @@ -34,7 +34,7 @@ namespace rocshmem { __host__ IPCContext::IPCContext(Backend *b, unsigned int ctx_id) - : Context(b, false) { + : Context(b) { IPCBackend *backend{static_cast(b)}; ipcImpl_.ipc_bases = b->ipcImpl.ipc_bases; ipcImpl_.shm_size = b->ipcImpl.shm_size; diff --git a/projects/rocshmem/src/ipc/context_ipc_host.cpp b/projects/rocshmem/src/ipc/context_ipc_host.cpp index e30fa2c379..5714aa4dd8 100644 --- a/projects/rocshmem/src/ipc/context_ipc_host.cpp +++ b/projects/rocshmem/src/ipc/context_ipc_host.cpp @@ -34,7 +34,7 @@ namespace rocshmem { __host__ IPCHostContext::IPCHostContext(Backend *backend, [[maybe_unused]] int64_t options) - : Context(backend, true) { + : Context(backend) { IPCBackend *b{static_cast(backend)}; host_interface = b->host_interface; diff --git a/projects/rocshmem/src/reverse_offload/context_ro_device.cpp b/projects/rocshmem/src/reverse_offload/context_ro_device.cpp index 7d3e8a9a6c..a47243a958 100644 --- a/projects/rocshmem/src/reverse_offload/context_ro_device.cpp +++ b/projects/rocshmem/src/reverse_offload/context_ro_device.cpp @@ -46,7 +46,7 @@ namespace rocshmem { __host__ ROContext::ROContext(Backend *b, size_t block_id, bool default_ctx) - : Context(b, false), + : Context(b), is_default_ctx{default_ctx} { ROBackend *backend{static_cast(b)}; diff --git a/projects/rocshmem/src/reverse_offload/context_ro_host.cpp b/projects/rocshmem/src/reverse_offload/context_ro_host.cpp index f72399341b..87ef5cc768 100644 --- a/projects/rocshmem/src/reverse_offload/context_ro_host.cpp +++ b/projects/rocshmem/src/reverse_offload/context_ro_host.cpp @@ -34,7 +34,7 @@ namespace rocshmem { __host__ ROHostContext::ROHostContext(Backend *backend, long options) - : Context(backend, true) { + : Context(backend) { ROBackend *b{static_cast(backend)}; host_interface = b->host_interface; diff --git a/projects/rocshmem/src/rocshmem_gpu.cpp b/projects/rocshmem/src/rocshmem_gpu.cpp index 9ef0add371..24462c5531 100644 --- a/projects/rocshmem/src/rocshmem_gpu.cpp +++ b/projects/rocshmem/src/rocshmem_gpu.cpp @@ -323,10 +323,7 @@ __device__ int rocshmem_wg_ctx_create(long options, rocshmem_ctx_t *ctx) { if (get_flat_block_id() == 0) { ctx->team_opaque = reinterpret_cast(ROCSHMEM_CTX_DEFAULT.team_opaque); result = device_backend_proxy->create_ctx(options, ctx); - if(result) { - reinterpret_cast(ctx->ctx_opaque)->setFence(options); - } - else { + if (!result) { *ctx = ROCSHMEM_CTX_INVALID; } } @@ -348,10 +345,7 @@ __device__ int rocshmem_wg_team_create_ctx(rocshmem_team_t team, long options, TeamInfo *info_wrt_world = team_obj->tinfo_wrt_world; ctx->team_opaque = info_wrt_world; result = device_backend_proxy->create_ctx(options, ctx); - if(result) { - reinterpret_cast(ctx->ctx_opaque)->setFence(options); - } - else { + if (!result) { *ctx = ROCSHMEM_CTX_INVALID; } } @@ -656,9 +650,9 @@ __global__ ATTR_NO_INLINE void rocshmem_alltoallmem_kernel(rocshmem_team_t team, // This allows parallel execution across multiple streams without serialization __shared__ rocshmem_ctx_t ctx; __shared__ int ctx_result; - + ctx_result = rocshmem_wg_team_create_ctx(team, 0, &ctx); - + // If context creation failed, fall back to default context if (ctx_result != 0) { ctx = ROCSHMEM_CTX_DEFAULT;