-
Notifications
You must be signed in to change notification settings - Fork 894
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Merge branch 'branch-24.10' into ipv4-uint32
- Loading branch information
Showing
33 changed files
with
548 additions
and
207 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,101 @@ | ||
/* | ||
* Copyright (c) 2024, NVIDIA CORPORATION. | ||
* | ||
* Licensed under the Apache License, Version 2.0 (the "License"); | ||
* you may not use this file except in compliance with the License. | ||
* You may obtain a copy of the License at | ||
* | ||
* http://www.apache.org/licenses/LICENSE-2.0 | ||
* | ||
* Unless required by applicable law or agreed to in writing, software | ||
* distributed under the License is distributed on an "AS IS" BASIS, | ||
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | ||
* See the License for the specific language governing permissions and | ||
* limitations under the License. | ||
*/ | ||
|
||
#include <benchmarks/common/generate_input.hpp> | ||
#include <benchmarks/fixture/benchmark_fixture.hpp> | ||
#include <benchmarks/io/cuio_common.hpp> | ||
#include <benchmarks/io/nvbench_helpers.hpp> | ||
|
||
#include <cudf/io/parquet.hpp> | ||
#include <cudf/utilities/default_stream.hpp> | ||
|
||
#include <nvbench/nvbench.cuh> | ||
|
||
// Size of the data in the benchmark dataframe; chosen to be low enough to allow benchmarks to | ||
// run on most GPUs, but large enough to allow highest throughput | ||
constexpr size_t data_size = 512 << 20; | ||
|
||
void parquet_read_common(cudf::size_type num_rows_to_read, | ||
cudf::size_type num_cols_to_read, | ||
cuio_source_sink_pair& source_sink, | ||
nvbench::state& state) | ||
{ | ||
cudf::io::parquet_reader_options read_opts = | ||
cudf::io::parquet_reader_options::builder(source_sink.make_source_info()); | ||
|
||
auto mem_stats_logger = cudf::memory_stats_logger(); | ||
state.set_cuda_stream(nvbench::make_cuda_stream_view(cudf::get_default_stream().value())); | ||
state.exec( | ||
nvbench::exec_tag::sync | nvbench::exec_tag::timer, [&](nvbench::launch& launch, auto& timer) { | ||
try_drop_l3_cache(); | ||
|
||
timer.start(); | ||
auto const result = cudf::io::read_parquet(read_opts); | ||
timer.stop(); | ||
|
||
CUDF_EXPECTS(result.tbl->num_columns() == num_cols_to_read, "Unexpected number of columns"); | ||
CUDF_EXPECTS(result.tbl->num_rows() == num_rows_to_read, "Unexpected number of rows"); | ||
}); | ||
|
||
auto const time = state.get_summary("nv/cold/time/gpu/mean").get_float64("value"); | ||
state.add_element_count(static_cast<double>(data_size) / time, "bytes_per_second"); | ||
state.add_buffer_size( | ||
mem_stats_logger.peak_memory_usage(), "peak_memory_usage", "peak_memory_usage"); | ||
state.add_buffer_size(source_sink.size(), "encoded_file_size", "encoded_file_size"); | ||
} | ||
|
||
template <data_type DataType> | ||
void bench_batched_memset(nvbench::state& state, nvbench::type_list<nvbench::enum_type<DataType>>) | ||
{ | ||
auto const d_type = get_type_or_group(static_cast<int32_t>(DataType)); | ||
auto const num_cols = static_cast<cudf::size_type>(state.get_int64("num_cols")); | ||
auto const cardinality = static_cast<cudf::size_type>(state.get_int64("cardinality")); | ||
auto const run_length = static_cast<cudf::size_type>(state.get_int64("run_length")); | ||
auto const source_type = retrieve_io_type_enum(state.get_string("io_type")); | ||
auto const compression = cudf::io::compression_type::NONE; | ||
cuio_source_sink_pair source_sink(source_type); | ||
auto const tbl = | ||
create_random_table(cycle_dtypes(d_type, num_cols), | ||
table_size_bytes{data_size}, | ||
data_profile_builder().cardinality(cardinality).avg_run_length(run_length)); | ||
auto const view = tbl->view(); | ||
|
||
cudf::io::parquet_writer_options write_opts = | ||
cudf::io::parquet_writer_options::builder(source_sink.make_sink_info(), view) | ||
.compression(compression); | ||
cudf::io::write_parquet(write_opts); | ||
auto const num_rows = view.num_rows(); | ||
|
||
parquet_read_common(num_rows, num_cols, source_sink, state); | ||
} | ||
|
||
using d_type_list = nvbench::enum_type_list<data_type::INTEGRAL, | ||
data_type::FLOAT, | ||
data_type::DECIMAL, | ||
data_type::TIMESTAMP, | ||
data_type::DURATION, | ||
data_type::STRING, | ||
data_type::LIST, | ||
data_type::STRUCT>; | ||
|
||
NVBENCH_BENCH_TYPES(bench_batched_memset, NVBENCH_TYPE_AXES(d_type_list)) | ||
.set_name("batched_memset") | ||
.set_type_axes_names({"data_type"}) | ||
.add_int64_axis("num_cols", {1000}) | ||
.add_string_axis("io_type", {"DEVICE_BUFFER"}) | ||
.set_min_samples(4) | ||
.add_int64_axis("cardinality", {0, 1000}) | ||
.add_int64_axis("run_length", {1, 32}); |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,82 @@ | ||
/* | ||
* Copyright (c) 2024, NVIDIA CORPORATION. | ||
* | ||
* Licensed under the Apache License, Version 2.0 (the "License"); | ||
* you may not use this file except in compliance with the License. | ||
* You may obtain a copy of the License at | ||
* | ||
* http://www.apache.org/licenses/LICENSE-2.0 | ||
* | ||
* Unless required by applicable law or agreed to in writing, software | ||
* distributed under the License is distributed on an "AS IS" BASIS, | ||
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | ||
* See the License for the specific language governing permissions and | ||
* limitations under the License. | ||
*/ | ||
|
||
#include <cudf/detail/iterator.cuh> | ||
#include <cudf/detail/utilities/vector_factories.hpp> | ||
|
||
#include <rmm/cuda_stream_view.hpp> | ||
#include <rmm/device_buffer.hpp> | ||
#include <rmm/resource_ref.hpp> | ||
|
||
#include <cub/device/device_copy.cuh> | ||
#include <cuda/functional> | ||
#include <thrust/iterator/counting_iterator.h> | ||
#include <thrust/iterator/transform_iterator.h> | ||
#include <thrust/transform.h> | ||
|
||
namespace CUDF_EXPORT cudf { | ||
namespace io::detail { | ||
|
||
/** | ||
* @brief A helper function that takes in a vector of device spans and memsets them to the | ||
* value provided using batches sent to the GPU. | ||
* | ||
* @param bufs Vector with device spans of data | ||
* @param value Value to memset all device spans to | ||
* @param _stream Stream used for device memory operations and kernel launches | ||
* | ||
* @return The data in device spans all set to value | ||
*/ | ||
template <typename T> | ||
void batched_memset(std::vector<cudf::device_span<T>> const& bufs, | ||
T const value, | ||
rmm::cuda_stream_view stream) | ||
{ | ||
// define task and bytes parameters | ||
auto const num_bufs = bufs.size(); | ||
|
||
// copy bufs into device memory and then get sizes | ||
auto gpu_bufs = | ||
cudf::detail::make_device_uvector_async(bufs, stream, rmm::mr::get_current_device_resource()); | ||
|
||
// get a vector with the sizes of all buffers | ||
auto sizes = cudf::detail::make_counting_transform_iterator( | ||
static_cast<std::size_t>(0), | ||
cuda::proclaim_return_type<std::size_t>( | ||
[gpu_bufs = gpu_bufs.data()] __device__(std::size_t i) { return gpu_bufs[i].size(); })); | ||
|
||
// get an iterator with a constant value to memset | ||
auto iter_in = thrust::make_constant_iterator(thrust::make_constant_iterator(value)); | ||
|
||
// get an iterator pointing to each device span | ||
auto iter_out = thrust::make_transform_iterator( | ||
thrust::counting_iterator<std::size_t>(0), | ||
cuda::proclaim_return_type<T*>( | ||
[gpu_bufs = gpu_bufs.data()] __device__(std::size_t i) { return gpu_bufs[i].data(); })); | ||
|
||
size_t temp_storage_bytes = 0; | ||
|
||
cub::DeviceCopy::Batched(nullptr, temp_storage_bytes, iter_in, iter_out, sizes, num_bufs, stream); | ||
|
||
rmm::device_buffer d_temp_storage( | ||
temp_storage_bytes, stream, rmm::mr::get_current_device_resource()); | ||
|
||
cub::DeviceCopy::Batched( | ||
d_temp_storage.data(), temp_storage_bytes, iter_in, iter_out, sizes, num_bufs, stream); | ||
} | ||
|
||
} // namespace io::detail | ||
} // namespace CUDF_EXPORT cudf |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Oops, something went wrong.