https://github.com/RiverDave updated https://github.com/llvm/llvm-project/pull/186562
>From 51fc1ffd97eecc57b933d6a5f28f8c6b606f2d62 Mon Sep 17 00:00:00 2001 From: David Rivera <[email protected]> Date: Sat, 14 Mar 2026 01:43:49 -0400 Subject: [PATCH 1/3] [CIR][NVPTX] NVPTX lowering info skeleton and target AS mapping --- .../Transforms/TargetLowering/CMakeLists.txt | 1 + .../Transforms/TargetLowering/LowerModule.cpp | 3 ++ .../TargetLowering/TargetLoweringInfo.h | 2 + .../TargetLowering/Targets/NVPTX.cpp | 39 +++++++++++++++ clang/test/CIR/CodeGenCUDA/address-spaces.cu | 47 ++++++++++++------- 5 files changed, 75 insertions(+), 17 deletions(-) create mode 100644 clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CMakeLists.txt b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CMakeLists.txt index 07e3a67f97859..86502b7f5dd4e 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CMakeLists.txt +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CMakeLists.txt @@ -4,6 +4,7 @@ add_clang_library(MLIRCIRTargetLowering LowerItaniumCXXABI.cpp TargetLoweringInfo.cpp Targets/AMDGPU.cpp + Targets/NVPTX.cpp DEPENDS clangBasic diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp index 26e63b3b676ae..6b6eec473ec89 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp @@ -50,6 +50,9 @@ createTargetLoweringInfo(LowerModule &lm) { switch (triple.getArch()) { case llvm::Triple::amdgcn: return createAMDGPUTargetLoweringInfo(); + case llvm::Triple::nvptx: + case llvm::Triple::nvptx64: + return createNVPTXTargetLoweringInfo(); 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 a307bcb373dec..2f778d8302f02 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h @@ -36,6 +36,8 @@ class TargetLoweringInfo { // Target-specific factory functions. std::unique_ptr<TargetLoweringInfo> createAMDGPUTargetLoweringInfo(); +std::unique_ptr<TargetLoweringInfo> createNVPTXTargetLoweringInfo(); + } // namespace cir #endif diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp new file mode 100644 index 0000000000000..f38d2b8bfa32d --- /dev/null +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp @@ -0,0 +1,39 @@ +//===- NVPTX.cpp ----------------------------------------------------------===// +// +// 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" +#include "llvm/Support/NVPTXAddrSpace.h" + +namespace cir { + +namespace { + +constexpr unsigned NVPTXAddrSpaceMap[] = { + llvm::NVPTXAS::ADDRESS_SPACE_GENERIC, llvm::NVPTXAS::ADDRESS_SPACE_GENERIC, + llvm::NVPTXAS::ADDRESS_SPACE_SHARED, llvm::NVPTXAS::ADDRESS_SPACE_GLOBAL, + llvm::NVPTXAS::ADDRESS_SPACE_CONST, llvm::NVPTXAS::ADDRESS_SPACE_GENERIC, +}; + +class NVPTXTargetLoweringInfo : public TargetLoweringInfo { +public: + unsigned getTargetAddrSpaceFromCIRAddrSpace( + cir::LangAddressSpace addrSpace) const override { + + auto idx = static_cast<unsigned>(addrSpace); + assert(idx < std::size(NVPTXAddrSpaceMap) && + "Unknown CIR address space for NVPTX target"); + return NVPTXAddrSpaceMap[idx]; + } +}; + +} // namespace + +std::unique_ptr<TargetLoweringInfo> createNVPTXTargetLoweringInfo() { + return std::make_unique<NVPTXTargetLoweringInfo>(); +} +} // namespace cir diff --git a/clang/test/CIR/CodeGenCUDA/address-spaces.cu b/clang/test/CIR/CodeGenCUDA/address-spaces.cu index a47a2867e7111..0053b888a19af 100644 --- a/clang/test/CIR/CodeGenCUDA/address-spaces.cu +++ b/clang/test/CIR/CodeGenCUDA/address-spaces.cu @@ -10,8 +10,9 @@ // RUN: -mmlir -mlir-print-ir-before=cir-target-lowering %s -o %t.cir 2> %t-pre.cir // RUN: FileCheck --check-prefix=CIR-PRE --input-file=%t-pre.cir %s -// TODO: Add CIR (post target lowering) and LLVM checks once NVPTX TargetLoweringInfo -// is implemented. +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda \ +// RUN: -fcuda-is-device -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR-POST --input-file=%t.cir %s // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir \ // RUN: -fcuda-is-device -emit-llvm -target-sdk-version=12.3 \ @@ -32,26 +33,38 @@ // CIR-DEVICE: cir.global "private" internal dso_local @_ZZ2fnvE1j = #cir.undef : !s32i {alignment = 4 : i64} // LLVM-DEVICE: @_ZZ2fnvE1j = internal global i32 undef, align 4 -__device__ int a; -// CIR-PRE: cir.global external lang_address_space(offload_global) @a = #cir.int<0> : !s32i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized} -// LLVM-DEVICE: @[[DEV_LD:.*]] = externally_initialized global i32 0, align 4 -// OGCG-DAG: @a = addrspace(1) externally_initialized global i32 0, align 4 -// OGCG-DEVICE: @[[DEV_OD:.*]] = addrspace(1) externally_initialized global i32 0, align 4 +// CIR-PRE: cir.global external lang_address_space(offload_global) @i = #cir.int<0> : !s32i +// CIR-POST: cir.global external target_address_space(1) @i = #cir.int<0> : !s32i +// OGCG-DAG: @i = addrspace(1) externally_initialized global i32 0, align 4 +__device__ int i; + +// CIR-PRE: cir.global constant external lang_address_space(offload_constant) @j = #cir.int<0> : !s32i +// CIR-POST: cir.global constant external target_address_space(4) @j = #cir.int<0> : !s32i +// OGCG-DAG: @j = addrspace(4) externally_initialized constant i32 0, align 4 +__constant__ int j; -__constant__ int c; -// CIR-PRE: cir.global constant external lang_address_space(offload_constant) @c = #cir.int<0> : !s32i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized} -// LLVM-DEVICE: @[[CONST_LL:.*]] = externally_initialized constant i32 0, align 4 -// OGCG-DAG: @c = addrspace(4) externally_initialized constant i32 0, align 4 -// OGCG-DEVICE: @[[CONST_OD:.*]] = addrspace(4) externally_initialized constant i32 0, align 4 +// CIR-PRE: cir.global external lang_address_space(offload_local) @k = #cir.poison : !s32i +// CIR-POST: cir.global external target_address_space(3) @k = #cir.poison : !s32i +// OGCG-DAG: @k = addrspace(3) global i32 undef, align 4 +__shared__ int k; -// OGCG-DEVICE: @_ZZ2fnvE1j = internal addrspace(3) global i32 undef, align 4 +// CIR-PRE: cir.global external lang_address_space(offload_local) @b = #cir.poison : !cir.float +// CIR-POST: cir.global external target_address_space(3) @b = #cir.poison : !cir.float +// OGCG-DAG: @b = addrspace(3) global float undef, align 4 +__shared__ float b; __device__ void foo() { - // CIR-PRE: cir.get_global @a : !cir.ptr<!s32i, lang_address_space(offload_global)> - a++; + // CIR-PRE: cir.get_global @i : !cir.ptr<!s32i, lang_address_space(offload_global)> + // CIR-POST: cir.get_global @i : !cir.ptr<!s32i, target_address_space(1)> + i++; + + // CIR-PRE: cir.get_global @j : !cir.ptr<!s32i, lang_address_space(offload_constant)> + // CIR-POST: cir.get_global @j : !cir.ptr<!s32i, target_address_space(4)> + j++; - // CIR-PRE: cir.get_global @c : !cir.ptr<!s32i, lang_address_space(offload_constant)> - c++; + // CIR-PRE: cir.get_global @k : !cir.ptr<!s32i, lang_address_space(offload_local)> + // CIR-POST: cir.get_global @k : !cir.ptr<!s32i, target_address_space(3)> + k++; } __global__ void fn() { >From d1cee04031cd2c6224a727a13b9334bee2eb62f9 Mon Sep 17 00:00:00 2001 From: David Rivera <[email protected]> Date: Thu, 2 Apr 2026 03:50:46 -0400 Subject: [PATCH 2/3] Poison attr lowering and llvm `__shared__` lowering. --- .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 18 +++++++++++++----- clang/test/CIR/CodeGenCUDA/address-spaces.cu | 4 ++++ 2 files changed, 17 insertions(+), 5 deletions(-) diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index b622fa1ef3205..b8b6e20e6fad7 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -391,7 +391,7 @@ class CIRAttrToValue { .Case<cir::BoolAttr, cir::IntAttr, cir::FPAttr, cir::ConstComplexAttr, cir::ConstArrayAttr, cir::ConstRecordAttr, cir::ConstVectorAttr, cir::ConstPtrAttr, cir::GlobalViewAttr, cir::TypeInfoAttr, - cir::UndefAttr, cir::VTableAttr, cir::ZeroAttr>( + cir::UndefAttr, cir::PoisonAttr, cir::VTableAttr, cir::ZeroAttr>( [&](auto attrT) { return visitCirAttr(attrT); }) .Default([&](auto attrT) { return mlir::Value(); }); } @@ -407,6 +407,7 @@ class CIRAttrToValue { mlir::Value visitCirAttr(cir::GlobalViewAttr attr); mlir::Value visitCirAttr(cir::TypeInfoAttr attr); mlir::Value visitCirAttr(cir::UndefAttr attr); + mlir::Value visitCirAttr(cir::PoisonAttr attr); mlir::Value visitCirAttr(cir::VTableAttr attr); mlir::Value visitCirAttr(cir::ZeroAttr attr); @@ -768,6 +769,13 @@ mlir::Value CIRAttrToValue::visitCirAttr(cir::UndefAttr undefAttr) { rewriter, loc, converter->convertType(undefAttr.getType())); } +/// PoisonAttr visitor. +mlir::Value CIRAttrToValue::visitCirAttr(cir::PoisonAttr poisonAttr) { + mlir::Location loc = parentOp->getLoc(); + return mlir::LLVM::PoisonOp::create( + rewriter, loc, converter->convertType(poisonAttr.getType())); +} + // VTableAttr visitor. mlir::Value CIRAttrToValue::visitCirAttr(cir::VTableAttr vtableArr) { mlir::Type llvmTy = converter->convertType(vtableArr.getType()); @@ -2629,8 +2637,8 @@ CIRToLLVMGlobalOpLowering::matchAndRewriteRegionInitializedGlobal( assert( (isa<cir::ConstArrayAttr, cir::ConstRecordAttr, cir::ConstVectorAttr, cir::ConstPtrAttr, cir::ConstComplexAttr, cir::GlobalViewAttr, - cir::TypeInfoAttr, cir::UndefAttr, cir::VTableAttr, cir::ZeroAttr>( - init))); + cir::TypeInfoAttr, cir::UndefAttr, cir::PoisonAttr, cir::VTableAttr, + cir::ZeroAttr>(init))); // TODO(cir): once LLVM's dialect has proper equivalent attributes this // should be updated. For now, we use a custom op to initialize globals @@ -2691,8 +2699,8 @@ mlir::LogicalResult CIRToLLVMGlobalOpLowering::matchAndRewrite( } else if (mlir::isa<cir::ConstArrayAttr, cir::ConstVectorAttr, cir::ConstRecordAttr, cir::ConstPtrAttr, cir::ConstComplexAttr, cir::GlobalViewAttr, - cir::TypeInfoAttr, cir::UndefAttr, cir::VTableAttr, - cir::ZeroAttr>(init.value())) { + cir::TypeInfoAttr, cir::UndefAttr, cir::PoisonAttr, + cir::VTableAttr, cir::ZeroAttr>(init.value())) { // TODO(cir): once LLVM's dialect has proper equivalent attributes this // should be updated. For now, we use a custom op to initialize globals // to the appropriate value. diff --git a/clang/test/CIR/CodeGenCUDA/address-spaces.cu b/clang/test/CIR/CodeGenCUDA/address-spaces.cu index 0053b888a19af..cc1791a8f2244 100644 --- a/clang/test/CIR/CodeGenCUDA/address-spaces.cu +++ b/clang/test/CIR/CodeGenCUDA/address-spaces.cu @@ -35,21 +35,25 @@ // CIR-PRE: cir.global external lang_address_space(offload_global) @i = #cir.int<0> : !s32i // CIR-POST: cir.global external target_address_space(1) @i = #cir.int<0> : !s32i +// LLVM-DEVICE-DAG: @i = addrspace(1) {{.*}}global i32 0, align 4 // OGCG-DAG: @i = addrspace(1) externally_initialized global i32 0, align 4 __device__ int i; // CIR-PRE: cir.global constant external lang_address_space(offload_constant) @j = #cir.int<0> : !s32i // CIR-POST: cir.global constant external target_address_space(4) @j = #cir.int<0> : !s32i +// LLVM-DEVICE-DAG: @j = addrspace(4) {{.*}}constant i32 0, align 4 // OGCG-DAG: @j = addrspace(4) externally_initialized constant i32 0, align 4 __constant__ int j; // CIR-PRE: cir.global external lang_address_space(offload_local) @k = #cir.poison : !s32i // CIR-POST: cir.global external target_address_space(3) @k = #cir.poison : !s32i +// LLVM-DEVICE-DAG: @k = addrspace(3) global i32 {{undef|poison}}, align 4 // OGCG-DAG: @k = addrspace(3) global i32 undef, align 4 __shared__ int k; // CIR-PRE: cir.global external lang_address_space(offload_local) @b = #cir.poison : !cir.float // CIR-POST: cir.global external target_address_space(3) @b = #cir.poison : !cir.float +// LLVM-DEVICE-DAG: @b = addrspace(3) global float {{undef|poison}}, align 4 // OGCG-DAG: @b = addrspace(3) global float undef, align 4 __shared__ float b; >From f95d97373e4c4a03c78b642669a8eaf63759efbe Mon Sep 17 00:00:00 2001 From: David Rivera <[email protected]> Date: Thu, 2 Apr 2026 04:01:41 -0400 Subject: [PATCH 3/3] fix fmt --- clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index b8b6e20e6fad7..c60f1276cf5f0 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -2634,11 +2634,10 @@ CIRToLLVMGlobalOpLowering::matchAndRewriteRegionInitializedGlobal( cir::GlobalOp op, mlir::Attribute init, mlir::ConversionPatternRewriter &rewriter) const { // TODO: Generalize this handling when more types are needed here. - assert( - (isa<cir::ConstArrayAttr, cir::ConstRecordAttr, cir::ConstVectorAttr, - cir::ConstPtrAttr, cir::ConstComplexAttr, cir::GlobalViewAttr, - cir::TypeInfoAttr, cir::UndefAttr, cir::PoisonAttr, cir::VTableAttr, - cir::ZeroAttr>(init))); + assert((isa<cir::ConstArrayAttr, cir::ConstRecordAttr, cir::ConstVectorAttr, + cir::ConstPtrAttr, cir::ConstComplexAttr, cir::GlobalViewAttr, + cir::TypeInfoAttr, cir::UndefAttr, cir::PoisonAttr, + cir::VTableAttr, cir::ZeroAttr>(init))); // TODO(cir): once LLVM's dialect has proper equivalent attributes this // should be updated. For now, we use a custom op to initialize globals _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
