Add transform c parallel implementation

c/parallel/CMakeLists.txt

@@ -40,8 +40,9 @@ target_link_libraries(cccl.c.parallel PRIVATE
   CUB::CUB
   Thrust::Thrust
 )
-target_compile_definitions(cccl.c.parallel PUBLIC CCCL_C_EXPERIMENTAL=1)
+target_compile_definitions(cccl.c.parallel PUBLIC CCCL_C_EXPERIMENTAL=1 _CUB_HAS_TRANSFORM_UBLKCP=0)
 target_compile_definitions(cccl.c.parallel PRIVATE NVRTC_GET_TYPE_NAME=1)
+target_compile_options(cccl.c.parallel PRIVATE $<$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>:--extended-lambda>)
 
 target_include_directories(cccl.c.parallel PUBLIC "include")
 target_include_directories(cccl.c.parallel PRIVATE "src")

c/parallel/include/cccl/c/transform.h

@@ -0,0 +1,78 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of CUDA Experimental in CUDA Core Compute Libraries,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#ifndef CCCL_C_EXPERIMENTAL
+#  error "C exposure is experimental and subject to change. Define CCCL_C_EXPERIMENTAL to acknowledge this notice."
+#endif // !CCCL_C_EXPERIMENTAL
+
+#include <cuda.h>
+
+#include <cccl/c/extern_c.h>
+#include <cccl/c/types.h>
+
+CCCL_C_EXTERN_C_BEGIN
+
+typedef struct cccl_device_transform_build_result_t
+{
+  int cc;
+  void* cubin;
+  size_t cubin_size;
+  CUlibrary library;
+  CUkernel transform_kernel;
+  int loaded_bytes_per_iteration;
+} cccl_device_transform_build_result_t;
+
+CCCL_C_API CUresult cccl_device_unary_transform_build(
+  cccl_device_transform_build_result_t* build_ptr,
+  cccl_iterator_t d_in,
+  cccl_iterator_t d_out,
+  cccl_op_t op,
+  int cc_major,
+  int cc_minor,
+  const char* cub_path,
+  const char* thrust_path,
+  const char* libcudacxx_path,
+  const char* ctk_path);
+
+CCCL_C_API CUresult cccl_device_unary_transform(
+  cccl_device_transform_build_result_t build,
+  cccl_iterator_t d_in,
+  cccl_iterator_t d_out,
+  unsigned long long num_items,
+  cccl_op_t op,
+  CUstream stream);
+
+CCCL_C_API CUresult cccl_device_binary_transform_build(
+  cccl_device_transform_build_result_t* build_ptr,
+  cccl_iterator_t d_in1,
+  cccl_iterator_t d_in2,
+  cccl_iterator_t d_out,
+  cccl_op_t op,
+  int cc_major,
+  int cc_minor,
+  const char* cub_path,
+  const char* thrust_path,
+  const char* libcudacxx_path,
+  const char* ctk_path);
+
+CCCL_C_API CUresult cccl_device_binary_transform(
+  cccl_device_transform_build_result_t build,
+  cccl_iterator_t d_in1,
+  cccl_iterator_t d_in2,
+  cccl_iterator_t d_out,
+  unsigned long long num_items,
+  cccl_op_t op,
+  CUstream stream);
+
+CCCL_C_API CUresult cccl_device_transform_cleanup(cccl_device_transform_build_result_t* bld_ptr);
+
+CCCL_C_EXTERN_C_END

c/parallel/src/kernels/operators.cpp

@@ -16,25 +16,27 @@
 #include <util/errors.h>
 #include <util/types.h>
 
-constexpr std::string_view op_template = R"XXX(
-#define VALUE_T {0}
-#define OP_NAME {1}
-#define OP_ALIGNMENT {2}
-#define OP_SIZE {3}
+constexpr std::string_view binary_op_template = R"XXX(
+#define LHS_T {0}
+#define RHS_T {1}
+#define OP_NAME {2}
+#define OP_ALIGNMENT {3}
+#define OP_SIZE {4}
 
 // Source
