Set calling convention for CUDA kernel

[yaxunl]

Summary

This patch sets target specific calling convention for CUDA kernels in IR.

Patch by Greg Rodgers.
Modified and lit test added by Yaxun Liu.

[rjmccall]

Is there a reason for this to be done as a special case in IRGen instead of just implicitly applying the calling convention in Sema?

[yaxunl]

Upload diff with full context.

[yaxunl]

In D44747#1044893, @rjmccall wrote:

Is there a reason for this to be done as a special case in IRGen instead of just implicitly applying the calling convention in Sema?

The calling convention is not used in Sema, therefore it seems simpler to do it in codegen. I could try doing this in Sema too. Is there any advantage of doing this in Sema?

Thanks.

[rjmccall]

In D44747#1044916, @yaxunl wrote:

In D44747#1044893, @rjmccall wrote:

Is there a reason for this to be done as a special case in IRGen instead of just implicitly applying the calling convention in Sema?

The calling convention is not used in Sema, therefore it seems simpler to do it in codegen. I could try doing this in Sema too. Is there any advantage of doing this in Sema?

In IRGen, it's a special case for your specific language mode on your specific target. In Sema, it can be done as part of the special checking for kernel functions.

Also, it looks like CUDA allows you to take the address of a __global__ function, and indirect calls to such functions presumably still follow the normal CUDA restrictions, so there must be *some* reflection of this in Sema.

[yaxunl]

Revised by John's comments. Introduce CC_CUDAKernel calling convention in AST, which is translated to target calling convention in IR.

[rjmccall]

lib/AST/Type.cpp

2762

For consistency with the rest of this switch, please put the return on the same line as its case.

lib/AST/TypePrinter.cpp

781

I think the spelling for this is __global__. You might need to adjust printing because this isn't the right place to print it, of course.

lib/CodeGen/CGCall.cpp

68

For consistency with the rest of this switch, please put the return on the same line as its case.

lib/Sema/SemaOverload.cpp

1492

It's cheaper not to check the CUDA language mode here; pulling the CC out of the FPT is easy.

Why is this necessary, anyway? From the spec, it doesn't look to me like kernel function pointers can be converted to ordinary function pointers. A kernel function pointer is supposed to be declared with something like __global__ void (*fn)(void). You'll need to change your patch to SemaType to apply the CC even when compiling for the host, of course.

I was going to say that you should use this CC in your validation that calls with execution configurations go to kernel functions, but... I can't actually find where you do that validation.

Do you need these function pointers to be a different size from the host function pointer?

tools/libclang/CXType.cpp

630

Formatting.

[yaxunl]

lib/Sema/SemaOverload.cpp

1492

In CUDA, __global__ can only be used with function declaration or definition. Using it in function pointer declaration will result in a warning: 'global' attribute only applies to functions.

Also, there is this lit test in SemaCUDA:

__global__ void kernel() {}

typedef void (*fn_ptr_t)();

__host__ fn_ptr_t get_ptr_h() {
  return kernel;
}

It allows implicit conversion of __global__ void() to void(*)(), therefore I need the above change to drop the CUDA kernel calling convention in such implicit conversion.

[rjmccall]

lib/Sema/SemaOverload.cpp

1492

I see. I must have mis-read the specification, but I see that the code samples I can find online agree with that test case. So __global__ function pointers are just treated as function pointers, and it's simply undefined behavior if you try to call a function pointer that happens to be a kernel without an execution configuration, or contrariwise if you use an execution configuration to call a function pointer that isn't a kernel.

In that case, I think the best solution is to just immediately strip __global__ from the type of a DRE to a kernel function, since __global__ isn't supposed to be part of the user-facing type system.

[yaxunl]

Revised by John's comments. Drop CUDA kernel calling convention in DRE.

[rjmccall]

LGTM.

If __global__ is supported in C++ structures, you might also need to make sure that member function constants (&A::kernel_function) drop the CC. And it might be a good idea to make sure that decltype(kernel_function) doesn't have a problem with it, either, since that does do some special-case work.

lib/Sema/SemaExpr.cpp

