[openmp][AMDGPU] - Correct getNumberOfBlocks calculation.

[estewart08]

Summary

This patch fixes the 6 amdgpu buildbot lit test failures
introduced by https://reviews.llvm.org/D135444.

libomptarget :: amdgcn-amd-amdhsa :: mapping/reduction_implicit_map.cpp
libomptarget :: amdgcn-amd-amdhsa :: offloading/cuda_no_devices.c
libomptarget :: amdgcn-amd-amdhsa :: offloading/target-teams-atomic.c
libomptarget :: amdgcn-amd-amdhsa-LTO :: mapping/reduction_implicit_map.cpp
libomptarget :: amdgcn-amd-amdhsa-LTO :: offloading/cuda_no_devices.c
libomptarget :: amdgcn-amd-amdhsa-LTO :: offloading/target-teams-atomic.c

[estewart08]

Adding clang-format to PATH.

[estewart08]

Updated summary.

openmp/libomptarget/DeviceRTL/src/Mapping.cpp

 //===------- Mapping.cpp - OpenMP device runtime mapping helpers -- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //

 }
 
 uint32_t getThreadIdInBlock() { return __builtin_amdgcn_workitem_id_x(); }
 
 uint32_t getKernelSize() { return __builtin_amdgcn_grid_size_x(); }
 
 uint32_t getBlockId() { return __builtin_amdgcn_workgroup_id_x(); }
 
-uint32_t getNumberOfBlocks() { return __builtin_amdgcn_grid_size_x(); }
+uint32_t getNumberOfBlocks() {
+  return __builtin_amdgcn_grid_size_x() / __builtin_amdgcn_workgroup_size_x();
+}
 
 uint32_t getWarpId() {
   return impl::getThreadIdInBlock() / mapping::getWarpSize();
 }
 
 uint32_t getNumberOfWarpsInBlock() {
   return mapping::getBlockSize() / mapping::getWarpSize();
 }