AMDGPU: Set amdgpu_kernel calling convention for OpenCL kernels.

[nhaustov]

Summary

Summary:
Change Clang calling convention SpirKernel to OpenCLKernel.
Set calling convention OpenCLKernel for amdgcn as well.
Add virtual method .getOpenCLKernelCallingConv() to TargetCodeGenInfo
and use it to set target calling convention for AMDGPU and SPIR.
Update tests.

[• tstellarAMD]

Does this new patch fix the OpenCL regression?

[nhaustov]

In D21367#461480, @tstellarAMD wrote:

Does this new patch fix the OpenCL regression?

Yes, it fixes the problem with calling kernel from kernel. I'll add a test too.

[nhaustov]

Add test for calling OpenCL kernel from kernel.

[• tstellarAMD]

This seems OK to me.

[yaxunl]

I am wondering what's the difference between this calling convention and spir_kernel.

[nhaustov]

In D21367#467894, @yaxunl wrote:

I am wondering what's the difference between this calling convention and spir_kernel.

spir_kernel has only effect and was created only for SPIR target. In theory it could perhaps be reused for AMDGPU target, however it seems not very clear to me.

[yaxunl]

In D21367#467982, @nhaustov wrote:

In D21367#467894, @yaxunl wrote:

I am wondering what's the difference between this calling convention and spir_kernel.

spir_kernel has only effect and was created only for SPIR target. In theory it could perhaps be reused for AMDGPU target, however it seems not very clear to me.

If the new calling convention is to indicate a function is an OpenCL kernel, then it has exactly the same meaning as spir_kernel. Can we just use it in AMDGPU target?

[nhaustov]

In D21367#467894, @yaxunl wrote:

I am wondering what's the difference between this calling convention and spir_kernel.

spir_kernel has only effect and was created only for SPIR target. In theory it could perhaps be reused for AMDGPU target, however it seems not very clear to me.

In D21367#468057, @yaxunl wrote:

In D21367#467982, @nhaustov wrote:

In D21367#467894, @yaxunl wrote:

I am wondering what's the difference between this calling convention and spir_kernel.

spir_kernel has only effect and was created only for SPIR target. In theory it could perhaps be reused for AMDGPU target, however it seems not very clear to me.

If the new calling convention is to indicate a function is an OpenCL kernel, then it has exactly the same meaning as spir_kernel. Can we just use it in AMDGPU target?

This looks like a hack to me. Calling conventions in LLVM IR are really target specific. spir_kernel really doesn't make sense for anything other than SPIR.

Maybe we can use the fact that calling conventions are separate in AST and in LLVM IR. So in AST maybe we can have opencl_kernel and in LLVM it can map to spir_kernel for SPIR and amdgpu_kernel for AMDGPU.

[Anastasia]

lib/Basic/Targets.cpp

2122 ↗

(On Diff #61247)

One extra line?

[yaxunl]

Maybe we can use the fact that calling conventions are separate in AST and in LLVM IR. So in AST maybe we can have opencl_kernel and in LLVM it can map to spir_kernel for SPIR and amdgpu_kernel for AMDGPU.

This approach looks good to me. Thanks.

[nhaustov]

Update after discussion with Sam and review from Anastasia.

[yaxunl]

LGTM. Thanks!

cfe/trunk/include/clang/Basic/Specifiers.h

enum CallingConv {
     CC_X86VectorCall, // __attribute__((vectorcall))
     CC_X86Pascal,   // __attribute__((pascal))
     CC_X86_64Win64, // __attribute__((ms_abi))
     CC_X86_64SysV,  // __attribute__((sysv_abi))
     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_SpirKernel,   // inferred for OpenCL kernels on SPIR target
+    CC_OpenCLKernel, // inferred for OpenCL kernels
     CC_Swift,        // __attribute__((swiftcall))
     CC_PreserveMost, // __attribute__((preserve_most))
     CC_PreserveAll,  // __attribute__((preserve_all))
   };
 
   /// \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:
     case CC_X86ThisCall:
     case CC_X86Pascal:
     case CC_X86VectorCall:
     case CC_SpirFunction:
-    case CC_SpirKernel:
+    case CC_OpenCLKernel:
     case CC_Swift:
       return false;
     default:
       return true;
     }
   }
 
   /// \brief The storage duration for an object (per C++ [basic.stc]).