1669

You should use getAs<FunctionType> here (no const necessary). It's possible to declare a function with a typedef of function type, you just can't define it that way.

[yaxunl]

Matt, are you OK with the change from amdgcn backend point of view? Thanks.

lib/Sema/SemaExpr.cpp

1669

will do.

[yaxunl]

Use getAs instead of dyn_cast as John suggested.

[rjmccall]

LGTM.

[tra]

lib/Sema/SemaType.cpp

3319–3330

This apparently breaks compilation of some CUDA code in our internal tests. I'm working on minimizing a reproduction case. Should this code be enabled for AMD GPUs only?

[tra]

lib/Sema/SemaType.cpp

3319–3330

Here's a small snippet of code that previously used to compile and work:

template <typename T>
__global__ void EmptyKernel(void) { }

struct Dummy {
  /// Type definition of the EmptyKernel kernel entry point
  typedef void (*EmptyKernelPtr)();
  EmptyKernelPtr Empty() { return EmptyKernel<void>; }
};

AFAICT, it's currently impossible to apply global to pointers, so there's no way to make the code above work with this patch applied.

[yaxunl]

I will try fixing that.

The CUDA kernel calling convention should be dropped in all DRE's since it is invisible to the user.

Sam

[tra]

In D44747#1055877, @yaxunl wrote:

I will try fixing that.

The CUDA kernel calling convention should be dropped in all DRE's since it is invisible to the user.

Sam

Apparently it's not always the case.
Here's a bit more elaborate example demonstrating inconsistency in this behavior. Calling convention is ignored for regular functions, but not for function templates.

__global__ void EmptyKernel(void) { }

template <typename T>
__global__ void EmptyKernelT(void) { }

struct Dummy {
  /// Type definition of the EmptyKernel kernel entry point
  typedef void (*EmptyKernelPtr)();
  EmptyKernelPtr Empty() { return EmptyKernel; } // this one works
  EmptyKernelPtr EmptyT() { return EmptyKernelT<void>; } // this one errors out.
};

Do you think this is something you can fix quickly or do you want to unroll the change while you figure out what's going on?

[yaxunl]

Let's revert it for now. I will create a review after fixing it and commit it again with the fix.

Thanks.

Sam

[tra]

In D44747#1055894, @yaxunl wrote:

Let's revert it for now. I will create a review after fixing it and commit it again with the fix.

Thanks.

Sam

reverted in r329099.

include/clang/Basic/Specifiers.h

