From 174a400318242fff4750b061b0f7b0e9e3d88e77 Mon Sep 17 00:00:00 2001 From: Ken Raffenetti Date: Tue, 10 Mar 2026 13:12:31 -0500 Subject: [PATCH 1/3] yaksa/cuda: Template pack/unpack kernels on operation type MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Replace the per-operation __global__ kernel functions with a single template kernel per direction (pack/unpack) per layout×type combination. Add yaksuri_cudai_ops.cuh defining one C++ functor struct per operation. --- .../typerep/yaksa/src/backend/cuda/genpup.py | 47 ++++++------ .../backend/cuda/pup/yaksuri_cudai_ops.cuh | 71 +++++++++++++++++++ 2 files changed, 97 insertions(+), 21 deletions(-) create mode 100644 src/mpi/datatype/typerep/yaksa/src/backend/cuda/pup/yaksuri_cudai_ops.cuh diff --git a/src/mpi/datatype/typerep/yaksa/src/backend/cuda/genpup.py b/src/mpi/datatype/typerep/yaksa/src/backend/cuda/genpup.py index e7a8766094c..acbfd216408 100755 --- a/src/mpi/datatype/typerep/yaksa/src/backend/cuda/genpup.py +++ b/src/mpi/datatype/typerep/yaksa/src/backend/cuda/genpup.py @@ -24,6 +24,19 @@ builtin_maps = { "YAKSA_TYPE__LONG_DOUBLE": "double", } +op_functor_names = { + 'REPLACE': 'YaksuriOpReplace', + 'SUM': 'YaksuriOpSum', + 'PROD': 'YaksuriOpProd', + 'MIN': 'YaksuriOpMin', + 'MAX': 'YaksuriOpMax', + 'LAND': 'YaksuriOpLand', + 'LOR': 'YaksuriOpLor', + 'LXOR': 'YaksuriOpLxor', + 'BAND': 'YaksuriOpBand', + 'BOR': 'YaksuriOpBor', + 'BXOR': 'YaksuriOpBxor', +} ######################################################################################## @@ -103,19 +116,20 @@ def resized(suffix, dtp, b, last): ######################################################################################## ##### Core kernels ######################################################################################## -def generate_kernels(b, darray, op): +def generate_kernels(b, darray): global need_extent global s global idx for func in "pack","unpack": ##### figure out the function name to use - funcprefix = "%s_%s_" % (func, op) + funcprefix = "%s_" % func for d in darray: funcprefix = funcprefix + "%s_" % d funcprefix = funcprefix + b.replace(" ", "_") ##### generate the CUDA kernel + yutils.display(OUTFILE, "template\n") yutils.display(OUTFILE, "__global__ void yaksuri_cudai_kernel_%s(const void *inbuf, void *outbuf, uintptr_t count, const yaksuri_cudai_md_s *__restrict__ md)\n" % funcprefix) yutils.display(OUTFILE, "{\n") yutils.display(OUTFILE, "const char *__restrict__ sbuf = (const char *) inbuf;\n"); @@ -184,15 +198,9 @@ def generate_kernels(b, darray, op): dtp = dtp + "->u.%s.child" % d if (func == "pack"): - if ((b == "float" or b == "double") and (op == "MAX" or op == "MIN")): - yutils.display(OUTFILE, "YAKSURI_CUDAI_OP_%s_FLOAT(%s, *((const %s *) (const void *) (sbuf + %s)), *((%s *) (void *) (dbuf + idx * sizeof(%s))));\n" % (op, b, b, s, b, b)) - else: - yutils.display(OUTFILE, "YAKSURI_CUDAI_OP_%s(*((const %s *) (const void *) (sbuf + %s)), *((%s *) (void *) (dbuf + idx * sizeof(%s))));\n" % (op, b, s, b, b)) + yutils.display(OUTFILE, "Op::apply(*((const %s *) (const void *) (sbuf + %s)), *((%s *) (void *) (dbuf + idx * sizeof(%s))));\n" % (b, s, b, b)) else: - if ((b == "float" or b == "double") and (op == "MAX" or op == "MIN")): - yutils.display(OUTFILE, "YAKSURI_CUDAI_OP_%s_FLOAT(%s, *((const %s *) (const void *) (sbuf + idx * sizeof(%s))), *((%s *) (void *) (dbuf + %s)));\n" % (op, b, b, b, b, s)) - else: - yutils.display(OUTFILE, "YAKSURI_CUDAI_OP_%s(*((const %s *) (const void *) (sbuf + idx * sizeof(%s))), *((%s *) (void *) (dbuf + %s)));\n" % (op, b, b, b, s)) + yutils.display(OUTFILE, "Op::apply(*((const %s *) (const void *) (sbuf + idx * sizeof(%s))), *((%s *) (void *) (dbuf + %s)));\n" % (b, b, b, s)) yutils.display(OUTFILE, "}\n\n") @@ -211,12 +219,9 @@ def generate_host_function(b, darray): yutils.display(OUTFILE, "cudaError_t cerr;\n") yutils.display(OUTFILE, "switch (op) {\n") for op in gencomm.type_ops[b]: - funcprefix = "%s_%s_" % (func, op) - for d in darray: - funcprefix = funcprefix + "%s_" % d - funcprefix = funcprefix + b.replace(" ", "_") + functor = "%s<%s>" % (op_functor_names[op], b.replace(" ", "_")) yutils.display(OUTFILE, "case YAKSA_OP__%s:\n" % op) - yutils.display(OUTFILE, "cerr = cudaLaunchKernel((const void *) yaksuri_cudai_kernel_%s,\n" % funcprefix) + yutils.display(OUTFILE, "cerr = cudaLaunchKernel((const void *) yaksuri_cudai_kernel_%s<%s>,\n" % (funcprefix, functor)) yutils.display(OUTFILE, " dim3(n_blocks_x, n_blocks_y, n_blocks_z), dim3(n_threads), args, 0, stream);\n") yutils.display(OUTFILE, "YAKSURI_CUDAI_CUDA_ERR_CHECK(cerr);\n") yutils.display(OUTFILE, "break;\n\n") @@ -252,11 +257,11 @@ def generate_host_function(b, darray): yutils.display(OUTFILE, "#include \n") yutils.display(OUTFILE, "#include \"yaksuri_cudai_base.h\"\n") yutils.display(OUTFILE, "#include \"yaksuri_cudai_pup.h\"\n") + yutils.display(OUTFILE, "#include \"yaksuri_cudai_ops.cuh\"\n") yutils.display(OUTFILE, "\n") emptylist = [ ] - for op in gencomm.type_ops[b]: - generate_kernels(b, emptylist, op) + generate_kernels(b, emptylist) generate_host_function(b, emptylist) OUTFILE.close() @@ -275,12 +280,12 @@ def generate_host_function(b, darray): yutils.display(OUTFILE, "#include \n") yutils.display(OUTFILE, "#include \"yaksuri_cudai_base.h\"\n") yutils.display(OUTFILE, "#include \"yaksuri_cudai_pup.h\"\n") + yutils.display(OUTFILE, "#include \"yaksuri_cudai_ops.cuh\"\n") yutils.display(OUTFILE, "\n") emptylist = [ ] emptylist.append(d) - for op in gencomm.type_ops[b]: - generate_kernels(b, emptylist, op) + generate_kernels(b, emptylist) generate_host_function(b, emptylist) emptylist.pop() @@ -304,13 +309,13 @@ def generate_host_function(b, darray): yutils.display(OUTFILE, "#include \n") yutils.display(OUTFILE, "#include \"yaksuri_cudai_base.h\"\n") yutils.display(OUTFILE, "#include \"yaksuri_cudai_pup.h\"\n") + yutils.display(OUTFILE, "#include \"yaksuri_cudai_ops.cuh\"\n") yutils.display(OUTFILE, "\n") for darray in darraylist: darray.append(d1) darray.append(d2) - for op in gencomm.type_ops[b]: - generate_kernels(b, darray, op) + generate_kernels(b, darray) generate_host_function(b, darray) darray.pop() darray.pop() diff --git a/src/mpi/datatype/typerep/yaksa/src/backend/cuda/pup/yaksuri_cudai_ops.cuh b/src/mpi/datatype/typerep/yaksa/src/backend/cuda/pup/yaksuri_cudai_ops.cuh new file mode 100644 index 00000000000..72bfd8e92dd --- /dev/null +++ b/src/mpi/datatype/typerep/yaksa/src/backend/cuda/pup/yaksuri_cudai_ops.cuh @@ -0,0 +1,71 @@ +/* + * Copyright (C) by Argonne National Laboratory + * See COPYRIGHT in top-level directory + */ + +#ifndef YAKSURI_CUDAI_OPS_CUH_INCLUDED +#define YAKSURI_CUDAI_OPS_CUH_INCLUDED + +template struct YaksuriOpReplace { + __device__ static void apply(const T& in, T& out) { out = in; } +}; +template struct YaksuriOpSum { + __device__ static void apply(const T& in, T& out) { out += in; } +}; +template struct YaksuriOpProd { + __device__ static void apply(const T& in, T& out) { out *= in; } +}; +template struct YaksuriOpLand { + __device__ static void apply(const T& in, T& out) { out = out && in; } +}; +template struct YaksuriOpBand { + __device__ static void apply(const T& in, T& out) { out &= in; } +}; +template struct YaksuriOpLor { + __device__ static void apply(const T& in, T& out) { out = out || in; } +}; +template struct YaksuriOpBor { + __device__ static void apply(const T& in, T& out) { out |= in; } +}; +template struct YaksuriOpLxor { + __device__ static void apply(const T& in, T& out) { out = !out != !in; } +}; +template struct YaksuriOpBxor { + __device__ static void apply(const T& in, T& out) { out ^= in; } +}; + +/* Integer MAX/MIN: branchless bitwise trick */ +template struct YaksuriOpMax { + __device__ static void apply(const T& in, T& out) { + out = in ^ ((in ^ out) & -(in < out)); + } +}; +template struct YaksuriOpMin { + __device__ static void apply(const T& in, T& out) { + out = out ^ ((in ^ out) & -(in < out)); + } +}; + +/* Float/double MAX/MIN */ +template<> struct YaksuriOpMax { + __device__ static void apply(const float& in, float& out) { + if (in > out) out = in; + } +}; +template<> struct YaksuriOpMax { + __device__ static void apply(const double& in, double& out) { + if (in > out) out = in; + } +}; +template<> struct YaksuriOpMin { + __device__ static void apply(const float& in, float& out) { + if (in < out) out = in; + } +}; +template<> struct YaksuriOpMin { + __device__ static void apply(const double& in, double& out) { + if (in < out) out = in; + } +}; + +#endif /* YAKSURI_CUDAI_OPS_CUH_INCLUDED */ From 2fae308b4f6d35f6b15d6aa518f402efb8e91adc Mon Sep 17 00:00:00 2001 From: Ken Raffenetti Date: Thu, 5 Mar 2026 23:00:27 -0600 Subject: [PATCH 2/3] yaksa/cuda: Collapse to one .cu file per layout MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Template pack/unpack kernels on both operation and element type using template class Op, typename T>, allowing all 9 builtin types to share a single kernel template per layout×direction. In a 2-level nesting configuration, the files (1 builtin, 5 single-level + 25 double-level) now contains 2 kernel templates plus 9×2 host functions that instantiate them as kernel. --- .../typerep/yaksa/src/backend/cuda/genpup.py | 168 +++++++++--------- 1 file changed, 87 insertions(+), 81 deletions(-) diff --git a/src/mpi/datatype/typerep/yaksa/src/backend/cuda/genpup.py b/src/mpi/datatype/typerep/yaksa/src/backend/cuda/genpup.py index acbfd216408..219275ce732 100755 --- a/src/mpi/datatype/typerep/yaksa/src/backend/cuda/genpup.py +++ b/src/mpi/datatype/typerep/yaksa/src/backend/cuda/genpup.py @@ -116,20 +116,19 @@ def resized(suffix, dtp, b, last): ######################################################################################## ##### Core kernels ######################################################################################## -def generate_kernels(b, darray): +def generate_kernels(darray): global need_extent global s global idx for func in "pack","unpack": - ##### figure out the function name to use - funcprefix = "%s_" % func + ##### figure out the function name to use (layout only, no type) + funcprefix = func for d in darray: - funcprefix = funcprefix + "%s_" % d - funcprefix = funcprefix + b.replace(" ", "_") + funcprefix = funcprefix + "_%s" % d ##### generate the CUDA kernel - yutils.display(OUTFILE, "template\n") + yutils.display(OUTFILE, "template class Op, typename T>\n") yutils.display(OUTFILE, "__global__ void yaksuri_cudai_kernel_%s(const void *inbuf, void *outbuf, uintptr_t count, const yaksuri_cudai_md_s *__restrict__ md)\n" % funcprefix) yutils.display(OUTFILE, "{\n") yutils.display(OUTFILE, "const char *__restrict__ sbuf = (const char *) inbuf;\n"); @@ -193,35 +192,40 @@ def generate_kernels(b, darray): last = 1 else: last = 0 - getattr(sys.modules[__name__], d)(x, dtp, b, last) + getattr(sys.modules[__name__], d)(x, dtp, "T", last) x = x + 1 dtp = dtp + "->u.%s.child" % d if (func == "pack"): - yutils.display(OUTFILE, "Op::apply(*((const %s *) (const void *) (sbuf + %s)), *((%s *) (void *) (dbuf + idx * sizeof(%s))));\n" % (b, s, b, b)) + yutils.display(OUTFILE, "Op::apply(*((const T *) (const void *) (sbuf + %s)), *((T *) (void *) (dbuf + idx * sizeof(T))));\n" % s) else: - yutils.display(OUTFILE, "Op::apply(*((const %s *) (const void *) (sbuf + idx * sizeof(%s))), *((%s *) (void *) (dbuf + %s)));\n" % (b, b, b, s)) + yutils.display(OUTFILE, "Op::apply(*((const T *) (const void *) (sbuf + idx * sizeof(T))), *((T *) (void *) (dbuf + %s)));\n" % s) yutils.display(OUTFILE, "}\n\n") def generate_host_function(b, darray): for func in "pack","unpack": - funcprefix = "%s_" % func + # Host function name includes the type (maintains C ABI) + host_funcprefix = "%s_" % func for d in darray: - funcprefix = funcprefix + "%s_" % d - funcprefix = funcprefix + b.replace(" ", "_") + host_funcprefix = host_funcprefix + "%s_" % d + host_funcprefix = host_funcprefix + b.replace(" ", "_") - yutils.display(OUTFILE, "void yaksuri_cudai_%s(const void *inbuf, void *outbuf, uintptr_t count, yaksa_op_t op, yaksuri_cudai_md_s *md, int n_threads, int n_blocks_x, int n_blocks_y, int n_blocks_z, cudaStream_t stream)\n" % funcprefix) + # Kernel name is layout-only (no type) — matches generate_kernels output + kernel_funcprefix = func + for d in darray: + kernel_funcprefix = kernel_funcprefix + "_%s" % d + + yutils.display(OUTFILE, "void yaksuri_cudai_%s(const void *inbuf, void *outbuf, uintptr_t count, yaksa_op_t op, yaksuri_cudai_md_s *md, int n_threads, int n_blocks_x, int n_blocks_y, int n_blocks_z, cudaStream_t stream)\n" % host_funcprefix) yutils.display(OUTFILE, "{\n") yutils.display(OUTFILE, "void *args[] = { &inbuf, &outbuf, &count, &md };\n") yutils.display(OUTFILE, "cudaError_t cerr;\n") yutils.display(OUTFILE, "switch (op) {\n") for op in gencomm.type_ops[b]: - functor = "%s<%s>" % (op_functor_names[op], b.replace(" ", "_")) yutils.display(OUTFILE, "case YAKSA_OP__%s:\n" % op) - yutils.display(OUTFILE, "cerr = cudaLaunchKernel((const void *) yaksuri_cudai_kernel_%s<%s>,\n" % (funcprefix, functor)) + yutils.display(OUTFILE, "cerr = cudaLaunchKernel((const void *) yaksuri_cudai_kernel_%s<%s, %s>,\n" % (kernel_funcprefix, op_functor_names[op], b.replace(" ", "_"))) yutils.display(OUTFILE, " dim3(n_blocks_x, n_blocks_y, n_blocks_z), dim3(n_threads), args, 0, stream);\n") yutils.display(OUTFILE, "YAKSURI_CUDAI_CUDA_ERR_CHECK(cerr);\n") yutils.display(OUTFILE, "break;\n\n") @@ -245,82 +249,86 @@ def generate_host_function(b, darray): ##### generate the core pack/unpack kernels (zero levels) if args.pup_max_nesting > 0: + filename = "src/backend/cuda/pup/yaksuri_cudai_pup_builtin.cu" + yutils.copyright_c(filename) + OUTFILE = open(filename, "a") + yutils.display(OUTFILE, "#include \n") + yutils.display(OUTFILE, "#include \n") + yutils.display(OUTFILE, "#include \n") + yutils.display(OUTFILE, "#include \n") + yutils.display(OUTFILE, "#include \n") + yutils.display(OUTFILE, "#include \n") + yutils.display(OUTFILE, "#include \"yaksuri_cudai_base.h\"\n") + yutils.display(OUTFILE, "#include \"yaksuri_cudai_pup.h\"\n") + yutils.display(OUTFILE, "#include \"yaksuri_cudai_ops.cuh\"\n") + yutils.display(OUTFILE, "\n") + + emptylist = [ ] + generate_kernels(emptylist) for b in builtin_types: - filename = "src/backend/cuda/pup/yaksuri_cudai_pup_%s.cu" % b.replace(" ","_") - yutils.copyright_c(filename) - OUTFILE = open(filename, "a") - yutils.display(OUTFILE, "#include \n") - yutils.display(OUTFILE, "#include \n") - yutils.display(OUTFILE, "#include \n") - yutils.display(OUTFILE, "#include \n") - yutils.display(OUTFILE, "#include \n") - yutils.display(OUTFILE, "#include \n") - yutils.display(OUTFILE, "#include \"yaksuri_cudai_base.h\"\n") - yutils.display(OUTFILE, "#include \"yaksuri_cudai_pup.h\"\n") - yutils.display(OUTFILE, "#include \"yaksuri_cudai_ops.cuh\"\n") - yutils.display(OUTFILE, "\n") - - emptylist = [ ] - generate_kernels(b, emptylist) generate_host_function(b, emptylist) - OUTFILE.close() + OUTFILE.close() ##### generate the core pack/unpack kernels (single level) - for b in builtin_types: - for d in gencomm.derived_types: - filename = "src/backend/cuda/pup/yaksuri_cudai_pup_%s_%s.cu" % (d, b.replace(" ","_")) - yutils.copyright_c(filename) - OUTFILE = open(filename, "a") - yutils.display(OUTFILE, "#include \n") - yutils.display(OUTFILE, "#include \n") - yutils.display(OUTFILE, "#include \n") - yutils.display(OUTFILE, "#include \n") - yutils.display(OUTFILE, "#include \n") - yutils.display(OUTFILE, "#include \n") - yutils.display(OUTFILE, "#include \"yaksuri_cudai_base.h\"\n") - yutils.display(OUTFILE, "#include \"yaksuri_cudai_pup.h\"\n") - yutils.display(OUTFILE, "#include \"yaksuri_cudai_ops.cuh\"\n") - yutils.display(OUTFILE, "\n") - - emptylist = [ ] - emptylist.append(d) - generate_kernels(b, emptylist) + for d in gencomm.derived_types: + filename = "src/backend/cuda/pup/yaksuri_cudai_pup_%s.cu" % d + yutils.copyright_c(filename) + OUTFILE = open(filename, "a") + yutils.display(OUTFILE, "#include \n") + yutils.display(OUTFILE, "#include \n") + yutils.display(OUTFILE, "#include \n") + yutils.display(OUTFILE, "#include \n") + yutils.display(OUTFILE, "#include \n") + yutils.display(OUTFILE, "#include \n") + yutils.display(OUTFILE, "#include \"yaksuri_cudai_base.h\"\n") + yutils.display(OUTFILE, "#include \"yaksuri_cudai_pup.h\"\n") + yutils.display(OUTFILE, "#include \"yaksuri_cudai_ops.cuh\"\n") + yutils.display(OUTFILE, "\n") + + emptylist = [ d ] + generate_kernels(emptylist) + for b in builtin_types: generate_host_function(b, emptylist) - emptylist.pop() - OUTFILE.close() + OUTFILE.close() ##### generate the core pack/unpack kernels (more than one level) if args.pup_max_nesting > 1: darraylist = [ ] yutils.generate_darrays(gencomm.derived_types, darraylist, args.pup_max_nesting - 2) - for b in builtin_types: - for d1 in gencomm.derived_types: - for d2 in gencomm.derived_types: - filename = "src/backend/cuda/pup/yaksuri_cudai_pup_%s_%s_%s.cu" % (d1, d2, b.replace(" ","_")) - yutils.copyright_c(filename) - OUTFILE = open(filename, "a") - yutils.display(OUTFILE, "#include \n") - yutils.display(OUTFILE, "#include \n") - yutils.display(OUTFILE, "#include \n") - yutils.display(OUTFILE, "#include \n") - yutils.display(OUTFILE, "#include \n") - yutils.display(OUTFILE, "#include \n") - yutils.display(OUTFILE, "#include \"yaksuri_cudai_base.h\"\n") - yutils.display(OUTFILE, "#include \"yaksuri_cudai_pup.h\"\n") - yutils.display(OUTFILE, "#include \"yaksuri_cudai_ops.cuh\"\n") - yutils.display(OUTFILE, "\n") + for d1 in gencomm.derived_types: + for d2 in gencomm.derived_types: + filename = "src/backend/cuda/pup/yaksuri_cudai_pup_%s_%s.cu" % (d1, d2) + yutils.copyright_c(filename) + OUTFILE = open(filename, "a") + yutils.display(OUTFILE, "#include \n") + yutils.display(OUTFILE, "#include \n") + yutils.display(OUTFILE, "#include \n") + yutils.display(OUTFILE, "#include \n") + yutils.display(OUTFILE, "#include \n") + yutils.display(OUTFILE, "#include \n") + yutils.display(OUTFILE, "#include \"yaksuri_cudai_base.h\"\n") + yutils.display(OUTFILE, "#include \"yaksuri_cudai_pup.h\"\n") + yutils.display(OUTFILE, "#include \"yaksuri_cudai_ops.cuh\"\n") + yutils.display(OUTFILE, "\n") + + for darray in darraylist: + darray.append(d1) + darray.append(d2) + generate_kernels(darray) + darray.pop() + darray.pop() + for b in builtin_types: for darray in darraylist: darray.append(d1) darray.append(d2) - generate_kernels(b, darray) generate_host_function(b, darray) darray.pop() darray.pop() - OUTFILE.close() + OUTFILE.close() ##### generate the core pack/unpack kernel declarations filename = "src/backend/cuda/pup/yaksuri_cudai_pup.h" @@ -417,16 +425,14 @@ def generate_host_function(b, darray): yutils.copyright_makefile(filename) OUTFILE = open(filename, "a") yutils.display(OUTFILE, "libyaksa_la_SOURCES += \\\n") - for b in builtin_types: - yutils.display(OUTFILE, "\tsrc/backend/cuda/pup/yaksuri_cudai_pup_%s.cu \\\n" % b.replace(" ","_")) - if args.pup_max_nesting > 0: - for d1 in gencomm.derived_types: - yutils.display(OUTFILE, "\tsrc/backend/cuda/pup/yaksuri_cudai_pup_%s_%s.cu \\\n" % \ - (d1, b.replace(" ","_"))) - if args.pup_max_nesting > 1: - for d2 in gencomm.derived_types: - yutils.display(OUTFILE, "\tsrc/backend/cuda/pup/yaksuri_cudai_pup_%s_%s_%s.cu \\\n" % \ - (d1, d2, b.replace(" ","_"))) + if args.pup_max_nesting > 0: + yutils.display(OUTFILE, "\tsrc/backend/cuda/pup/yaksuri_cudai_pup_builtin.cu \\\n") + for d1 in gencomm.derived_types: + yutils.display(OUTFILE, "\tsrc/backend/cuda/pup/yaksuri_cudai_pup_%s.cu \\\n" % d1) + if args.pup_max_nesting > 1: + for d2 in gencomm.derived_types: + yutils.display(OUTFILE, "\tsrc/backend/cuda/pup/yaksuri_cudai_pup_%s_%s.cu \\\n" % \ + (d1, d2)) yutils.display(OUTFILE, "\tsrc/backend/cuda/pup/yaksuri_cudai_pup.c\n") yutils.display(OUTFILE, "\n") yutils.display(OUTFILE, "noinst_HEADERS += \\\n") From d7b1d2ef89bbfa37228be6598073ea139b6aa7df Mon Sep 17 00:00:00 2001 From: Ken Raffenetti Date: Tue, 10 Mar 2026 13:12:49 -0500 Subject: [PATCH 3/3] yaksa/cuda: Use CUDA math functions for Min/Max kernels Replace the if-based specializations and integer bitwise trick with CUDA's overloaded max()/min(), which resolve to hardware MAX/MIN instructions for integer types and FMAX/FMIN for float/double. No type-specific specializations are needed. Note that CUDA's max()/min() for floating-point follow IEEE 754-2008 and return the non-NaN operand when one input is NaN. --- .../backend/cuda/pup/yaksuri_cudai_ops.cuh | 32 ++----------------- 1 file changed, 3 insertions(+), 29 deletions(-) diff --git a/src/mpi/datatype/typerep/yaksa/src/backend/cuda/pup/yaksuri_cudai_ops.cuh b/src/mpi/datatype/typerep/yaksa/src/backend/cuda/pup/yaksuri_cudai_ops.cuh index 72bfd8e92dd..4a2d92ffb37 100644 --- a/src/mpi/datatype/typerep/yaksa/src/backend/cuda/pup/yaksuri_cudai_ops.cuh +++ b/src/mpi/datatype/typerep/yaksa/src/backend/cuda/pup/yaksuri_cudai_ops.cuh @@ -34,38 +34,12 @@ template struct YaksuriOpBxor { __device__ static void apply(const T& in, T& out) { out ^= in; } }; -/* Integer MAX/MIN: branchless bitwise trick */ +/* MAX/MIN: CUDA overloaded max()/min() map to hardware instructions for all types */ template struct YaksuriOpMax { - __device__ static void apply(const T& in, T& out) { - out = in ^ ((in ^ out) & -(in < out)); - } + __device__ static void apply(const T& in, T& out) { out = max(in, out); } }; template struct YaksuriOpMin { - __device__ static void apply(const T& in, T& out) { - out = out ^ ((in ^ out) & -(in < out)); - } -}; - -/* Float/double MAX/MIN */ -template<> struct YaksuriOpMax { - __device__ static void apply(const float& in, float& out) { - if (in > out) out = in; - } -}; -template<> struct YaksuriOpMax { - __device__ static void apply(const double& in, double& out) { - if (in > out) out = in; - } -}; -template<> struct YaksuriOpMin { - __device__ static void apply(const float& in, float& out) { - if (in < out) out = in; - } -}; -template<> struct YaksuriOpMin { - __device__ static void apply(const double& in, double& out) { - if (in < out) out = in; - } + __device__ static void apply(const T& in, T& out) { out = min(in, out); } }; #endif /* YAKSURI_CUDAI_OPS_CUH_INCLUDED */