Add Python scripts to generate header files based on sections in OpenSHMEM specifications

utlis/header_files_gen/AMO.py

@@ -0,0 +1,639 @@
+"""
+******************************************************************************
+ * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *****************************************************************************
+"""
+
+import os
+
+types = [
+    ("int", "int"),
+    ("long", "long"),
+    ("long long", "longlong"),
+    ("unsigned int", "uint"),
+    ("unsigned long", "ulong"),
+    ("unsigned long long", "ulonglong"),
+    ("int32_t", "int32"),
+    ("int64_t", "int64"),
+    ("uint32_t", "uint32"),
+    ("uint64_t", "uint64"),
+    ("size_t", "size"),
+    ("ptrdiff_t", "ptrdiff"),
+]
+
+
+float_types = [
+    ("float", "float"),
+    ("double", "double"),
+]
+
+bitwise_types = types[3:10]
+
+
+def atomic_fetch_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE {T} rocshmem_ctx_{TNAME}_atomic_fetch(\n"
+        f"    rocshmem_ctx_t ctx, {T} *source, int pe);\n"
+        f"__device__ ATTR_NO_INLINE {T} rocshmem_{TNAME}_atomic_fetch(\n"
+        f"    {T} *source, int pe);\n"
+        f"__host__ {T} rocshmem_ctx_{TNAME}_atomic_fetch(\n"
+        f"    rocshmem_ctx_t ctx, {T} *source, int pe);\n"
+        f"__host__ {T} rocshmem_{TNAME}_atomic_fetch(\n"
+        f"    {T} *source, int pe);\n\n"
+    )
+
+
+def generate_atomic_fetch_api():
+    expanded_code = """
+/**
+ * @name SHMEM_ATOMIC_FETCH
+ * @brief Atomically return the value of \p dest to the calling PE.
+ *
+ * The operation is blocking.
+ *
+ * This function can be called from divergent control paths at per-thread
+ * granularity.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+                      heap.
+ * @param[in] val     The value to be atomically added.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return            The value of \p dest.
+ */\n"""
+
+    for type_, tname_ in float_types:
+        expanded_code += atomic_fetch_api(type_, tname_)
+
+    for type_, tname_ in types:
+        expanded_code += atomic_fetch_api(type_, tname_)
+
+    return expanded_code
+
+
+def atomic_set_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_atomic_set(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__device__ ATTR_NO_INLINE void rocshmem_{TNAME}_atomic_set(\n"
+        f"    {T} *dest, {T} value, int pe);\n"
+        f"__host__ void rocshmem_ctx_{TNAME}_atomic_set(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__host__ void rocshmem_{TNAME}_atomic_set(\n"
+        f"    {T} *dest, {T} value, int pe);\n\n"
+    )
+
+
+def generate_atomic_set_api():
+    expanded_code = """
+/**
+ * @name SHMEM_ATOMIC_SET
+ * @brief Atomically set the value \p val to \p dest on \p pe.
+ *
+ * The operation is blocking.
+ *
+ * This function can be called from divergent control paths at per-thread
+ * granularity.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+                      heap.
+ * @param[in] val     The value to be atomically added.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return void
+ */\n"""
+
+    for type_, tname_ in float_types:
+        expanded_code += atomic_set_api(type_, tname_)
+
+    for type_, tname_ in types:
+        expanded_code += atomic_set_api(type_, tname_)
+
+    return expanded_code
+
+
+def atomic_compare_swap_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE {T} rocshmem_ctx_{TNAME}_atomic_compare_swap(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} cond, {T} value, int pe);\n"
+        f"__device__ ATTR_NO_INLINE {T} rocshmem_{TNAME}_atomic_compare_swap(\n"
+        f"    {T} *dest, {T} cond, {T} value, int pe);\n"
+        f"__host__ {T} rocshmem_ctx_{TNAME}_atomic_compare_swap(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} cond, {T} value, int pe);\n"
+        f"__host__ {T} rocshmem_{TNAME}_atomic_compare_swap(\n"
+        f"    {T} *dest, {T} cond, {T} value, int pe);\n\n"
+    )
+
+
+def generate_atomic_compare_swap_api():
+    expanded_code = """
+/**
+ * @name SHMEM_ATOMIC_COMPARE_SWAP
+ * @brief Atomically compares if the value in \p dest with \p cond is equal
+ * then put \p val in \p dest. The operation returns the older value of \p dest
+ * to the calling PE.
+ *
+ * The operation is blocking.
+ *
+ * This function can be called from divergent control paths at per-thread
+ * granularity.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+                      heap.
+ * @param[in] cond    The value to be compare with.
+ * @param[in] val     The value to be atomically swapped.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return            The old value of \p dest.
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += atomic_compare_swap_api(type_, tname_)
+
+    return expanded_code
+
+
+def atomic_swap_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE {T} rocshmem_ctx_{TNAME}_atomic_swap(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__device__ ATTR_NO_INLINE {T} rocshmem_{TNAME}_atomic_swap(\n"
+        f"    {T} *dest, {T} value, int pe);\n"
+        f"__host__ {T} rocshmem_ctx_{TNAME}_atomic_swap(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__host__ {T} rocshmem_{TNAME}_atomic_swap(\n"
+        f"    {T} *dest, {T} value, int pe);\n\n"
+    )
+
+
+def generate_atomic_swap_api():
+    expanded_code = """
+/**
+ * @name SHMEM_ATOMIC_SWAP
+ * @brief Atomically swap the value \p val to \p dest on \p pe.
+ *
+ * The operation is blocking.
+ *
+ * This function can be called from divergent control paths at per-thread
+ * granularity.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+                      heap.
+ * @param[in] val     The value to be atomically added.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return original value
+ */\n"""
+
+    for type_, tname_ in float_types:
+        expanded_code += atomic_swap_api(type_, tname_)
+
+    for type_, tname_ in types:
+        expanded_code += atomic_swap_api(type_, tname_)
+
+    return expanded_code
+
+
+def atomic_fetch_inc_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE {T} rocshmem_ctx_{TNAME}_atomic_fetch_inc(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, int pe);\n"
+        f"__device__ ATTR_NO_INLINE {T} rocshmem_{TNAME}_atomic_fetch_inc(\n"
+        f"    {T} *dest, int pe);\n"
+        f"__host__ {T} rocshmem_ctx_{TNAME}_atomic_fetch_inc(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, int pe);\n"
+        f"__host__ {T} rocshmem_{TNAME}_atomic_fetch_inc(\n"
+        f"    {T} *dest, int pe);\n\n"
+    )
+
+
+def generate_atomic_fetch_inc_api():
+    expanded_code = """
+/**
+ * @name SHMEM_ATOMIC_FETCH_INC
+ * @brief Atomically add 1 to \p dest on \p pe. The operation
+ * returns the older value of \p dest to the calling PE.
+ *
+ * The operation is blocking.
+ *
+ * This function can be called from divergent control paths at per-thread
+ * granularity.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+                      heap.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return            The old value of \p dest before it was incremented by 1.
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += atomic_fetch_inc_api(type_, tname_)
+
+    return expanded_code
+
+
+def atomic_inc_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_atomic_inc(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, int pe);\n"
+        f"__device__ ATTR_NO_INLINE void rocshmem_{TNAME}_atomic_inc(\n"
+        f"    {T} *dest, int pe);\n"
+        f"__host__ void rocshmem_ctx_{TNAME}_atomic_inc(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, int pe);\n"
+        f"__host__ void rocshmem_{TNAME}_atomic_inc(\n"
+        f"    {T} *dest, int pe);\n\n"
+    )
+
+
+def generate_atomic_inc_api():
+    expanded_code = """
+/**
+ * @name SHMEM_ATOMIC_INC
+ * @brief Atomically add 1 to \p dest on \p pe.
+ *
+ * The operation is blocking.
+ *
+ * This function can be called from divergent control paths at per-thread
+ * granularity.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+                      heap.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return void
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += atomic_inc_api(type_, tname_)
+
+    return expanded_code
+
+
+def atomic_fetch_add_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE {T} rocshmem_ctx_{TNAME}_atomic_fetch_add(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__device__ ATTR_NO_INLINE {T} rocshmem_{TNAME}_atomic_fetch_add(\n"
+        f"    {T} *dest, {T} value, int pe);\n"
+        f"__host__ {T} rocshmem_ctx_{TNAME}_atomic_fetch_add(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__host__ {T} rocshmem_{TNAME}_atomic_fetch_add(\n"
+        f"    {T} *dest, {T} value, int pe);\n\n"
+    )
+
+
+def generate_atomic_fetch_add_api():
+    expanded_code = """
+/**
+ * @name SHMEM_ATOMIC_FETCH_ADD
+ * @brief Atomically add the value \p val to \p dest on \p pe. The operation
+ * returns the older value of \p dest to the calling PE.
+ *
+ * The operation is blocking.
+ *
+ * This function can be called from divergent control paths at per-thread
+ * granularity.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+                      heap.
+ * @param[in] val     The value to be atomically added.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return            The old value of \p dest before the \p val was added.
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += atomic_fetch_add_api(type_, tname_)
+
+    return expanded_code
+
+
+def atomic_add_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_atomic_add(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__device__ ATTR_NO_INLINE void rocshmem_{TNAME}_atomic_add(\n"
+        f"    {T} *dest, {T} value, int pe);\n"
+        f"__host__ void rocshmem_ctx_{TNAME}_atomic_add(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__host__ void rocshmem_{TNAME}_atomic_add(\n"
+        f"    {T} *dest, {T} value, int pe);\n\n"
+    )
+
+
+def generate_atomic_add_api():
+    expanded_code = """
+/**
+ * @name SHMEM_ATOMIC_ADD
+ * @brief Atomically add the value \p val to \p dest on \p pe.
+ *
+ * The operation is blocking.
+ *
+ * This function can be called from divergent control paths at per-thread
+ * granularity.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+                      heap.
+ * @param[in] val     The value to be atomically added.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return void
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += atomic_add_api(type_, tname_)
+
+    return expanded_code
+
+
+def atomic_fetch_and_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE {T} rocshmem_ctx_{TNAME}_atomic_fetch_and(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__device__ ATTR_NO_INLINE {T} rocshmem_{TNAME}_atomic_fetch_and(\n"
+        f"    {T} *dest, {T} value, int pe);\n"
+        f"__host__ {T} rocshmem_ctx_{TNAME}_atomic_fetch_and(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__host__ {T} rocshmem_{TNAME}_atomic_fetch_and(\n"
+        f"    {T} *dest, {T} value, int pe);\n\n"
+    )
+
+
+def generate_atomic_fetch_and_api():
+    expanded_code = """
+/**
+ * @name SHMEM_ATOMIC_FETCH_AND
+ * @brief Atomically bitwise-and the value \p val to \p dest on \p pe.
+ *
+ * The operation is blocking.
+ *
+ * This function can be called from divergent control paths at per-thread
+ * granularity.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+                      heap.
+ * @param[in] val     The value to be atomically added.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return original value
+ */\n"""
+    for type_, tname_ in bitwise_types:
+        expanded_code += atomic_fetch_and_api(type_, tname_)
+
+    return expanded_code
+
+
+def atomic_and_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_atomic_and(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__device__ ATTR_NO_INLINE void rocshmem_{TNAME}_atomic_and(\n"
+        f"    {T} *dest, {T} value, int pe);\n"
+        f"__host__ void rocshmem_ctx_{TNAME}_atomic_and(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__host__ void rocshmem_{TNAME}_atomic_and(\n"
+        f"    {T} *dest, {T} value, int pe);\n\n"
+    )
+
+
+def generate_atomic_and_api():
+    expanded_code = """
+/**
+ * @name SHMEM_ATOMIC_AND
+ * @brief Atomically bitwise-and the value \p val to \p dest on \p pe.
+ *
+ * The operation is blocking.
+ *
+ * This function can be called from divergent control paths at per-thread
+ * granularity.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+                      heap.
+ * @param[in] val     The value to be atomically added.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return void
+ */\n"""
+    for type_, tname_ in bitwise_types:
+        expanded_code += atomic_and_api(type_, tname_)
+
+    return expanded_code
+
+
+def atomic_fetch_or_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE {T} rocshmem_ctx_{TNAME}_atomic_fetch_or(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__device__ ATTR_NO_INLINE {T} rocshmem_{TNAME}_atomic_fetch_or(\n"
+        f"    {T} *dest, {T} value, int pe);\n"
+        f"__host__ {T} rocshmem_ctx_{TNAME}_atomic_fetch_or(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__host__ {T} rocshmem_{TNAME}_atomic_fetch_or(\n"
+        f"    {T} *dest, {T} value, int pe);\n\n"
+    )
+
+
+def generate_atomic_fetch_or_api():
+    expanded_code = """
+/**
+ * @name SHMEM_ATOMIC_FETCH_OR
+ * @brief Atomically bitwise-or the value \p val to \p dest on \p pe.
+ *
+ * The operation is blocking.
+ *
+ * This function can be called from divergent control paths at per-thread
+ * granularity.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+                      heap.
+ * @param[in] val     The value to be atomically added.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return original value
+ */\n"""
+    for type_, tname_ in bitwise_types:
+        expanded_code += atomic_fetch_or_api(type_, tname_)
+
+    return expanded_code
+
+
+def atomic_or_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_atomic_or(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__device__ ATTR_NO_INLINE void rocshmem_{TNAME}_atomic_or(\n"
+        f"    {T} *dest, {T} value, int pe);\n"
+        f"__host__ void rocshmem_ctx_{TNAME}_atomic_or(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__host__ void rocshmem_{TNAME}_atomic_or(\n"
+        f"    {T} *dest, {T} value, int pe);\n\n"
+    )
+
+
+def generate_atomic_or_api():
+    expanded_code = """
+/**
+ * @name SHMEM_ATOMIC_OR
+ * @brief Atomically bitwise-or the value \p val to \p dest on \p pe.
+ *
+ * The operation is blocking.
+ *
+ * This function can be called from divergent control paths at per-thread
+ * granularity.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+                      heap.
+ * @param[in] val     The value to be atomically added.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return void
+ */\n"""
+    for type_, tname_ in bitwise_types:
+        expanded_code += atomic_or_api(type_, tname_)
+
+    return expanded_code
+
+
+def atomic_fetch_xor_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE {T} rocshmem_ctx_{TNAME}_atomic_fetch_xor(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__device__ ATTR_NO_INLINE {T} rocshmem_{TNAME}_atomic_fetch_xor(\n"
+        f"    {T} *dest, {T} value, int pe);\n"
+        f"__host__ {T} rocshmem_ctx_{TNAME}_atomic_fetch_xor(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__host__ {T} rocshmem_{TNAME}_atomic_fetch_xor(\n"
+        f"    {T} *dest, {T} value, int pe);\n\n"
+    )
+
+
+def generate_atomic_fetch_xor_api():
+    expanded_code = """
+/**
+ * @name SHMEM_ATOMIC_FETCH_XOR
+ * @brief Atomically bitwise-xor the value \p val to \p dest on \p pe.
+ *
+ * The operation is blocking.
+ *
+ * This function can be called from divergent control paths at per-thread
+ * granularity.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+                      heap.
+ * @param[in] val     The value to be atomically added.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return original value
+ */\n"""
+    for type_, tname_ in bitwise_types:
+        expanded_code += atomic_fetch_xor_api(type_, tname_)
+
+    return expanded_code
+
+
+def atomic_xor_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_atomic_xor(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__device__ ATTR_NO_INLINE void rocshmem_{TNAME}_atomic_xor(\n"
+        f"    {T} *dest, {T} value, int pe);\n"
+        f"__host__ void rocshmem_ctx_{TNAME}_atomic_xor(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value, int pe);\n"
+        f"__host__ void rocshmem_{TNAME}_atomic_xor(\n"
+        f"    {T} *dest, {T} value, int pe);\n\n"
+    )
+
+
+def generate_atomic_xor_api():
+    expanded_code = """
+/**
+ * @name SHMEM_ATOMIC_XOR
+ * @brief Atomically bitwise-xor the value \p val to \p dest on \p pe.
+ *
+ * The operation is blocking.
+ *
+ * This function can be called from divergent control paths at per-thread
+ * granularity.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+                      heap.
+ * @param[in] val     The value to be atomically added.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return void
+ */\n"""
+    for type_, tname_ in bitwise_types:
+        expanded_code += atomic_xor_api(type_, tname_)
+
+    return expanded_code
+
+def write_to_file(filename, content):
+    with open(filename, 'w') as file:
+        file.write(content)
+
+
+def generate_AMO_header(output_dir, copyright):
+    expanded_code = copyright
+
+    expanded_code += """
+#ifndef LIBRARY_INCLUDE_ROCSHMEM_AMO_HPP
+#define LIBRARY_INCLUDE_ROCSHMEM_AMO_HPP
+
+namespace rocshmem {
+"""
+
+    expanded_code += (
+        generate_atomic_fetch_api() +
+        generate_atomic_set_api() +
+        generate_atomic_compare_swap_api() +
+        generate_atomic_swap_api() +
+        generate_atomic_fetch_inc_api() +
+        generate_atomic_inc_api() +
+        generate_atomic_fetch_add_api() +
+        generate_atomic_add_api() +
+        generate_atomic_fetch_and_api() +
+        generate_atomic_and_api() +
+        generate_atomic_fetch_or_api() +
+        generate_atomic_or_api() +
+        generate_atomic_fetch_xor_api() +
+        generate_atomic_xor_api()
+    )
+
+    expanded_code += """
+}  // namespace rocshmem
+
+#endif  // LIBRARY_INCLUDE_ROCSHMEM_AMO_HPP
+"""
+
+    output_file = os.path.join(
+        output_dir, 'rocshmem_AMO.hpp'
+    )
+
+    write_to_file(output_file, expanded_code)

utlis/header_files_gen/COLL.py

@@ -0,0 +1,246 @@
+"""
+******************************************************************************
+ * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *****************************************************************************
+"""
+
+import os
+
+types = [
+    ("float", "float"),
+    ("double", "double"),
+    ("char", "char"),
+    ("signed char", "schar"),
+    ("short", "short"),
+    ("int", "int"),
+    ("long", "long"),
+    ("long long", "longlong"),
+    ("unsigned char", "uchar"),
+    ("unsigned short", "ushort"),
+    ("unsigned int", "uint"),
+    ("unsigned long", "ulong"),
+    ("unsigned long long", "ulonglong"),
+]
+
+
+def alltoall_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_wg_alltoall(\n"
+        f"    rocshmem_ctx_t ctx, rocshmem_team_t team, {T} *dest,\n"
+        f"    const {T} *source, int nelems);\n\n"
+    )
+
+
+def generate_alltoall_api():
+    expanded_code = """
+/**
+ * @name SHMEM_ALLTOALL
+ * @brief Exchanges a fixed amount of contiguous data blocks between all pairs
+ * of PEs participating in the collective routine.
+ *
+ * This function must be called as a work-group collective.
+ *
+ * @param[in] team         The team participating in the collective.
+ * @param[in] dest         Destination address. Must be an address on the
+ *                         symmetric heap.
+ * @param[in] source       Source address. Must be an address on the symmetric
+                           heap.
+ * @param[in] nelems       Number of data blocks transferred per pair of PEs.
+ *
+ * @return void
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += alltoall_api(type_, tname_)
+
+    return expanded_code
+
+
+def broadcast_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_wg_broadcast(\n"
+        f"    rocshmem_ctx_t ctx, rocshmem_team_t team, {T} *dest,\n"
+        f"    const {T} *source, int nelems, int pe_root);\n"
+        f"__host__ void rocshmem_ctx_{TNAME}_broadcast(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, const {T} *source,\n"
+        f"    int nelems, int pe_root, int pe_start, int log_pe_stride,\n"
+        f"    int pe_size, long *p_sync);\n"
+        f"__host__ void rocshmem_ctx_{TNAME}_broadcast(\n"
+        f"    rocshmem_ctx_t ctx, rocshmem_team_t team, {T} *dest,\n"
+        f"    const {T} *source, int nelems, int pe_root);\n\n"
+    )
+
+
+def generate_broadcast_api():
+    expanded_code = """
+/**
+ * @name SHMEM_BROADCAST
+ * @brief Perform a broadcast between PEs in the active set. The caller
+ * is blocked until the broadcase completes.
+ *
+ * This function must be called as a work-group collective.
+ *
+ * @param[in] dest         Destination address. Must be an address on the
+ *                         symmetric heap.
+ * @param[in] source       Source address. Must be an address on the symmetric
+                           heap.
+ * @param[in] nelement     Size of the buffer to participate in the broadcast.
+ * @param[in] PE_root      Zero-based ordinal of the PE, with respect to the
+                           active set, from which the data is copied
+ * @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 ROCSHMEM. Must
+                           be of size at least ROCSHMEM_REDUCE_SYNC_SIZE.
+ *
+ * @return void
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += broadcast_api(type_, tname_)
+
+    return expanded_code
+
+
+def fcollect_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_wg_fcollect(\n"
+        f"    rocshmem_ctx_t ctx, rocshmem_team_t team, {T} *dest,\n"
+        f"    const {T} *source, int nelems);\n\n"
+    )
+
+
+def generate_fcollect_api():
+    expanded_code = """
+/**
+ * @name SHMEM_FCOLLECT
+ * @brief Concatenates blocks of data from multiple PEs to an array in every
+ * PE participating in the collective routine.
+ *
+ * This function must be called as a work-group collective.
+ *
+ * @param[in] team         The team participating in the collective.
+ * @param[in] dest         Destination address. Must be an address on the
+ *                         symmetric heap.
+ * @param[in] source       Source address. Must be an address on the symmetric
+                           heap.
+ * @param[in] nelems       Number of data blocks in source array.
+ *
+ * @return void
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += fcollect_api(type_, tname_)
+
+    return expanded_code
+
+
+def reduction_api(T, TNAME, Op_API):
+    return (
+        f"__device__ ATTR_NO_INLINE int rocshmem_ctx_{TNAME}_{Op_API}_wg_reduce(\n"
+        f"    rocshmem_ctx_t ctx, rocshmem_team_t team, {T} *dest, const {T} *source,\n"
+        f"    int nreduce);\n"
+        f"__host__ int rocshmem_ctx_{TNAME}_{Op_API}_reduce(\n"
+        f"    rocshmem_ctx_t ctx, rocshmem_team_t team, {T} *dest, const {T} *source,\n"
+        f"    int nreduce);\n\n"
+    )
+
+
+def arith_reduction_api(T, TNAME):
+    operations = ["sum", "min", "max", "prod"]
+    return "".join([reduction_api(T, TNAME, op) for op in operations])
+
+def bitwise_reduction_api(T, TNAME):
+    operations = ["or", "and", "xor"]
+    return "".join([reduction_api(T, TNAME, op) for op in operations])
+
+
+def generate_reduction_api():
+    expanded_code = """
+/**
+ * @name SHMEM_REDUCTIONS
+ * @brief Perform an allreduce between PEs in the active set. The caller
+ * is blocked until the reduction completes.
+ *
+ * This function must be called as a work-group collective.
+ *
+ * @param[in] team         The team participating in the collective.
+ * @param[in] dest         Destination address. Must be an address on the
+ *                         symmetric heap.
+ * @param[in] source       Source address. Must be an address on the symmetric
+                           heap.
+ * @param[in] nreduce      Size of the buffer to participate in the reduction.
+ *
+ * @return int (Zero on successful local completion. Nonzero otherwise.)
+ */\n"""
+
+    int_types = [
+        ("short", "short"),
+        ("int", "int"),
+        ("long", "long"),
+        ("long long", "longlong")
+    ]
+
+    float_types = [
+        ("float", "float"),
+        ("double", "double")
+    ]
+
+    for type_, tname_ in int_types:
+        expanded_code += arith_reduction_api(type_, tname_)
+        expanded_code += bitwise_reduction_api(type_, tname_)
+
+    for type_, tname_ in float_types:
+        expanded_code += arith_reduction_api(type_, tname_)
+
+    return expanded_code
+
+
+def write_to_file(filename, content):
+    with open(filename, 'w') as file:
+        file.write(content)
+
+
+def generate_COLL_header(output_dir, copyright):
+    expanded_code = copyright
+
+    expanded_code += """
+#ifndef LIBRARY_INCLUDE_ROCSHMEM_COLL_HPP
+#define LIBRARY_INCLUDE_ROCSHMEM_COLL_HPP
+
+namespace rocshmem {
+"""
+
+    expanded_code += (
+        generate_alltoall_api() +
+        generate_broadcast_api() +
+        generate_fcollect_api() +
+        generate_reduction_api()
+    )
+
+    expanded_code += """
+}  // namespace rocshmem
+
+#endif  // LIBRARY_INCLUDE_ROCSHMEM_COLL_HPP
+"""
+
+    output_file = os.path.join(
+        output_dir, 'rocshmem_COLL.hpp'
+    )
+
+    write_to_file(output_file, expanded_code)

utlis/header_files_gen/P2P_SYNC.py

@@ -0,0 +1,176 @@
+"""
+******************************************************************************
+ * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *****************************************************************************
+"""
+
+import os
+
+types = [
+    ("float", "float"),
+    ("double", "double"),
+    ("char", "char"),
+    ("signed char", "schar"),
+    ("short", "short"),
+    ("int", "int"),
+    ("long", "long"),
+    ("long long", "longlong"),
+    ("unsigned char", "uchar"),
+    ("unsigned short", "ushort"),
+    ("unsigned int", "uint"),
+    ("unsigned long", "ulong"),
+    ("unsigned long long", "ulonglong"),
+]
+
+
+def wait_until_api(T, TNAME):
+    return (
+        f"__device__ void rocshmem_{TNAME}_wait_until(\n"
+        f"    {T} *ivars, int cmp, {T} val);\n"
+        f"__device__ size_t rocshmem_{TNAME}_wait_until_any(\n"
+        f"    {T} *ivars, size_t nelems, const int* status,\n"
+        f"    int cmp, {T} val);\n"
+        f"__device__ void rocshmem_{TNAME}_wait_until_all(\n"
+        f"    {T} *ivars, size_t nelems, const int* status,\n"
+        f"    int cmp, {T} val);\n"
+        f"__device__ size_t rocshmem_{TNAME}_wait_until_some(\n"
+        f"    {T} *ivars, size_t nelems, size_t* indices, const int* status,\n"
+        f"    int cmp, {T} val);\n"
+        f"__device__ size_t rocshmem_{TNAME}_wait_until_any_vector(\n"
+        f"    {T} *ivars, size_t nelems, const int* status,\n"
+        f"    int cmp, {T} val);\n"
+        f"__device__ void rocshmem_{TNAME}_wait_until_all_vector(\n"
+        f"    {T} *ivars, size_t nelems, const int* status,\n"
+        f"    int cmp, {T} val);\n"
+        f"__device__ size_t rocshmem_{TNAME}_wait_until_some_vector(\n"
+        f"    {T} *ivars, size_t nelems, size_t* indices, const int* status,\n"
+        f"    int cmp, {T} val);\n"
+        f"__host__ void rocshmem_{TNAME}_wait_until(\n"
+        f"    {T} *ivars, int cmp, {T} val);\n"
+        f"__host__ size_t rocshmem_{TNAME}_wait_until_any(\n"
+        f"    {T} *ivars, size_t nelems, const int* status,\n"
+        f"    int cmp, {T} val);\n"
+        f"__host__ void rocshmem_{TNAME}_wait_until_all(\n"
+        f"    {T} *ivars, size_t nelems, const int* status,\n"
+        f"    int cmp, {T} val);\n"
+        f"__host__ size_t rocshmem_{TNAME}_wait_until_some(\n"
+        f"    {T} *ivars, size_t nelems, size_t* indices, const int* status,\n"
+        f"    int cmp, {T} val);\n"
+        f"__host__ size_t rocshmem_{TNAME}_wait_until_any_vector(\n"
+        f"    {T} *ivars, size_t nelems, const int* status,\n"
+        f"    int cmp, {T} val);\n"
+        f"__host__ void rocshmem_{TNAME}_wait_until_all_vector(\n"
+        f"    {T} *ivars, size_t nelems, const int* status,\n"
+        f"    int cmp, {T} val);\n"
+        f"__host__ size_t rocshmem_{TNAME}_wait_until_some_vector(\n"
+        f"    {T} *ivars, size_t nelems, size_t* indices, const int* status,\n"
+        f"    int cmp, {T} val);\n\n"
+    )
+
+
+def generate_wait_until_api():
+    expanded_code = """
+/**
+ * @name SHMEM_WAIT_UNTIL
+ * @brief Block the caller until the condition (* \p ptr \p cmps \p val) is
+ * true.
+ *
+ * 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] ivars Pointer to memory on the symmetric heap to wait for.
+ * @param[in] cmp Operation for the comparison.
+ * @param[in] val Value to compare the memory at \p ptr to.
+ *
+ * @return void
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += wait_until_api(type_, tname_)
+
+    return expanded_code
+
+
+def test_api(T, TNAME):
+    return (
+        f"__device__ int rocshmem_{TNAME}_test(\n"
+        f"    {T} *ivars, int cmp, {T} val);\n"
+        f"__host__ int rocshmem_{TNAME}_test(\n"
+        f"    {T} *ivars, int cmp, {T} val);\n\n"
+    )
+
+
+def generate_test_api():
+    expanded_code = """
+/**
+ * @name SHMEM_TEST
+ * @brief test if the condition (* \p ptr \p cmps \p val) is
+ * true.
+ *
+ * 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] ivars Pointer to memory on the symmetric heap to wait for.
+ * @param[in] cmp Operation for the comparison.
+ * @param[in] val Value to compare the memory at \p ptr to.
+ *
+ * @return 1 if the evaluation is true else 0
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += test_api(type_, tname_)
+
+    return expanded_code
+
+
+def write_to_file(filename, content):
+    with open(filename, 'w') as file:
+        file.write(content)
+
+
+def generate_P2P_SYNC_header(output_dir, copyright):
+    expanded_code = copyright
+
+    expanded_code += """
+#ifndef LIBRARY_INCLUDE_ROCSHMEM_P2P_SYNC_HPP
+#define LIBRARY_INCLUDE_ROCSHMEM_P2P_SYNC_HPP
+
+namespace rocshmem {
+"""
+
+    expanded_code += (
+        generate_wait_until_api() +
+        generate_test_api()
+    )
+
+    expanded_code += """
+}  // namespace rocshmem
+
+#endif  // LIBRARY_INCLUDE_ROCSHMEM_P2P_SYNC_HPP
+"""
+
+    output_file = os.path.join(
+        output_dir, 'rocshmem_P2P_SYNC.hpp'
+    )
+
+    write_to_file(output_file, expanded_code)

utlis/header_files_gen/RMA.py

@@ -0,0 +1,335 @@
+"""
+******************************************************************************
+ * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *****************************************************************************
+"""
+
+import os
+
+types = [
+    ("float", "float"),
+    ("double", "double"),
+    ("char", "char"),
+    ("signed char", "schar"),
+    ("short", "short"),
+    ("int", "int"),
+    ("long", "long"),
+    ("long long", "longlong"),
+    ("unsigned char", "uchar"),
+    ("unsigned short", "ushort"),
+    ("unsigned int", "uint"),
+    ("unsigned long", "ulong"),
+    ("unsigned long long", "ulonglong"),
+]
+
+
+def put_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_put(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, const {T} *source,\n"
+        f"    size_t nelems, int pe);\n"
+        f"__device__ ATTR_NO_INLINE void rocshmem_{TNAME}_put(\n"
+        f"    {T} *dest, const {T} *source, size_t nelems, int pe);\n"
+        f"__host__ void rocshmem_ctx_{TNAME}_put(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, const {T} *source,\n"
+        f"    size_t nelems, int pe);\n"
+        f"__host__ void rocshmem_{TNAME}_put({T} *dest,\n"
+        f"    const {T} *source, size_t nelems, int pe);\n\n"
+    )
+
+
+def generate_put_api():
+    expanded_code = """
+/**
+ * @name SHMEM_PUT
+ * @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 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
+ * ROCSHMEM function.
+ *
+ * @param[in] ctx    Context with which to perform this operation.
+ * @param[in] dest   Destination address. Must be an address on the symmetric
+ *                   heap.
+ * @param[in] source Source address. Must be an address on the symmetric heap.
+ * @param[in] nelems Size of the transfer in number of elements.
+ * @param[in] pe     PE of the remote process.
+ *
+ * @return void.
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += put_api(type_, tname_)
+
+    return expanded_code
+
+
+def get_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_get(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, const {T} *source,\n"
+        f"    size_t nelems, int pe);\n"
+        f"__device__ ATTR_NO_INLINE void rocshmem_{TNAME}_get(\n"
+        f"    {T} *dest, const {T} *source, size_t nelems, int pe);\n"
+        f"__host__ void rocshmem_ctx_{TNAME}_get(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, const {T} *source,\n"
+        f"    size_t nelems, int pe);\n"
+        f"__host__ void rocshmem_{TNAME}_get({T} *dest,\n"
+        f"    const {T} *source, size_t nelems, int pe);\n\n"
+    )
+
+
+def generate_get_api():
+    expanded_code = """
+/**
+ * @name SHMEM_GET
+ * @brief Reads contiguous data of \p nelems elements from \p source on \p pe
+ * to \p dest on the calling PE. The calling work-group will block until the
+ * operation completes (data has been placed in \p dest).
+ *
+ * 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.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+ *                    heap.
+ * @param[in] source  Source address. Must be an address on the symmetric heap.
+ * @param[in] nelems  Size of the transfer in bytes.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return void.
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += get_api(type_, tname_)
+
+    return expanded_code
+
+
+def p_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_p(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value,\n"
+        f"    int pe);\n"
+        f"__device__ ATTR_NO_INLINE void rocshmem_{TNAME}_p(\n"
+        f"    {T} *dest, {T} value, int pe);\n"
+        f"__host__ void rocshmem_ctx_{TNAME}_p(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, {T} value,\n"
+        f"    int pe);\n"
+        f"__host__ void rocshmem_{TNAME}_p(\n"
+        f"    {T} *dest, {T} value, int pe);\n\n"
+    )
+
+
+def generate_p_api():
+    expanded_code = """
+/**
+ * @name SHMEM_P
+ * @brief Writes a single value to \p dest at \p pe PE to \p dst at \p pe.
+ * 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
+ * ROCSHMEM function.
+ *
+ * @param[in] ctx    Context with which to perform this operation.
+ * @param[in] dest   Destination address. Must be an address on the symmetric
+ *                   heap.
+ * @param[in] value  Value to write to dest at \p pe.
+ * @param[in] pe     PE of the remote process.
+ *
+ * @return void.
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += p_api(type_, tname_)
+
+    return expanded_code
+
+
+def g_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE {T} rocshmem_ctx_{TNAME}_g(\n"
+        f"    rocshmem_ctx_t ctx, const {T} *source, int pe);\n"
+        f"__device__ ATTR_NO_INLINE {T} rocshmem_{TNAME}_g(\n"
+        f"    const {T} *source, int pe);\n"
+        f"__host__ {T} rocshmem_ctx_{TNAME}_g(\n"
+        f"    rocshmem_ctx_t ctx, const {T} *source, int pe);\n"
+        f"__host__ {T} rocshmem_{TNAME}_g(\n"
+        f"    const {T} *source, int pe);\n\n"
+    )
+
+
+def generate_g_api():
+    expanded_code = """
+/**
+ * @name SHMEM_G
+ * @brief reads and returns single value from \p source at \p pe.
+ * The calling work-group/thread will block until the operation completes.
+ *
+ * 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.
+ * @param[in] source Source address. Must be an address on the symmetric
+ *                   heap.
+ * @param[in] pe     PE of the remote process.
+ *
+ * @return the value read from remote \p source at \p pe.
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += g_api(type_, tname_)
+
+    return expanded_code
+
+
+def put_nbi_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_put_nbi(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, const {T} *source,\n"
+        f"    size_t nelems, int pe);\n"
+        f"__device__ ATTR_NO_INLINE void rocshmem_{TNAME}_put_nbi(\n"
+        f"    {T} *dest, const {T} *source, size_t nelems, int pe);\n"
+        f"__host__ void rocshmem_ctx_{TNAME}_put_nbi(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, const {T} *source,\n"
+        f"    size_t nelems, int pe);\n"
+        f"__host__ void rocshmem_{TNAME}_put_nbi(\n"
+        f"    {T} *dest, const {T} *source, size_t nelems, int pe);\n\n"
+    )    
+
+
+def generate_put_nbi_api():
+    expanded_code = """
+/**
+ * @name SHMEM_PUT_NBI
+ * @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
+ * 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
+ * ROCSHMEM function.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+                      heap.
+ * @param[in] source  Source address. Must be an address on the symmetric heap.
+ * @param[in] nelems  Size of the transfer in bytes.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return void.
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += put_nbi_api(type_, tname_)
+
+    return expanded_code
+
+
+def get_nbi_api(T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_get_nbi(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, const {T} *source,\n"
+        f"    size_t nelems, int pe);\n"
+        f"__device__ ATTR_NO_INLINE void rocshmem_{TNAME}_get_nbi(\n"
+        f"    {T} *dest, const {T} *source, size_t nelems, int pe);\n"
+        f"__host__ void rocshmem_ctx_{TNAME}_get_nbi(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, const {T} *source,\n"
+        f"    size_t nelems, int pe);\n"
+        f"__host__ void rocshmem_{TNAME}_get_nbi({T} *dest,\n"
+        f"    const {T} *source, size_t nelems, int pe);\n\n"
+    )
+
+
+def generate_get_nbi_api():
+    expanded_code = """
+/**
+ * @name SHMEM_GET_NBI
+ * @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
+ * 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
+ * ROCSHMEM function.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+ *                    heap.
+ * @param[in] source  Source address. Must be an address on the symmetric heap.
+ * @param[in] nelems  Size of the transfer in bytes.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return void.
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += get_nbi_api(type_, tname_)
+
+    return expanded_code
+
+
+def write_to_file(filename, content):
+    with open(filename, 'w') as file:
+        file.write(content)
+
+
+def generate_RMA_header(output_dir, copyright):
+    expanded_code = copyright
+
+    expanded_code += """
+#ifndef LIBRARY_INCLUDE_ROCSHMEM_RMA_HPP
+#define LIBRARY_INCLUDE_ROCSHMEM_RMA_HPP
+
+namespace rocshmem {
+"""
+
+    expanded_code += (
+        generate_put_api() +
+        generate_p_api() +
+        generate_get_api() +
+        generate_g_api() +
+        generate_put_nbi_api() +
+        generate_get_nbi_api()
+    )
+
+    expanded_code += """
+}  // namespace rocshmem
+
+#endif  // LIBRARY_INCLUDE_ROCSHMEM_RMA_HPP
+"""
+
+    output_file = os.path.join(
+        output_dir, 'rocshmem_RMA.hpp'
+    )
+
+    write_to_file(output_file, expanded_code)

utlis/header_files_gen/RMA_X.py

@@ -0,0 +1,318 @@
+"""
+******************************************************************************
+ * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *****************************************************************************
+"""
+
+import os
+
+types = [
+    ("float", "float"),
+    ("double", "double"),
+    ("char", "char"),
+    ("signed char", "schar"),
+    ("short", "short"),
+    ("int", "int"),
+    ("long", "long"),
+    ("long long", "longlong"),
+    ("unsigned char", "uchar"),
+    ("unsigned short", "ushort"),
+    ("unsigned int", "uint"),
+    ("unsigned long", "ulong"),
+    ("unsigned long long", "ulonglong"),
+]
+
+
+def put_api_x(GRAN, T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_put_{GRAN}(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, const {T} *source,\n"
+        f"    size_t nelems, int pe);\n"
+        f"__device__ ATTR_NO_INLINE void rocshmem_{TNAME}_put_{GRAN}(\n"
+        f"    {T} *dest, const {T} *source, size_t nelems, int pe);\n\n"
+    )
+
+
+def generate_put_api_x():
+    expanded_code = """
+/**
+ * @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 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 collectively participate
+ * in the call using the same arguments
+ *
+ * @param[in] ctx    Context with which to perform this operation.
+ * @param[in] dest   Destination address. Must be an address on the symmetric
+ *                   heap.
+ * @param[in] source Source address. Must be an address on the symmetric heap.
+ * @param[in] nelems Size of the transfer in number of elements.
+ * @param[in] pe     PE of the remote process.
+ *
+ * @return void.
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += put_api_x("wave", type_, tname_)
+    
+    expanded_code += """
+/**
+ * @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 rocshmem_quiet() if remote completion is required.
+ *
+ * This function can be called from divergent control paths at per-workgroup
+ * (WG) granularity. However, All threads in a WG must collectively participate
+ * in the call using the same arguments.
+ *
+ * @param[in] ctx    Context with which to perform this operation.
+ * @param[in] dest   Destination address. Must be an address on the symmetric
+ *                   heap.
+ * @param[in] source Source address. Must be an address on the symmetric heap.
+ * @param[in] nelems Size of the transfer in number of elements.
+ * @param[in] pe     PE of the remote process.
+ *
+ * @return void.
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += put_api_x("wg", type_, tname_)
+
+    return expanded_code
+
+
+def get_api_x(GRAN, T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_get_{GRAN}(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, const {T} *source,\n"
+        f"    size_t nelems, int pe);\n"
+        f"__device__ ATTR_NO_INLINE void rocshmem_{TNAME}_get_{GRAN}(\n"
+        f"    {T} *dest, const {T} *source, size_t nelems, int pe);\n\n"
+    )
+
+
+def generate_get_api_x():
+    expanded_code = """
+/**
+ * @brief Reads contiguous data of \p nelems elements from \p source on \p pe
+ * to \p dest on the calling PE. The calling work-group will block until the
+ * operation completes (data has been placed in \p dest).
+ *
+ * This function can be called from divergent control paths at per-wave
+ * granularity. However,  all threads in the wave must participate in the
+ * call using the same parameters
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+ *                    heap.
+ * @param[in] source  Source address. Must be an address on the symmetric heap.
+ * @param[in] nelems  Size of the transfer in bytes.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return void.
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += get_api_x("wave", type_, tname_)
+    
+    expanded_code += """
+/**
+ * @brief Reads contiguous data of \p nelems elements from \p source on \p pe
+ * to \p dest on the calling PE. The calling work-group will block until the
+ * operation completes (data has been placed in \p dest).
+ *
+ * This function can be called from divergent control paths at per-workgroup
+ * granularity. However,  all threads in the workgroup must participate in
+ * the call using the same parameters
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+ *                    heap.
+ * @param[in] source  Source address. Must be an address on the symmetric heap.
+ * @param[in] nelems  Size of the transfer in bytes.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return void.
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += get_api_x("wg", type_, tname_)
+
+    return expanded_code
+
+
+def put_nbi_api_x(GRAN, T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_put_nbi_{GRAN}(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, const {T} *source,\n"
+        f"    size_t nelems, int pe);\n"
+        f"__device__ ATTR_NO_INLINE void rocshmem_{TNAME}_put_nbi_{GRAN}(\n"
+        f"    {T} *dest, const {T} *source, size_t nelems, int pe);\n\n"
+    )    
+
+
+def generate_put_nbi_api_x():
+    expanded_code = """
+/**
+ * @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
+ * rocshmem_quiet() on the same context if completion notification is
+ * required.
+ *
+ * This function can be called from divergent control paths at per-wave
+ * granularity. However, all threads in the wave must call in with the same
+ * arguments.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+                      heap.
+ * @param[in] source  Source address. Must be an address on the symmetric heap.
+ * @param[in] nelems  Size of the transfer in bytes.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return void.
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += put_nbi_api_x("wave", type_, tname_)
+
+    expanded_code += """
+/**
+ * @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
+ * rocshmem_quiet() on the same context if completion notification is
+ * required.
+ *
+ * This function can be called from divergent control paths at per-workgroup
+ * granularity. However, all threads in the WG must call in with the sameo
+ * arguments.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+                      heap.
+ * @param[in] source  Source address. Must be an address on the symmetric heap.
+ * @param[in] nelems  Size of the transfer in bytes.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return void.
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += put_nbi_api_x("wg", type_, tname_)
+
+    return expanded_code
+
+
+def get_nbi_api_x(GRAN, T, TNAME):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_get_nbi_{GRAN}(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, const {T} *source,\n"
+        f"    size_t nelems, int pe);\n"
+        f"__device__ ATTR_NO_INLINE void rocshmem_{TNAME}_get_nbi_{GRAN}(\n"
+        f"    {T} *dest, const {T} *source, size_t nelems, int pe);\n\n"
+    )
+
+
+def generate_get_nbi_api_x():
+    expanded_code = """
+/**
+ * @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
+ * rocshmem_quiet() on the same context if completion notification is
+ * required.
+ *
+ * This function can be called from divergent control paths at per-wave
+ * granularity. However, all threads in the wave must call in with the same
+ * arguments.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+ *                    heap.
+ * @param[in] source  Source address. Must be an address on the symmetric heap.
+ * @param[in] nelems  Size of the transfer in bytes.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return void.
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += get_nbi_api_x("wave", type_, tname_)
+
+    expanded_code += """
+/**
+ * @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
+ * rocshmem_quiet() on the same context if completion notification is
+ * required.
+ *
+ * This function can be called from divergent control paths at per-workgroup
+ * granularity. However, all threads in the WG must call in with the same
+ * arguments.
+ *
+ * @param[in] ctx     Context with which to perform this operation.
+ * @param[in] dest    Destination address. Must be an address on the symmetric
+ *                    heap.
+ * @param[in] source  Source address. Must be an address on the symmetric heap.
+ * @param[in] nelems  Size of the transfer in bytes.
+ * @param[in] pe      PE of the remote process.
+ *
+ * @return void.
+ */\n"""
+    for type_, tname_ in types:
+        expanded_code += get_nbi_api_x("wg", type_, tname_)
+
+    return expanded_code
+
+
+def write_to_file(filename, content):
+    with open(filename, 'w') as file:
+        file.write(content)
+
+
+def generate_RMA_X_header(output_dir, copyright):
+    expanded_code = copyright
+
+    expanded_code += """
+#ifndef LIBRARY_INCLUDE_ROCSHMEM_RMA_X_HPP
+#define LIBRARY_INCLUDE_ROCSHMEM_RMA_X_HPP
+
+namespace rocshmem {
+"""
+
+    expanded_code += (
+        generate_put_api_x() +
+        generate_get_api_x() +
+        generate_put_nbi_api_x() +
+        generate_get_nbi_api_x()
+    )
+
+    expanded_code += """
+}  // namespace rocshmem
+
+#endif  // LIBRARY_INCLUDE_ROCSHMEM_RMA_X_HPP
+"""
+
+    output_file = os.path.join(
+        output_dir, 'rocshmem_RMA_X.hpp'
+    )
+
+    write_to_file(output_file, expanded_code)

utlis/header_files_gen/SIG_OP.py

@@ -0,0 +1,108 @@
+"""
+******************************************************************************
+ * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *****************************************************************************
+"""
+
+import os
+
+types = [
+    ("float", "float"),
+    ("double", "double"),
+    ("char", "char"),
+    ("signed char", "schar"),
+    ("short", "short"),
+    ("int", "int"),
+    ("long", "long"),
+    ("long long", "longlong"),
+    ("unsigned char", "uchar"),
+    ("unsigned short", "ushort"),
+    ("unsigned int", "uint"),
+    ("unsigned long", "ulong"),
+    ("unsigned long long", "ulonglong"),
+]
+
+
+def putmem_signal_dec(SUFFIX):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_putmem_signal{SUFFIX}(\n"
+        f"    void *dest, const void *source, size_t nelems, uint64_t *sig_addr,\n"
+        f"    uint64_t signal, int sig_op, int pe);\n"
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_putmem_signal{SUFFIX}(\n"
+        f"    rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems,\n"
+        f"    uint64_t *sig_addr, uint64_t signal, int sig_op, int pe);\n\n"
+    )
+
+
+def put_signal_typed_dec(T, TNAME, SUFFIX):
+    return (
+        f"__device__ ATTR_NO_INLINE void rocshmem_ctx_{TNAME}_put_signal{SUFFIX}(\n"
+        f"    rocshmem_ctx_t ctx, {T} *dest, const {T} *source, size_t nelems,\n"
+        f"    uint64_t *sig_addr, uint64_t signal, int sig_op, int pe);\n"
+        f"__device__ ATTR_NO_INLINE void rocshmem_{TNAME}_put_signal{SUFFIX}(\n"
+        f"    {T} *dest, const {T} *source, size_t nelems, uint64_t *sig_addr,\n"
+        f"    uint64_t signal, int sig_op, int pe);\n\n"
+    )
+
+
+def put_signal_dec(SUFFIX):
+    return "".join([put_signal_typed_dec(T, TNAME, SUFFIX) for T, TNAME in types])
+
+
+def signaling_api_dec(SUFFIX):
+    return (putmem_signal_dec(SUFFIX) + put_signal_dec(SUFFIX))
+
+
+def generate_signal_api():
+
+    suffixes = ["", "_wg", "_wave", "_nbi", "_nbi_wg", "_nbi_wave"]
+
+    return "".join([signaling_api_dec(suffix) for suffix in suffixes])
+
+
+def write_to_file(filename, content):
+    with open(filename, 'w') as file:
+        file.write(content)
+
+
+def generate_SIG_OP_header(output_dir, copyright):
+    expanded_code = copyright
+
+    expanded_code += """
+#ifndef LIBRARY_INCLUDE_ROCSHMEM_SIG_OP_HPP
+#define LIBRARY_INCLUDE_ROCSHMEM_SIG_OP_HPP
+
+namespace rocshmem {
+"""
+
+    expanded_code += generate_signal_api()
+
+    expanded_code += """
+}  // namespace rocshmem
+
+#endif  // LIBRARY_INCLUDE_ROCSHMEM_SIG_OP_HPP
+"""
+
+    output_file = os.path.join(
+        output_dir, 'rocshmem_SIG_OP.hpp'
+    )
+
+    write_to_file(output_file, expanded_code)

utlis/header_files_gen/header_files_gen.py

@@ -0,0 +1,77 @@
+"""
+******************************************************************************
+ * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *****************************************************************************
+ """
+
+import argparse
+from RMA import generate_RMA_header
+from AMO import generate_AMO_header
+from SIG_OP import generate_SIG_OP_header
+from COLL import generate_COLL_header
+from P2P_SYNC import generate_P2P_SYNC_header
+from RMA_X import generate_RMA_X_header
+
+
+copyright = """
+/******************************************************************************
+ * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *****************************************************************************/
+"""
+
+
+def main():
+    parser = argparse.ArgumentParser(
+        description='Generate an expanded header files.'
+    )
+    parser.add_argument(
+        'directory', type=str, nargs='?', default='.',
+        help='Directory to write the header files to (default: current directory)'
+    )
+    args = parser.parse_args()
+
+    generate_RMA_header(args.directory, copyright)
+    generate_AMO_header(args.directory, copyright)
+    generate_SIG_OP_header(args.directory, copyright)
+    generate_COLL_header(args.directory, copyright)
+    generate_P2P_SYNC_header(args.directory, copyright)
+    generate_RMA_X_header(args.directory, copyright)
+
+if __name__ == "__main__":
+    main()