#include "clang/Basic/OpenCLImageTypes.def"
   enum InClassInitStyle {
     ICIS_NoInit,   ///< No in-class initializer.
     ICIS_CopyInit, ///< Copy initialization.
     ICIS_ListInit  ///< Direct list-initialization.
   };
 
   /// \brief CallingConv - Specifies the calling convention that a function uses.
   enum CallingConv {
-    CC_C,           // __attribute__((cdecl))
-    CC_X86StdCall,  // __attribute__((stdcall))
-    CC_X86FastCall, // __attribute__((fastcall))
-    CC_X86ThisCall, // __attribute__((thiscall))
+    CC_C,             // __attribute__((cdecl))
+    CC_X86StdCall,    // __attribute__((stdcall))
+    CC_X86FastCall,   // __attribute__((fastcall))
+    CC_X86ThisCall,   // __attribute__((thiscall))
     CC_X86VectorCall, // __attribute__((vectorcall))
-    CC_X86Pascal,   // __attribute__((pascal))
-    CC_Win64,       // __attribute__((ms_abi))
-    CC_X86_64SysV,  // __attribute__((sysv_abi))
-    CC_X86RegCall, // __attribute__((regcall))
-    CC_AAPCS,       // __attribute__((pcs("aapcs")))
-    CC_AAPCS_VFP,   // __attribute__((pcs("aapcs-vfp")))
-    CC_IntelOclBicc, // __attribute__((intel_ocl_bicc))
-    CC_SpirFunction, // default for OpenCL functions on SPIR target
-    CC_OpenCLKernel, // inferred for OpenCL kernels
-    CC_Swift,        // __attribute__((swiftcall))
-    CC_PreserveMost, // __attribute__((preserve_most))
-    CC_PreserveAll,  // __attribute__((preserve_all))
+    CC_X86Pascal,     // __attribute__((pascal))
+    CC_Win64,         // __attribute__((ms_abi))
+    CC_X86_64SysV,    // __attribute__((sysv_abi))
+    CC_X86RegCall,    // __attribute__((regcall))
+    CC_AAPCS,         // __attribute__((pcs("aapcs")))
+    CC_AAPCS_VFP,     // __attribute__((pcs("aapcs-vfp")))
+    CC_IntelOclBicc,  // __attribute__((intel_ocl_bicc))
+    CC_SpirFunction,  // default for OpenCL functions on SPIR target
+    CC_OpenCLKernel,  // inferred for OpenCL kernels
+    CC_Swift,         // __attribute__((swiftcall))
+    CC_PreserveMost,  // __attribute__((preserve_most))
+    CC_PreserveAll,   // __attribute__((preserve_all))
+    CC_CUDAKernel,    // inferred for CUDA kernels
   };
 
   /// \brief Checks whether the given calling convention supports variadic
   /// calls. Unprototyped calls also use the variadic call rules.
   inline bool supportsVariadicCall(CallingConv CC) {
     switch (CC) {
     case CC_X86StdCall:
     case CC_X86FastCall:

lib/AST/ItaniumMangle.cpp

   case CC_X86RegCall:
   case CC_AAPCS:
   case CC_AAPCS_VFP:
   case CC_IntelOclBicc:
   case CC_SpirFunction:
   case CC_OpenCLKernel:
   case CC_PreserveMost:
   case CC_PreserveAll:
+  case CC_CUDAKernel:
     // FIXME: we should be mangling all of the above.
     return "";
 
   case CC_Swift:
     return "swiftcall";
   }
   llvm_unreachable("bad calling convention");
 }

lib/AST/Type.cpp

   case CC_AAPCS: return "aapcs";
   case CC_AAPCS_VFP: return "aapcs-vfp";
   case CC_IntelOclBicc: return "intel_ocl_bicc";
   case CC_SpirFunction: return "spir_function";
   case CC_OpenCLKernel: return "opencl_kernel";
   case CC_Swift: return "swiftcall";
   case CC_PreserveMost: return "preserve_most";
   case CC_PreserveAll: return "preserve_all";
