Use thread_index_type to avoid index overflow in grid-stride loops

[PointKernel]

Description

This PR checks all related files under src/hash, src/bitmask and src/transform folders and fixes potential index overflow issues by using thread_index_type.

Checklist

• I am familiar with the Contributing Guidelines.
• New or existing tests cover these changes.
• The documentation is up to date with these changes.

[PointKernel]

Question for reviewers: I don't think we have tests exercising an input that could cause overflow. I hesitated to add those tests since it requires relatively large memory and may take a long time to execute. What do you think?

[ttnghia]

Probably we also need to upgrade first_ and last_ indices above.

[PointKernel]

We can reply on implicit conversions for < or <= since they are always cast to the larger data type for comparison. first_word_index and last_word_index are explicitly used as size_type later in the loop thus setting them to thread_index_type seems not worth the effort.

[mythrocks]

I think I understand, but I'd better confirm:

1. first_word_index and last_word_index are int32_t.
2. tid is uint32_t. So thread_word_index is eventually uint32_t.
3. while (thread_word_idx <= last_word_index) compares the two as uint32_t.
4. The place where thread_word_index might overflow is at line 275 (thread_word_index += stride). But both thread_word_index and stride are uint32_t already.

Yeah, it looks like this should work.

[karthikeyann]

Question for reviewers: I don't think we have tests exercising an input that could cause overflow. I hesitated to add those tests since it requires relatively large memory and may take a long time to execute. What do you think?

The for-loop comparison happens against shorter-width type size_type. So, is overflow possible to test?

[PointKernel]

The for-loop comparison happens against shorter-width type size_type. So, is overflow possible to test?

It should be possible if the input size is smaller than INT_MAX and larger than INT_MAX - stride, then the "last" iteration with int32_t index will produce negative indices when doing tid += stride; which is smaller than the end condition thus causing illegal memory access in the iteration right after.

[karthikeyann]

Benchmarks could be a good place for very large sizes (not for testing, but to catch illegal accesses). But Only unit tests are run through memcheck regularly. So, it won't be caught regularly unless it's added in unit tests. Besides, memcheck will be very slow to run on very large inputs.

Almost all of our algorithms aren't tested for INT32 max sizes. So, it may be okay to not add unit tests for this.
@PointKernel Were you able to verify with a unit test case locally? (How much time does it take to run?)

[PointKernel]

@PointKernel Were you able to verify with a unit test case locally? (How much time does it take to run?)

I wrote a test as below and realized bitmask is a special case: the loop end condition is no smaller than the number of bitmask words which is total_bits / num_bits_per_word. Since total_bits is size_type, the thread index would almost never overflow (unless we change the "word" type to something 1 bit only).

TEST_F(SetBitmaskTest, index_overflow)
{
  auto const begin = 0;
  auto const end   = INT_MAX - 10;
  auto const valid = true;
  auto const size  = end - begin;

  thrust::host_vector<bool> expected(size, valid);
  rmm::device_buffer mask = create_null_mask(size, cudf::mask_state::UNINITIALIZED);

  auto bitmask = static_cast<cudf::bitmask_type*>(mask.data());
  cudf::set_null_mask(bitmask, begin, end, valid);

  auto stream = cudf::get_default_stream();

  rmm::device_uvector<bool> output(size, stream);
  auto counting_iter = thrust::counting_iterator<cudf::size_type>{0};
  thrust::transform(rmm::exec_policy(stream),
                    counting_iter,
                    counting_iter + size,
                    output.begin(),
                    valid_bit_functor{bitmask});

  auto const result = thrust::all_of(
    rmm::exec_policy(stream), output.begin(), output.end(), thrust::identity<bool>{});

  EXPECT_EQ(result, valid);
}

Not sure if it's still relevant, to answer your question about runtime, the test took about 700 ms to run.

(base) yunsongw@yunsongw-dt:~/dev/rapids/cudf/cpp/build/release/gtests$ ./BITMASK_TEST --gtest_filter=SetBitmaskTest.index_overflow
Note: Google Test filter = SetBitmaskTest.index_overflow
[==========] Running 1 test from 1 test suite.
[----------] Global test environment set-up.
[----------] 1 test from SetBitmaskTest
[ RUN      ] SetBitmaskTest.index_overflow
[       OK ] SetBitmaskTest.index_overflow (708 ms)
[----------] 1 test from SetBitmaskTest (708 ms total)

[----------] Global test environment tear-down
[==========] 1 test from 1 test suite ran. (708 ms total)
[  PASSED  ] 1 test.

[mythrocks]

LGTM!

[PointKernel]

/merge