Remove unused fence policy (#348)

[ROCm/rocshmem commit: ecd4c9f561]

projects/rocshmem/src/context.hpp

@@ -28,7 +28,6 @@
 #include <hip/hip_runtime.h>
 
 #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

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

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) {
 }
 

projects/rocshmem/src/fence_policy.hpp

@@ -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_

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<GDABackend *>(b)};
   base_heap = backend->heap.get_heap_bases().data();
 

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<GDABackend *>(backend)};
 
   host_interface = b->host_interface;

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<IPCBackend *>(b)};
   ipcImpl_.ipc_bases = b->ipcImpl.ipc_bases;
   ipcImpl_.shm_size = b->ipcImpl.shm_size;

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<IPCBackend *>(backend)};
 
   host_interface = b->host_interface;

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<ROBackend *>(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<ROBackend *>(backend)};
 
   host_interface = b->host_interface;

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<TeamInfo *>(ROCSHMEM_CTX_DEFAULT.team_opaque);
     result = device_backend_proxy->create_ctx(options, ctx);
-    if(result) {
-      reinterpret_cast<Context *>(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<Context *>(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;