Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Reduces memory requirements in JSON parser and adds bytes/s and peak memory usage to benchmarks #11732

Merged
merged 5 commits into from
Sep 22, 2022
Merged

Reduces memory requirements in JSON parser and adds bytes/s and peak memory usage to benchmarks #11732

Show file tree
Hide file tree
Changes from 4 commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
232c9b4
adds memory requirements and bytes/sec to bm
elstehle Sep 21, 2022
01d6605
lowers memory requirements in json parser
elstehle Sep 21, 2022
c2492e1
Merge remote-tracking branch 'upstream/branch-22.10' into feature/jso…
elstehle Sep 21, 2022
9d19481
Merge remote-tracking branch 'upstream/branch-22.10' into feature/jso…
elstehle Sep 22, 2022
a9d155a
addresses review comments
elstehle Sep 22, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
9 changes: 8 additions & 1 deletion cpp/benchmarks/io/json/nested_json.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,6 +15,7 @@
*/

#include <benchmarks/common/generate_input.hpp>
#include <benchmarks/fixture/benchmark_fixture.hpp>
#include <benchmarks/fixture/rmm_pool_raii.hpp>

#include <nvbench/nvbench.cuh>
Expand Down Expand Up @@ -71,15 +72,21 @@ void BM_NESTED_JSON(nvbench::state& state)
state.add_element_count(input.size());

// Run algorithm
auto mem_stats_logger = cudf::memory_stats_logger();
elstehle marked this conversation as resolved.
Show resolved Hide resolved
state.set_cuda_stream(nvbench::make_cuda_stream_view(cudf::default_stream_value.value()));
state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
// Allocate device-side temporary storage & run algorithm
cudf::io::json::detail::parse_nested_json(input, default_options, cudf::default_stream_value);
});

auto const time = state.get_summary("nv/cold/time/gpu/mean").get_float64("value");
state.add_element_count(static_cast<double>(string_size) / time, "bytes_per_second");
state.add_buffer_size(
mem_stats_logger.peak_memory_usage(), "peak_memory_usage", "peak_memory_usage");
}

NVBENCH_BENCH(BM_NESTED_JSON)
.set_name("nested_json_gpu_parser")
.add_int64_power_of_two_axis("string_size", nvbench::range(20, 31, 1));
.add_int64_power_of_two_axis("string_size", nvbench::range(20, 30, 1));

} // namespace cudf
38 changes: 20 additions & 18 deletions cpp/src/io/json/nested_json_gpu.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1079,26 +1079,25 @@ std::pair<rmm::device_uvector<PdaTokenT>, rmm::device_uvector<SymbolOffsetT>> ge
// Range of encapsulating function that parses to internal columnar data representation
CUDF_FUNC_RANGE();

rmm::device_uvector<PdaTokenT> tokens{json_in.size(), stream, mr};
rmm::device_uvector<SymbolOffsetT> tokens_indices{json_in.size(), stream, mr};
rmm::device_scalar<SymbolOffsetT> num_written_tokens{stream, mr};

auto const new_line_delimited_json = options.is_enabled_lines();

// Memory holding the top-of-stack stack context for the input
rmm::device_uvector<StackSymbolT> stack_op_indices{json_in.size(), stream};

// Identify what is the stack context for each input character (is it: JSON-root, struct, or list)
get_stack_context(json_in, stack_op_indices.data(), stream);

// Prepare for PDA transducer pass, merging input symbols with stack symbols
rmm::device_uvector<PdaSymbolGroupIdT> pda_sgids{json_in.size(), stream};
auto zip_in = thrust::make_zip_iterator(json_in.data(), stack_op_indices.data());
thrust::transform(rmm::exec_policy(stream),
zip_in,
zip_in + json_in.size(),
pda_sgids.data(),
tokenizer_pda::PdaSymbolToSymbolGroupId{});
rmm::device_uvector<PdaSymbolGroupIdT> pda_sgids = [json_in, stream, mr]() {
rmm::device_uvector<PdaSymbolGroupIdT> pda_sgids{json_in.size(), stream, mr};
// Memory holding the top-of-stack stack context for the input
rmm::device_uvector<StackSymbolT> stack_op_indices{json_in.size(), stream, mr};

// Identify what is the stack context for each input character (JSON-root, struct, or list)
get_stack_context(json_in, stack_op_indices.data(), stream);

auto zip_in = thrust::make_zip_iterator(json_in.data(), stack_op_indices.data());
thrust::transform(rmm::exec_policy(stream),
zip_in,
zip_in + json_in.size(),
pda_sgids.data(),
tokenizer_pda::PdaSymbolToSymbolGroupId{});
return pda_sgids;
}();

// PDA transducer alias
using ToTokenStreamFstT =
Expand All @@ -1118,6 +1117,9 @@ std::pair<rmm::device_uvector<PdaTokenT>, rmm::device_uvector<SymbolOffsetT>> ge
stream};

// Perform a PDA-transducer pass
rmm::device_scalar<SymbolOffsetT> num_written_tokens{stream, mr};
elstehle marked this conversation as resolved.
Show resolved Hide resolved
rmm::device_uvector<PdaTokenT> tokens{json_in.size(), stream, mr};
rmm::device_uvector<SymbolOffsetT> tokens_indices{json_in.size(), stream, mr};
json_to_tokens_fst.Transduce(pda_sgids.begin(),
static_cast<SymbolOffsetT>(json_in.size()),
tokens.data(),
Expand All @@ -1126,7 +1128,7 @@ std::pair<rmm::device_uvector<PdaTokenT>, rmm::device_uvector<SymbolOffsetT>> ge
tokenizer_pda::start_state,
stream);

auto num_total_tokens = num_written_tokens.value(stream);
auto const num_total_tokens = num_written_tokens.value(stream);
tokens.resize(num_total_tokens, stream);
tokens_indices.resize(num_total_tokens, stream);

Expand Down