Skip to content

Commit f205cbc

Browse files
PointKernelPointKernel
authored
Merge branch 'main' into host-span-cuda-std-span
2 parents f874272 + 35e7cd6 commit f205cbc

171 files changed

Lines changed: 3206 additions & 999 deletions

File tree

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

‎.clang-format‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,5 @@
1+
# SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
2+
# SPDX-License-Identifier: Apache-2.0
13
---
24
# Refer to the following link for the explanation of each params:
35
# http://releases.llvm.org/8.0.0/tools/clang/docs/ClangFormatStyleOptions.html

‎.pre-commit-config.yaml‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -242,6 +242,7 @@ repos:
242242
recipe[.]yaml$|
243243
dependencies[.]yaml$|
244244
pytest[.]ini$|
245+
^[.]clang-format$|
245246
^[.]pre-commit-config[.]yaml$|
246247
Makefile$
247248
# TODO: Remove FindCUDAToolkit once we require CMake 4.4

‎CONTRIBUTING.md‎

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -54,7 +54,10 @@ documentation docs](https://docs.rapids.ai/api/cudf/stable/cudf/developer_guide/
5454
8. Verify that CI passes all [status checks](https://docs.github.com/en/pull-requests/collaborating-with-pull-requests/collaborating-on-repositories-with-code-quality-features/about-status-checks).
5555
Fix if needed.
5656
9. Wait for other developers to review your code and update code as needed.
57-
Changes to any C++ files require at least 2 approvals from the cudf-cpp-codeowners before merging.
57+
Changes to libcudf C++ files require at least 2 approvals from the cudf-cpp-codeowners before
58+
merging.
59+
Changes limited to libcudf_streaming C++ files require at least 1 approval from the
60+
rapidsmpf-cpp-codeowners before merging.
5861
10. Once reviewed and approved, a RAPIDS developer will merge your pull request.
5962

6063
If you are unsure about anything, don't hesitate to comment on issues and ask for clarification!

‎cpp/CMakeLists.txt‎

Lines changed: 107 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -413,6 +413,9 @@ include(cmake/thirdparty/get_nanoarrow.cmake)
413413
# find thread_pool
414414
include(cmake/thirdparty/get_thread_pool.cmake)
415415

416+
# find xxhash
417+
include(cmake/thirdparty/get_xxhash.cmake)
418+
416419
# find zstd
417420
include(cmake/thirdparty/get_zstd.cmake)
418421

@@ -422,6 +425,9 @@ add_subdirectory(librtcx)
422425
# JIT Embedding helper functions
423426
include(librtcx/embed.cmake)
424427

428+
# Pre-compiled fragment management helper functions
429+
include(cmake/Modules/AddFragment.cmake)
430+
425431
# Workaround until https://github.com/rapidsai/rapids-cmake/issues/176 is resolved
426432
if(NOT BUILD_SHARED_LIBS)
427433
include("${rapids-cmake-dir}/export/find_package_file.cmake")
@@ -444,39 +450,39 @@ if(NOT BUILD_SHARED_LIBS)
444450
)
445451
endif()
446452

447-
add_embed(cudf_cuda_embed)
453+
rtcx_add_embed(cudf_cuda_embed)
448454

449-
embed_includes(
455+
rtcx_embed_includes(
450456
cudf_cuda_embed SOURCE_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/librtcx/libcxx DEST_DIRECTORY
451457
librtcx/libcxx INCLUDE_DIRECTORIES librtcx/libcxx
452458
)
453459

454-
embed_includes(
460+
rtcx_embed_includes(
455461
cudf_cuda_embed SOURCE_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include/cudf DEST_DIRECTORY
456462
cudf/cpp/include/cudf INCLUDE_DIRECTORIES cudf/cpp/include
457463
)
458464

459-
embed_includes(
465+
rtcx_embed_includes(
460466
cudf_cuda_embed SOURCE_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/src/jit DEST_DIRECTORY
461467
cudf/cpp/src/jit INCLUDE_DIRECTORIES cudf/cpp/src
462468
)
463469

464-
embed_includes(
470+
rtcx_embed_includes(
465471
cudf_cuda_embed SOURCE_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/src/binaryop/jit DEST_DIRECTORY
466472
cudf/cpp/src/binaryop/jit INCLUDE_DIRECTORIES cudf/cpp/src
467473
)
468474

469-
embed_includes(
475+
rtcx_embed_includes(
470476
cudf_cuda_embed SOURCE_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/src/join/jit DEST_DIRECTORY
471477
cudf/cpp/src/join/jit INCLUDE_DIRECTORIES cudf/cpp/src
472478
)
473479

474-
embed_includes(
480+
rtcx_embed_includes(
475481
cudf_cuda_embed SOURCE_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/src/rolling DEST_DIRECTORY
476482
cudf/cpp/src/rolling INCLUDE_DIRECTORIES cudf/cpp/src
477483
)
478484

479-
embed_includes(
485+
rtcx_embed_includes(
480486
cudf_cuda_embed SOURCE_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/src/transform/jit DEST_DIRECTORY
481487
cudf/cpp/src/transform/jit INCLUDE_DIRECTORIES cudf/cpp/src
482488
)
@@ -486,13 +492,91 @@ get_target_property(LIBCUDACXX_RAW_INCLUDE_DIRS CCCL::libcudacxx INTERFACE_INCLU
486492
foreach(INC_DIR IN LISTS LIBCUDACXX_RAW_INCLUDE_DIRS)
487493
cmake_path(GET INC_DIR FILENAME INC_DIR_NAME)
488494

489-
embed_includes(
495+
rtcx_embed_includes(
490496
cudf_cuda_embed SOURCE_DIRECTORY ${INC_DIR} DEST_DIRECTORY CCCL/libcudacxx/${INC_DIR_NAME}
491497
INCLUDE_DIRECTORIES CCCL/libcudacxx/${INC_DIR_NAME}
492498
)
493499
endforeach()
494500

495-
embed(cudf_cuda_embed COMPRESSION zstd OUTPUT_DIRECTORY "${CUDF_GENERATED_INCLUDE_DIR}/rtcx_embed")
501+
rtcx_embed(
502+
cudf_cuda_embed COMPRESSION zstd OUTPUT_DIRECTORY "${CUDF_GENERATED_INCLUDE_DIR}/rtcx_embed"
503+
)
504+
505+
rtcx_add_embed(cudf_fragments)
506+
507+
list(APPEND CUDF_PRECOMPILE_PHYSICAL_TYPES uint8_t uint16_t uint32_t uint64_t numeric::decimal32
508+
numeric::decimal64 numeric::decimal128
509+
)
510+
511+
foreach(TYPE IN ITEMS ${CUDF_PRECOMPILE_PHYSICAL_TYPES})
512+
set(FRAGMENT_NAME transform_kernel)
513+
get_property(
514+
FILE_INDEX
515+
TARGET cudf_fragments__embed_props
516+
PROPERTY EMBED_FILE_INDEX
517+
)
518+
set(VARIANT_NAME transform_kernel_${FILE_INDEX})
519+
set(INSTANCE
520+
"cudf::jit::transform_kernel<false, false, cudf::jit::type_list<cudf::jit::column_accessor<0ULL, cudf::column_device_view_core, ${TYPE}, false, 0>>, cudf::jit::type_list<cudf::jit::column_accessor<0ULL, cudf::mutable_column_device_view_core, ${TYPE}, false, 0>>>"
521+
)
522+
add_fragment(
523+
cudf_fragments
524+
FRAGMENT
525+
${VARIANT_NAME}
526+
SOURCE
527+
src/transform/jit/kernel.cu
528+
KERNEL_INSTANCE
529+
${INSTANCE}
530+
UDF_TYPE
531+
"int(${TYPE} *, ${TYPE})"
532+
DEFINITIONS
533+
CUDF_LTO_MODE
534+
ARRAY_IDS
535+
${FRAGMENT_NAME}_FILE_INDEX
536+
${FRAGMENT_NAME}_INSTANCE
537+
ARRAY_VALUES
538+
${FILE_INDEX}
539+
"${INSTANCE}"
540+
)
541+
endforeach()
542+
543+
foreach(TYPE IN ITEMS ${CUDF_PRECOMPILE_PHYSICAL_TYPES})
544+
foreach(RHS_IS_SCALAR IN ITEMS "false" "true")
545+
set(FRAGMENT_NAME transform_kernel)
546+
get_property(
547+
FILE_INDEX
548+
TARGET cudf_fragments__embed_props
549+
PROPERTY EMBED_FILE_INDEX
550+
)
551+
set(VARIANT_NAME transform_kernel_${FILE_INDEX})
552+
set(INSTANCE
553+
"cudf::jit::transform_kernel<false, false, cudf::jit::type_list<cudf::jit::column_accessor<0ULL, cudf::column_device_view_core, ${TYPE}, false, 0>, cudf::jit::column_accessor<1ULL, cudf::column_device_view_core, ${TYPE}, ${RHS_IS_SCALAR}, 0>>, cudf::jit::type_list<cudf::jit::column_accessor<0ULL, cudf::mutable_column_device_view_core, ${TYPE}, false, 0>>>"
554+
)
555+
add_fragment(
556+
cudf_fragments
557+
FRAGMENT
558+
${VARIANT_NAME}
559+
SOURCE
560+
src/transform/jit/kernel.cu
561+
KERNEL_INSTANCE
562+
${INSTANCE}
563+
UDF_TYPE
564+
"int(${TYPE} *, ${TYPE}, ${TYPE})"
565+
DEFINITIONS
566+
CUDF_LTO_MODE
567+
ARRAY_IDS
568+
${FRAGMENT_NAME}_FILE_INDEX
569+
${FRAGMENT_NAME}_INSTANCE
570+
ARRAY_VALUES
571+
${FILE_INDEX}
572+
"${INSTANCE}"
573+
)
574+
endforeach()
575+
endforeach()
576+
577+
rtcx_embed(
578+
cudf_fragments COMPRESSION none OUTPUT_DIRECTORY "${CUDF_GENERATED_INCLUDE_DIR}/rtcx_embed"
579+
)
496580

497581
# ##################################################################################################
498582
# * library targets -------------------------------------------------------------------------------
@@ -1058,9 +1142,10 @@ add_library(
10581142
src/utilities/type_checks.cpp
10591143
src/utilities/type_dispatcher.cpp
10601144
${cudf_cuda_embed_SOURCE_DIR}/cudf_cuda_embed.s
1145+
${cudf_fragments_SOURCE_DIR}/cudf_fragments.s
10611146
)
10621147

1063-
add_dependencies(cudf cudf_cuda_embed)
1148+
add_dependencies(cudf cudf_cuda_embed cudf_fragments)
10641149

10651150
set_property(
10661151
SOURCE src/io/parquet/writer_impl.cu
@@ -1130,7 +1215,9 @@ target_include_directories(
11301215
"$<BUILD_INTERFACE:${nanoarrow_SOURCE_DIR}/src>"
11311216
"$<BUILD_INTERFACE:${FlatBuffers_SOURCE_DIR}/include>"
11321217
"$<BUILD_INTERFACE:${ZSTD_INCLUDE_DIR}>"
1218+
"$<BUILD_INTERFACE:${CMAKE_CURRENT_LIST_DIR}/librtcx>"
11331219
"$<BUILD_INTERFACE:${cudf_cuda_embed_INCLUDE_DIRS}>"
1220+
"$<BUILD_INTERFACE:${cudf_fragments_INCLUDE_DIRS}>"
11341221
INTERFACE "$<INSTALL_INTERFACE:include>"
11351222
)
11361223

@@ -1172,8 +1259,15 @@ target_compile_definitions(cudf PRIVATE THRUST_FORCE_32_BIT_OFFSET_TYPE=1 CCCL_A
11721259
target_link_libraries(
11731260
cudf
11741261
PUBLIC CCCL::CCCL $<BUILD_LOCAL_INTERFACE:BS::thread_pool>
1175-
PRIVATE $<BUILD_LOCAL_INTERFACE:nvtx3::nvtx3-cpp> $<BUILD_LOCAL_INTERFACE:cuco::cuco> ZLIB::ZLIB
1176-
${CUDF_nvcomp_TARGET} kvikio::kvikio ${CUDF_nanoarrow_TARGET} zstd rtcx::rtcx
1262+
PRIVATE $<BUILD_LOCAL_INTERFACE:nvtx3::nvtx3-cpp>
1263+
$<BUILD_LOCAL_INTERFACE:cuco::cuco>
1264+
ZLIB::ZLIB
1265+
${CUDF_nvcomp_TARGET}
1266+
kvikio::kvikio
1267+
${CUDF_nanoarrow_TARGET}
1268+
zstd
1269+
$<BUILD_LOCAL_INTERFACE:xxhash>
1270+
rtcx::rtcx
11771271
)
11781272

11791273
# When rmm is a static library being absorbed via whole-archive, strip nvtx3 from its public

‎cpp/benchmarks/CMakeLists.txt‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -389,11 +389,24 @@ ConfigureNVBench(CSV_WRITER_NVBENCH io/csv/csv_writer.cpp)
389389
# * ast benchmark ---------------------------------------------------------------------------------
390390
ConfigureNVBench(AST_NVBENCH ast/polynomials.cpp ast/transform.cpp)
391391

392+
# ##################################################################################################
393+
# * LTO Fragments ----------------------------------------------------------------------------
394+
rtcx_add_embed(cudf_benchmark_fragments)
395+
add_fragment(cudf_benchmark_fragments FRAGMENT add_f32 SOURCE binaryop/fragments/add_f32.cu)
396+
add_fragment(cudf_benchmark_fragments FRAGMENT mul_f32 SOURCE binaryop/fragments/mul_f32.cu)
397+
rtcx_embed(
398+
cudf_benchmark_fragments COMPRESSION none OUTPUT_DIRECTORY
399+
"${CUDF_GENERATED_INCLUDE_DIR}/rtcx_embed"
400+
)
401+
392402
# ##################################################################################################
393403
# * binaryop benchmark ----------------------------------------------------------------------------
394404
ConfigureNVBench(
395405
BINARYOP_NVBENCH binaryop/binaryop.cpp binaryop/compiled_binaryop.cpp binaryop/polynomials.cpp
406+
${cudf_benchmark_fragments_SOURCE_DIR}/cudf_benchmark_fragments.s
396407
)
408+
target_include_directories(BINARYOP_NVBENCH PRIVATE ${cudf_benchmark_fragments_SOURCE_DIR})
409+
add_dependencies(BINARYOP_NVBENCH cudf_benchmark_fragments)
397410

398411
# ##################################################################################################
399412
# * transform benchmark

‎cpp/benchmarks/binaryop/compiled_binaryop.cpp‎

Lines changed: 116 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,13 +1,15 @@
11
/*
2-
* SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION.
2+
* SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
33
* SPDX-License-Identifier: Apache-2.0
44
*/
55

66
#include <benchmarks/common/generate_input.hpp>
77
#include <benchmarks/common/memory_stats.hpp>
88

99
#include <cudf/binaryop.hpp>
10+
#include <cudf/transform.hpp>
1011

12+
#include <cudf_benchmark_fragments.hpp>
1113
#include <nvbench/nvbench.cuh>
1214

1315
template <typename TypeLhs, typename TypeRhs, typename TypeOut>
@@ -72,6 +74,7 @@ BINARYOP_BENCHMARK_DEFINE(timestamp_s, duration_s, ADD, time
7274
BINARYOP_BENCHMARK_DEFINE(duration_s, duration_D, SUB, duration_ms);
7375
BINARYOP_BENCHMARK_DEFINE(int64_t, int64_t, SUB, int64_t);
7476
BINARYOP_BENCHMARK_DEFINE(float, float, MUL, int64_t);
77+
BINARYOP_BENCHMARK_DEFINE(float, float, MUL, float);
7578
BINARYOP_BENCHMARK_DEFINE(duration_s, int64_t, MUL, duration_s);
7679
BINARYOP_BENCHMARK_DEFINE(int64_t, int64_t, DIV, int64_t);
7780
BINARYOP_BENCHMARK_DEFINE(duration_ms, int32_t, DIV, duration_ms);
@@ -101,3 +104,115 @@ BINARYOP_BENCHMARK_DEFINE(duration_ms, duration_ns, NULL_EQUALS, bool
101104
BINARYOP_BENCHMARK_DEFINE(duration_ms, duration_ns, NULL_NOT_EQUALS, bool);
102105
BINARYOP_BENCHMARK_DEFINE(decimal32, decimal32, NULL_MAX, decimal32);
103106
BINARYOP_BENCHMARK_DEFINE(timestamp_D, timestamp_s, NULL_MIN, timestamp_s);
107+
// clang-format on
108+
109+
template <typename TypeLhs, typename TypeRhs, typename TypeOut>
110+
void BM_jit_binaryop(nvbench::state& state, cudf::binary_operator binop)
111+
{
112+
constexpr auto const jit_mul_cuda = R"***(
113+
__device__ void transform(float* out, float a, float b) {
114+
*out = a * b;
115+
}
116+
)***";
117+
118+
constexpr auto const jit_add_cuda = R"***(
119+
__device__ void transform(float* out, float a, float b) {
120+
*out = a + b;
121+
}
122+
)***";
123+
124+
auto const num_rows = static_cast<cudf::size_type>(state.get_int64("num_rows"));
125+
auto const use_lto = state.get_string("use_lto") == "true";
126+
static_assert(std::is_same_v<TypeLhs, TypeRhs> && std::is_same_v<TypeRhs, TypeOut>);
127+
static_assert(std::is_same_v<TypeLhs, float>);
128+
129+
auto const source_table = create_random_table(
130+
{cudf::type_to_id<TypeLhs>(), cudf::type_to_id<TypeRhs>()}, row_count{num_rows});
131+
132+
auto lhs = cudf::column_view(source_table->get_column(0));
133+
auto rhs = cudf::column_view(source_table->get_column(1));
134+
135+
size_t fragment_id = 0;
136+
char const* cuda = nullptr;
137+
138+
switch (binop) {
139+
case cudf::binary_operator::ADD: {
140+
fragment_id = cudf_benchmark_fragments::add_f32;
141+
cuda = jit_add_cuda;
142+
} break;
143+
case cudf::binary_operator::MUL: {
144+
fragment_id = cudf_benchmark_fragments::mul_f32;
145+
cuda = jit_mul_cuda;
146+
} break;
147+
default: throw std::runtime_error("Unsupported binary operator for JIT benchmark");
148+
}
149+
150+
// Call once for hot cache.
151+
cudf::transform_input inputs[] = {lhs, rhs};
152+
cudf::transform_output outputs[] = {
153+
{cudf::data_type{cudf::type_to_id<TypeOut>()}, cudf::output_nullability::ALL_VALID}};
154+
155+
auto const range = cudf_benchmark_fragments::file_ranges[fragment_id];
156+
std::span<uint8_t const> udf{cudf_benchmark_fragments::files.subspan(range[0], range[1])};
157+
158+
auto result = use_lto ? cudf::transform_lto(udf,
159+
cudf::lto_binary_type::FATBIN,
160+
cudf::null_aware::NO,
161+
std::nullopt,
162+
inputs,
163+
outputs,
164+
{},
165+
std::nullopt)
166+
: cudf::multi_transform(cuda,
167+
cudf::udf_source_type::CUDA,
168+
cudf::null_aware::NO,
169+
std::nullopt,
170+
inputs,
171+
outputs,
172+
{},
173+
std::nullopt);
174+
175+
// use number of bytes read and written to global memory
176+
state.add_global_memory_reads<TypeLhs>(num_rows);
177+
state.add_global_memory_reads<TypeRhs>(num_rows);
178+
state.add_global_memory_writes<TypeOut>(num_rows);
179+
180+
state.exec(nvbench::exec_tag::sync, [&](nvbench::launch&) {
181+
[[maybe_unused]] auto result = use_lto ? cudf::transform_lto(udf,
182+
cudf::lto_binary_type::FATBIN,
183+
cudf::null_aware::NO,
184+
std::nullopt,
185+
inputs,
186+
outputs,
187+
{},
188+
std::nullopt)
189+
: cudf::multi_transform(cuda,
190+
cudf::udf_source_type::CUDA,
191+
cudf::null_aware::NO,
192+
std::nullopt,
193+
inputs,
194+
outputs,
195+
{},
196+
std::nullopt);
197+
});
198+
}
199+
200+
#define BM_JIT_BINARYOP_BENCHMARK_DEFINE(name, lhs, rhs, bop, tout) \
201+
static void name(::nvbench::state& st) \
202+
{ \
203+
::BM_jit_binaryop<lhs, rhs, tout>(st, ::cudf::binary_operator::bop); \
204+
} \
205+
NVBENCH_BENCH(name) \
206+
.set_name("jit_binary_op_" BM_STRINGIFY(name)) \
207+
.add_int64_axis("num_rows", {10'000, 100'000, 1'000'000, 10'000'000, 100'000'000}) \
208+
.add_string_axis("use_lto", {"true", "false"})
209+
210+
#define build_name_jit(a, b, c, d) a##_##b##_##c##_##d##_jit
211+
212+
#define JIT_BINARYOP_BENCHMARK_DEFINE(lhs, rhs, bop, tout) \
213+
BM_JIT_BINARYOP_BENCHMARK_DEFINE(build_name_jit(bop, lhs, rhs, tout), lhs, rhs, bop, tout)
214+
215+
// clang-format off
216+
JIT_BINARYOP_BENCHMARK_DEFINE(float, float, ADD, float);
217+
JIT_BINARYOP_BENCHMARK_DEFINE(float, float, MUL, float);
218+
// clang-format on

0 commit comments

Comments
 (0)