cfe/trunk/lib/AST/ItaniumMangle.cpp

StringRef CXXNameMangler::getCallingConvQualifierName(CallingConv CC) {
   case CC_X86VectorCall:
   case CC_X86Pascal:
   case CC_X86_64Win64:
   case CC_X86_64SysV:
   case CC_AAPCS:
   case CC_AAPCS_VFP:
   case CC_IntelOclBicc:
   case CC_SpirFunction:
-  case CC_SpirKernel:
+  case CC_OpenCLKernel:
   case CC_PreserveMost:
   case CC_PreserveAll:
     // FIXME: we should be mangling all of the above.
     return "";
 
   case CC_Swift:
     return "swiftcall";
   }

cfe/trunk/lib/AST/Type.cpp

StringRef FunctionType::getNameForCallConv(CallingConv CC) {
   case CC_X86Pascal: return "pascal";
   case CC_X86VectorCall: return "vectorcall";
   case CC_X86_64Win64: return "ms_abi";
   case CC_X86_64SysV: return "sysv_abi";
   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_SpirKernel: return "spir_kernel";
+  case CC_OpenCLKernel: return "opencl_kernel";
   case CC_Swift: return "swiftcall";
   case CC_PreserveMost: return "preserve_most";
   case CC_PreserveAll: return "preserve_all";
   }
 
   llvm_unreachable("Invalid calling convention.");
 }
 

cfe/trunk/lib/AST/TypePrinter.cpp

case CC_IntelOclBicc:
       break;
     case CC_X86_64Win64:
       OS << " __attribute__((ms_abi))";
       break;
     case CC_X86_64SysV:
       OS << " __attribute__((sysv_abi))";
       break;
     case CC_SpirFunction:
-    case CC_SpirKernel:
+    case CC_OpenCLKernel:
       // Do nothing. These CCs are not available as attributes.
       break;
     case CC_Swift:
       OS << " __attribute__((swiftcall))";
       break;
     case CC_PreserveMost:
       OS << " __attribute__((preserve_most))";
       break;

cfe/trunk/lib/Basic/Targets.cpp

void setSupportedOpenCLOpts() override {
     }
     if (GPU >= GK_SOUTHERN_ISLANDS) {
       Opts.cl_khr_fp16 = 1;
       Opts.cl_khr_int64_base_atomics = 1;
       Opts.cl_khr_int64_extended_atomics = 1;
       Opts.cl_khr_3d_image_writes = 1;
     }
   }
