From 6c10a66ec7732e8b7afd8ac0260c2e572fbe0a31 Mon Sep 17 00:00:00 2001 From: Yaxun Liu Date: Tue, 12 Jun 2018 00:16:33 +0000 Subject: [PATCH] [CUDA][HIP] Set kernel calling convention before arrange function Currently clang set kernel calling convention for CUDA/HIP after arranging function, which causes incorrect kernel function type since it depends on calling convention. This patch moves setting kernel convention before arranging function. Differential Revision: https://reviews.llvm.org/D47733 llvm-svn: 334457 --- clang/lib/CodeGen/CGCall.cpp | 15 ++++++++++- clang/lib/CodeGen/CodeGenModule.cpp | 2 -- clang/lib/CodeGen/TargetInfo.cpp | 7 ++--- clang/lib/CodeGen/TargetInfo.h | 2 +- clang/test/CodeGenCUDA/kernel-args.cu | 39 +++++++++++++++++++++++++++ 5 files changed, 58 insertions(+), 7 deletions(-) create mode 100644 clang/test/CodeGenCUDA/kernel-args.cu diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index 2d8339ef6970..064b5fcb102f 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -255,6 +255,16 @@ CodeGenTypes::arrangeCXXMethodType(const CXXRecordDecl *RD, FTP->getCanonicalTypeUnqualified().getAs(), MD); } +/// Set calling convention for CUDA/HIP kernel. +static void setCUDAKernelCallingConvention(CanQualType &FTy, CodeGenModule &CGM, + const FunctionDecl *FD) { + if (FD->hasAttr()) { + const FunctionType *FT = FTy->getAs(); + CGM.getTargetCodeGenInfo().setCUDAKernelCallingConvention(FT); + FTy = FT->getCanonicalTypeUnqualified(); + } +} + /// Arrange the argument and result information for a declaration or /// definition of the given C++ non-static member function. The /// member function must be an ordinary function, i.e. not a @@ -264,7 +274,9 @@ CodeGenTypes::arrangeCXXMethodDeclaration(const CXXMethodDecl *MD) { assert(!isa(MD) && "wrong method for constructors!"); assert(!isa(MD) && "wrong method for destructors!"); - CanQual prototype = GetFormalType(MD); + CanQualType FT = GetFormalType(MD).getAs(); + setCUDAKernelCallingConvention(FT, CGM, MD); + auto prototype = FT.getAs(); if (MD->isInstance()) { // The abstract case is perfectly fine. @@ -424,6 +436,7 @@ CodeGenTypes::arrangeFunctionDeclaration(const FunctionDecl *FD) { CanQualType FTy = FD->getType()->getCanonicalTypeUnqualified(); assert(isa(FTy)); + setCUDAKernelCallingConvention(FTy, CGM, FD); // When declaring a function without a prototype, always use a // non-variadic type. diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 814eda4381b5..6595694ebd53 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -3671,8 +3671,6 @@ void CodeGenModule::EmitGlobalFunctionDefinition(GlobalDecl GD, MaybeHandleStaticInExternC(D, Fn); - if (D->hasAttr()) - getTargetCodeGenInfo().setCUDAKernelCallingConvention(Fn); maybeSetTrivialComdat(*D, *Fn); diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp index 3ec00553607c..b29bcce237cf 100644 --- a/clang/lib/CodeGen/TargetInfo.cpp +++ b/clang/lib/CodeGen/TargetInfo.cpp @@ -7646,7 +7646,7 @@ public: llvm::Function *BlockInvokeFunc, llvm::Value *BlockLiteral) const override; bool shouldEmitStaticExternCAliases() const override; - void setCUDAKernelCallingConvention(llvm::Function *F) const override; + void setCUDAKernelCallingConvention(const FunctionType *&FT) const override; }; } @@ -7783,8 +7783,9 @@ bool AMDGPUTargetCodeGenInfo::shouldEmitStaticExternCAliases() const { } void AMDGPUTargetCodeGenInfo::setCUDAKernelCallingConvention( - llvm::Function *F) const { - F->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL); + const FunctionType *&FT) const { + FT = getABIInfo().getContext().adjustFunctionType( + FT, FT->getExtInfo().withCallingConv(CC_OpenCLKernel)); } //===----------------------------------------------------------------------===// diff --git a/clang/lib/CodeGen/TargetInfo.h b/clang/lib/CodeGen/TargetInfo.h index 5c19c7141243..b530260ea48f 100644 --- a/clang/lib/CodeGen/TargetInfo.h +++ b/clang/lib/CodeGen/TargetInfo.h @@ -302,7 +302,7 @@ public: /// as 'used', and having internal linkage. virtual bool shouldEmitStaticExternCAliases() const { return true; } - virtual void setCUDAKernelCallingConvention(llvm::Function *F) const {} + virtual void setCUDAKernelCallingConvention(const FunctionType *&FT) const {} }; } // namespace CodeGen diff --git a/clang/test/CodeGenCUDA/kernel-args.cu b/clang/test/CodeGenCUDA/kernel-args.cu new file mode 100644 index 000000000000..d0986629f8d4 --- /dev/null +++ b/clang/test/CodeGenCUDA/kernel-args.cu @@ -0,0 +1,39 @@ +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \ +// RUN: -emit-llvm %s -o - | FileCheck -check-prefix=AMDGCN %s +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda- -fcuda-is-device \ +// RUN: -emit-llvm %s -o - | FileCheck -check-prefix=NVPTX %s +#include "Inputs/cuda.h" + +struct A { + int a[32]; +}; + +// AMDGCN: define amdgpu_kernel void @_Z6kernel1A(%struct.A %x.coerce) +// NVPTX: define void @_Z6kernel1A(%struct.A* byval align 4 %x) +__global__ void kernel(A x) { +} + +class Kernel { +public: + // AMDGCN: define amdgpu_kernel void @_ZN6Kernel12memberKernelE1A(%struct.A %x.coerce) + // NVPTX: define void @_ZN6Kernel12memberKernelE1A(%struct.A* byval align 4 %x) + static __global__ void memberKernel(A x){} + template static __global__ void templateMemberKernel(T x) {} +}; + + +template +__global__ void templateKernel(T x) {} + +void launch(void*); + +void test() { + Kernel K; + // AMDGCN: define amdgpu_kernel void @_Z14templateKernelI1AEvT_(%struct.A %x.coerce) + // NVPTX: define void @_Z14templateKernelI1AEvT_(%struct.A* byval align 4 %x) + launch((void*)templateKernel); + + // AMDGCN: define amdgpu_kernel void @_ZN6Kernel20templateMemberKernelI1AEEvT_(%struct.A %x.coerce) + // NVPTX: define void @_ZN6Kernel20templateMemberKernelI1AEEvT_(%struct.A* byval align 4 %x) + launch((void*)Kernel::templateMemberKernel); +}