https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/196920
>From 841d4412c187d27f9c66707f50bcb1614c04aeb9 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Mon, 11 May 2026 12:34:23 +0200 Subject: [PATCH 1/4] [CIR][SPIR-V] Add initial SPIR-V target CodeGen support Connect SPIR/SPIR-V triples in CIRGenModule and the TargetLowering pipeline, set proper `spir_kernel` calling convention on OpenCL kernel --- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 7 +++ clang/lib/CIR/CodeGen/CMakeLists.txt | 1 + clang/lib/CIR/CodeGen/TargetInfo.h | 2 + clang/lib/CIR/CodeGen/Targets/SPIRV.cpp | 58 +++++++++++++++++++ .../Transforms/TargetLowering/CMakeLists.txt | 1 + .../Transforms/TargetLowering/LowerModule.cpp | 6 ++ .../TargetLowering/TargetLoweringInfo.h | 2 + .../TargetLowering/Targets/SPIRV.cpp | 43 ++++++++++++++ clang/test/CIR/CodeGenOpenCL/spirv-kernel.cl | 27 +++++++++ 9 files changed, 147 insertions(+) create mode 100644 clang/lib/CIR/CodeGen/Targets/SPIRV.cpp create mode 100644 clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp create mode 100644 clang/test/CIR/CodeGenOpenCL/spirv-kernel.cl diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index 04e413aa916ec..53c9e068da06b 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -312,6 +312,13 @@ const TargetCIRGenInfo &CIRGenModule::getTargetCIRGenInfo() { theTargetCIRGenInfo = createAMDGPUTargetCIRGenInfo(genTypes); return *theTargetCIRGenInfo; } + case llvm::Triple::spir: + case llvm::Triple::spir64: + case llvm::Triple::spirv: + case llvm::Triple::spirv32: + case llvm::Triple::spirv64: + theTargetCIRGenInfo = createSPIRVTargetCIRGenInfo(genTypes); + return *theTargetCIRGenInfo; } } diff --git a/clang/lib/CIR/CodeGen/CMakeLists.txt b/clang/lib/CIR/CodeGen/CMakeLists.txt index 7320afe626bdc..f3f523ab3c21f 100644 --- a/clang/lib/CIR/CodeGen/CMakeLists.txt +++ b/clang/lib/CIR/CodeGen/CMakeLists.txt @@ -55,6 +55,7 @@ add_clang_library(clangCIR CIRGenVTables.cpp TargetInfo.cpp Targets/AMDGPU.cpp + Targets/SPIRV.cpp DEPENDS MLIRCIR diff --git a/clang/lib/CIR/CodeGen/TargetInfo.h b/clang/lib/CIR/CodeGen/TargetInfo.h index ecdfb7cb42c0e..5c5659056d696 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.h +++ b/clang/lib/CIR/CodeGen/TargetInfo.h @@ -151,6 +151,8 @@ std::unique_ptr<TargetCIRGenInfo> createX8664TargetCIRGenInfo(CIRGenTypes &cgt); std::unique_ptr<TargetCIRGenInfo> createNVPTXTargetCIRGenInfo(CIRGenTypes &cgt); +std::unique_ptr<TargetCIRGenInfo> createSPIRVTargetCIRGenInfo(CIRGenTypes &cgt); + } // namespace clang::CIRGen #endif // LLVM_CLANG_LIB_CIR_TARGETINFO_H diff --git a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp new file mode 100644 index 0000000000000..84094df5850b9 --- /dev/null +++ b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp @@ -0,0 +1,58 @@ +//===---- SPIRV.cpp - SPIR/SPIR-V-specific CIR CodeGen --------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This provides SPIR/SPIR-V-specific CIR CodeGen logic for function attributes. +// +//===----------------------------------------------------------------------===// + +#include "../CIRGenModule.h" +#include "../TargetInfo.h" + +#include "clang/AST/Attr.h" +#include "clang/AST/Decl.h" +#include "clang/CIR/Dialect/IR/CIRDialect.h" + +using namespace clang; +using namespace clang::CIRGen; + +namespace { + +class SPIRVABIInfo : public ABIInfo { +public: + SPIRVABIInfo(CIRGenTypes &cgt) : ABIInfo(cgt) {} +}; + +class SPIRVTargetCIRGenInfo : public TargetCIRGenInfo { +public: + 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::dyn_cast<cir::CIRGlobalValueInterface>(global); + if (globalValue && globalValue.isDeclaration()) + return; + + const auto *fd = dyn_cast_or_null<FunctionDecl>(decl); + if (!fd) + return; + + if (cgm.getLangOpts().OpenCL && fd->hasAttr<DeviceKernelAttr>()) { + auto func = mlir::cast<cir::FuncOp>(global); + func.setCallingConv(cir::CallingConv::SpirKernel); + func.setInlineKind(cir::InlineKind::NoInline); + } + } +}; + +} // namespace + +std::unique_ptr<TargetCIRGenInfo> +clang::CIRGen::createSPIRVTargetCIRGenInfo(CIRGenTypes &cgt) { + return std::make_unique<SPIRVTargetCIRGenInfo>(cgt); +} diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CMakeLists.txt b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CMakeLists.txt index 86502b7f5dd4e..d86fb9eb2aa05 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CMakeLists.txt +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CMakeLists.txt @@ -5,6 +5,7 @@ add_clang_library(MLIRCIRTargetLowering TargetLoweringInfo.cpp Targets/AMDGPU.cpp Targets/NVPTX.cpp + Targets/SPIRV.cpp DEPENDS clangBasic diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp index 6b6eec473ec89..69f1c66f6f119 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp @@ -53,6 +53,12 @@ createTargetLoweringInfo(LowerModule &lm) { case llvm::Triple::nvptx: case llvm::Triple::nvptx64: return createNVPTXTargetLoweringInfo(); + case llvm::Triple::spir: + case llvm::Triple::spir64: + case llvm::Triple::spirv: + case llvm::Triple::spirv32: + case llvm::Triple::spirv64: + return createSPIRVTargetLoweringInfo(); default: assert(!cir::MissingFeatures::targetLoweringInfo()); return std::make_unique<TargetLoweringInfo>(); diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h index 2f778d8302f02..ecba391c0016f 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h @@ -38,6 +38,8 @@ std::unique_ptr<TargetLoweringInfo> createAMDGPUTargetLoweringInfo(); std::unique_ptr<TargetLoweringInfo> createNVPTXTargetLoweringInfo(); +std::unique_ptr<TargetLoweringInfo> createSPIRVTargetLoweringInfo(); + } // namespace cir #endif diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp new file mode 100644 index 0000000000000..b759acccd1ac6 --- /dev/null +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp @@ -0,0 +1,43 @@ +//===- SPIRV.cpp - Emit CIR for SPIR/SPIR-V -------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "../TargetLoweringInfo.h" +#include "clang/CIR/Dialect/IR/CIROpsEnums.h" + +namespace cir { + +namespace { + +// SPIR-V OpenCL storage classes, indexed by cir::LangAddressSpace. +constexpr unsigned SPIRVAddrSpaceMap[] = { + 0, // Function + 0, // Function + 3, // Workgroup + 1, // CrossWorkgroup + 2, // UniformConstant + 4, // Generic +}; + +class SPIRVTargetLoweringInfo : public TargetLoweringInfo { +public: + unsigned getTargetAddrSpaceFromCIRAddrSpace( + cir::LangAddressSpace addrSpace) const override { + auto idx = static_cast<unsigned>(addrSpace); + assert(idx < std::size(SPIRVAddrSpaceMap) && + "Unknown CIR address space for SPIR-V target"); + return SPIRVAddrSpaceMap[idx]; + } +}; + +} // namespace + +std::unique_ptr<TargetLoweringInfo> createSPIRVTargetLoweringInfo() { + return std::make_unique<SPIRVTargetLoweringInfo>(); +} + +} // namespace cir diff --git a/clang/test/CIR/CodeGenOpenCL/spirv-kernel.cl b/clang/test/CIR/CodeGenOpenCL/spirv-kernel.cl new file mode 100644 index 0000000000000..0357d52b079ab --- /dev/null +++ b/clang/test/CIR/CodeGenOpenCL/spirv-kernel.cl @@ -0,0 +1,27 @@ +// RUN: %clang_cc1 %s -fclangir -emit-cir -triple spirv64-unknown-unknown -o %t.cir +// RUN: FileCheck %s --input-file=%t.cir --check-prefix=CIR + +// RUN: %clang_cc1 %s -fclangir -emit-llvm -triple spirv64-unknown-unknown -o %t.ll +// RUN: FileCheck %s --input-file=%t.ll --check-prefix=LLVM + +// RUN: %clang_cc1 %s -emit-llvm -triple spirv64-unknown-unknown -o %t.ll +// RUN: FileCheck %s --input-file=%t.ll --check-prefix=OGCG + +__kernel void inc(__global int *x) { *x = *x + 1; } + +__kernel void add(__global int *a, __global int *b, __global int *out) { + *out = *a + *b; +} + +void plain(int x) {} + +// CIR: cir.func {{.*}} @inc{{.*}} cc(spir_kernel) +// CIR: cir.func {{.*}} @add{{.*}} cc(spir_kernel) +// CIR: cir.func{{.*}} @plain +// CIR-NOT: cc(spir_kernel) + +// LLVM: define spir_kernel void @inc +// LLVM: define spir_kernel void @add + +// OGCG: define spir_kernel void @inc +// OGCG: define spir_kernel void @add >From 484ed42f65ab84599380d2744b92a8a5f40f5c85 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Mon, 18 May 2026 07:04:12 +0200 Subject: [PATCH 2/4] Address review comments --- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 2 +- clang/lib/CIR/CodeGen/TargetInfo.h | 2 +- clang/lib/CIR/CodeGen/Targets/SPIRV.cpp | 22 ++++++++++---------- clang/test/CIR/CodeGenOpenCL/spirv-kernel.cl | 4 ++-- 4 files changed, 15 insertions(+), 15 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index 53c9e068da06b..f1e664390561c 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -317,7 +317,7 @@ const TargetCIRGenInfo &CIRGenModule::getTargetCIRGenInfo() { case llvm::Triple::spirv: case llvm::Triple::spirv32: case llvm::Triple::spirv64: - theTargetCIRGenInfo = createSPIRVTargetCIRGenInfo(genTypes); + theTargetCIRGenInfo = createSPIRTargetCIRGenInfo(genTypes); return *theTargetCIRGenInfo; } } diff --git a/clang/lib/CIR/CodeGen/TargetInfo.h b/clang/lib/CIR/CodeGen/TargetInfo.h index 5c5659056d696..0d1a7fca12cf2 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.h +++ b/clang/lib/CIR/CodeGen/TargetInfo.h @@ -151,7 +151,7 @@ std::unique_ptr<TargetCIRGenInfo> createX8664TargetCIRGenInfo(CIRGenTypes &cgt); std::unique_ptr<TargetCIRGenInfo> createNVPTXTargetCIRGenInfo(CIRGenTypes &cgt); -std::unique_ptr<TargetCIRGenInfo> createSPIRVTargetCIRGenInfo(CIRGenTypes &cgt); +std::unique_ptr<TargetCIRGenInfo> createSPIRTargetCIRGenInfo(CIRGenTypes &cgt); } // namespace clang::CIRGen diff --git a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp index 84094df5850b9..d36a9ee6c9c46 100644 --- a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp +++ b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp @@ -22,30 +22,30 @@ using namespace clang::CIRGen; namespace { -class SPIRVABIInfo : public ABIInfo { +class SPIRABIInfo : public ABIInfo { public: - SPIRVABIInfo(CIRGenTypes &cgt) : ABIInfo(cgt) {} + SPIRABIInfo(CIRGenTypes &cgt) : ABIInfo(cgt) {} }; -class SPIRVTargetCIRGenInfo : public TargetCIRGenInfo { +class SPIRTargetCIRGenInfo : public TargetCIRGenInfo { public: - SPIRVTargetCIRGenInfo(CIRGenTypes &cgt) - : TargetCIRGenInfo(std::make_unique<SPIRVABIInfo>(cgt)) {} + SPIRTargetCIRGenInfo(CIRGenTypes &cgt) + : TargetCIRGenInfo(std::make_unique<SPIRABIInfo>(cgt)) {} void setTargetAttributes(const clang::Decl *decl, mlir::Operation *global, CIRGenModule &cgm) const override { - auto globalValue = mlir::dyn_cast<cir::CIRGlobalValueInterface>(global); - if (globalValue && globalValue.isDeclaration()) + auto globalValue = mlir::cast<cir::CIRGlobalValueInterface>(global); + if (globalValue.isDeclaration()) return; const auto *fd = dyn_cast_or_null<FunctionDecl>(decl); if (!fd) return; - if (cgm.getLangOpts().OpenCL && fd->hasAttr<DeviceKernelAttr>()) { + if (cgm.getLangOpts().OpenCL && + DeviceKernelAttr::isOpenCLSpelling(fd->getAttr<DeviceKernelAttr>())) { auto func = mlir::cast<cir::FuncOp>(global); func.setCallingConv(cir::CallingConv::SpirKernel); - func.setInlineKind(cir::InlineKind::NoInline); } } }; @@ -53,6 +53,6 @@ class SPIRVTargetCIRGenInfo : public TargetCIRGenInfo { } // namespace std::unique_ptr<TargetCIRGenInfo> -clang::CIRGen::createSPIRVTargetCIRGenInfo(CIRGenTypes &cgt) { - return std::make_unique<SPIRVTargetCIRGenInfo>(cgt); +clang::CIRGen::createSPIRTargetCIRGenInfo(CIRGenTypes &cgt) { + return std::make_unique<SPIRTargetCIRGenInfo>(cgt); } diff --git a/clang/test/CIR/CodeGenOpenCL/spirv-kernel.cl b/clang/test/CIR/CodeGenOpenCL/spirv-kernel.cl index 0357d52b079ab..7e52ae832f041 100644 --- a/clang/test/CIR/CodeGenOpenCL/spirv-kernel.cl +++ b/clang/test/CIR/CodeGenOpenCL/spirv-kernel.cl @@ -15,8 +15,8 @@ __kernel void add(__global int *a, __global int *b, __global int *out) { void plain(int x) {} -// CIR: cir.func {{.*}} @inc{{.*}} cc(spir_kernel) -// CIR: cir.func {{.*}} @add{{.*}} cc(spir_kernel) +// CIR: cir.func{{.*}} @inc{{.*}} cc(spir_kernel) +// CIR: cir.func{{.*}} @add{{.*}} cc(spir_kernel) // CIR: cir.func{{.*}} @plain // CIR-NOT: cc(spir_kernel) >From 23cfc24ba7270d24dae95674619151121f70a4fc Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Mon, 18 May 2026 17:17:15 +0200 Subject: [PATCH 3/4] Address review comments --- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 4 +--- clang/lib/CIR/CodeGen/TargetInfo.h | 2 +- clang/lib/CIR/CodeGen/Targets/SPIRV.cpp | 18 +++++++++--------- 3 files changed, 11 insertions(+), 13 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index f1e664390561c..2dc8af8b3f295 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -312,12 +312,10 @@ const TargetCIRGenInfo &CIRGenModule::getTargetCIRGenInfo() { theTargetCIRGenInfo = createAMDGPUTargetCIRGenInfo(genTypes); return *theTargetCIRGenInfo; } - case llvm::Triple::spir: - case llvm::Triple::spir64: case llvm::Triple::spirv: case llvm::Triple::spirv32: case llvm::Triple::spirv64: - theTargetCIRGenInfo = createSPIRTargetCIRGenInfo(genTypes); + theTargetCIRGenInfo = createSPIRVTargetCIRGenInfo(genTypes); return *theTargetCIRGenInfo; } } diff --git a/clang/lib/CIR/CodeGen/TargetInfo.h b/clang/lib/CIR/CodeGen/TargetInfo.h index 0d1a7fca12cf2..5c5659056d696 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.h +++ b/clang/lib/CIR/CodeGen/TargetInfo.h @@ -151,7 +151,7 @@ std::unique_ptr<TargetCIRGenInfo> createX8664TargetCIRGenInfo(CIRGenTypes &cgt); std::unique_ptr<TargetCIRGenInfo> createNVPTXTargetCIRGenInfo(CIRGenTypes &cgt); -std::unique_ptr<TargetCIRGenInfo> createSPIRTargetCIRGenInfo(CIRGenTypes &cgt); +std::unique_ptr<TargetCIRGenInfo> createSPIRVTargetCIRGenInfo(CIRGenTypes &cgt); } // namespace clang::CIRGen diff --git a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp index d36a9ee6c9c46..643c635128d09 100644 --- a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp +++ b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp @@ -1,4 +1,4 @@ -//===---- SPIRV.cpp - SPIR/SPIR-V-specific CIR CodeGen --------------------===// +//===---- SPIRV.cpp - SPIR-V-specific CIR CodeGen -------------------------===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// // -// This provides SPIR/SPIR-V-specific CIR CodeGen logic for function attributes. +// This provides SPIR-V-specific CIR CodeGen logic for function attributes. // //===----------------------------------------------------------------------===// @@ -22,15 +22,15 @@ using namespace clang::CIRGen; namespace { -class SPIRABIInfo : public ABIInfo { +class SPIRVABIInfo : public ABIInfo { public: - SPIRABIInfo(CIRGenTypes &cgt) : ABIInfo(cgt) {} + SPIRVABIInfo(CIRGenTypes &cgt) : ABIInfo(cgt) {} }; -class SPIRTargetCIRGenInfo : public TargetCIRGenInfo { +class SPIRVTargetCIRGenInfo : public TargetCIRGenInfo { public: - SPIRTargetCIRGenInfo(CIRGenTypes &cgt) - : TargetCIRGenInfo(std::make_unique<SPIRABIInfo>(cgt)) {} + SPIRVTargetCIRGenInfo(CIRGenTypes &cgt) + : TargetCIRGenInfo(std::make_unique<SPIRVABIInfo>(cgt)) {} void setTargetAttributes(const clang::Decl *decl, mlir::Operation *global, CIRGenModule &cgm) const override { @@ -53,6 +53,6 @@ class SPIRTargetCIRGenInfo : public TargetCIRGenInfo { } // namespace std::unique_ptr<TargetCIRGenInfo> -clang::CIRGen::createSPIRTargetCIRGenInfo(CIRGenTypes &cgt) { - return std::make_unique<SPIRTargetCIRGenInfo>(cgt); +clang::CIRGen::createSPIRVTargetCIRGenInfo(CIRGenTypes &cgt) { + return std::make_unique<SPIRVTargetCIRGenInfo>(cgt); } >From 53bfcbf8d99c09b4cd1520a365a9d793b25d7295 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Mon, 18 May 2026 17:37:40 +0200 Subject: [PATCH 4/4] cleanup --- clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp index 69f1c66f6f119..cccbfdde3bfdf 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp @@ -53,8 +53,6 @@ createTargetLoweringInfo(LowerModule &lm) { case llvm::Triple::nvptx: case llvm::Triple::nvptx64: return createNVPTXTargetLoweringInfo(); - case llvm::Triple::spir: - case llvm::Triple::spir64: case llvm::Triple::spirv: case llvm::Triple::spirv32: case llvm::Triple::spirv64: _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
