[OpenMP][NVPTX] Added forward declaration for atomic operations

[tianshilei1992]

Summary

Pretty similar to D95058, this patch added forward declaration for
CUDA atomic functions. We already have definitions with right mangled names in
internal CUDA headers so the forward declaration here can work properly.

[JonChesterfield]

Possibly obsolete?

Sorry, misread. I have a nearly-ready patch that turns these into non-template functions and moves them into the .cu

Fine by me too

[jdoerfert]

LGTM

[JonChesterfield]

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

136

I think these all need to be marked DEVICE

[tianshilei1992]

Mark those forward declarations as DEVICE

[JonChesterfield]

Thanks. If you land this I'll rebase D95093 manually

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

 //===------------ target_impl.h - NVPTX OpenMP GPU options ------- CUDA -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Definitions of target specific functions
 //
 //===----------------------------------------------------------------------===//
 #ifndef _TARGET_IMPL_H_
 #define _TARGET_IMPL_H_
 
 #include <assert.h>
 #include <cuda.h>
 #include <inttypes.h>
 #include <stdio.h>
 #include <stdlib.h>
 
 #include "nvptx_interface.h"
 
 #define DEVICE __device__
 #define INLINE __forceinline__ DEVICE

 // Calls to the NVPTX layer (assuming 1D layout)
 DEVICE int GetThreadIdInBlock();
 DEVICE int GetBlockIdInKernel();
 DEVICE int GetNumberOfBlocksInKernel();
 DEVICE int GetNumberOfThreadsInBlock();
 DEVICE unsigned GetWarpId();
 DEVICE unsigned GetLaneId();
 
+// Forward declaration of atomics. Although they're template functions, we
+// already have definitions for different types in CUDA internal headers with
+// the right mangled names.
+template <typename T> DEVICE T atomicAdd(T *address, T val);
+template <typename T> DEVICE T atomicInc(T *address, T val);
+template <typename T> DEVICE T atomicMax(T *address, T val);
+template <typename T> DEVICE T atomicExch(T *address, T val);
+template <typename T> DEVICE T atomicCAS(T *address, T compare, T val);
+
 // Atomics
 template <typename T> INLINE T __kmpc_atomic_add(T *address, T val) {
   return atomicAdd(address, val);
 }
 
 template <typename T> INLINE T __kmpc_atomic_inc(T *address, T val) {
   return atomicInc(address, val);
 }