https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/214246
>From e0289b4451833784512fa7f6f3702ea10802e536 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Wed, 5 Aug 2026 17:05:13 +0200 Subject: [PATCH 1/5] [CIR][SPIR-V] Set spir_kernel calling convention for AMDGCN-flavored SPIR-V HIP kernels --- clang/lib/CIR/CodeGen/Targets/SPIRV.cpp | 9 +++++++- .../CIR/CodeGenHIP/amdgcnspirv-kernel.hip | 23 +++++++++++++++++++ 2 files changed, 31 insertions(+), 1 deletion(-) create mode 100644 clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip diff --git a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp index 643c635128d09..b56bbfe41e87c 100644 --- a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp +++ b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp @@ -42,11 +42,18 @@ class SPIRVTargetCIRGenInfo : public TargetCIRGenInfo { if (!fd) return; + auto func = mlir::cast<cir::FuncOp>(global); + if (cgm.getLangOpts().OpenCL && DeviceKernelAttr::isOpenCLSpelling(fd->getAttr<DeviceKernelAttr>())) { - auto func = mlir::cast<cir::FuncOp>(global); func.setCallingConv(cir::CallingConv::SpirKernel); + return; } + + if (cgm.getLangOpts().HIP && + cgm.getTriple().getVendor() == llvm::Triple::AMD && + fd->hasAttr<CUDAGlobalAttr>()) + func.setCallingConv(cir::CallingConv::SpirKernel); } }; diff --git a/clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip b/clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip new file mode 100644 index 0000000000000..039ab35f1c906 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip @@ -0,0 +1,23 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fclangir \ +// RUN: -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR %s --input-file=%t.cir + +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fclangir \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.ll + +// Test that HIP kernels on AMDGCN-flavored SPIR-V get the spir_kernel +// calling convention. + +#define __global__ __attribute__((global)) +#define __device__ __attribute__((device)) + +// CIR: cir.func{{.*}} @_Z13kernel_scalari{{.*}} cc(spir_kernel) +// LLVM: define spir_kernel void @_Z13kernel_scalari +__global__ void kernel_scalar(int a) {} + +// CIR: cir.func{{.*}} @_Z9device_fni +// CIR-NOT: cc(spir_kernel) +// LLVM: define{{.*}} void @_Z9device_fni +__device__ void device_fn(int a) {} >From 5a17e67fd7786430df8330a9b1c36535984a3ec2 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Wed, 5 Aug 2026 17:42:59 +0200 Subject: [PATCH 2/5] fix integerattr The program address space DLTI entry is stored as an unsigned integer attribute, but was read with getInt(), which asserts on signless types. --- clang/lib/CIR/Dialect/IR/CIRDataLayout.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/CIR/Dialect/IR/CIRDataLayout.cpp b/clang/lib/CIR/Dialect/IR/CIRDataLayout.cpp index d7d563d79a0fa..5ce0c8f8e7a2c 100644 --- a/clang/lib/CIR/Dialect/IR/CIRDataLayout.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRDataLayout.cpp @@ -26,7 +26,7 @@ void CIRDataLayout::reset(mlir::DataLayoutSpecInterface spec) { if (mlir::DataLayoutEntryInterface entry = spec.getSpecForIdentifier(addrSpKey)) if (auto val = llvm::dyn_cast<mlir::IntegerAttr>(entry.getValue())) - programAddrSpace = val.getInt(); + programAddrSpace = val.getUInt(); } } >From 93eb6b0d66627648bc57bbee5e9d49c58e5cc2da Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Mon, 24 Aug 2026 12:43:49 +0200 Subject: [PATCH 3/5] rm AMD triple check --- clang/lib/CIR/CodeGen/Targets/SPIRV.cpp | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp index b56bbfe41e87c..8a19fa1f8a9c0 100644 --- a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp +++ b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp @@ -50,9 +50,7 @@ class SPIRVTargetCIRGenInfo : public TargetCIRGenInfo { return; } - if (cgm.getLangOpts().HIP && - cgm.getTriple().getVendor() == llvm::Triple::AMD && - fd->hasAttr<CUDAGlobalAttr>()) + if (cgm.getLangOpts().HIP && fd->hasAttr<CUDAGlobalAttr>()) func.setCallingConv(cir::CallingConv::SpirKernel); } }; >From d22d3b3b8b0487ac731153ddda2dc2df2485c4a4 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Wed, 26 Aug 2026 08:36:39 +0200 Subject: [PATCH 4/5] Align CI and classic codegen --- clang/lib/CIR/CodeGen/CIRGenCall.cpp | 32 ++++++++++++++++------ clang/lib/CIR/CodeGen/CIRGenFunctionInfo.h | 21 ++++++++++++-- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 6 +--- clang/lib/CIR/CodeGen/CIRGenTypes.cpp | 19 +++++++++++-- clang/lib/CIR/CodeGen/CIRGenTypes.h | 3 ++ clang/lib/CIR/CodeGen/TargetInfo.cpp | 8 ++++++ clang/lib/CIR/CodeGen/TargetInfo.h | 6 ++++ clang/lib/CIR/CodeGen/Targets/SPIRV.cpp | 29 ++++++-------------- 8 files changed, 85 insertions(+), 39 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenCall.cpp b/clang/lib/CIR/CodeGen/CIRGenCall.cpp index 28670cf31694a..bf94d0fcbeb24 100644 --- a/clang/lib/CIR/CodeGen/CIRGenCall.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenCall.cpp @@ -26,9 +26,11 @@ using namespace clang; using namespace clang::CIRGen; -CIRGenFunctionInfo *CIRGenFunctionInfo::create( - FunctionType::ExtInfo info, bool isInstanceMethod, CanQualType resultType, - llvm::ArrayRef<CanQualType> argTypes, RequiredArgs required) { +CIRGenFunctionInfo * +CIRGenFunctionInfo::create(cir::CallingConv cirCC, FunctionType::ExtInfo info, + bool isInstanceMethod, CanQualType resultType, + llvm::ArrayRef<CanQualType> argTypes, + RequiredArgs required) { // The first slot allocated for arg type slot is for the return value. void *buffer = operator new( totalSizeToAlloc<CanQualType>(argTypes.size() + 1)); @@ -37,6 +39,8 @@ CIRGenFunctionInfo *CIRGenFunctionInfo::create( CIRGenFunctionInfo *fi = new (buffer) CIRGenFunctionInfo(); + fi->callingConvention = llvm::to_underlying(cirCC); + fi->astCallingConvention = info.getCC(); fi->noReturn = info.getNoReturn(); fi->instanceMethod = isInstanceMethod; @@ -311,7 +315,7 @@ void CIRGenModule::constructAttributeList( llvm::MutableArrayRef<mlir::NamedAttrList> argAttrs, mlir::NamedAttrList &retAttrs, cir::CallingConv &callingConv, cir::SideEffect &sideEffect, bool attrOnCallSite, bool isThunk) { - assert(!cir::MissingFeatures::opCallCallConv()); + callingConv = info.getCallingConvention(); sideEffect = cir::SideEffect::All; auto addUnitAttr = [&](llvm::StringRef name) { @@ -1028,6 +1032,17 @@ CIRGenTypes::arrangeBuiltinFunctionCall(QualType resultType, FunctionType::ExtInfo(), RequiredArgs::All); } +/// Set calling convention for CUDA/HIP kernel. +static void setCUDAKernelCallingConvention(CanQualType &funcTy, + CIRGenModule &cgm, + const FunctionDecl *fd) { + if (fd->hasAttr<CUDAGlobalAttr>()) { + const FunctionType *ft = funcTy->getAs<FunctionType>(); + cgm.getTargetCIRGenInfo().setCUDAKernelCallingConvention(ft); + funcTy = 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 constructor or destructor. @@ -1036,9 +1051,9 @@ CIRGenTypes::arrangeCXXMethodDeclaration(const CXXMethodDecl *md) { assert(!isa<CXXConstructorDecl>(md) && "wrong method for constructors!"); assert(!isa<CXXDestructorDecl>(md) && "wrong method for destructors!"); - auto prototype = - md->getType()->getCanonicalTypeUnqualified().getAs<FunctionProtoType>(); - assert(!cir::MissingFeatures::cudaSupport()); + CanQualType funcTy = md->getType()->getCanonicalTypeUnqualified(); + setCUDAKernelCallingConvention(funcTy, cgm, md); + auto prototype = funcTy.getAs<FunctionProtoType>(); // Mirrors classic CodeGen's check at CGCall.cpp. C++23 explicit-object // member functions (P0847R7, `void f(this Self&&)`) do not receive an @@ -1088,8 +1103,7 @@ CIRGenTypes::arrangeFunctionDeclaration(const FunctionDecl *fd) { CanQualType funcTy = fd->getType()->getCanonicalTypeUnqualified(); assert(isa<FunctionType>(funcTy)); - // TODO: setCUDAKernelCallingConvention - assert(!cir::MissingFeatures::cudaSupport()); + setCUDAKernelCallingConvention(funcTy, cgm, fd); // When declaring a function without a prototype, always use a non-variadic // type. diff --git a/clang/lib/CIR/CodeGen/CIRGenFunctionInfo.h b/clang/lib/CIR/CodeGen/CIRGenFunctionInfo.h index d37a6149bcafa..2957293c83623 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunctionInfo.h +++ b/clang/lib/CIR/CodeGen/CIRGenFunctionInfo.h @@ -17,6 +17,7 @@ #include "clang/AST/CanonicalType.h" #include "clang/CIR/ABIArgInfo.h" +#include "clang/CIR/Dialect/IR/CIROpsEnums.h" #include "clang/CIR/MissingFeatures.h" #include "llvm/ADT/FoldingSet.h" #include "llvm/Support/TrailingObjects.h" @@ -84,6 +85,12 @@ class RequiredArgs { class CIRGenFunctionInfo final : public llvm::FoldingSetNode, private llvm::TrailingObjects<CIRGenFunctionInfo, CanQualType> { + /// The CIR-level calling convention to use for this function. + unsigned callingConvention : 8; + + /// The AST-level calling convention this function was declared with. + unsigned astCallingConvention : 8; + // Whether this function has noreturn. LLVM_PREFERRED_TYPE(bool) unsigned noReturn : 1; @@ -107,13 +114,14 @@ class CIRGenFunctionInfo final // here instead of explicit false/0. return FunctionType::ExtInfo( isNoReturn(), /*getHasRegParm=*/false, /*getRegParm=*/false, - /*getASTCallingConvention=*/CallingConv(0), /*isReturnsRetained=*/false, + getASTCallingConvention(), /*isReturnsRetained=*/false, /*isNoCallerSavedRegs=*/false, /*isNoCfCheck=*/false, /*isCmseNSCall=*/false); } public: - static CIRGenFunctionInfo *create(FunctionType::ExtInfo info, + static CIRGenFunctionInfo *create(cir::CallingConv cirCC, + FunctionType::ExtInfo info, bool instanceMethod, CanQualType resultType, llvm::ArrayRef<CanQualType> argTypes, RequiredArgs required); @@ -133,6 +141,7 @@ class CIRGenFunctionInfo final FunctionType::ExtInfo info, RequiredArgs required, CanQualType resultType, llvm::ArrayRef<CanQualType> argTypes) { + id.AddInteger(info.getCC()); id.AddBoolean(instanceMethod); id.AddBoolean(info.getNoReturn()); id.AddInteger(required.getOpaqueData()); @@ -192,6 +201,14 @@ class CIRGenFunctionInfo final bool isNoReturn() const { return noReturn; } bool isInstanceMethod() const { return instanceMethod; } + + cir::CallingConv getCallingConvention() const { + return static_cast<cir::CallingConv>(callingConvention); + } + + CallingConv getASTCallingConvention() const { + return static_cast<CallingConv>(astCallingConvention); + } }; } // namespace clang::CIRGen diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index a7143408c9bee..c1b5f1faf16d9 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -3125,11 +3125,7 @@ void CIRGenModule::setCIRFunctionAttributes(GlobalDecl globalDecl, // TODO(cir): Check X86_VectorCall incompatibility wiht WinARM64EC - // TODO(cir): Set the calling convention computed by constructAttributeList - // on the function. FuncOp supports calling_conv, but target-specific - // CodeGen is needed to set it correctly (e.g., AMDGPU kernel functions - // should be marked with AMDGPUKernel). - assert(!cir::MissingFeatures::opFuncCallingConv()); + func.setCallingConv(callingConv); } void CIRGenModule::setFunctionAttributes(GlobalDecl globalDecl, diff --git a/clang/lib/CIR/CodeGen/CIRGenTypes.cpp b/clang/lib/CIR/CodeGen/CIRGenTypes.cpp index e5af4eec7720f..469f052c66780 100644 --- a/clang/lib/CIR/CodeGen/CIRGenTypes.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenTypes.cpp @@ -750,6 +750,19 @@ bool CIRGenTypes::isZeroInitializable(const RecordDecl *rd) { return getCIRGenRecordLayout(rd).isZeroInitializable(); } +cir::CallingConv +CIRGenTypes::clangCallConvToCIRCallConv(clang::CallingConv cc) { + switch (cc) { + default: + // TODO(cir): Support the remaining target-specific calling conventions. + return cir::CallingConv::C; + case CC_SpirFunction: + return cir::CallingConv::SpirFunction; + case CC_DeviceKernel: + return cgm.getTargetCIRGenInfo().getDeviceKernelCallingConv(); + } +} + const CIRGenFunctionInfo &CIRGenTypes::arrangeCIRFunctionInfo( CanQualType returnType, bool isInstanceMethod, llvm::ArrayRef<CanQualType> argTypes, FunctionType::ExtInfo info, @@ -773,11 +786,11 @@ const CIRGenFunctionInfo &CIRGenTypes::arrangeCIRFunctionInfo( return *fi; } - assert(!cir::MissingFeatures::opCallCallConv()); + cir::CallingConv cirCC = clangCallConvToCIRCallConv(info.getCC()); // Construction the function info. We co-allocate the ArgInfos. - fi = CIRGenFunctionInfo::create(info, isInstanceMethod, returnType, argTypes, - required); + fi = CIRGenFunctionInfo::create(cirCC, info, isInstanceMethod, returnType, + argTypes, required); functionInfos.InsertNode(fi, insertPos); return *fi; diff --git a/clang/lib/CIR/CodeGen/CIRGenTypes.h b/clang/lib/CIR/CodeGen/CIRGenTypes.h index a7827f76bd5f2..4c9280876882d 100644 --- a/clang/lib/CIR/CodeGen/CIRGenTypes.h +++ b/clang/lib/CIR/CodeGen/CIRGenTypes.h @@ -219,6 +219,9 @@ class CIRGenTypes { const CIRGenFunctionInfo &arrangeFreeFunctionCall(const CallArgList &args, const FunctionType *fnType); + /// Convert a clang calling convention to a CIR calling convention. + cir::CallingConv clangCallConvToCIRCallConv(clang::CallingConv cc); + const CIRGenFunctionInfo & arrangeCIRFunctionInfo(CanQualType returnType, bool isInstanceMethod, llvm::ArrayRef<CanQualType> argTypes, diff --git a/clang/lib/CIR/CodeGen/TargetInfo.cpp b/clang/lib/CIR/CodeGen/TargetInfo.cpp index ba7eeb29dd252..42906a5f68c12 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.cpp +++ b/clang/lib/CIR/CodeGen/TargetInfo.cpp @@ -143,6 +143,14 @@ bool TargetCIRGenInfo::isNoProtoCallVariadic( return false; } +cir::CallingConv TargetCIRGenInfo::getDeviceKernelCallingConv() const { + // Device kernels are entered through a runtime API, not called as normal + // sub-functions, so a modified C calling convention is used. + assert(getABIInfo().cgt.getASTContext().getLangOpts().OpenCL && + "Kernel calling convention only defined for OpenCL"); + return cir::CallingConv::C; +} + clang::LangAS TargetCIRGenInfo::getGlobalVarAddressSpace(CIRGenModule &cgm, const clang::VarDecl *d) const { diff --git a/clang/lib/CIR/CodeGen/TargetInfo.h b/clang/lib/CIR/CodeGen/TargetInfo.h index 308d472234f99..326a05bf22e86 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.h +++ b/clang/lib/CIR/CodeGen/TargetInfo.h @@ -128,6 +128,12 @@ class TargetCIRGenInfo { return false; } + /// Returns the calling convention used for device kernels on this target. + virtual cir::CallingConv getDeviceKernelCallingConv() const; + + virtual void + setCUDAKernelCallingConvention(const clang::FunctionType *&ft) const {} + /// Corrects the MLIR type for a given constraint and "usual" /// type. /// diff --git a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp index 8a19fa1f8a9c0..598cd693d6d26 100644 --- a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp +++ b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp @@ -32,26 +32,15 @@ class SPIRVTargetCIRGenInfo : public TargetCIRGenInfo { SPIRVTargetCIRGenInfo(CIRGenTypes &cgt) : TargetCIRGenInfo(std::make_unique<SPIRVABIInfo>(cgt)) {} - void setTargetAttributes(const clang::Decl *decl, mlir::Operation *global, - CIRGenModule &cgm) const override { - auto globalValue = mlir::cast<cir::CIRGlobalValueInterface>(global); - if (globalValue.isDeclaration()) - return; - - const auto *fd = dyn_cast_or_null<FunctionDecl>(decl); - if (!fd) - return; - - auto func = mlir::cast<cir::FuncOp>(global); - - if (cgm.getLangOpts().OpenCL && - DeviceKernelAttr::isOpenCLSpelling(fd->getAttr<DeviceKernelAttr>())) { - func.setCallingConv(cir::CallingConv::SpirKernel); - return; - } - - if (cgm.getLangOpts().HIP && fd->hasAttr<CUDAGlobalAttr>()) - func.setCallingConv(cir::CallingConv::SpirKernel); + cir::CallingConv getDeviceKernelCallingConv() const override { + return cir::CallingConv::SpirKernel; + } + + void setCUDAKernelCallingConvention(const FunctionType *&ft) const override { + // Convert HIP kernels to SPIR-V kernels. + if (getABIInfo().cgt.getASTContext().getLangOpts().HIP) + ft = getABIInfo().cgt.getASTContext().adjustFunctionType( + ft, ft->getExtInfo().withCallingConv(CC_DeviceKernel)); } }; >From b8313e894cde477279e7100b14e6654979f6b882 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Wed, 26 Aug 2026 12:01:46 +0200 Subject: [PATCH 5/5] fix build --- clang/lib/CIR/CodeGen/CIRGenTypes.cpp | 11 +++++++---- clang/lib/CIR/CodeGen/TargetInfo.h | 6 ------ clang/lib/CIR/CodeGen/Targets/SPIRV.cpp | 4 ---- 3 files changed, 7 insertions(+), 14 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenTypes.cpp b/clang/lib/CIR/CodeGen/CIRGenTypes.cpp index a4982ff707cf1..f9fa018d70540 100644 --- a/clang/lib/CIR/CodeGen/CIRGenTypes.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenTypes.cpp @@ -756,13 +756,16 @@ bool CIRGenTypes::isZeroInitializable(const RecordDecl *rd) { cir::CallingConv CIRGenTypes::clangCallConvToCIRCallConv(clang::CallingConv cc) { switch (cc) { - default: - // TODO(cir): Support the remaining target-specific calling conventions. + case CC_C: + // SPIR/SPIR-V lowers the default CC to spir_func, not plain C. + if (cgm.getTriple().isSPIROrSPIRV()) + return cir::CallingConv::SpirFunction; return cir::CallingConv::C; - case CC_SpirFunction: - return cir::CallingConv::SpirFunction; case CC_DeviceKernel: return cgm.getTargetCIRGenInfo().getDeviceKernelCallingConv(); + default: + // TODO(cir): Support the remaining target-specific calling conventions. + return cir::CallingConv::C; } } diff --git a/clang/lib/CIR/CodeGen/TargetInfo.h b/clang/lib/CIR/CodeGen/TargetInfo.h index 54cb256bcb774..0e91f4e95a4a9 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.h +++ b/clang/lib/CIR/CodeGen/TargetInfo.h @@ -153,12 +153,6 @@ class TargetCIRGenInfo { mlir::Operation *global, CIRGenModule &module) const {} - /// Get the CIR calling convention to use for a device kernel entry point - /// (e.g. an OpenCL/SYCL or CUDA/HIP kernel) on this target. - virtual cir::CallingConv getDeviceKernelCallingConv() const { - return cir::CallingConv::C; - } - virtual bool isScalarizableAsmOperand(CIRGenFunction &cgf, mlir::Type ty) const { return false; diff --git a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp index e88086242e99e..598cd693d6d26 100644 --- a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp +++ b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp @@ -42,10 +42,6 @@ class SPIRVTargetCIRGenInfo : public TargetCIRGenInfo { ft = getABIInfo().cgt.getASTContext().adjustFunctionType( ft, ft->getExtInfo().withCallingConv(CC_DeviceKernel)); } - - cir::CallingConv getDeviceKernelCallingConv() const override { - return cir::CallingConv::SpirKernel; - } }; } // namespace _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