+  case CC_CUDAKernel: return "cuda_kernel";
   }
 
   llvm_unreachable("Invalid calling convention.");
 }
 
 FunctionProtoType::FunctionProtoType(QualType result, ArrayRef<QualType> params,
                                      QualType canonical,
                                      const ExtProtoInfo &epi)
     : FunctionType(FunctionProto, result, canonical,

lib/AST/TypePrinter.cpp

if (!InsideCCAttribute) {
     case CC_X86_64SysV:
       OS << " __attribute__((sysv_abi))";
       break;
     case CC_X86RegCall:
       OS << " __attribute__((regcall))";
       break;
     case CC_SpirFunction:
     case CC_OpenCLKernel:
       // Do nothing. These CCs are not available as attributes.
       break;
+    case CC_CUDAKernel:
+      // ToDo: print this before the function.
+      OS << " __global__";
+      break;
     case CC_Swift:
       OS << " __attribute__((swiftcall))";
       break;
     case CC_PreserveMost:
       OS << " __attribute__((preserve_most))";
       break;
     case CC_PreserveAll:
       OS << " __attribute__((preserve_all))";

lib/CodeGen/CGCall.cpp

unsigned CodeGenTypes::ClangCallConvToLLVMCallConv(CallingConv CC) {
   case CC_X86Pascal: return llvm::CallingConv::C;
   // TODO: Add support for __vectorcall to LLVM.
   case CC_X86VectorCall: return llvm::CallingConv::X86_VectorCall;
   case CC_SpirFunction: return llvm::CallingConv::SPIR_FUNC;
   case CC_OpenCLKernel: return CGM.getTargetCodeGenInfo().getOpenCLKernelCallingConv();
   case CC_PreserveMost: return llvm::CallingConv::PreserveMost;
   case CC_PreserveAll: return llvm::CallingConv::PreserveAll;
   case CC_Swift: return llvm::CallingConv::Swift;
+  case CC_CUDAKernel: return CGM.getTargetCodeGenInfo().getCUDAKernelCallingConv();
   }
 }
 
 /// Derives the 'this' type for codegen purposes, i.e. ignoring method
 /// qualification.
 /// FIXME: address space qualification?
 static CanQualType GetThisType(ASTContext &Context, const CXXRecordDecl *RD) {
   QualType RecTy = Context.getTagDeclType(RD)->getCanonicalTypeInternal();
   return Context.getPointerType(CanQualType::CreateUnsafe(RecTy));

lib/CodeGen/CGDebugInfo.cpp

   case CC_Swift:
     return llvm::dwarf::DW_CC_LLVM_Swift;
   case CC_PreserveMost:
     return llvm::dwarf::DW_CC_LLVM_PreserveMost;
   case CC_PreserveAll:
     return llvm::dwarf::DW_CC_LLVM_PreserveAll;
   case CC_X86RegCall:
     return llvm::dwarf::DW_CC_LLVM_X86RegCall;
+  case CC_CUDAKernel:
+    // ToDo: Add llvm::dwarf::DW_CC_LLVM_CUDAKernel;
+    return 0;
   }
   return 0;
 }
 
 llvm::DIType *CGDebugInfo::CreateType(const FunctionType *Ty,
                                       llvm::DIFile *Unit) {
   SmallVector<llvm::Metadata *, 16> EltTys;
 

lib/CodeGen/TargetInfo.h

public:
   /// this platform.
   virtual void getDetectMismatchOption(llvm::StringRef Name,
                                        llvm::StringRef Value,
                                        llvm::SmallString<32> &Opt) const {}
 
   /// Get LLVM calling convention for OpenCL kernel.
   virtual unsigned getOpenCLKernelCallingConv() const;
 
+  /// Get LLVM calling convention for CUDA kernel.
+  virtual unsigned getCUDAKernelCallingConv() const;
+
   /// Get target specific null pointer.
   /// \param T is the LLVM type of the null pointer.
   /// \param QT is the clang QualType of the null pointer.
   /// \return ConstantPointerNull with the given type \p T.
   /// Each target can override it to return its own desired constant value.
   virtual llvm::Constant *getNullPointer(const CodeGen::CodeGenModule &CGM,
       llvm::PointerType *T, QualType QT) const;
 

lib/CodeGen/TargetInfo.cpp

   // list to enable feasible implementation of clSetKernelArg() with
   // aggregates etc. In case we would use the default C calling conv here,
   // clSetKernelArg() might break depending on the target-specific
   // conventions; different targets might split structs passed as values
   // to multiple function arguments etc.
   return llvm::CallingConv::SPIR_KERNEL;
 }
 
+unsigned TargetCodeGenInfo::getCUDAKernelCallingConv() const {
+  return llvm::CallingConv::C;
+}
+
 llvm::Constant *TargetCodeGenInfo::getNullPointer(const CodeGen::CodeGenModule &CGM,
     llvm::PointerType *T, QualType QT) const {
   return llvm::ConstantPointerNull::get(T);
 }
 
 LangAS TargetCodeGenInfo::getGlobalVarAddressSpace(CodeGenModule &CGM,
                                                    const VarDecl *D) const {
   assert(!CGM.getLangOpts().OpenCL &&

 
 class AMDGPUTargetCodeGenInfo : public TargetCodeGenInfo {
 public:
   AMDGPUTargetCodeGenInfo(CodeGenTypes &CGT)
     : TargetCodeGenInfo(new AMDGPUABIInfo(CGT)) {}
   void setTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
                            CodeGen::CodeGenModule &M) const override;
   unsigned getOpenCLKernelCallingConv() const override;
+  unsigned getCUDAKernelCallingConv() const override;
 
   llvm::Constant *getNullPointer(const CodeGen::CodeGenModule &CGM,
       llvm::PointerType *T, QualType QT) const override;
 
   LangAS getASTAllocaAddressSpace() const override {
     return getLangASFromTargetAS(
         getABIInfo().getDataLayout().getAllocaAddrSpace());
   }

       F->addFnAttr("amdgpu-num-vgpr", llvm::utostr(NumVGPR));
   }
 }
 
 unsigned AMDGPUTargetCodeGenInfo::getOpenCLKernelCallingConv() const {
   return llvm::CallingConv::AMDGPU_KERNEL;
 }
 
