From 455f9330e9842c860049c558f1fdd57adc6d844d Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Thu, 5 Dec 2024 23:53:10 +0000 Subject: [PATCH] Add Python scripts to generate header files based on sections in OpenSHMEM specifications --- utlis/header_files_gen/AMO.py | 639 +++++++++++++++++++++ utlis/header_files_gen/COLL.py | 246 ++++++++ utlis/header_files_gen/P2P_SYNC.py | 176 ++++++ utlis/header_files_gen/RMA.py | 335 +++++++++++ utlis/header_files_gen/RMA_X.py | 318 ++++++++++ utlis/header_files_gen/SIG_OP.py | 108 ++++ utlis/header_files_gen/header_files_gen.py | 77 +++ 7 files changed, 1899 insertions(+) create mode 100644 utlis/header_files_gen/AMO.py create mode 100644 utlis/header_files_gen/COLL.py create mode 100644 utlis/header_files_gen/P2P_SYNC.py create mode 100644 utlis/header_files_gen/RMA.py create mode 100644 utlis/header_files_gen/RMA_X.py create mode 100644 utlis/header_files_gen/SIG_OP.py create mode 100644 utlis/header_files_gen/header_files_gen.py diff --git a/utlis/header_files_gen/AMO.py b/utlis/header_files_gen/AMO.py new file mode 100644 index 0000000000..afee8d0f43 --- /dev/null +++ b/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) diff --git a/utlis/header_files_gen/COLL.py b/utlis/header_files_gen/COLL.py new file mode 100644 index 0000000000..2755f8b26f --- /dev/null +++ b/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) diff --git a/utlis/header_files_gen/P2P_SYNC.py b/utlis/header_files_gen/P2P_SYNC.py new file mode 100644 index 0000000000..430e0b8aed --- /dev/null +++ b/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) diff --git a/utlis/header_files_gen/RMA.py b/utlis/header_files_gen/RMA.py new file mode 100644 index 0000000000..cf12a3c3fa --- /dev/null +++ b/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) diff --git a/utlis/header_files_gen/RMA_X.py b/utlis/header_files_gen/RMA_X.py new file mode 100644 index 0000000000..705731fa53 --- /dev/null +++ b/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) diff --git a/utlis/header_files_gen/SIG_OP.py b/utlis/header_files_gen/SIG_OP.py new file mode 100644 index 0000000000..b9880c0d0a --- /dev/null +++ b/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) diff --git a/utlis/header_files_gen/header_files_gen.py b/utlis/header_files_gen/header_files_gen.py new file mode 100644 index 0000000000..ca7268298a --- /dev/null +++ b/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()