cuda-memcheck: initcheck and racecheck failed

Hi,
I have run cuda-memcheck on this project, but it showed errors on initcheck while compressing data and warnings on racecheck while decompressing data. I was running on **Tesla V100-SXM2(32GB) and A100-SXM4(40GB)** with **cuda-11.2.** I'm wondering whether it is a false alarm or not. Below are the examples of the errors and warnings and how I removed them.

The initcheck error can be removed by adding ```CHECKED_CUDA_CALL(cudaMemset, _memory, 0,size * sizeof(T));``` in **cuda_bits.cuh** (line 196 and 216).
Here is an example of initcheck error:
```
========= Uninitialized __global__ memory read of size 4
========= at 0x000000e0 in void ndzip::detail::gpu_cuda::hierarchical_inclusive_scan_reduce>(unsigned int*, unsigned int, unsigned int)
========= by thread (64,0,0) in block (0,0,0)
========= Address 0x7fffcfe20d00
========= Device Frame:void ndzip::detail::gpu_cuda::hierarchical_inclusive_scan_reduce>(unsigned int*, unsigned int, unsigned int) (void ndzip::detail::gpu_cuda::hierarchical_inclusive_scan_reduce>(unsigned int*, unsigned int, unsigned int) : 0xe0)
========= Saved host backtrace up to driver entry point
========= Host Frame:/lib64/libcuda.so.1 (cuLaunchKernel + 0x2b8) [0x222dc8]
========= Host Frame:${prefix_path}/ndzip/build/libndzip-cuda.so [0x277bb]
========= Host Frame:${prefix_path}/ndzip/build/libndzip-cuda.so [0x74270]
========= Host Frame:${prefix_path}/ndzip/build/libndzip-cuda.so (_ZN5ndzip6detail8gpu_cuda27hierarchical_inclusive_scanIjNS1_4plusIjEEEEvPT_RSt6vectorINS1_11cuda_bufferIS5_EESaIS9_EEjT0_P11CUstream_st + 0x1ed) [0xc91d]
========= Host Frame:${prefix_path}/ndzip/build/libndzip-cuda.so (_ZN5ndzip4cuda14compress_asyncIfLj1EEEvNS_5sliceIKT_XT0_EEEPvPmRNS0_25compressor_scratch_memoryIS3_XT0_EEEP11CUstream_st + 0x14b) [0xcc6b]
========= Host Frame:${prefix_path}/ndzip/build/libndzip-cuda.so (_ZNK5ndzip12cuda_encoderIfLj1EE8compressERKNS_5sliceIKfLj1EEEPvPNSt6chrono8durationImSt5ratioILl1ELl1000000000EEEE + 0x3ee) [0xefae]
========= Host Frame:./build/compress [0x11741]
========= Host Frame:./build/compress [0x12451]
========= Host Frame:./build/compress [0xb1b6]
========= Host Frame:./build/compress [0xab27]
========= Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22555]
========= Host Frame:./build/compress [0xaec3]
=========
```
Here is the code change
```
--- ndzip/src/ndzip/cuda_bits.cuh
+++ ndzip/src/ndzip/cuda_bits.cuh
@@ -193,6 +193,7 @@

explicit cuda_buffer(index_type size) : _size(size) {
CHECKED_CUDA_CALL(cudaMalloc, &_memory, size * sizeof(T));
+ CHECKED_CUDA_CALL(cudaMemset, _memory, 0,size * sizeof(T));//<---ADD
}

cuda_buffer(cuda_buffer &&other) noexcept {
@@ -212,6 +213,7 @@
void allocate(index_type size) {
reset();
CHECKED_CUDA_CALL(cudaMalloc, &_memory, size * sizeof(T));
+ CHECKED_CUDA_CALL(cudaMemset, _memory, 0,size * sizeof(T));//<---ADD
_size = size;
}

```

For the racecheck, warnings can be removed by adding ```__syncwarp()``` in **cuda_encoder.inl** after line 355. I think the warnings are because that some threads might enter line 323 in the second iteration while the other threads are still at line 345 in the first iteration after the __syncwarp() in line 326 ?

Here is the example of racecheck warning:

```
========= WARN: Race reported between Read access at 0x00000f20 in void ndzip::detail::gpu_cuda::decompress_block>(floatbits_type const *, ndzip::slice::data_type>, __scope__(dimensions)>)
========= and Write access at 0x00000b30 in void ndzip::detail::gpu_cuda::decompress_block>(floatbits_type const *, ndzip::slice::data_type>, __scope__(dimensions)>) [21864 hazards]
=========
```
Here is the code change
```
--- ndzip/src/ndzip/cuda_encoder.inl
+++ ndzip/src/ndzip/cuda_encoder.inl
@@ -353,6 +353,7 @@
__builtin_memcpy(&row_bits, row, sizeof row_bits);
hc.store(item, row_bits);
}
+ __syncwarp();//<-------ADD
} else {
// TODO duplication of the `item` calculation above. The term can be simplified!
for (index_type w = 0; w < warps_per_col_chunk; ++w) {
```
Thanks!

Start with src/ndzip/cuda_bits.cuh and src/ndzip/cuda_encoder.inl, then reproduce the reported initcheck and racecheck findings while compressing and decompressing. Compare the diagnostics with the proposed changes and determine whether they indicate real synchronization or initialization defects. Done means the findings are resolved or confirmed as false alarms without changing compression behavior.

Written by the indexing model from the issue text.