Fix racecheck in the gpu_debrotli_kernel

[davidwendt]

Description

Fixes a racecheck reported by compute-sanitizer in the gpu_debrotli_kernel code found by the COMPRESSION_TEST BrotliDecompressTest/BrotliDecompressTest.HelloWorld/0 gtest:

[ RUN      ] BrotliDecompressTest/BrotliDecompressTest.HelloWorld/0
========= Error: Race reported between Read access at cudf::io::detail::gpu_debrotli_kernel(cudf::device_span<const cudf::device_span<const unsigned char, (unsigned long)18446744073709551615>, (unsigned long)18446744073709551615>, cudf::device_span<const cudf::device_span<unsigned char, (unsigned long)18446744073709551615>, (unsigned long)18446744073709551615>, cudf::device_span<cudf::io::detail::codec_exec_result, (unsigned long)18446744073709551615>, unsigned char *, unsigned int)+0x1900 in debrotli.cu:1956
=========     and Write access at cudf::io::detail::initbits(cudf::io::detail::debrotli_state_s *, const unsigned char *, unsigned long, unsigned long)+0x1c20 in debrotli.cu:200 [8 hazards]

The s->cur variable is read by all threads but may be set by thread 0 in a race condition.
Adding a __syncthreads() right after the s->cur is read ensures it is not reset by any threads before all threads in the block have read it.

Checklist

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

[copy-pr-bot]

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

[davidwendt]

/ok to test

[davidwendt]

/merge