Intel(R) SHMEM (ISHMEM) 1.5.1 (#14)

sys-shmem · web-flow · commit 3a0c65869e92 · 2025-11-12T13:08:30.000-08:00
Signed-off-by: sys_shmem &lt;sys_shmem@intel.com&gt;
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -13,7 +13,7 @@ set(PROJECT_FULL_NAME "Intel® SHMEM")
 
 set(ISHMEM_MAJOR_VERSION "1")
 set(ISHMEM_MINOR_VERSION "5")
-set(ISHMEM_PATCH_VERSION "0")
+set(ISHMEM_PATCH_VERSION "1")
 set(PROJECT_VERSION "${ISHMEM_MAJOR_VERSION}.${ISHMEM_MINOR_VERSION}.${ISHMEM_PATCH_VERSION}")
 
 project(${PROJECT_NAME} VERSION ${PROJECT_VERSION} LANGUAGES C CXX)
diff --git a/RELEASE_NOTES.md b/RELEASE_NOTES.md
@@ -1,6 +1,23 @@
 # Release Notes <!-- omit in toc -->
 This document contains a list of new features and known limitations of Intel® SHMEM releases.
 
+## Release 1.5.1
+
+### New Features and Enhancements
+- Fix a compilation issue due to missing files
+
+### Known Limitations
+- Only [Sandia OpenSHMEM](https://github.com/Sandia-OpenSHMEM/SOS) and [Intel® MPI Library](https://www.intel.com/content/www/us/en/developer/tools/oneapi/mpi-library.html) are currently supported as the host back-end.
+- Not all APIs from OpenSHMEM standard are supported. Please refer to [Supported/Unsupported Features](https://oneapi-src.github.io/ishmem/supported_features.html) to get a complete view.
+- Intel® SHMEM requires a one-to-one mapping of PEs to SYCL devices. This implies that Intel® SHMEM executions must launch with a number of processes on each compute node that is no more than the number of available SYCL devices on each one of those nodes. By default, the Intel® SHMEM runtime considers each individual device tile to make up a single SYCL device and assigns a tile per PE.
+- All collective operations within a kernel must complete before invoking subsequent kernel-initiated collective operation.
+- To run Intel® SHMEM with SOS enabling the Slingshot provider in OFI, environment variable `FI_CXI_OPTIMIZED_MRS=0` must be used. It is also recommended to use `FI_CXI_DEFAULT_CQ_SIZE=131072`.
+- To run Intel® SHMEM with SOS enabling the verbs provider, environment variable `MLX5_SCATTER_TO_CQE=0` must be used.
+- To run Intel® SHMEM with Intel® MPI Library, environment variable `I_MPI_OFFLOAD=1` must be used. Additionally, `I_MPI_OFFLOAD_RDMA=1` may be necessary for GPU RDMA depending on the OFI provider. Please refer to the [reference guide](https://www.intel.com/content/www/us/en/docs/mpi-library/developer-reference-linux/2021-16/gpu-buffers-support.html) for further details.
+- Inter-node communication in Intel® SHMEM requires [dma-buf](https://www.kernel.org/doc/html/latest/driver-api/dma-buf.html) support in the Linux kernel. Inter-node functionality in Intel® SHMEM Release 1.5.0 is tested with SUSE Linux Enterprise Server 15 SP4.
+- Support for Intel® Arc™ B-Series GPUs is preliminary. As such, not all APIs are currently supported.
+- When using Intel® Arc™ B-Series GPUs, environment variable `RenderCompressedBuffersEnabled=0` is required. This is automatically set when running with the launcher script `ishmrun`.
+
 ## Release 1.5.0
 
 ### New Features and Enhancements
diff --git a/docs/source/conf.py b/docs/source/conf.py
@@ -16,7 +16,7 @@
 project = u'Intel® SHMEM'
 copyright = u'2024 Intel Corporation licensed under Creative Commons BY 4.0'
 author = u'Intel Corporation'
-release = u'1.5.0'
+release = u'1.5.1'
 version = release
 
 # -- General configuration ---------------------------------------------------
diff --git a/src/collectives/scan.cpp b/src/collectives/scan.cpp
@@ -0,0 +1,74 @@
+/* Copyright (C) 2025 Intel Corporation
+ * SPDX-License-Identifier: BSD-3-Clause
+ */
+
+#include "collectives/scan_impl.h"
+
+/* clang-format off */
+#define ISHMEMI_API_IMPL_INSCAN(TYPENAME, TYPE)                                                                                                                                             \
+    int ishmem_##TYPENAME##_sum_inscan(TYPE *dest, const TYPE *src, size_t nelems) { return ishmem_sum_inscan(ISHMEM_TEAM_WORLD, dest, src, nelems); }                                      \
+    int ishmem_##TYPENAME##_sum_inscan(ishmem_team_t team, TYPE *dest, const TYPE *src, size_t nelems) { return ishmem_sum_inscan(team, dest, src, nelems); }                               \
+    sycl::event ishmemx_##TYPENAME##_sum_inscan_on_queue(TYPE *dest, const TYPE *src, size_t nelems, int *ret, sycl::queue &q, const std::vector<sycl::event> &deps) {                      \
+        return ishmemx_sum_inscan_on_queue(ISHMEM_TEAM_WORLD, dest, src, nelems, ret, q, deps);                                                                                             \
+    }                                                                                                                                                                                       \
+    sycl::event ishmemx_##TYPENAME##_sum_inscan_on_queue(ishmem_team_t team, TYPE *dest, const TYPE *src, size_t nelems, int *ret, sycl::queue &q, const std::vector<sycl::event> &deps) {  \
+        return ishmemx_sum_inscan_on_queue(team, dest, src, nelems, ret, q, deps);                                                                                                          \
+    }
+#define ISHMEMI_API_IMPL_EXSCAN(TYPENAME, TYPE)                                                                                                                                             \
+    int ishmem_##TYPENAME##_sum_exscan(TYPE *dest, const TYPE *src, size_t nelems) { return ishmem_sum_exscan(ISHMEM_TEAM_WORLD, dest, src, nelems); }                                      \
+    int ishmem_##TYPENAME##_sum_exscan(ishmem_team_t team, TYPE *dest, const TYPE *src, size_t nelems) { return ishmem_sum_exscan(team, dest, src, nelems); }                               \
+    sycl::event ishmemx_##TYPENAME##_sum_exscan_on_queue(TYPE *dest, const TYPE *src, size_t nelems, int *ret, sycl::queue &q, const std::vector<sycl::event> &deps) {                      \
+        return ishmemx_sum_exscan_on_queue(ISHMEM_TEAM_WORLD, dest, src, nelems, ret, q, deps);                                                                                             \
+    }                                                                                                                                                                                       \
+    sycl::event ishmemx_##TYPENAME##_sum_exscan_on_queue(ishmem_team_t team, TYPE *dest, const TYPE *src, size_t nelems, int *ret, sycl::queue &q, const std::vector<sycl::event> &deps) {  \
+        return ishmemx_sum_exscan_on_queue(team, dest, src, nelems, ret, q, deps);                                                                                                          \
+    }
+/* clang-format on */
+
+ISHMEMI_API_IMPL_INSCAN(float, float)
+ISHMEMI_API_IMPL_INSCAN(double, double)
+ISHMEMI_API_IMPL_INSCAN(char, char)
+ISHMEMI_API_IMPL_INSCAN(schar, signed char)
+ISHMEMI_API_IMPL_INSCAN(short, short)
+ISHMEMI_API_IMPL_INSCAN(int, int)
+ISHMEMI_API_IMPL_INSCAN(long, long)
+ISHMEMI_API_IMPL_INSCAN(longlong, long long)
+ISHMEMI_API_IMPL_INSCAN(uchar, unsigned char)
+ISHMEMI_API_IMPL_INSCAN(ushort, unsigned short)
+ISHMEMI_API_IMPL_INSCAN(uint, unsigned int)
+ISHMEMI_API_IMPL_INSCAN(ulong, unsigned long)
+ISHMEMI_API_IMPL_INSCAN(ulonglong, unsigned long long)
+ISHMEMI_API_IMPL_INSCAN(int8, int8_t)
+ISHMEMI_API_IMPL_INSCAN(int16, int16_t)
+ISHMEMI_API_IMPL_INSCAN(int32, int32_t)
+ISHMEMI_API_IMPL_INSCAN(int64, int64_t)
+ISHMEMI_API_IMPL_INSCAN(uint8, uint8_t)
+ISHMEMI_API_IMPL_INSCAN(uint16, uint16_t)
+ISHMEMI_API_IMPL_INSCAN(uint32, uint32_t)
+ISHMEMI_API_IMPL_INSCAN(uint64, uint64_t)
+ISHMEMI_API_IMPL_INSCAN(size, size_t)
+ISHMEMI_API_IMPL_INSCAN(ptrdiff, ptrdiff_t)
+
+ISHMEMI_API_IMPL_EXSCAN(float, float)
+ISHMEMI_API_IMPL_EXSCAN(double, double)
+ISHMEMI_API_IMPL_EXSCAN(char, char)
+ISHMEMI_API_IMPL_EXSCAN(schar, signed char)
+ISHMEMI_API_IMPL_EXSCAN(short, short)
+ISHMEMI_API_IMPL_EXSCAN(int, int)
+ISHMEMI_API_IMPL_EXSCAN(long, long)
+ISHMEMI_API_IMPL_EXSCAN(longlong, long long)
+ISHMEMI_API_IMPL_EXSCAN(uchar, unsigned char)
+ISHMEMI_API_IMPL_EXSCAN(ushort, unsigned short)
+ISHMEMI_API_IMPL_EXSCAN(uint, unsigned int)
+ISHMEMI_API_IMPL_EXSCAN(ulong, unsigned long)
+ISHMEMI_API_IMPL_EXSCAN(ulonglong, unsigned long long)
+ISHMEMI_API_IMPL_EXSCAN(int8, int8_t)
+ISHMEMI_API_IMPL_EXSCAN(int16, int16_t)
+ISHMEMI_API_IMPL_EXSCAN(int32, int32_t)
+ISHMEMI_API_IMPL_EXSCAN(int64, int64_t)
+ISHMEMI_API_IMPL_EXSCAN(uint8, uint8_t)
+ISHMEMI_API_IMPL_EXSCAN(uint16, uint16_t)
+ISHMEMI_API_IMPL_EXSCAN(uint32, uint32_t)
+ISHMEMI_API_IMPL_EXSCAN(uint64, uint64_t)
+ISHMEMI_API_IMPL_EXSCAN(size, size_t)
+ISHMEMI_API_IMPL_EXSCAN(ptrdiff, ptrdiff_t)
diff --git a/src/collectives/scan_impl.h b/src/collectives/scan_impl.h
@@ -0,0 +1,98 @@
+/* Copyright (C) 2025 Intel Corporation
+ * SPDX-License-Identifier: BSD-3-Clause
+ */
+
+#ifndef COLLECTIVES_SCAN_IMPL_H
+#define COLLECTIVES_SCAN_IMPL_H
+
+#include "collectives.h"
+#include "runtime.h"
+#include "proxy_impl.h"
+#include "on_queue.h"
+
+template <typename T, ishmemi_op_t OP>
+int scan_impl(ishmem_team_t team, T *dest, const T *src, size_t nelems)
+{
+    if constexpr (enable_error_checking) {
+        validate_parameters((void *) dest, (void *) src, nelems * sizeof(T));
+    }
+
+    ishmemi_request_t req;
+    req.src = src;
+    req.dst = dest;
+    req.nelems = nelems;
+    req.op = OP;
+    req.type = ishmemi_union_get_base_type<T, OP>();
+    req.team = team;
+
+#ifdef __SYCL_DEVICE_ONLY__
+    return ishmemi_proxy_blocking_request_status(req);
+#else
+    ishmemi_ringcompletion_t comp;
+    ishmemi_runtime->proxy_funcs[req.op][req.type](&req, &comp);
+    return ishmemi_proxy_get_status(comp.completion.ret);
+#endif
+}
+
+template <typename T>
+int ishmem_sum_inscan(ishmem_team_t team, T *dest, const T *src, size_t nelems)
+{
+    return scan_impl<T, INSCAN>(team, dest, src, nelems);
+}
+
+template <typename T>
+int ishmem_sum_exscan(ishmem_team_t team, T *dest, const T *src, size_t nelems)
+{
+    return scan_impl<T, EXSCAN>(team, dest, src, nelems);
+}
+
+template <typename T, ishmemi_op_t OP>
+sycl::event scan_on_queue_impl(ishmem_team_t team, T *dest, const T *src, size_t nelems, int *ret,
+                               sycl::queue &q, const std::vector<sycl::event> &deps)
+{
+    bool entry_already_exists = true;
+    const std::lock_guard<std::mutex> lock(ishmemi_on_queue_events_map.map_mtx);
+    auto iter = ishmemi_on_queue_events_map.get_entry_info(q, entry_already_exists);
+
+    auto e = q.submit([&](sycl::handler &cgh) {
+        set_cmd_grp_dependencies(cgh, entry_already_exists, iter->second->event, deps);
+        cgh.single_task([=]() {
+            int tmp_ret = scan_impl<T, OP>(team, dest, src, nelems);
+            if (ret) *ret = tmp_ret;
+        });
+    });
+    ishmemi_on_queue_events_map[&q]->event = e;
+    return e;
+}
+
+template <typename T>
+sycl::event ishmemx_sum_inscan_on_queue(ishmem_team_t team, T *dest, const T *src, size_t nelems,
+                                        int *ret, sycl::queue &q,
+                                        const std::vector<sycl::event> &deps)
+{
+    return scan_on_queue_impl<T, INSCAN>(team, dest, src, nelems, ret, q, deps);
+}
+
+template <typename T>
+sycl::event ishmemx_sum_inscan_on_queue(T *dest, const T *src, size_t nelems, int *ret,
+                                        sycl::queue &q, const std::vector<sycl::event> &deps)
+{
+    return scan_on_queue_impl<T, INSCAN>(ISHMEM_TEAM_WORLD, dest, src, nelems, ret, q, deps);
+}
+
+template <typename T>
+sycl::event ishmemx_sum_exscan_on_queue(ishmem_team_t team, T *dest, const T *src, size_t nelems,
+                                        int *ret, sycl::queue &q,
+                                        const std::vector<sycl::event> &deps)
+{
+    return scan_on_queue_impl<T, EXSCAN>(team, dest, src, nelems, ret, q, deps);
+}
+
+template <typename T>
+sycl::event ishmemx_sum_exscan_on_queue(T *dest, const T *src, size_t nelems, int *ret,
+                                        sycl::queue &q, const std::vector<sycl::event> &deps)
+{
+    return scan_on_queue_impl<T, EXSCAN>(ISHMEM_TEAM_WORLD, dest, src, nelems, ret, q, deps);
+}
+
+#endif
diff --git a/src/collectives/sync_impl.h b/src/collectives/sync_impl.h
@@ -0,0 +1,71 @@
+/* Copyright (C) 2025 Intel Corporation
+ * SPDX-License-Identifier: BSD-3-Clause
+ */
+
+#ifndef COLLECTIVES_SYNC_IMPL_H
+#define COLLECTIVES_SYNC_IMPL_H
+
+#include "ishmem/types.h"
+#include "proxy_impl.h"
+#include "collectives.h"
+#include "runtime.h"
+#include "teams.h"
+
+static inline void sync_team_fallback(ishmem_team_t team)
+{
+    ishmemi_request_t req;
+    ishmemi_ringcompletion_t comp __attribute__((unused));
+    req.op = TEAM_SYNC;
+    req.type = NONE;
+    req.team = team;
+
+#ifdef __SYCL_DEVICE_ONLY__
+    ishmemi_proxy_blocking_request(req);
+    atomic_fence(sycl::memory_order::seq_cst, sycl::memory_scope::system);
+#else
+    ishmemi_runtime->proxy_funcs[req.op][req.type](&req, &comp);
+#endif
+}
+
+ISHMEM_DEVICE_ATTRIBUTES inline void ishmemi_team_sync(ishmem_team_t team)
+{
+    /* Node-local, on-device implementation */
+    if constexpr (ishmemi_is_device) {
+        ishmemi_info_t *info = global_info;
+        ishmemi_team_device_t *team_ptr = &info->team_device_pool[team];
+        if (team_ptr->only_intra) {
+            sycl::atomic_fence(sycl::memory_order::seq_cst, sycl::memory_scope::system);
+
+            int index = team_ptr->psync_idx;
+            long *my_psync = &team_ptr->psync[team_ptr->psync_idx];
+            team_ptr->psync_idx = (index + 1) % N_PSYNCS_PER_TEAM;
+
+            /* This atomic has to be seq_cst because we definitely want it to happen in order */
+            sycl::atomic_ref<long, sycl::memory_order::seq_cst, sycl::memory_scope::system,
+                             sycl::access::address_space::global_space>
+                atomic_psync(*my_psync);
+
+            for (int i = team_ptr->start; i <= team_ptr->last_pe; i += team_ptr->stride) {
+                uint8_t local_index = ISHMEMI_LOCAL_PES[i];
+                long *remote_psync = ISHMEMI_FAST_ADJUST(long, info, local_index, my_psync);
+
+                /* These atomics can be relaxed because we don't care about their ordering */
+                sycl::atomic_ref<long, sycl::memory_order::relaxed, sycl::memory_scope::system,
+                                 sycl::access::address_space::global_space>
+                    atomic_psync(*remote_psync);
+                atomic_psync += 1L;
+            }
+
+            while (atomic_psync.load() != team_ptr->size)
+                ;
+            atomic_psync.store(0);
+
+            return;
+        }
+    }
+
+    /* Otherwise */
+    sync_team_fallback(team);
+}
+
+#endif
diff --git a/src/ishmem.h b/src/ishmem.h
@@ -16,7 +16,7 @@
 
 #define ISHMEM_MAJOR_VERSION 1
 #define ISHMEM_MINOR_VERSION 5
-#define ISHMEM_PATCH_VERSION 0
+#define ISHMEM_PATCH_VERSION 1
 #define ISHMEM_MAX_NAME_LEN  256
 #define ISHMEM_VENDOR_STRING "Intel® SHMEM"
 
diff --git a/test/unit/exscan.cpp b/test/unit/exscan.cpp
@@ -0,0 +1,78 @@
+/* Copyright (C) 2025 Intel Corporation
+ * SPDX-License-Identifier: BSD-3-Clause
+ */
+
+#include <ishmem_tester.h>
+#include "rma_test.h"
+
+#undef TEST_BRANCH_ON_QUEUE
+
+#define TEST_BRANCH_SINGLE(testname, typeenum, typename, type, op, opname)                         \
+    *res = ishmem_##typename##_##testname(ISHMEM_TEAM_WORLD, (type *) dest, (type *) src, nelems);
+
+#define TEST_BRANCH_ON_QUEUE(testname, typeenum, typename, type, op, opname)                       \
+    ishmemx_##typename##_##testname##_on_queue((type *) dest, (type *) src, nelems, res, q);
+
+GEN_HOST_FNS(sum_exscan, host, NOP, nop)
+GEN_ON_QUEUE_FNS(sum_exscan, on_queue, NOP, nop)
+GEN_SINGLE_FNS(sum_exscan, single, NOP, nop)
+
+class sum_exscan_tester : public ishmem_tester {
+  public:
+    sum_exscan_tester(int argc, char *argv[]) : ishmem_tester(argc, argv, true) {}
+    virtual size_t create_source_pattern(ishmemi_type_t t, ishmemi_op_t op, testmode_t mode,
+                                         size_t nelems);
+    virtual size_t create_check_pattern(ishmemi_type_t t, ishmemi_op_t op, testmode_t mode,
+                                        size_t nelems);
+};
+
+/* result written into aligned_source, using host_source as a temp buffer if needed */
+size_t sum_exscan_tester::create_source_pattern(ishmemi_type_t t, ishmemi_op_t op, testmode_t mode,
+                                                size_t nelems)
+{
+    size_t test_size = nelems * typesize(t);
+    size_t from_pe = (size_t) my_pe;
+    for (size_t idx = 0; idx < ((test_size / sizeof(long)) + 1); idx += 1) {
+        aligned_source[idx] = (long) from_pe + (long) idx;
+        if (patterndebugflag && (idx < 16)) {
+            printf("[%d] source pattern idx %lu val %016lx\n", my_pe, idx, aligned_source[idx]);
+        }
+    }
+
+    return (test_size);
+}
+
+/* check pattern written into host_check, using host_source as a temp buffer */
+size_t sum_exscan_tester::create_check_pattern(ishmemi_type_t t, ishmemi_op_t op, testmode_t mode,
+                                               size_t nelems)
+{
+    size_t test_size = nelems * typesize(t);
+    size_t to_pe = (size_t) my_pe;
+    for (size_t idx = 0; idx < ((test_size / sizeof(long)) + 1); idx += 1) {
+        host_check[idx] = ((to_pe + idx) * (to_pe + idx - 1) / 2 - (idx * (idx - 1)) / 2);
+        if (patterndebugflag && (idx < 16)) {
+            printf("[%d] check pattern idx %lu val %016lx\n", my_pe, idx, host_check[idx]);
+        }
+    }
+
+    return (test_size);
+}
+
+int main(int argc, char **argv)
+{
+    class sum_exscan_tester t(argc, argv);
+
+    size_t bufsize = (t.max_nelems * t.typesize(SIZE128) * (size_t) t.n_pes) + 4096L;
+    t.alloc_memory(bufsize);
+    size_t errors = 0;
+
+    GEN_FN_TABLE(sum_exscan, host, NOP, nop)
+    GEN_FN_TABLE(sum_exscan, on_queue, NOP, nop)
+    GEN_FN_TABLE(sum_exscan, single, NOP, nop)
+
+    if (!t.test_types_set) t.add_test_type_list(scan_types);
+    errors += t.run_aligned_tests(NOP);
+    errors += t.run_offset_tests(NOP);
+
+    return (t.finalize_and_report(errors));
+}
diff --git a/test/unit/inscan.cpp b/test/unit/inscan.cpp
