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

Jump to bottom

[FIX] Fix incorrect nullCount in get_json_object #11633

Merged
merged 3 commits into from
Sep 2, 2022
Merged

[FIX] Fix incorrect nullCount in get_json_object #11633

Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
1aa4ec8
[FIX] Fix incorrect `nullCount` in `get_json_object`
trxcllnt Aug 31, 2022
a633da3
Initialize `out_valid_count` before it is filled.
thomcom Sep 1, 2022
1abd0d2
Revert previous proposal for correcting the initialization concern.
thomcom Sep 1, 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
1 change: 0 additions & 1 deletion cpp/src/strings/json/json_path.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -909,7 +909,6 @@ __launch_bounds__(block_size) __global__
size_type tid = threadIdx.x + (blockDim.x * blockIdx.x);
size_type stride = blockDim.x * gridDim.x;

if (out_valid_count.has_value()) { *(out_valid_count.value()) = 0; }
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

So this reset is no longer needed? I see that out_valid_count is then updated by atomicAdd, but what if out_valid_count was never been initialized before being passed into this function?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It is initialized here:

cudf/cpp/src/strings/json/json_path.cu

Line 1020 in 1aa4ec8

rmm::device_scalar<size_type> d_valid_count{0, stream};

I don't know why this if-statement was here. Maybe it was for some debug purpose.

Copy link
Contributor

@thomcom thomcom Sep 1, 2022 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It is initialized here but I see your point. It can probably be safely initialized here, by threadIdx.x = 0, the thread that is responsible for loading the correct value.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If so then this looks somewhat unsafe IMO because this needs to rely on the caller to initialize the variable without any guarantee from anywhere.

If we decided to not have the initialization here, it is better to have a comment line in the function doxygen clearly stressing/clarifying that.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I moved the initialization @ttnghia @davidwendt please re-review.

Copy link
Contributor Author

@trxcllnt trxcllnt Sep 1, 2022 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If so then this looks somewhat unsafe IMO because this needs to rely on the caller to initialize the variable without any guarantee from anywhere.

Passing a pointer to an uninitialized device_scalar seems anti-RAII. If this were regular host code passing a pointer to a function and expecting that function to initialize it, without a very good reason (like that function has knowledge that the caller doesn't about how it should be initialized) I'd probably flag it as a code smell.

size_type warp_valid_count{0};

auto active_threads = __ballot_sync(0xffff'ffffu, tid < col.size());
Expand Down