[OpenMP][AMDGPU] Use AMDGPU_KERNEL calling convention for entry function

[pdhaliwal]

Summary

AMDGPU backend requires entry functions/kernels to have AMDGPU_KERNEL
calling convention for proper linking.

[JonChesterfield]

LGTM. This is the same hook hip & opencl use to indicate that a function is a kernel entry point. We may need to create the metadata openmp nvptx uses as well, but that's separate.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
   // the offloading entry and launched by the runtime library. We also mark the
   // outlined function to have external linkage in case we are emitting code for
   // the device, because these functions will be entry points to the device.
 
   if (CGM.getLangOpts().OpenMPIsDevice) {
     OutlinedFnID = llvm::ConstantExpr::getBitCast(OutlinedFn, CGM.Int8PtrTy);
     OutlinedFn->setLinkage(llvm::GlobalValue::WeakAnyLinkage);
     OutlinedFn->setDSOLocal(false);
+    if (CGM.getTriple().isAMDGCN())
+      OutlinedFn->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL);
   } else {
     std::string Name = getName({EntryFnName, "region_id"});
     OutlinedFnID = new llvm::GlobalVariable(
         CGM.getModule(), CGM.Int8Ty, /*isConstant=*/true,
         llvm::GlobalValue::WeakAnyLinkage,
         llvm::Constant::getNullValue(CGM.Int8Ty), Name);
   }
 

clang/test/OpenMP/amdgcn_target_codegen.cpp

 // REQUIRES: amdgpu-registered-target
 
 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc
 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
 #define N 1000
 
 int test_amdgcn_target_tid_threads() {
-// CHECK-LABEL: define weak void @{{.*}}test_amdgcn_target_tid_threads
+// CHECK-LABEL: define weak amdgpu_kernel void @{{.*}}test_amdgcn_target_tid_threads
 
   int arr[N];
 
 // CHECK: [[NUM_THREADS:%.+]] = call i32 @__kmpc_amdgcn_gpu_num_threads()
 // CHECK: sub nuw i32 [[NUM_THREADS]], 64
 // CHECK: call i32 @llvm.amdgcn.workitem.id.x()
 #pragma omp target
   for (int i = 0; i < N; i++) {
     arr[i] = 1;
   }
 
   return arr[0];
 }
 
 int test_amdgcn_target_tid_threads_simd() {
-// CHECK-LABEL: define weak void @{{.*}}test_amdgcn_target_tid_threads_simd
+// CHECK-LABEL: define weak amdgpu_kernel void @{{.*}}test_amdgcn_target_tid_threads_simd
 
   int arr[N];
 
 // CHECK: [[NUM_THREADS:%.+]] = call i32 @__kmpc_amdgcn_gpu_num_threads()
 // CHECK: call void @__kmpc_spmd_kernel_init(i32 [[NUM_THREADS]], i16 0)
 #pragma omp target simd
   for (int i = 0; i < N; i++) {
     arr[i] = 1;
   }
   return arr[0];
 }
 
 #endif