-{4}
+{5}
 
-#undef VALUE_T
+#undef LHS_T
+#undef RHS_T
 #undef OP_NAME
 #undef OP_ALIGNMENT
 #undef OP_SIZE
 )XXX";
 
 constexpr std::string_view stateless_binary_op_template = R"XXX(
-extern "C" __device__ {0} OP_NAME(VALUE_T lhs, VALUE_T rhs);
+extern "C" __device__ {0} OP_NAME(LHS_T lhs, RHS_T rhs);
 struct op_wrapper {{
-  __device__ {0} operator()(VALUE_T lhs, VALUE_T rhs) const {{
+  __device__ {0} operator()(LHS_T lhs, RHS_T rhs) const {{
     return OP_NAME(lhs, rhs);
   }}
 }};
@@ -44,25 +46,26 @@ constexpr std::string_view stateful_binary_op_template = R"XXX(
 struct __align__(OP_ALIGNMENT) op_state {{
   char data[OP_SIZE];
 }};
-extern "C" __device__ {0} OP_NAME(op_state *state, VALUE_T lhs, VALUE_T rhs);
+extern "C" __device__ {0} OP_NAME(op_state *state, LHS_T lhs, RHS_T rhs);
 struct op_wrapper {{
   op_state state;
-  __device__ {0} operator()(VALUE_T lhs, VALUE_T rhs) {{
+  __device__ {0} operator()(LHS_T lhs, RHS_T rhs) {{
     return OP_NAME(&state, lhs, rhs);
   }}
 }};
 )XXX";
 
-std::string
-make_kernel_binary_operator_full_source(std::string_view input_t, cccl_op_t operation, std::string_view return_type)
+std::string make_kernel_binary_operator_full_source(
+  std::string_view lhs_t, std::string_view rhs_t, cccl_op_t operation, std::string_view return_type)
 {
   const std::string op_alignment =
     operation.type == cccl_op_kind_t::CCCL_STATELESS ? "" : std::format("{}", operation.alignment);
   const std::string op_size = operation.type == cccl_op_kind_t::CCCL_STATELESS ? "" : std::format("{}", operation.size);
 
   return std::format(
-    op_template,
-    input_t,
+    binary_op_template,
+    lhs_t,
+    rhs_t,
     operation.name,
     op_alignment,
     op_size,
@@ -71,22 +74,40 @@ make_kernel_binary_operator_full_source(std::string_view input_t, cccl_op_t oper
       : std::format(stateful_binary_op_template, return_type));
 }
 
-std::string make_kernel_user_binary_operator(std::string_view input_t, cccl_op_t operation)
+std::string make_kernel_user_binary_operator(
+  std::string_view lhs_t, std::string_view rhs_t, std::string_view output_t, cccl_op_t operation)
 {
-  return make_kernel_binary_operator_full_source(input_t, operation, "VALUE_T");
+  return make_kernel_binary_operator_full_source(lhs_t, rhs_t, operation, output_t);
 }
 
 std::string make_kernel_user_comparison_operator(std::string_view input_t, cccl_op_t operation)
 {
-  return make_kernel_binary_operator_full_source(input_t, operation, "bool");
+  return make_kernel_binary_operator_full_source(input_t, input_t, operation, "bool");
 }
 
-std::string make_kernel_user_unary_operator(std::string_view input_t, cccl_op_t operation)
+std::string make_kernel_user_unary_operator(std::string_view input_t, std::string_view output_t, cccl_op_t operation)
 {
+  constexpr std::string_view unary_op_template = R"XXX(
+#define INPUT_T {0}
+#define OUTPUT_T {1}
+#define OP_NAME {2}
+#define OP_ALIGNMENT {3}
+#define OP_SIZE {4}
+
+// Source
+{5}
+
+#undef INPUT_T
+#undef OUTPUT_T
+#undef OP_NAME
+#undef OP_ALIGNMENT
+#undef OP_SIZE
+)XXX";
+
   constexpr std::string_view stateless_op = R"XXX(
-extern "C" __device__ VALUE_T OP_NAME(VALUE_T val);
+extern "C" __device__ OUTPUT_T OP_NAME(INPUT_T val);
 struct op_wrapper {
-  __device__ VALUE_T operator()(VALUE_T val) const {
+  __device__ OUTPUT_T operator()(INPUT_T val) const {
     return OP_NAME(val);
   }
 };
@@ -96,16 +117,19 @@ struct op_wrapper {
 struct __align__(OP_ALIGNMENT) op_state {
   char data[OP_SIZE];
 };
-extern "C" __device__ VALUE_T OP_NAME(op_state *state, VALUE_T val);
+extern "C" __device__ OUPUT_T OP_NAME(op_state *state, INPUT_T val);
 struct op_wrapper {
   op_state state;
-  __device__ VALUE_T operator()(VALUE_T val) {
+  __device__ OUTPUT_T operator()(INPUT_T val) {
     return OP_NAME(&state, val);
   }
 };
+
+
 )XXX";
 
   return (operation.type == cccl_op_kind_t::CCCL_STATELESS)
-         ? std::format(op_template, input_t, operation.name, "", "", stateless_op)
-         : std::format(op_template, input_t, operation.name, operation.alignment, operation.size, stateful_op);
+         ? std::format(unary_op_template, input_t, output_t, operation.name, "", "", stateless_op)
+         : std::format(
+             unary_op_template, input_t, output_t, operation.name, operation.alignment, operation.size, stateful_op);
 }

c/parallel/src/kernels/operators.h

@@ -14,6 +14,10 @@
 
 #include <cccl/c/types.h>
 
-std::string make_kernel_user_binary_operator(std::string_view input_value_t, cccl_op_t operation);
+std::string make_kernel_user_binary_operator(
+  std::string_view lhs_value_t, std::string_view rhs_value_t, std::string_view output_value_t, cccl_op_t operation);
+
+std::string
+make_kernel_user_unary_operator(std::string_view input_value_t, std::string_view output_value_t, cccl_op_t operation);
 
 std::string make_kernel_user_comparison_operator(std::string_view input_value_t, cccl_op_t operation);

c/parallel/src/reduce.cu

@@ -281,7 +281,7 @@ CUresult cccl_device_reduce_build(
     const std::string output_iterator_src =
       make_kernel_output_iterator(offset_t, output_iterator_typename, accum_cpp, output_it);
 
-    const std::string op_src = make_kernel_user_binary_operator(accum_cpp, op);
+    const std::string op_src = make_kernel_user_binary_operator(accum_cpp, accum_cpp, accum_cpp, op);
 
     const std::string src = std::format(
       R"XXX(

c/parallel/src/scan.cu

@@ -239,7 +239,7 @@ CUresult cccl_device_scan_build(
     const std::string output_iterator_src =
       make_kernel_output_iterator(offset_t, "output_iterator_t", accum_cpp, output_it);
 
-    const std::string op_src = make_kernel_user_binary_operator(accum_cpp, op);
+    const std::string op_src = make_kernel_user_binary_operator(accum_cpp, accum_cpp, accum_cpp, op);
 
     constexpr std::string_view src_template = R"XXX(
 #include <cub/block/block_scan.cuh>

c/parallel/src/segmented_reduce.cu

@@ -283,7 +283,7 @@ CUresult cccl_device_segmented_reduce_build(
     const std::string end_offset_iterator_src =
       make_kernel_input_iterator(offset_t, "end_offset_iterator_t", end_offset_it_value_t, end_offset_it);
 
-    const std::string op_src = make_kernel_user_binary_operator(accum_cpp, op);
+    const std::string op_src = make_kernel_user_binary_operator(accum_cpp, accum_cpp, accum_cpp, op);
 
     // agent_policy_t is to specify parameters like policy_hub does in dispatch_reduce.cuh
     constexpr std::string_view program_preamble_template = R"XXX(