Fix bitmask_tests.cpp host accessing device memory (#8370)

This fixes an error in the `cleanEndWord` utility in the `bitmask_tests.cpp` source file. The code is trying to set the last remaining bits of a bitmask word by accessing the byte through a device pointer in host code. The end result was not actually being used so it appears the release compile optimized out the entire function. The expected segfault does occur in a debug build.
This PR corrects the logic to copy the last word, set the appropriate bits, and then copy the word back to the device buffer.

Authors:
  - David Wendt (https://github.com/davidwendt)

Approvers:
  - Mark Harris (https://github.com/harrism)
  - Nghia Truong (https://github.com/ttnghia)

URL: #8370

cpp/tests/bitmask/bitmask_tests.cpp

@@ -384,8 +384,12 @@ void cleanEndWord(rmm::device_buffer &mask, int begin_bit, int end_bit)
   auto number_of_mask_words = cudf::num_bitmask_words(static_cast<size_t>(end_bit - begin_bit));
   auto number_of_bits       = end_bit - begin_bit;
   if (number_of_bits % 32 != 0) {
-    auto end_mask = ptr[number_of_mask_words - 1];
-    end_mask      = end_mask & ((1 << (number_of_bits % 32)) - 1);
+    cudf::bitmask_type end_mask = 0;
+    CUDA_TRY(cudaMemcpy(
+      &end_mask, ptr + number_of_mask_words - 1, sizeof(end_mask), cudaMemcpyDeviceToHost));
+    end_mask = end_mask & ((1 << (number_of_bits % 32)) - 1);
+    CUDA_TRY(cudaMemcpy(
+      ptr + number_of_mask_words - 1, &end_mask, sizeof(end_mask), cudaMemcpyHostToDevice));
   }
 }