[OpenMP][NVPTX] Replaced CUDA builtin vars with LLVM intrinsics

[tianshilei1992]

Summary

Replaced CUDA builtin vars with LLVM intrinsics such that we don't need
definitions of those intrinsics.

[tianshilei1992]

Fixed clang-format

[JonChesterfield]

Nice. Particularly like getting rid of these because they use field access syntax to make a function call, which clang does with a Microsoft extension.

openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.cu

asm volatile("bar.sync %0, %1;"
                : "memory");
 }
 
 DEVICE void __kmpc_impl_threadfence() { __threadfence(); }
 DEVICE void __kmpc_impl_threadfence_block() { __threadfence_block(); }
 DEVICE void __kmpc_impl_threadfence_system() { __threadfence_system(); }
 
 // Calls to the NVPTX layer (assuming 1D layout)
-DEVICE int GetThreadIdInBlock() { return threadIdx.x; }
-DEVICE int GetBlockIdInKernel() { return blockIdx.x; }
-DEVICE int GetNumberOfBlocksInKernel() { return gridDim.x; }
-DEVICE int GetNumberOfThreadsInBlock() { return blockDim.x; }
+DEVICE int GetThreadIdInBlock() { return __nvvm_read_ptx_sreg_tid_x(); }
+DEVICE int GetBlockIdInKernel() { return __nvvm_read_ptx_sreg_ctaid_x(); }
+DEVICE int GetNumberOfBlocksInKernel() { return __nvvm_read_ptx_sreg_nctaid_x(); }
+DEVICE int GetNumberOfThreadsInBlock() { return __nvvm_read_ptx_sreg_ntid_x(); }
 DEVICE unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; }
 DEVICE unsigned GetLaneId() { return GetThreadIdInBlock() & (WARPSIZE - 1); }
 
 #define __OMP_SPIN 1000
 #define UNSET 0u
 #define SET 1u
 
 DEVICE void __kmpc_impl_init_lock(omp_lock_t *lock) {