+unsigned AMDGPUTargetCodeGenInfo::getCUDAKernelCallingConv() const {
+  return llvm::CallingConv::AMDGPU_KERNEL;
+}
+
 // Currently LLVM assumes null pointers always have value 0,
 // which results in incorrectly transformed IR. Therefore, instead of
 // emitting null pointers in private and local address spaces, a null
 // pointer in generic address space is emitted which is casted to a
 // pointer in local or private address space.
 llvm::Constant *AMDGPUTargetCodeGenInfo::getNullPointer(
     const CodeGen::CodeGenModule &CGM, llvm::PointerType *PT,
     QualType QT) const {

lib/Sema/SemaExpr.cpp

 #include "clang/AST/DeclObjC.h"
 #include "clang/AST/DeclTemplate.h"
 #include "clang/AST/EvaluatedExprVisitor.h"
 #include "clang/AST/Expr.h"
 #include "clang/AST/ExprCXX.h"
 #include "clang/AST/ExprObjC.h"
 #include "clang/AST/ExprOpenMP.h"
 #include "clang/AST/RecursiveASTVisitor.h"
+#include "clang/AST/Type.h"
 #include "clang/AST/TypeLoc.h"
 #include "clang/Basic/PartialDiagnostic.h"
 #include "clang/Basic/SourceManager.h"
 #include "clang/Basic/TargetInfo.h"
 #include "clang/Lex/LiteralSupport.h"
 #include "clang/Lex/Preprocessor.h"
 #include "clang/Sema/AnalysisBasedWarnings.h"
 #include "clang/Sema/DeclSpec.h"

 Sema::BuildDeclRefExpr(ValueDecl *D, QualType Ty, ExprValueKind VK,
                        const DeclarationNameInfo &NameInfo,
                        const CXXScopeSpec *SS, NamedDecl *FoundD,
                        const TemplateArgumentListInfo *TemplateArgs) {
   bool RefersToCapturedVariable =
       isa<VarDecl>(D) &&
       NeedToCaptureVariable(cast<VarDecl>(D), NameInfo.getLoc());
 
+  // Drop CUDA kernel calling convention since it is invisible to the user
+  // in DRE.
+  if (const auto *FT = Ty->getAs<FunctionType>()) {
+    if (FT->getCallConv() == CC_CUDAKernel) {
+      FT = Context.adjustFunctionType(FT,
+                                      FT->getExtInfo().withCallingConv(CC_C));
+      Ty = QualType(FT, Ty.getQualifiers().getAsOpaqueValue());
+    }
+  }
+
   DeclRefExpr *E;
   if (isa<VarTemplateSpecializationDecl>(D)) {
     VarTemplateSpecializationDecl *VarSpec =
         cast<VarTemplateSpecializationDecl>(D);
 
     E = DeclRefExpr::Create(Context, SS ? SS->getWithLocInContext(Context)
                                         : NestedNameSpecifierLoc(),
                             VarSpec->getTemplateKeywordLoc(), D,

lib/Sema/SemaOverload.cpp

     const auto *ToFPT = cast<FunctionProtoType>(ToFn);
     if (FromFPT->isNothrow(Context) && !ToFPT->isNothrow(Context)) {
       FromFn = cast<FunctionType>(
           Context.getFunctionTypeWithExceptionSpec(QualType(FromFPT, 0),
                                                    EST_None)
                  .getTypePtr());
       Changed = true;
     }
-
     // Convert FromFPT's ExtParameterInfo if necessary. The conversion is valid
     // only if the ExtParameterInfo lists of the two function prototypes can be
     // merged and the merged list is identical to ToFPT's ExtParameterInfo list.
     SmallVector<FunctionProtoType::ExtParameterInfo, 4> NewParamInfos;
     bool CanUseToFPT, CanUseFromFPT;
     if (Context.mergeExtParameterInfo(ToFPT, FromFPT, CanUseToFPT,
                                       CanUseFromFPT, NewParamInfos) &&
         CanUseToFPT && !CanUseFromFPT) {
       FunctionProtoType::ExtProtoInfo ExtInfo = FromFPT->getExtProtoInfo();
       ExtInfo.ExtParameterInfos =
           NewParamInfos.empty() ? nullptr : NewParamInfos.data();
       QualType QT = Context.getFunctionType(FromFPT->getReturnType(),
                                             FromFPT->getParamTypes(), ExtInfo);
       FromFn = QT->getAs<FunctionType>();
       Changed = true;
     }
   }

lib/Sema/SemaType.cpp

           D.getDeclSpec().getStorageClassSpec() != DeclSpec::SCS_typedef &&
           !D.isStaticMember();
     }
   }
 
   CallingConv CC = S.Context.getDefaultCallingConvention(FTI.isVariadic,
                                                          IsCXXInstanceMethod);
 
+  // Attribute AT_CUDAGlobal affects the calling convention for AMDGPU targets.
+  // This is the simplest place to infer calling convention for CUDA kernels.
+  if (S.getLangOpts().CUDA && S.getLangOpts().CUDAIsDevice) {
+    for (const AttributeList *Attr = D.getDeclSpec().getAttributes().getList();
+         Attr; Attr = Attr->getNext()) {
+      if (Attr->getKind() == AttributeList::AT_CUDAGlobal) {
+        CC = CC_CUDAKernel;
+        break;
+      }
+    }
+  }
+
   // Attribute AT_OpenCLKernel affects the calling convention for SPIR
   // and AMDGPU targets, hence it cannot be treated as a calling
   // convention attribute. This is the simplest place to infer
   // calling convention for OpenCL kernels.
   if (S.getLangOpts().OpenCL) {
     for (const AttributeList *Attr = D.getDeclSpec().getAttributes().getList();
          Attr; Attr = Attr->getNext()) {
       if (Attr->getKind() == AttributeList::AT_OpenCLKernel) {

test/CodeGenCUDA/kernel-amdgcn.cu

+// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm %s -o - | FileCheck %s
+#include "Inputs/cuda.h"
+
+// CHECK: define amdgpu_kernel void @_ZN1A6kernelEv
+class A {
+public:
+  static __global__ void kernel(){}
+};
+
+// CHECK: define void @_Z10non_kernelv
+__device__ void non_kernel(){}
+
+// CHECK: define amdgpu_kernel void @_Z6kerneli
+__global__ void kernel(int x) {
+  non_kernel();
+}
+
+// CHECK: define amdgpu_kernel void @_Z15template_kernelI1AEvT_
+template<class T>
+__global__ void template_kernel(T x) {}
+
+void launch(void *f);
+
+int main() {
+  launch((void*)A::kernel);
+  launch((void*)kernel);
+  launch((void*)template_kernel<A>);
+  return 0;
+}

tools/libclang/CXType.cpp

switch (FD->getCallConv()) {
       TCALLINGCONV(AAPCS);
       TCALLINGCONV(AAPCS_VFP);
       TCALLINGCONV(IntelOclBicc);
       TCALLINGCONV(Swift);
       TCALLINGCONV(PreserveMost);
       TCALLINGCONV(PreserveAll);
     case CC_SpirFunction: return CXCallingConv_Unexposed;
     case CC_OpenCLKernel: return CXCallingConv_Unexposed;
+    case CC_CUDAKernel: return CXCallingConv_Unexposed;
       break;
     }
 #undef TCALLINGCONV
   }
   
   return CXCallingConv_Invalid;
 }
 
 int clang_getNumArgTypes(CXType X) {