Skip to content

Commit

Permalink
delete some debug printing
Browse files Browse the repository at this point in the history
  • Loading branch information
@pmattione-nvidia
pmattione-nvidia committed Oct 7, 2024
1 parent 4b7d1df commit b36b3b2
Show file tree
Hide file tree
Showing 5 changed files with 1 addition and 58 deletions.
14 changes: 1 addition & 13 deletions cpp/include/cudf/table/experimental/row_operators.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1429,30 +1429,18 @@ class device_row_comparator {
__device__ bool operator()(size_type const lhs_element_index,
size_type const rhs_element_index) const noexcept
{
static constexpr bool enable_print = false;
if (check_nulls) {
bool const lhs_is_null{lhs.is_null(lhs_element_index)};
bool const rhs_is_null{rhs.is_null(rhs_element_index)};
if (lhs_is_null and rhs_is_null) {
return nulls_are_equal == null_equality::EQUAL;
} else if (lhs_is_null != rhs_is_null) {
if constexpr (enable_print) {
printf("NULLS UNEQUAL AT %d, %d; values: %d %d\n",
lhs_element_index, rhs_element_index, int(lhs_is_null), int(rhs_is_null));
}
return false;
}
}

bool result = comparator(lhs.element<Element>(lhs_element_index),
return comparator(lhs.element<Element>(lhs_element_index),
rhs.element<Element>(rhs_element_index));
if constexpr (enable_print && cuda::std::is_integral_v<Element>) {
if(!result) {
printf("VALUES UNEQUAL: AT %d, %d, VALUES %d, %d\n", lhs_element_index, rhs_element_index,
(int)lhs.element<Element>(lhs_element_index), (int)rhs.element<Element>(rhs_element_index));
}
}
return result;
}

template <typename Element,
Expand Down
18 changes: 0 additions & 18 deletions cpp/src/io/parquet/page_data.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -89,14 +89,6 @@ inline __device__ void gpuStoreOutput(uint32_t* dst,
bytebuf = 0;
}
*dst = bytebuf;

static constexpr bool enable_print = false;
if constexpr (enable_print) {
if (threadIdx.x == 0) {
printf("STORE VALUE %u at %p, src8 %p, dict_pos %u, dict_size %u, ofs %u\n",
bytebuf, dst, src8, dict_pos, dict_size, ofs);
}
}
}

/**
Expand Down Expand Up @@ -347,16 +339,6 @@ inline __device__ void gpuOutputFast(page_state_s* s, state_buf* sb, int src_pos
dict = s->data_start;
}
dict_pos *= (uint32_t)s->dtype_len_in;

static constexpr bool enable_print = false;
if constexpr (enable_print) {
if (threadIdx.x == 0) {
auto dict_lookup_idx = rolling_index<state_buf::dict_buf_size>(src_pos);
printf("PREP OUTPUT VALUE at dst %p, dict %p, dict_pos %u, dict_size %u, dict_base %p, dict_bits %d, dict_lookup_idx %d, dtype_len_in %d\n",
dst, dict, dict_pos, dict_size, s->dict_base, s->dict_bits, dict_lookup_idx, s->dtype_len_in);
}
}

gpuStoreOutput(dst, dict, dict_pos, dict_size);
}

Expand Down
9 changes: 0 additions & 9 deletions cpp/src/io/parquet/reader_impl.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -50,8 +50,6 @@ void reader::impl::decode_page_data(read_mode mode, size_t skip_rows, size_t num
{
auto& pass = *_pass_itm_data;
auto& subpass = *pass.subpass;
//printf("PREP LAUNCH: decode_page_data: mode %d, skip_rows %lu, num_rows %lu, #pages %lu\n",
// (int)mode, skip_rows, num_rows, subpass.pages.size());

auto& page_nesting = subpass.page_nesting_info;
auto& page_nesting_decode = subpass.page_nesting_decode_info;
Expand Down Expand Up @@ -223,11 +221,6 @@ void reader::impl::decode_page_data(read_mode mode, size_t skip_rows, size_t num
int const nkernels = std::bitset<32>(kernel_mask).count();
auto streams = cudf::detail::fork_streams(_stream, nkernels);

static constexpr bool enable_print = false;
if constexpr (enable_print) {
printf("PAGE DATA DECODE MASK: %d\n", kernel_mask);
}

// launch string decoder
int s_idx = 0;
if (BitAnd(kernel_mask, decode_kernel_mask::STRING) != 0) {
Expand Down Expand Up @@ -419,11 +412,9 @@ void reader::impl::decode_page_data(read_mode mode, size_t skip_rows, size_t num
page_nesting.device_to_host_async(_stream);
page_nesting_decode.device_to_host_async(_stream);

//printf("SYNC ERROR CODE\n");
if (auto const error = error_code.value_sync(_stream); error != 0) {
CUDF_FAIL("Parquet data decode failed with code(s) " + kernel_error::to_string(error));
}
//printf("ERROR CODE SUNK\n");

// for list columns, add the final offset to every offset buffer.
// TODO : make this happen in more efficiently. Maybe use thrust::for_each
Expand Down
12 changes: 0 additions & 12 deletions cpp/src/io/parquet/reader_impl_preprocess.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -138,12 +138,7 @@ void generate_depth_remappings(
// depth.
//

static constexpr bool enable_print = false;

// compute "X" from above
if constexpr (enable_print) {
printf("REMAPPING: max def %d, max rep %d\n", schema.max_definition_level, schema.max_repetition_level);
}
for (int s_idx = schema.max_repetition_level; s_idx >= 0; s_idx--) {
auto find_shallowest = [&](int r) {
int shallowest = -1;
Expand All @@ -162,9 +157,6 @@ void generate_depth_remappings(
if (!cur_schema.is_stub()) { cur_depth--; }
schema_idx = cur_schema.parent_idx;
}
if constexpr (enable_print) {
printf("REMAPPING: s_idx / r %d, shallowest %d\n", r, shallowest);
}
return shallowest;
};
rep_depth_remap[s_idx] = find_shallowest(s_idx);
Expand Down Expand Up @@ -203,10 +195,6 @@ void generate_depth_remappings(
prev_schema = cur_schema;
schema_idx = cur_schema.parent_idx;
}

if constexpr (enable_print) {
printf("REMAPPING: s_idx %d, r1 %d, end_depth %d\n", s_idx, r1, depth);
}
return depth;
};
def_depth_remap[s_idx] = find_deepest(s_idx);
Expand Down
6 changes: 0 additions & 6 deletions cpp/src/io/parquet/rle_stream.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -252,8 +252,6 @@ struct rle_stream {
run.level_run = level_run;
run.remaining = run.size;
cur += run_bytes;
//printf("STORE RUN: decode_index %d, fill_index %d, output_pos %d, run.size %d\n",
//decode_index, fill_index, output_pos, run.size);
output_pos += run.size;
fill_index++;
}
Expand Down Expand Up @@ -355,8 +353,6 @@ struct rle_stream {
// this is the last batch we will process this iteration if:
// - either this run still has remaining values
// - or it is consumed fully and its last index corresponds to output_count
//printf("STATUS: run_index %d, batch_len %d, remaining %d, at_end %d, last_run_pos %d, cur_values %d\n",
//run_index, batch_len, remaining, at_end, last_run_pos, cur_values);
if (remaining > 0 || at_end) { values_processed_shared = output_count; }
if (remaining == 0 && (at_end || is_last_decode_warp(warp_id))) {
decode_index_shared = run_index + 1;
Expand Down Expand Up @@ -401,15 +397,13 @@ struct rle_stream {
}

if((output_pos + run_size) > target_count) {
//printf("SKIPPING: target_count %d, run_size %d, output_pos %d\n", target_count, run_size, output_pos);
return output_pos; //bail! we've reached the starting one
}

output_pos += run_size;
cur += run_bytes;
}

//printf("SKIPPING: target_count %d, output_pos %d\n", target_count, output_pos);
return output_pos; //we skipped everything
}

Expand Down

0 comments on commit b36b3b2

Please sign in to comment.