+
+  CallingConvCheckResult checkCallingConvention(CallingConv CC) const override {
+    switch (CC) {
+      default:
+        return CCCR_Warning;
+      case CC_C:
+      case CC_OpenCLKernel:
+        return CCCR_OK;
+    }
+  }
 };
 
 const Builtin::Info AMDGPUTargetInfo::BuiltinInfo[] = {
 #define BUILTIN(ID, TYPE, ATTRS)                \
   { #ID, TYPE, ATTRS, nullptr, ALL_LANGUAGES, nullptr },
 #define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE)                               \
   { #ID, TYPE, ATTRS, nullptr, ALL_LANGUAGES, FEATURE },
 #include "clang/Basic/BuiltinsAMDGPU.def"
public:
   ArrayRef<TargetInfo::GCCRegAlias> getGCCRegAliases() const override {
     return None;
   }
   BuiltinVaListKind getBuiltinVaListKind() const override {
     return TargetInfo::VoidPtrBuiltinVaList;
   }
 
   CallingConvCheckResult checkCallingConvention(CallingConv CC) const override {
-    return (CC == CC_SpirFunction || CC == CC_SpirKernel) ? CCCR_OK
+    return (CC == CC_SpirFunction || CC == CC_OpenCLKernel) ? CCCR_OK
                                                             : CCCR_Warning;
   }
 
   CallingConv getDefaultCallingConv(CallingConvMethodType MT) const override {
     return CC_SpirFunction;
   }
 
   void setSupportedOpenCLOpts() override {
     // Assume all OpenCL extensions and optional core features are supported

cfe/trunk/lib/CodeGen/CGCall.cpp

 #include "clang/AST/DeclObjC.h"
 #include "clang/Basic/TargetBuiltins.h"
 #include "clang/Basic/TargetInfo.h"
 #include "clang/CodeGen/CGFunctionInfo.h"
 #include "clang/CodeGen/SwiftCallingConv.h"
 #include "clang/Frontend/CodeGenOptions.h"
 #include "llvm/ADT/StringExtras.h"
 #include "llvm/IR/Attributes.h"
+#include "llvm/IR/CallingConv.h"
 #include "llvm/IR/CallSite.h"
 #include "llvm/IR/DataLayout.h"
 #include "llvm/IR/InlineAsm.h"
 #include "llvm/IR/Intrinsics.h"
 #include "llvm/IR/IntrinsicInst.h"
 #include "llvm/Transforms/Utils/Local.h"
 using namespace clang;
 using namespace CodeGen;
 
 /***/
 
-static unsigned ClangCallConvToLLVMCallConv(CallingConv CC) {
+unsigned CodeGenTypes::ClangCallConvToLLVMCallConv(CallingConv CC) {
   switch (CC) {
   default: return llvm::CallingConv::C;
   case CC_X86StdCall: return llvm::CallingConv::X86_StdCall;
   case CC_X86FastCall: return llvm::CallingConv::X86_FastCall;
   case CC_X86ThisCall: return llvm::CallingConv::X86_ThisCall;
   case CC_X86_64Win64: return llvm::CallingConv::X86_64_Win64;
   case CC_X86_64SysV: return llvm::CallingConv::X86_64_SysV;
   case CC_AAPCS: return llvm::CallingConv::ARM_AAPCS;
   case CC_AAPCS_VFP: return llvm::CallingConv::ARM_AAPCS_VFP;
   case CC_IntelOclBicc: return llvm::CallingConv::Intel_OCL_BI;
   // TODO: Add support for __pascal to LLVM.
   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_SpirKernel: return llvm::CallingConv::SPIR_KERNEL;
+  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;
   }
 }
 
 /// Derives the 'this' type for codegen purposes, i.e. ignoring method
 /// qualification.

cfe/trunk/lib/CodeGen/CGDebugInfo.cpp

static unsigned getDwarfCC(CallingConv CC) {
 
   // FIXME: Create new DW_CC_ codes for these calling conventions.
   case CC_X86_64Win64:
   case CC_X86_64SysV:
   case CC_AAPCS:
   case CC_AAPCS_VFP:
   case CC_IntelOclBicc:
   case CC_SpirFunction:
-  case CC_SpirKernel:
+  case CC_OpenCLKernel:
   case CC_Swift:
   case CC_PreserveMost:
   case CC_PreserveAll:
     return 0;
   }
   return 0;
 }
 

cfe/trunk/lib/CodeGen/CodeGenTypes.h

class CodeGenTypes {
   SmallVector<const RecordDecl *, 8> DeferredRecords;
   
   /// This map keeps cache of llvm::Types and maps clang::Type to
   /// corresponding llvm::Type.
   llvm::DenseMap<const Type *, llvm::Type *> TypeCache;
 
   llvm::SmallSet<const Type *, 8> RecordsWithOpaqueMemberPointers;
 
+  unsigned ClangCallConvToLLVMCallConv(CallingConv CC);
+
 public:
   CodeGenTypes(CodeGenModule &cgm);
   ~CodeGenTypes();
 
   const llvm::DataLayout &getDataLayout() const {
     return TheModule.getDataLayout();
   }
   ASTContext &getContext() const { return Context; }

cfe/trunk/lib/CodeGen/TargetInfo.h

public:
   virtual void getDependentLibraryOption(llvm::StringRef Lib,
                                          llvm::SmallString<24> &Opt) const;
 
   /// Gets the linker options necessary to detect object file mismatches on
   /// 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;
 };
 
 } // namespace CodeGen
 } // namespace clang
 
 #endif // LLVM_CLANG_LIB_CODEGEN_TARGETINFO_H

cfe/trunk/lib/CodeGen/TargetInfo.cpp

TargetCodeGenInfo::getDependentLibraryOption(llvm::StringRef Lib,
                                              llvm::SmallString<24> &Opt) const {
   // This assumes the user is passing a library name like "rt" instead of a
   // filename like "librt.a/so", and that they don't care whether it's static or
   // dynamic.
   Opt = "-l";
   Opt += Lib;
 }
 
+unsigned TargetCodeGenInfo::getOpenCLKernelCallingConv() const {
+  return llvm::CallingConv::C;
+}
 static bool isEmptyRecord(ASTContext &Context, QualType T, bool AllowArrays);
 
 /// isEmptyField - Return true iff a the field is "empty", that is it
 /// is an unnamed bit-field or an (array of) empty record(s).
 static bool isEmptyField(ASTContext &Context, const FieldDecl *FD,
                          bool AllowArrays) {
   if (FD->isUnnamedBitfield())
     return true;

 namespace {
 
 class AMDGPUTargetCodeGenInfo : public TargetCodeGenInfo {
 public:
   AMDGPUTargetCodeGenInfo(CodeGenTypes &CGT)
     : TargetCodeGenInfo(new DefaultABIInfo(CGT)) {}
   void setTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
                            CodeGen::CodeGenModule &M) const override;
+  unsigned getOpenCLKernelCallingConv() const override;
 };
 
 }
 
 void AMDGPUTargetCodeGenInfo::setTargetAttributes(
   const Decl *D,
   llvm::GlobalValue *GV,
   CodeGen::CodeGenModule &M) const {
if (const auto Attr = FD->getAttr<AMDGPUNumSGPRAttr>()) {
     llvm::Function *F = cast<llvm::Function>(GV);
     unsigned NumSGPR = Attr->getNumSGPR();
     if (NumSGPR != 0)
       F->addFnAttr("amdgpu_num_sgpr", llvm::utostr(NumSGPR));
   }
 }
 
 
+unsigned AMDGPUTargetCodeGenInfo::getOpenCLKernelCallingConv() const {
+  return llvm::CallingConv::AMDGPU_KERNEL;
+}
+
 //===----------------------------------------------------------------------===//
 // SPARC v8 ABI Implementation.
 // Based on the SPARC Compliance Definition version 2.4.1.
 //
 // Ensures that complex values are passed in registers.
 //
 namespace {
 class SparcV8ABIInfo : public DefaultABIInfo {

 
 namespace {
 class SPIRTargetCodeGenInfo : public TargetCodeGenInfo {
 public:
   SPIRTargetCodeGenInfo(CodeGen::CodeGenTypes &CGT)
     : TargetCodeGenInfo(new DefaultABIInfo(CGT)) {}
   void emitTargetMD(const Decl *D, llvm::GlobalValue *GV,
                     CodeGen::CodeGenModule &M) const override;
+  unsigned getOpenCLKernelCallingConv() const override;
 };
 } // End anonymous namespace.
 
 /// Emit SPIR specific metadata: OpenCL and SPIR version.
 void SPIRTargetCodeGenInfo::emitTargetMD(const Decl *D, llvm::GlobalValue *GV,
                                          CodeGen::CodeGenModule &CGM) const {
   llvm::LLVMContext &Ctx = CGM.getModule().getContext();
   llvm::Type *Int32Ty = llvm::Type::getInt32Ty(Ctx);
llvm::Metadata *OCLVerElts[] = {
           Int32Ty, CGM.getLangOpts().OpenCLVersion / 100)),
       llvm::ConstantAsMetadata::get(llvm::ConstantInt::get(
           Int32Ty, (CGM.getLangOpts().OpenCLVersion % 100) / 10))};
   llvm::NamedMDNode *OCLVerMD =
       M.getOrInsertNamedMetadata("opencl.ocl.version");
   OCLVerMD->addOperand(llvm::MDNode::get(Ctx, OCLVerElts));
 }
 
+unsigned SPIRTargetCodeGenInfo::getOpenCLKernelCallingConv() const {
+  return llvm::CallingConv::SPIR_KERNEL;
+}
+
 static bool appendType(SmallStringEnc &Enc, QualType QType,
                        const CodeGen::CodeGenModule &CGM,
                        TypeStringCache &TSC);
 
 /// Helper function for appendRecordType().
 /// Builds a SmallVector containing the encoded field types in declaration
 /// order.
 static bool extractFieldType(SmallVectorImpl<FieldEncoding> &FE,

cfe/trunk/lib/Sema/SemaType.cpp

if (FoundNonParen) {
           D.getDeclSpec().getStorageClassSpec() != DeclSpec::SCS_typedef &&
           !D.isStaticMember();
     }
   }
 
   CallingConv CC = S.Context.getDefaultCallingConvention(FTI.isVariadic,
                                                          IsCXXInstanceMethod);
 
-  // Attribute AT_OpenCLKernel affects the calling convention only on
-  // the SPIR target, hence it cannot be treated as a calling
+  // 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
-  // "spir_kernel" for OpenCL kernels on SPIR.
-  if (CC == CC_SpirFunction) {
+  // 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) {
-        CC = CC_SpirKernel;
+        llvm::Triple::ArchType arch = S.Context.getTargetInfo().getTriple().getArch();
+        if (arch == llvm::Triple::spir || arch == llvm::Triple::spir64 ||
+            arch == llvm::Triple::amdgcn) {
+          CC = CC_OpenCLKernel;
+        }
         break;
       }
     }
   }
 
   return CC;
 }
 

