[libomptarget][amdgpu] Fix truncation error for partial wavefront

[JonChesterfield]

Summary

[libomptarget][amdgpu] Fix truncation error for partial wavefront

The partial barrier implementation involves one wavefront resetting and N-1
waiting. This change future proofs against launching with a number of threads
that is not a multiple of the wavefront size.

openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip

 EXTERN __kmpc_impl_lanemask_t __kmpc_impl_activemask() {
   return __builtin_amdgcn_read_exec();
 }
 
 static void pteam_mem_barrier(uint32_t num_threads, uint32_t * barrier_state)
 {
   __atomic_thread_fence(__ATOMIC_ACQUIRE);
 
-  uint32_t num_waves = num_threads / WARPSIZE;
+  uint32_t num_waves = (num_threads + WARPSIZE - 1) / WARPSIZE;
 
   // Partial barrier implementation for amdgcn.
   // Uses two 16 bit unsigned counters. One for the number of waves to have
   // reached the barrier, and one to count how many times the barrier has been
   // passed. These are packed in a single atomically accessed 32 bit integer.
   // Low bits for the number of waves, assumed zero before this call.
   // High bits to count the number of times the barrier has been passed.