[libc] Silence integer shortening warnings on NVPTX masks

[jhuber6]

Summary

Nvidia uses a 32-bit mask, but we store it in a common 64-bit integer to
provide it with a compatible ABI with the AMD implementaiton which may
use a 64-bit mask. Silence these warnings by explicitly casting to the
smaller value, we know this is always legal as the result will always
fit into the smaller value if it was generated on NVPTX.

[sivachandra]

libc/src/__support/GPU/nvptx/utils.h

119

May be cast the return value of this builtin also?

[jhuber6]

libc/src/__support/GPU/nvptx/utils.h

119

This should return a 32-bit value on the target.

libc/src/__support/GPU/nvptx/utils.h

[[clang::convergent]] LIBC_INLINE uint64_t get_lane_mask() {
   return mask;
 }
 
 /// Copies the value from the first active thread in the warp to the rest.
 [[clang::convergent]] LIBC_INLINE uint32_t broadcast_value(uint32_t x) {
   // NOTE: This is not sufficient in all cases on Volta hardware or later. The
   // lane mask returned here is not always the true lane mask used by the
   // intrinsics in cases of incedental or enforced divergence by the user.
-  uint64_t lane_mask = get_lane_mask();
-  uint64_t id = __builtin_ffsl(lane_mask) - 1;
+  uint32_t lane_mask = static_cast<uint32_t>(get_lane_mask());
+  uint32_t id = __builtin_ffs(lane_mask) - 1;
 #if __CUDA_ARCH__ >= 600
   return __nvvm_shfl_sync_idx_i32(lane_mask, x, id, get_lane_size() - 1);
 #else
   return __nvvm_shfl_idx_i32(x, id, get_lane_size() - 1);
 #endif
 }
 
 /// Returns a bitmask of threads in the current lane for which \p x is true.
 [[clang::convergent]] LIBC_INLINE uint64_t ballot(uint64_t lane_mask, bool x) {
 #if __CUDA_ARCH__ >= 600
-  return __nvvm_vote_ballot_sync(lane_mask, x);
+  return __nvvm_vote_ballot_sync(static_cast<uint32_t>(lane_mask), x);
 #else
-  return lane_mask & __nvvm_vote_ballot(x);
+  return static_cast<uint32_t>(lane_mask) & __nvvm_vote_ballot(x);
 #endif
 }
 /// Waits for all the threads in the block to converge and issues a fence.
 [[clang::convergent]] LIBC_INLINE void sync_threads() { __syncthreads(); }
 
 /// Waits for all threads in the warp to reconverge for independent scheduling.
 [[clang::convergent]] LIBC_INLINE void sync_lane(uint64_t mask) {
-  __nvvm_bar_warp_sync(mask);
+  __nvvm_bar_warp_sync(static_cast<uint32_t>(mask));
 }
 
 /// Returns the current value of the GPU's processor clock.
 LIBC_INLINE uint64_t processor_clock() {
   uint64_t timestamp;
   LIBC_INLINE_ASM("mov.u64  %0, %%clock64;" : "=l"(timestamp));
   return timestamp;
 }