cfe/trunk/test/CodeGenOpenCL/amdgpu-call-kernel.cl

+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s
+// CHECK: define amdgpu_kernel void @test_call_kernel(i32 addrspace(1)* nocapture %out)
+// CHECK: store i32 4, i32 addrspace(1)* %out, align 4
+
+kernel void test_kernel(global int *out)
+{
+  out[0] = 4;
+}
+
+__kernel void test_call_kernel(__global int *out)
+{
+  test_kernel(out);
+}

cfe/trunk/test/CodeGenOpenCL/amdgpu-calling-conv.cl

+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s
+
+// CHECK: define amdgpu_kernel void @calling_conv_amdgpu_kernel()
+kernel void calling_conv_amdgpu_kernel()
+{
+}
+
+// CHECK: define void @calling_conv_none()
+void calling_conv_none()
+{
+}

cfe/trunk/test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl

 // RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -O0 -emit-llvm -o - %s | FileCheck %s
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -verify -o - %s | FileCheck -check-prefix=X86 %s
 
 // Make sure this is silently accepted on other targets.
 
 __attribute__((amdgpu_num_vgpr(64))) // expected-no-diagnostics
 kernel void test_num_vgpr64() {
-// CHECK: define void @test_num_vgpr64() [[ATTR_VGPR64:#[0-9]+]]
+// CHECK: define amdgpu_kernel void @test_num_vgpr64() [[ATTR_VGPR64:#[0-9]+]]
 }
 
 __attribute__((amdgpu_num_sgpr(32))) // expected-no-diagnostics
 kernel void test_num_sgpr32() {
-// CHECK: define void @test_num_sgpr32() [[ATTR_SGPR32:#[0-9]+]]
+// CHECK: define amdgpu_kernel void @test_num_sgpr32() [[ATTR_SGPR32:#[0-9]+]]
 }
 
 __attribute__((amdgpu_num_vgpr(64), amdgpu_num_sgpr(32))) // expected-no-diagnostics
 kernel void test_num_vgpr64_sgpr32() {
-// CHECK: define void @test_num_vgpr64_sgpr32() [[ATTR_VGPR64_SGPR32:#[0-9]+]]
+// CHECK: define amdgpu_kernel void @test_num_vgpr64_sgpr32() [[ATTR_VGPR64_SGPR32:#[0-9]+]]
 
 }
 
 __attribute__((amdgpu_num_sgpr(20), amdgpu_num_vgpr(40))) // expected-no-diagnostics
 kernel void test_num_sgpr20_vgpr40() {
-// CHECK: define void @test_num_sgpr20_vgpr40() [[ATTR_SGPR20_VGPR40:#[0-9]+]]
+// CHECK: define amdgpu_kernel void @test_num_sgpr20_vgpr40() [[ATTR_SGPR20_VGPR40:#[0-9]+]]
 }
 
 __attribute__((amdgpu_num_vgpr(0))) // expected-no-diagnostics
 kernel void test_num_vgpr0() {
 }
 
 __attribute__((amdgpu_num_sgpr(0))) // expected-no-diagnostics
 kernel void test_num_sgpr0() {

cfe/trunk/tools/libclang/CXType.cpp

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