fix merge with upstream kernel attribute refactor (#18868)

Fix merge with my upstream change to refactor the device kernel
attribute
[here](llvm/llvm-project@3b9ebe9).

---------

Signed-off-by: Sarnie, Nick <[email protected]>

Author: sarnex

clang/include/clang/Basic/Attr.td

@@ -1650,7 +1650,7 @@ def DeviceKernel : DeclOrTypeAttr {
     // list, but here we have the same spelling with unscores and without,
     // so handle that case manually.
       return A.getAttributeSpellingListIndex() == Keyword_kernel ||
-      A.getAttrName()->getName() == "kernel";
+      (A.getAttrName() && A.getAttrName()->getName() == "kernel");
     }
     static inline bool isOpenCLSpelling(const AttributeCommonInfo* A) {
     if (!A) return false;

clang/lib/Basic/Targets/X86.h

@@ -796,7 +796,7 @@ class LLVM_LIBRARY_VISIBILITY X86_64TargetInfo : public X86TargetInfo {
     case CC_X86RegCall:
       return CCCR_OK;
     case CC_DeviceKernel:
-      return IsOpenCL ? CCCR_OK : CCCR_Warning;
+      return IsOpenCL || IsSYCLDevice ? CCCR_OK : CCCR_Warning;
     default:
       return CCCR_Warning;
     }
@@ -842,10 +842,12 @@ class LLVM_LIBRARY_VISIBILITY X86_64TargetInfo : public X86TargetInfo {
   void adjust(DiagnosticsEngine &Diags, LangOptions &Opts) override {
     TargetInfo::adjust(Diags, Opts);
     IsOpenCL = Opts.OpenCL;
+    IsSYCLDevice = Opts.SYCLIsDevice;
   }
 
 private:
   bool IsOpenCL = false;
+  bool IsSYCLDevice = false;
 };
 
 // x86-64 UEFI target

clang/lib/CodeGen/CGCall.cpp

@@ -94,6 +94,10 @@ unsigned CodeGenTypes::ClangCallConvToLLVMCallConv(CallingConv CC) {
       return llvm::CallingConv::AMDGPU_KERNEL;
     if (CGM.getTriple().isNVPTX())
       return llvm::CallingConv::PTX_Kernel;
+    if (CGM.getLangOpts().SYCLIsNativeCPU)
+      return CGM.getTarget().getDefaultCallingConv();
+    if (CGM.getLangOpts().SYCLIsDevice)
+      return CGM.getTarget().getDefaultCallingConv();
     llvm_unreachable("Unknown kernel calling convention");
   }
   case CC_PreserveMost:
@@ -2628,8 +2632,8 @@ void CodeGenModule::ConstructAttributeList(StringRef Name,
                                  NumElemsParam);
     }
 
-    if (DeviceKernelAttr::isOpenCLSpelling(
-            TargetDecl->getAttr<DeviceKernelAttr>()) &&
+    if (TargetDecl->hasAttr<DeviceKernelAttr>() &&
+        (getLangOpts().OpenCL || getLangOpts().SYCLIsDevice) &&
         CallingConv != CallingConv::CC_C &&
         CallingConv != CallingConv::CC_SpirFunction) {
       // Check CallingConv to avoid adding uniform-work-group-size attribute to
@@ -3013,8 +3017,9 @@ void CodeGenModule::ConstructAttributeList(StringRef Name,
     // > data type, the OpenCL compiler can assume that the pointee is always
     // > appropriately aligned as required by the data type.
     if (TargetDecl &&
-        DeviceKernelAttr::isOpenCLSpelling(
-            TargetDecl->getAttr<DeviceKernelAttr>()) &&
+
+        (TargetDecl->hasAttr<DeviceKernelAttr>() &&
+         (getLangOpts().OpenCL || getLangOpts().SYCLIsDevice)) &&
         ParamType->isPointerType()) {
       QualType PTy = ParamType->getPointeeType();
       if (!PTy->isIncompleteType() && PTy->isConstantSizeType()) {

clang/lib/Sema/SemaChecking.cpp

@@ -4014,7 +4014,7 @@ void Sema::checkCall(NamedDecl *FDecl, const FunctionProtoType *Proto,
   if (FD)
     diagnoseArgDependentDiagnoseIfAttrs(FD, ThisArg, Args, Loc);
 
-  if (FD && FD->hasAttr<DeviceKernelAttr>())
+  if (FD && FD->hasAttr<DeviceKernelAttr>() && getLangOpts().isSYCL())
     SYCL().CheckSYCLKernelCall(FD, Args);
 
   // Diagnose variadic calls in SYCL.

clang/lib/Sema/SemaDecl.cpp

@@ -20762,7 +20762,8 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(const FunctionDecl *FD,
     return FunctionEmissionStatus::TemplateDiscarded;
 
   if (LangOpts.SYCLIsDevice &&
-      (FD->hasAttr<SYCLDeviceAttr>() || FD->hasAttr<DeviceKernelAttr>()))
+      (FD->hasAttr<SYCLDeviceAttr>() || FD->hasAttr<DeviceKernelAttr>()) &&
+      !FD->hasAttr<ArtificialAttr>())
     return FunctionEmissionStatus::Emitted;
 
   // Check whether this function is an externally visible definition.

clang/lib/Sema/SemaDeclAttr.cpp

@@ -5355,7 +5355,7 @@ static void handleCallConvAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
 static void handleDeviceKernelAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   const auto *FD = dyn_cast_or_null<FunctionDecl>(D);
   bool IsFunctionTemplate = FD && FD->getDescribedFunctionTemplate();
-  if (S.getLangOpts().SYCLIsDevice) {
+  if (S.getLangOpts().isSYCL()) {
     if (!IsFunctionTemplate) {
       S.Diag(AL.getLoc(), diag::warn_attribute_wrong_decl_type_str)
           << AL << AL.isRegularKeywordAttribute() << "function templates";

clang/lib/Sema/SemaType.cpp

@@ -3785,12 +3785,11 @@ static CallingConv getCCForDeclaratorChunk(
       }
     }
   }
-  if (!S.getLangOpts().isSYCL()) {
-    for (const ParsedAttr &AL : D.getDeclSpec().getAttributes()) {
-      if (AL.getKind() == ParsedAttr::AT_DeviceKernel) {
-        CC = CC_DeviceKernel;
-        break;
-      }
+  for (const ParsedAttr &AL : D.getDeclSpec().getAttributes()) {
+    if (AL.getKind() == ParsedAttr::AT_DeviceKernel &&
+	!DeviceKernelAttr::isSYCLSpelling(AL)) {
+      CC = CC_DeviceKernel;
+      break;
     }
   }
   return CC;

clang/test/CodeGenSYCL/kernel-op-calls.cpp

@@ -53,7 +53,7 @@ int main() {
 
   Q.submit([&](sycl::handler& cgh) {
       ESIMDFunctor EF;
-      // CHECK: define dso_local spir_kernel void @_ZTS12ESIMDFunctor() {{.*}} !intel_reqd_sub_group_size !{{[0-9]+}} !work_group_size_hint !{{[0-9]+}} !kernel_arg_accessor_ptr !{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
+      // CHECK: define dso_local spir_kernel void @_ZTS12ESIMDFunctor() {{.*}} !sycl_explicit_simd !{{[0-9]+}} !intel_reqd_sub_group_size !{{[0-9]+}} !work_group_size_hint !{{[0-9]+}} !kernel_arg_accessor_ptr !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
       cgh.parallel_for(sycl::range<1>(10), EF);
     });
 

clang/test/CodeGenSYCL/usm-int-header.cpp

@@ -30,4 +30,4 @@ int main() {
     });
 }
 
-// CHECK: FunctionDecl {{.*}}usm_test{{.*}} 'void (__global int *, __global float *)'
+// CHECK: FunctionDecl {{.*}}usm_test{{.*}} 'void (__global int *, __global float *)  __attribute__((device_kernel))'

clang/test/SemaOpenCL/sampled_image_overload.cl

@@ -3,7 +3,7 @@
 void __attribute__((overloadable)) foo(__ocl_sampled_image1d_ro_t);
 void __attribute__((overloadable)) foo(__ocl_sampled_image2d_ro_t);
 
-// CHECK: FunctionDecl {{.*}} <{{.*}}> line:{{.*}} ker 'void (__private __ocl_sampled_image1d_ro_t, __private __ocl_sampled_image2d_ro_t)'
+// CHECK: FunctionDecl {{.*}} <{{.*}}> line:{{.*}} ker 'void (__private __ocl_sampled_image1d_ro_t, __private __ocl_sampled_image2d_ro_t) __attribute__((device_kernel))'
 void kernel ker(__ocl_sampled_image1d_ro_t src1, __ocl_sampled_image2d_ro_t src2) {
   // CHECK: CallExpr
   // CHECK-NEXT: ImplicitCastExpr {{.*}} <{{.*}}> 'void (*)(__private __ocl_sampled_image1d_ro_t)'