https://github.com/wenhu1024 updated https://github.com/llvm/llvm-project/pull/231567
>From eeaa28f6f2233ba4b3832c556808ca7886ff726c Mon Sep 17 00:00:00 2001 From: wenhu1024 <[email protected]> Date: Tue, 13 Oct 2026 23:58:33 +0800 Subject: [PATCH] [Clang][CUDA] Preserve addresses of scalar grid-constant parameters --- clang/include/clang/CodeGen/CGFunctionInfo.h | 29 +++- clang/lib/CodeGen/CGCall.cpp | 67 ++++++-- clang/lib/CodeGen/CodeGenTypes.h | 5 +- clang/lib/CodeGen/Targets/NVPTX.cpp | 15 +- clang/test/CodeGenCUDA/grid-constant-args.cu | 165 +++++++++++++++++++ clang/test/CodeGenCUDA/grid-constant.cu | 3 +- 6 files changed, 257 insertions(+), 27 deletions(-) create mode 100644 clang/test/CodeGenCUDA/grid-constant-args.cu diff --git a/clang/include/clang/CodeGen/CGFunctionInfo.h b/clang/include/clang/CodeGen/CGFunctionInfo.h index d9c449d78eb08f..11e36626ed43e6 100644 --- a/clang/include/clang/CodeGen/CGFunctionInfo.h +++ b/clang/include/clang/CodeGen/CGFunctionInfo.h @@ -19,6 +19,7 @@ #include "clang/AST/CharUnits.h" #include "clang/AST/Type.h" #include "llvm/ADT/FoldingSet.h" +#include "llvm/ADT/STLExtras.h" #include "llvm/IR/DerivedTypes.h" #include "llvm/Support/TrailingObjects.h" #include <cassert> @@ -598,7 +599,8 @@ struct CGFunctionInfoArgInfo { class CGFunctionInfo final : public llvm::FoldingSetNode, private llvm::TrailingObjects<CGFunctionInfo, CGFunctionInfoArgInfo, - FunctionProtoType::ExtParameterInfo> { + FunctionProtoType::ExtParameterInfo, + unsigned> { typedef CGFunctionInfoArgInfo ArgInfo; typedef FunctionProtoType::ExtParameterInfo ExtParameterInfo; @@ -670,6 +672,10 @@ class CGFunctionInfo final unsigned NumArgs; + /// Indices of CUDA grid-constant arguments. These affect ABI classification + /// but are not part of the source function type. + unsigned NumGridConstantArgs; + ArgInfo *getArgsBuffer() { return getTrailingObjects<ArgInfo>(); } @@ -692,7 +698,8 @@ class CGFunctionInfo final bool delegateCall, unsigned X86ABIAVXLevel, const FunctionType::ExtInfo &extInfo, ArrayRef<ExtParameterInfo> paramInfos, CanQualType resultType, - ArrayRef<CanQualType> argTypes, RequiredArgs required); + ArrayRef<CanQualType> argTypes, RequiredArgs required, + ArrayRef<unsigned> gridConstantArgs = {}); void operator delete(void *p) { ::operator delete(p); } // Friending class TrailingObjects is apparently not good enough for MSVC, @@ -792,6 +799,14 @@ class CGFunctionInfo final return getExtParameterInfos()[argIndex]; } + ArrayRef<unsigned> getGridConstantArgs() const { + return {getTrailingObjects<unsigned>(), NumGridConstantArgs}; + } + + bool isGridConstantArg(unsigned argIndex) const { + return llvm::is_contained(getGridConstantArgs(), argIndex); + } + /// Return true if this function uses inalloca arguments. bool usesInAlloca() const { return ArgStruct; } @@ -837,6 +852,9 @@ class CGFunctionInfo final for (auto paramInfo : getExtParameterInfos()) ID.AddInteger(paramInfo.getOpaqueValue()); } + ID.AddInteger(NumGridConstantArgs); + for (unsigned argIndex : getGridConstantArgs()) + ID.AddInteger(argIndex); getReturnType().Profile(ID); for (const auto &I : arguments()) I.type.Profile(ID); @@ -847,7 +865,8 @@ class CGFunctionInfo final const FunctionType::ExtInfo &info, ArrayRef<ExtParameterInfo> paramInfos, RequiredArgs required, CanQualType resultType, - ArrayRef<CanQualType> argTypes) { + ArrayRef<CanQualType> argTypes, + ArrayRef<unsigned> gridConstantArgs = {}) { ID.AddInteger(info.getCC()); ID.AddBoolean(InstanceMethod); ID.AddBoolean(ChainCall); @@ -866,6 +885,10 @@ class CGFunctionInfo final for (auto paramInfo : paramInfos) ID.AddInteger(paramInfo.getOpaqueValue()); } + // Match the unsigned NumGridConstantArgs used by the instance profile. + ID.AddInteger(static_cast<unsigned>(gridConstantArgs.size())); + for (unsigned argIndex : gridConstantArgs) + ID.AddInteger(argIndex); resultType.Profile(ID); for (const CanQualType &argType : argTypes) argType.Profile(ID); diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index 2193a89203f293..3c0ada9257b066 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -343,6 +343,28 @@ static CallingConv getCallingConventionForDecl(const ObjCMethodDecl *D, return CC_C; } +/// Collect grid-constant argument indices separately from the source function +/// type, so functions with different grid-constant parameters do not share ABI +/// information. Host stubs retain their normal argument classification. +static SmallVector<unsigned, 4> getGridConstantArgs(CodeGenModule &CGM, + const FunctionDecl *FD, + unsigned PrefixArgs = 0) { + SmallVector<unsigned, 4> Args; + if (!FD || !CGM.getLangOpts().CUDAIsDevice || + !CGM.getTarget().getTriple().isNVPTX() || !FD->hasAttr<CUDAGlobalAttr>()) + return Args; + + unsigned ArgIndex = PrefixArgs; + for (const ParmVarDecl *Param : FD->parameters()) { + if (Param->hasAttr<CUDAGridConstantAttr>()) + Args.push_back(ArgIndex); + ++ArgIndex; + if (Param->hasAttr<PassObjectSizeAttr>()) + ++ArgIndex; + } + return Args; +} + /// Arrange the argument and result information for a call to an /// unknown C++ non-static member function of the given abstract type. /// (A null RD means we don't have any meaningful "this" argument type, @@ -366,7 +388,7 @@ CodeGenTypes::arrangeCXXMethodType(const CXXRecordDecl *RD, return arrangeLLVMFunctionInfo( CanonicalFTP->getReturnType().getUnqualifiedType(), FnInfoOpts::IsInstanceMethod, argTypes, CanonicalFTP->getExtInfo(), - paramInfos, required, MD); + paramInfos, required, MD, getGridConstantArgs(CGM, MD, 1)); } /// Set calling convention for CUDA/HIP kernel. @@ -405,7 +427,8 @@ CodeGenTypes::arrangeCXXMethodDeclaration(const CXXMethodDecl *MD) { return arrangeLLVMFunctionInfo( prototype->getReturnType().getUnqualifiedType(), FnInfoOpts::None, argTypes, prototype->getExtInfo(), paramInfos, - RequiredArgs::forPrototypePlus(prototype.getTypePtr(), 0), MD); + RequiredArgs::forPrototypePlus(prototype.getTypePtr(), 0), MD, + getGridConstantArgs(CGM, MD)); } bool CodeGenTypes::inheritingCtorHasParams( @@ -572,10 +595,10 @@ CodeGenTypes::arrangeFunctionDeclaration(const GlobalDecl GD) { CanQualTypeList argTypes; ExtParameterInfoList paramInfos; appendParameterTypes(*this, argTypes, paramInfos, FTP); - return arrangeLLVMFunctionInfo(FTP->getReturnType().getUnqualifiedType(), - FnInfoOpts::None, argTypes, FTP->getExtInfo(), - paramInfos, - RequiredArgs::forPrototypePlus(FTP, 0), FD); + return arrangeLLVMFunctionInfo( + FTP->getReturnType().getUnqualifiedType(), FnInfoOpts::None, argTypes, + FTP->getExtInfo(), paramInfos, RequiredArgs::forPrototypePlus(FTP, 0), FD, + getGridConstantArgs(CGM, FD)); } /// Arrange the argument and result information for the declaration or @@ -859,7 +882,7 @@ const CGFunctionInfo &CodeGenTypes::arrangeCall(const CGFunctionInfo &signature, signature.isInstanceMethod(), signature.isChainCall(), signature.isDelegateCall(), X86ABIAVXLevel, signature.getExtInfo(), paramInfos, signature.getRequiredArgs(), signature.getReturnType(), - argTypes); + argTypes, signature.getGridConstantArgs()); return *newFI; } @@ -1075,7 +1098,8 @@ const CGFunctionInfo &CodeGenTypes::arrangeLLVMFunctionInfo( CanQualType resultType, FnInfoOpts opts, ArrayRef<CanQualType> argTypes, FunctionType::ExtInfo info, ArrayRef<FunctionProtoType::ExtParameterInfo> paramInfos, - RequiredArgs required, const FunctionDecl *ABIInfoFD) { + RequiredArgs required, const FunctionDecl *ABIInfoFD, + ArrayRef<unsigned> gridConstantArgs) { assert(llvm::all_of(argTypes, [](CanQualType T) { return T.isCanonicalAsParam(); })); @@ -1091,7 +1115,7 @@ const CGFunctionInfo &CodeGenTypes::arrangeLLVMFunctionInfo( const CGFunctionInfo *newFI = findOrInsertCGFunctionInfo( isInstanceMethod, isChainCall, isDelegateCall, X86ABIAVXLevel, info, - paramInfos, required, resultType, argTypes); + paramInfos, required, resultType, argTypes, gridConstantArgs); return *newFI; } @@ -1100,11 +1124,11 @@ CGFunctionInfo *CodeGenTypes::findOrInsertCGFunctionInfo( unsigned X86ABIAVXLevel, const FunctionType::ExtInfo &info, ArrayRef<FunctionProtoType::ExtParameterInfo> paramInfos, RequiredArgs required, CanQualType resultType, - ArrayRef<CanQualType> argTypes) { + ArrayRef<CanQualType> argTypes, ArrayRef<unsigned> gridConstantArgs) { llvm::FoldingSetNodeID ID; CGFunctionInfo::Profile(ID, isInstanceMethod, isChainCall, isDelegateCall, X86ABIAVXLevel, info, paramInfos, required, - resultType, argTypes); + resultType, argTypes, gridConstantArgs); llvm::FoldingSetInsertToken InsertToken; CGFunctionInfo *FI = FunctionInfos.lookup(ID, InsertToken); @@ -1116,7 +1140,7 @@ CGFunctionInfo *CodeGenTypes::findOrInsertCGFunctionInfo( // Construct the function info. We co-allocate the ArgInfos. FI = CGFunctionInfo::create(CC, isInstanceMethod, isChainCall, isDelegateCall, X86ABIAVXLevel, info, paramInfos, resultType, - argTypes, required); + argTypes, required, gridConstantArgs); FunctionInfos.insert(FI, InsertToken); bool inserted = FunctionsBeingProcessed.insert(FI).second; @@ -1162,13 +1186,18 @@ CGFunctionInfo *CGFunctionInfo::create( unsigned llvmCC, bool instanceMethod, bool chainCall, bool delegateCall, unsigned X86ABIAVXLevel, const FunctionType::ExtInfo &info, ArrayRef<ExtParameterInfo> paramInfos, CanQualType resultType, - ArrayRef<CanQualType> argTypes, RequiredArgs required) { + ArrayRef<CanQualType> argTypes, RequiredArgs required, + ArrayRef<unsigned> gridConstantArgs) { assert(paramInfos.empty() || paramInfos.size() == argTypes.size()); assert(!required.allowsOptionalArgs() || required.getNumRequiredArgs() <= argTypes.size()); + assert(llvm::all_of(gridConstantArgs, [&](unsigned ArgIndex) { + return ArgIndex < argTypes.size(); + })); - void *buffer = operator new(totalSizeToAlloc<ArgInfo, ExtParameterInfo>( - argTypes.size() + 1, paramInfos.size())); + void *buffer = operator new( + totalSizeToAlloc<ArgInfo, ExtParameterInfo, unsigned>( + argTypes.size() + 1, paramInfos.size(), gridConstantArgs.size())); CGFunctionInfo *FI = new (buffer) CGFunctionInfo(); FI->CallingConvention = llvmCC; @@ -1189,6 +1218,7 @@ CGFunctionInfo *CGFunctionInfo::create( FI->ArgStruct = nullptr; FI->ArgStructAlign = 0; FI->NumArgs = argTypes.size(); + FI->NumGridConstantArgs = gridConstantArgs.size(); FI->HasExtParameterInfos = !paramInfos.empty(); FI->getArgsBuffer()[0].type = resultType; FI->MaxVectorWidth = 0; @@ -1196,6 +1226,7 @@ CGFunctionInfo *CGFunctionInfo::create( FI->getArgsBuffer()[i + 1].type = argTypes[i]; for (unsigned i = 0, e = paramInfos.size(); i != e; ++i) FI->getExtParameterInfosBuffer()[i] = paramInfos[i]; + llvm::copy(gridConstantArgs, FI->getTrailingObjects<unsigned>()); return FI; } @@ -3561,7 +3592,11 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI, Fn->getArg(FirstIRArg), Ty, ArgI.getIndirectAlign(), false, nullptr, nullptr, KnownNonNull); - if (!hasScalarEvaluationKind(Ty)) { + if (FI.isGridConstantArg(ArgNo)) { + // Grid-constant parameters must retain their parameter-space address, + // including scalars that would otherwise be loaded into a local copy. + ArgVals.push_back(ParamValue::forIndirect(ParamAddr)); + } else if (!hasScalarEvaluationKind(Ty)) { // Aggregates and complex variables are accessed by reference. All we // need to do is realign the value, if requested. Also, if the address // may be aliased, copy it to ensure that the parameter variable is diff --git a/clang/lib/CodeGen/CodeGenTypes.h b/clang/lib/CodeGen/CodeGenTypes.h index 860f1f1ccdce28..b5acaa5bd3e05a 100644 --- a/clang/lib/CodeGen/CodeGenTypes.h +++ b/clang/lib/CodeGen/CodeGenTypes.h @@ -100,7 +100,7 @@ class CodeGenTypes { unsigned X86ABIAVXLevel, const FunctionType::ExtInfo &info, ArrayRef<FunctionProtoType::ExtParameterInfo> paramInfos, RequiredArgs required, CanQualType resultType, - ArrayRef<CanQualType> argTypes); + ArrayRef<CanQualType> argTypes, ArrayRef<unsigned> gridConstantArgs = {}); public: CodeGenTypes(CodeGenModule &cgm); @@ -294,7 +294,8 @@ class CodeGenTypes { CanQualType returnType, FnInfoOpts opts, ArrayRef<CanQualType> argTypes, FunctionType::ExtInfo info, ArrayRef<FunctionProtoType::ExtParameterInfo> paramInfos, - RequiredArgs args, const FunctionDecl *ABIInfoFD); + RequiredArgs args, const FunctionDecl *ABIInfoFD, + ArrayRef<unsigned> gridConstantArgs = {}); /// Compute a new LLVM record layout object for the given record. std::unique_ptr<CGRecordLayout> ComputeRecordLayout(const RecordDecl *D, diff --git a/clang/lib/CodeGen/Targets/NVPTX.cpp b/clang/lib/CodeGen/Targets/NVPTX.cpp index 00fe271c85b0b3..8f89c4cba4f1d1 100644 --- a/clang/lib/CodeGen/Targets/NVPTX.cpp +++ b/clang/lib/CodeGen/Targets/NVPTX.cpp @@ -219,10 +219,17 @@ void NVPTXABIInfo::computeInfo(CGFunctionInfo &FI) const { if (!getCXXABI().classifyReturnType(FI)) FI.getReturnInfo() = classifyReturnType(FI.getReturnType()); - for (auto &&[ArgumentsCount, I] : llvm::enumerate(FI.arguments())) - I.info = ArgumentsCount < FI.getNumRequiredArgs() - ? classifyArgumentType(I.type) - : ABIArgInfo::getDirect(); + for (auto &&[ArgumentsCount, I] : llvm::enumerate(FI.arguments())) { + if (FI.isGridConstantArg(ArgumentsCount)) + // NVVM represents grid-constant parameter objects as byval pointers, + // even when their source type is a scalar or vector. + I.info = getNaturalAlignIndirect( + I.type, getDataLayout().getAllocaAddrSpace(), /*byval=*/true); + else + I.info = ArgumentsCount < FI.getNumRequiredArgs() + ? classifyArgumentType(I.type) + : ABIArgInfo::getDirect(); + } // Always honor user-specified calling convention. if (FI.getCallingConvention() != llvm::CallingConv::C) diff --git a/clang/test/CodeGenCUDA/grid-constant-args.cu b/clang/test/CodeGenCUDA/grid-constant-args.cu new file mode 100644 index 00000000000000..61d020c0247ebc --- /dev/null +++ b/clang/test/CodeGenCUDA/grid-constant-args.cu @@ -0,0 +1,165 @@ +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -O0 -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s --check-prefixes=DEVICE,PTR64 +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -O2 -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s --check-prefixes=DEVICE,PTR64 +// RUN: %clang_cc1 -triple nvptx-nvidia-cuda -fcuda-is-device -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s --check-prefixes=DEVICE,PTR32 +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s | FileCheck %s --check-prefix=HOST + +#include "Inputs/cuda.h" + +extern "C" __device__ void escape(const void *); + +struct S { + int x, y; +}; + +// Functions of the same source type must not share ABI classifications when +// their parameters have different grid-constant attributes. +// DEVICE-LABEL: define{{.*}} ptx_kernel void @plain_before(i32 noundef %arg) +// DEVICE: %arg.addr = alloca i32, align 4 +// DEVICE: store i32 %arg, ptr %arg.addr, align 4 +// DEVICE: call void @escape(ptr noundef %arg.addr) +extern "C" __global__ void plain_before(const int arg) { + escape(&arg); +} + +// DEVICE-LABEL: define{{.*}} ptx_kernel void @grid_after(ptr noundef byval(i32) align 4 "nvvm.grid_constant" %arg) +// DEVICE-NOT: alloca +// DEVICE: call void @escape(ptr noundef %arg) +// DEVICE-NEXT: ret void +// HOST-LABEL: define{{.*}} void @__device_stub__grid_after(i32 noundef %arg) +extern "C" __global__ void grid_after(const __grid_constant__ int arg) { + escape(&arg); +} + +// DEVICE-LABEL: define{{.*}} ptx_kernel void @plain_after(i32 noundef %arg) +// DEVICE: %arg.addr = alloca i32, align 4 +// DEVICE: call void @escape(ptr noundef %arg.addr) +extern "C" __global__ void plain_after(const int arg) { + escape(&arg); +} + +// All three addresses must refer to the incoming byval parameter objects. +// DEVICE-LABEL: define{{.*}} ptx_kernel void @mixed(ptr noundef byval(%struct.S) align 4 "nvvm.grid_constant" %s, ptr noundef byval(i32) align 4 "nvvm.grid_constant" %scalar, ptr noundef byval(%struct.char4) align 1 "nvvm.grid_constant" %vec) +// DEVICE-NOT: alloca +// DEVICE: call void @escape(ptr noundef %s) +// DEVICE-NEXT: call void @escape(ptr noundef %scalar) +// DEVICE-NEXT: call void @escape(ptr noundef %vec) +// DEVICE-NEXT: ret void +extern "C" __global__ void mixed(const __grid_constant__ S s, + const __grid_constant__ int scalar, + const __grid_constant__ char4 vec) { + escape(&s); + escape(&scalar); + escape(&vec); +} + +// The grid-constant argument positions are part of the ABI cache key. +// DEVICE-LABEL: define{{.*}} ptx_kernel void @grid_first(ptr noundef byval(i32) align 4 "nvvm.grid_constant" %a, i32 noundef %b) +// DEVICE: %b.addr = alloca i32, align 4 +// DEVICE: call void @escape(ptr noundef %a) +// DEVICE-NEXT: call void @escape(ptr noundef %b.addr) +extern "C" __global__ void grid_first(const __grid_constant__ int a, const int b) { + escape(&a); + escape(&b); +} + +// DEVICE-LABEL: define{{.*}} ptx_kernel void @grid_second(i32 noundef %a, ptr noundef byval(i32) align 4 "nvvm.grid_constant" %b) +// DEVICE: %a.addr = alloca i32, align 4 +// DEVICE: call void @escape(ptr noundef %a.addr) +// DEVICE-NEXT: call void @escape(ptr noundef %b) +extern "C" __global__ void grid_second(const int a, const __grid_constant__ int b) { + escape(&a); + escape(&b); +} + +enum E : unsigned short { E0 }; +using V = int __attribute__((ext_vector_type(4))); + +// DEVICE-LABEL: define{{.*}} ptx_kernel void @scalar_types( +// DEVICE-SAME: ptr noundef byval(i8) align 1 "nvvm.grid_constant" %c, +// DEVICE-SAME: ptr noundef byval(i16) align 2 "nvvm.grid_constant" %h, +// DEVICE-SAME: ptr noundef byval(i8) align 1 "nvvm.grid_constant" %b, +// DEVICE-SAME: ptr noundef byval(float) align 4 "nvvm.grid_constant" %f, +// DEVICE-SAME: ptr noundef byval(double) align 8 "nvvm.grid_constant" %d, +// PTR64-SAME: ptr noundef byval(ptr) align 8 "nvvm.grid_constant" %p, +// PTR32-SAME: ptr noundef byval(ptr) align 4 "nvvm.grid_constant" %p, +// DEVICE-SAME: ptr noundef byval(i16) align 2 "nvvm.grid_constant" %e, +// DEVICE-SAME: ptr noundef byval(<4 x i32>) align 16 "nvvm.grid_constant" %v) +// DEVICE-NOT: alloca +// DEVICE: call void @escape(ptr noundef %c) +// DEVICE-NEXT: call void @escape(ptr noundef %h) +// DEVICE-NEXT: call void @escape(ptr noundef %b) +// DEVICE-NEXT: call void @escape(ptr noundef %f) +// DEVICE-NEXT: call void @escape(ptr noundef %d) +// DEVICE-NEXT: call void @escape(ptr noundef %p) +// DEVICE-NEXT: call void @escape(ptr noundef %e) +// DEVICE-NEXT: call void @escape(ptr noundef %v) +// DEVICE-NEXT: ret void +extern "C" __global__ void scalar_types( + const __grid_constant__ char c, const __grid_constant__ short h, + const __grid_constant__ bool b, const __grid_constant__ float f, + const __grid_constant__ double d, int *const __grid_constant__ p, + const __grid_constant__ E e, const __grid_constant__ V v) { + escape(&c); + escape(&h); + escape(&b); + escape(&f); + escape(&d); + escape(&p); + escape(&e); + escape(&v); +} + +// Values are loaded from the parameter objects, while a source-level copy is +// still a distinct local object. +// DEVICE-LABEL: define{{.*}} ptx_kernel void @values( +// DEVICE: %copy = alloca i32, align 4 +// DEVICE: load i32, ptr %s, align 4 +// DEVICE: store i32 {{.*}}, ptr %copy, align 4 +// DEVICE: call void @escape(ptr noundef %s) +// DEVICE-NEXT: call void @escape(ptr noundef %copy) +// DEVICE: load i8, ptr %b, align 1 +// DEVICE: load <4 x i32>, ptr %v, align 16 +// DEVICE: load ptr, ptr %p, align +extern "C" __global__ void values( + const __grid_constant__ int s, const __grid_constant__ bool b, + int *const __grid_constant__ p, const __grid_constant__ V v) { + const int copy = s; + escape(&s); + escape(©); + if (b) + *p = v.x; +} + +// A prior declaration and template instantiation must use the same byval ABI. +extern "C" __global__ void declared(const __grid_constant__ int arg); +void launch_declared() { + declared<<<1, 1>>>(1); +} +// DEVICE-LABEL: define{{.*}} ptx_kernel void @declared(ptr noundef byval(i32) align 4 "nvvm.grid_constant" %arg) +// DEVICE-NOT: alloca +// DEVICE: call void @escape(ptr noundef %arg) +extern "C" __global__ void declared(const __grid_constant__ int arg) { + escape(&arg); +} + +// DEVICE-LABEL: define{{.*}} ptx_kernel void @_ZN7Kernels6memberEi(ptr noundef byval(i32) align 4 "nvvm.grid_constant" %arg) +// DEVICE-NOT: alloca +// DEVICE: call void @escape(ptr noundef %arg) +struct Kernels { + static __global__ void member(const __grid_constant__ int arg) { + escape(&arg); + } +}; + +// DEVICE-LABEL: define{{.*}} ptx_kernel void @_Z15template_kernelIiEvT_(ptr noundef byval(i32) align 4 "nvvm.grid_constant" %arg) +// DEVICE-NOT: alloca +// DEVICE: call void @escape(ptr noundef %arg) +template <typename T> +__global__ void template_kernel(const __grid_constant__ T arg) { + escape(&arg); +} + +void launch() { + Kernels::member<<<1, 1>>>(1); + template_kernel<int><<<1, 1>>>(1); +} diff --git a/clang/test/CodeGenCUDA/grid-constant.cu b/clang/test/CodeGenCUDA/grid-constant.cu index 120b854e56746f..468001f8e05736 100644 --- a/clang/test/CodeGenCUDA/grid-constant.cu +++ b/clang/test/CodeGenCUDA/grid-constant.cu @@ -20,8 +20,7 @@ void foo() { tkernel<const S><<<1,1>>>(1, {}); } -// CHECK: define dso_local ptx_kernel void @_Z6kernel1Sii(ptr noundef byval(%struct.S) align 1 "nvvm.grid_constant" %gc_arg1, i32 noundef %arg2, i32 noundef "nvvm.grid_constant" %gc_arg3) +// CHECK: define dso_local ptx_kernel void @_Z6kernel1Sii(ptr noundef byval(%struct.S) align 1 "nvvm.grid_constant" %gc_arg1, i32 noundef %arg2, ptr noundef byval(i32) align 4 "nvvm.grid_constant" %gc_arg3) // CHECK: define ptx_kernel void @_Z13tkernel_constIK1SEvT_(ptr noundef byval(%struct.S) align 1 "nvvm.grid_constant" %arg) // CHECK: define ptx_kernel void @_Z13tkernel_constI1SEvT_(ptr noundef byval(%struct.S) align 1 "nvvm.grid_constant" %arg) // CHECK: define ptx_kernel void @_Z7tkernelIK1SEviT_(i32 noundef %dummy, ptr noundef byval(%struct.S) align 1 "nvvm.grid_constant" %arg) - _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
