[clang] d8ef5bc - [CIR] Emit target-cpu, target-features, and tune-cpu attrs on cir.func (#193458)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Apr 29 03:07:31 PDT 2026
Author: Chaitanya
Date: 2026-04-29T15:37:26+05:30
New Revision: d8ef5bcc87a8be8310279082c635ffef77e073ce
URL: https://github.com/llvm/llvm-project/commit/d8ef5bcc87a8be8310279082c635ffef77e073ce
DIFF: https://github.com/llvm/llvm-project/commit/d8ef5bcc87a8be8310279082c635ffef77e073ce.diff
LOG: [CIR] Emit target-cpu, target-features, and tune-cpu attrs on cir.func (#193458)
Add `getCPUAndFeaturesAttributes` to `CIRGenModule`, mirroring OGCG's
`GetCPUAndFeaturesAttributes`.
This sets `cir.target-cpu`, `cir.target-features` and `cir.tune-cpu`
string attributes on `cir.func`.
For AMDGPU, only features that differ from the target CPU's defaults are
emitted matching OGCG.
Added:
clang/test/CIR/CodeGen/function-target-features.c
clang/test/CIR/CodeGen/tune-cpu.c
clang/test/CIR/CodeGenHIP/target-features.hip
Modified:
clang/include/clang/CIR/MissingFeatures.h
clang/lib/CIR/CodeGen/CIRGenModule.cpp
clang/lib/CIR/CodeGen/CIRGenModule.h
clang/test/CIR/CodeGen/ctor-try-body.cpp
clang/test/CIR/CodeGen/lambda.cpp
clang/test/CIR/CodeGen/misc-attrs.cpp
clang/test/CIR/CodeGen/new-delete-deactivation.cpp
clang/test/CIR/CodeGen/new-delete.cpp
clang/test/CIR/CodeGen/noreturn.cpp
clang/test/CIR/CodeGen/optsize-func-attr.cpp
clang/test/CIR/CodeGen/paren-list-agg-init.cpp
clang/test/CIR/CodeGen/pointer-to-data-member.cpp
clang/test/CIR/CodeGen/pointer-to-member-func.cpp
clang/test/CIR/CodeGen/ret-attrs.cpp
clang/test/CIR/CodeGen/side-effect.cpp
clang/test/CIR/CodeGen/ternary.cpp
clang/test/CIR/CodeGen/try-catch.cpp
clang/test/CIR/CodeGenCUDA/kernel-stub-name.cu
clang/test/CIR/CodeGenHIP/simple.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/CIR/MissingFeatures.h b/clang/include/clang/CIR/MissingFeatures.h
index 57dc48c5484ea..3cb9dafc0f9ce 100644
--- a/clang/include/clang/CIR/MissingFeatures.h
+++ b/clang/include/clang/CIR/MissingFeatures.h
@@ -72,12 +72,12 @@ struct MissingFeatures {
static bool opFuncAstDeclAttr() { return false; }
static bool opFuncCallingConv() { return false; }
static bool opFuncColdHotAttr() { return false; }
- static bool opFuncCPUAndFeaturesAttributes() { return false; }
static bool opFuncExceptions() { return false; }
static bool opFuncExtraAttrs() { return false; }
static bool opFuncMaybeHandleStaticInExternC() { return false; }
static bool opFuncMinSizeAttr() { return false; }
static bool opFuncMultipleReturnVals() { return false; }
+ static bool opFuncMultiVersioning() { return false; }
static bool opFuncNakedAttr() { return false; }
static bool opFuncNoDuplicateAttr() { return false; }
static bool opFuncOpenCLKernelMetadata() { return false; }
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index ddf36e7fa9166..354fae59b3807 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -33,6 +33,7 @@
#include "clang/CIR/Dialect/IR/CIRTypes.h"
#include "clang/CIR/Interfaces/CIROpInterfaces.h"
#include "clang/CIR/MissingFeatures.h"
+#include "llvm/ADT/StringExtras.h"
#include "llvm/ADT/StringRef.h"
#include "CIRGenFunctionInfo.h"
@@ -724,6 +725,85 @@ void CIRGenModule::setCommonAttributes(GlobalDecl gd, mlir::Operation *gv) {
}
}
+/// Get the feature delta from the default feature map for the given target CPU.
+static std::vector<std::string>
+getFeatureDeltaFromDefault(const CIRGenModule &cgm, llvm::StringRef targetCPU,
+ llvm::StringMap<bool> &featureMap) {
+ llvm::StringMap<bool> defaultFeatureMap;
+ cgm.getTarget().initFeatureMap(
+ defaultFeatureMap, cgm.getASTContext().getDiagnostics(), targetCPU, {});
+
+ std::vector<std::string> delta;
+ for (const auto &[k, v] : featureMap) {
+ auto defaultIt = defaultFeatureMap.find(k);
+ if (defaultIt == defaultFeatureMap.end() || defaultIt->getValue() != v)
+ delta.push_back((v ? "+" : "-") + k.str());
+ }
+
+ return delta;
+}
+
+bool CIRGenModule::getCPUAndFeaturesAttributes(
+ GlobalDecl gd, llvm::StringMap<std::string> &attrs,
+ bool setTargetFeatures) {
+ // Add target-cpu and target-features attributes to functions. If
+ // we have a decl for the function and it has a target attribute then
+ // parse that and add it to the feature set.
+ llvm::StringRef targetCPU = getTarget().getTargetOpts().CPU;
+ llvm::StringRef tuneCPU = getTarget().getTargetOpts().TuneCPU;
+ std::vector<std::string> features;
+ // `fd` may be null when emitting attributes for globals that don't have a
+ // FunctionDecl. The AMDGPU branch below handles
+ // the null case via initFeatureMap.
+ const auto *fd = dyn_cast_or_null<FunctionDecl>(gd.getDecl());
+ fd = fd ? fd->getMostRecentDecl() : fd;
+ const auto *td = fd ? fd->getAttr<TargetAttr>() : nullptr;
+ const auto *tv = fd ? fd->getAttr<TargetVersionAttr>() : nullptr;
+ assert((!td || !tv) && "both target_version and target specified");
+ const auto *sd = fd ? fd->getAttr<CPUSpecificAttr>() : nullptr;
+ const auto *tc = fd ? fd->getAttr<TargetClonesAttr>() : nullptr;
+ bool addedAttr = false;
+ if (td || tv || sd || tc) {
+ assert(!cir::MissingFeatures::opFuncMultiVersioning());
+ } else {
+ // Just add the existing target cpu and target features to the function.
+ if (setTargetFeatures && getTarget().getTriple().isAMDGPU()) {
+ llvm::StringMap<bool> featureMap;
+ if (fd)
+ astContext.getFunctionFeatureMap(featureMap, gd);
+ else
+ getTarget().initFeatureMap(featureMap, astContext.getDiagnostics(),
+ targetCPU,
+ getTarget().getTargetOpts().Features);
+ features = getFeatureDeltaFromDefault(*this, targetCPU, featureMap);
+ } else {
+ features = getTarget().getTargetOpts().Features;
+ }
+ }
+
+ if (!targetCPU.empty()) {
+ attrs["cir.target-cpu"] = targetCPU.str();
+ addedAttr = true;
+ }
+ if (!tuneCPU.empty()) {
+ attrs["cir.tune-cpu"] = tuneCPU.str();
+ addedAttr = true;
+ }
+ if (!features.empty() && setTargetFeatures) {
+ llvm::erase_if(features, [&](const std::string &f) {
+ assert(!f.empty() && (f[0] == '+' || f[0] == '-') &&
+ "feature string must start with '+' or '-'");
+ return getTarget().isReadOnlyFeature(f.substr(1));
+ });
+ llvm::sort(features);
+ attrs["cir.target-features"] = llvm::join(features, ",");
+ addedAttr = true;
+ }
+ // TODO(cir): add metadata for AArch64 Function Multi Versioning.
+ assert(!cir::MissingFeatures::opFuncMultiVersioning());
+ return addedAttr;
+}
+
void CIRGenModule::setNonAliasAttributes(GlobalDecl gd, mlir::Operation *op) {
setCommonAttributes(gd, op);
@@ -734,12 +814,21 @@ void CIRGenModule::setNonAliasAttributes(GlobalDecl gd, mlir::Operation *op) {
gvi.setSection(builder.getStringAttr(sa->getName()));
if (d->hasAttr<RetainAttr>())
addUsedGlobal(gvi);
+
+ if (auto func = dyn_cast<cir::FuncOp>(op)) {
+ llvm::StringMap<std::string> attrs;
+ if (getCPUAndFeaturesAttributes(gd, attrs)) {
+ // TODO(cir): Classic codegen removes the existing target-cpu,
+ // target-features, tune-cpu and fmv-features attributes here
+ // before adding the new ones.
+ for (const auto &[key, val] : attrs)
+ func->setAttr(key, builder.getStringAttr(val));
+ }
+ }
}
}
assert(!cir::MissingFeatures::opGlobalPragmaClangSection());
- assert(!cir::MissingFeatures::opFuncCPUAndFeaturesAttributes());
-
getTargetCIRGenInfo().setTargetAttributes(gd.getDecl(), op, *this);
}
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
index d6f675841a227..2869411015bc5 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -897,6 +897,9 @@ class CIRGenModule : public CIRGenTypeCache {
/// Call replaceAllUsesWith on all pairs in replacements.
void applyReplacements();
+ bool getCPUAndFeaturesAttributes(GlobalDecl gd,
+ llvm::StringMap<std::string> &attrs,
+ bool setTargetFeatures = true);
void setNonAliasAttributes(GlobalDecl gd, mlir::Operation *op);
/// Map source language used to a CIR attribute.
diff --git a/clang/test/CIR/CodeGen/ctor-try-body.cpp b/clang/test/CIR/CodeGen/ctor-try-body.cpp
index befbd578e5537..567fe9d7c77ac 100644
--- a/clang/test/CIR/CodeGen/ctor-try-body.cpp
+++ b/clang/test/CIR/CodeGen/ctor-try-body.cpp
@@ -30,7 +30,7 @@ struct HasThings : Base {
side_effect2();
}
-// CIR: cir.func {{.*}}@_ZN9HasThingsC2ERK4Ctor(%[[THIS_ARG:.*]]: !cir.ptr<!rec_HasThings> {{.*}}, %[[C_ARG:.*]]: !cir.ptr<!rec_Ctor> {{.*}}) {{.*}}special_member<#cir.cxx_ctor<!rec_HasThings, custom>> {
+// CIR: cir.func {{.*}}@_ZN9HasThingsC2ERK4Ctor(%[[THIS_ARG:.*]]: !cir.ptr<!rec_HasThings> {{.*}}, %[[C_ARG:.*]]: !cir.ptr<!rec_Ctor> {{.*}}) {{.*}}special_member<#cir.cxx_ctor<!rec_HasThings, custom>>{{.*}} {
// CIR-NEXT: %[[THIS_ALLOC:.*]] = cir.alloca !cir.ptr<!rec_HasThings>, !cir.ptr<!cir.ptr<!rec_HasThings>>, ["this", init]
// CIR-NEXT: %[[C_ALLOC:.*]] = cir.alloca !cir.ptr<!rec_Ctor>, !cir.ptr<!cir.ptr<!rec_Ctor>>, ["c", init, const]
// CIR-NEXT: cir.store %[[THIS_ARG]], %[[THIS_ALLOC]] : !cir.ptr<!rec_HasThings>, !cir.ptr<!cir.ptr<!rec_HasThings>>
diff --git a/clang/test/CIR/CodeGen/function-target-features.c b/clang/test/CIR/CodeGen/function-target-features.c
new file mode 100644
index 0000000000000..6b5d134f92eb2
--- /dev/null
+++ b/clang/test/CIR/CodeGen/function-target-features.c
@@ -0,0 +1,27 @@
+// This test verifies that we produce cir.target-cpu and cir.target-features
+// attributes on functions when they're
diff erent from the standard cpu and
+// have written features.
+
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fclangir -emit-cir -o - %s -target-feature +avx | FileCheck %s -check-prefix=AVX-FEATURE
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fclangir -emit-cir -o - %s -target-feature +avx | FileCheck %s -check-prefix=AVX-NO-CPU
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fclangir -emit-cir -o - %s -target-feature +avx512f -target-feature +avx512bw | FileCheck %s -check-prefix=TWO-AVX
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fclangir -emit-cir -o - %s -target-cpu corei7 | FileCheck %s -check-prefix=CORE-CPU
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fclangir -emit-cir -o - %s -target-cpu corei7 -target-feature +avx | FileCheck %s -check-prefix=CORE-CPU-AND-FEATURES
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fclangir -emit-cir -o - %s -target-cpu x86-64 | FileCheck %s -check-prefix=X86-64-CPU
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fclangir -emit-cir -o - %s -target-cpu corei7-avx -target-feature -avx | FileCheck %s -check-prefix=AVX-MINUS-FEATURE
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fclangir -emit-llvm -o - %s -target-cpu corei7 -target-feature +avx | FileCheck %s -check-prefix=LLVM
+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -emit-llvm -o - %s -target-cpu corei7 -target-feature +avx | FileCheck %s -check-prefix=LLVM
+
+void foo(void) {}
+
+// AVX-FEATURE: "cir.target-features"{{.*}}+avx
+// AVX-NO-CPU-NOT: cir.target-cpu
+// TWO-AVX: "cir.target-features"{{.*}}+avx512bw{{.*}}+avx512f
+// CORE-CPU: "cir.target-cpu" = "corei7"
+// CORE-CPU-AND-FEATURES: "cir.target-cpu" = "corei7"
+// CORE-CPU-AND-FEATURES-SAME: "cir.target-features"{{.*}}+avx
+// X86-64-CPU: "cir.target-cpu" = "x86-64"
+// AVX-MINUS-FEATURE: "cir.target-features"{{.*}}-avx
+
+// LLVM: define{{.*}} void @foo(){{.*}} #[[ATTR:[0-9]+]]
+// LLVM: attributes #[[ATTR]] = {{.*}}"target-cpu"="corei7"{{.*}}"target-features"={{.*}}+avx
diff --git a/clang/test/CIR/CodeGen/lambda.cpp b/clang/test/CIR/CodeGen/lambda.cpp
index 7a0c0427959fe..0faa3bc9b8441 100644
--- a/clang/test/CIR/CodeGen/lambda.cpp
+++ b/clang/test/CIR/CodeGen/lambda.cpp
@@ -157,7 +157,7 @@ auto g() {
};
}
-// CIR: cir.func {{.*}} @_Z1gv() -> ![[REC_LAM_G:.*]] attributes {nothrow} {
+// CIR: cir.func {{.*}} @_Z1gv() -> ![[REC_LAM_G:.*]] attributes {{{.*}}nothrow} {
// CIR: %[[RETVAL:.*]] = cir.alloca ![[REC_LAM_G]], !cir.ptr<![[REC_LAM_G]]>, ["__retval"]
// CIR: %[[I_ADDR:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["i", init]
// CIR: %[[TWELVE:.*]] = cir.const #cir.int<12> : !s32i
@@ -199,7 +199,7 @@ auto g2() {
}
// Should be same as above because of NRVO
-// CIR: cir.func {{.*}} @_Z2g2v() -> ![[REC_LAM_G2:.*]] attributes {nothrow} {
+// CIR: cir.func {{.*}} @_Z2g2v() -> ![[REC_LAM_G2:.*]] attributes {{{.*}}nothrow} {
// CIR: %[[RETVAL:.*]] = cir.alloca ![[REC_LAM_G2]], !cir.ptr<![[REC_LAM_G2]]>, ["__retval", init]
// CIR: %[[I_ADDR:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["i", init]
// CIR: %[[TWELVE:.*]] = cir.const #cir.int<12> : !s32i
diff --git a/clang/test/CIR/CodeGen/misc-attrs.cpp b/clang/test/CIR/CodeGen/misc-attrs.cpp
index 4cfddc7c25666..d9e7ad261dd96 100644
--- a/clang/test/CIR/CodeGen/misc-attrs.cpp
+++ b/clang/test/CIR/CodeGen/misc-attrs.cpp
@@ -6,47 +6,47 @@
// RUN: FileCheck --input-file=%t.ll %s -check-prefix=LLVM
extern "C" {
- // CIR: cir.func{{.*}}@returns_twice() attributes {nothrow, returns_twice} {
+ // CIR: cir.func{{.*}}@returns_twice() attributes {{{.*}}nothrow, returns_twice} {
// LLVM: Function Attrs:{{.*}}returns_twice
// LLVM-NEXT: define{{.*}}@returns_twice() #[[RT_ATTR:.*]] {
__attribute__((returns_twice))
void returns_twice() {}
- // CIR: cir.func{{.*}}@cold() attributes {cold, nothrow} {
+ // CIR: cir.func{{.*}}@cold() attributes {{{.*}}cold, nothrow} {
// LLVM: Function Attrs:{{.*}}cold
// LLVM-NEXT: define{{.*}}@cold() #[[COLD_ATTR:.*]] {
__attribute__((cold))
void cold() {}
- // CIR: cir.func{{.*}}@hot() attributes {hot, nothrow} {
+ // CIR: cir.func{{.*}}@hot() attributes {{{.*}}hot, nothrow} {
// LLVM: Function Attrs:{{.*}}hot
// LLVM-NEXT: define{{.*}}@hot() #[[HOT_ATTR:.*]] {
__attribute__((hot))
void hot() {}
- // CIR: cir.func{{.*}}@nodupes() attributes {noduplicate, nothrow} {
+ // CIR: cir.func{{.*}}@nodupes() attributes {{{.*}}noduplicate, nothrow} {
// LLVM: Function Attrs:{{.*}}noduplicate
// LLVM-NEXT: define{{.*}}@nodupes() #[[ND_ATTR:.*]] {
__attribute__((noduplicate))
void nodupes() {}
- // CIR: cir.func{{.*}}@convergent() attributes {convergent, nothrow} {
+ // CIR: cir.func{{.*}}@convergent() attributes {{{.*}}convergent, nothrow} {
// LLVM: Function Attrs:{{.*}}convergent
// LLVM-NEXT: define{{.*}}@convergent() #[[CONV_ATTR:.*]] {
__attribute__((convergent))
void convergent() {}
- // CIR: cir.func{{.*}}@no_caller_saved_registers() attributes {no_caller_saved_registers, nothrow} {
+ // CIR: cir.func{{.*}}@no_caller_saved_registers() attributes {{{.*}}no_caller_saved_registers, nothrow} {
// LLVM: Function Attrs:
// LLVM-NOT: no_caller_saved_registers
// LLVM-NEXT: define{{.*}}@no_caller_saved_registers() #[[NCSR_ATTR:.*]] {
__attribute__((no_caller_saved_registers))
void no_caller_saved_registers() {}
- // CIR: cir.func{{.*}}@leaf() attributes {nocallback, nothrow} {
+ // CIR: cir.func{{.*}}@leaf() attributes {{{.*}}nocallback, nothrow} {
// LLVM: Function Attrs:
// LLVM-NOT: leaf
// LLVM-NEXT: define{{.*}}@leaf() #[[LEAF_ATTR:.*]] {
__attribute__((leaf))
void leaf() {}
- // CIR: cir.func{{.*}}@modular_format({{.*}}) attributes {modular_format = "kprintf,1,2,someIdent,someStr,aspect,aspect2", nothrow} {
+ // CIR: cir.func{{.*}}@modular_format({{.*}}) attributes {{{.*}}modular_format = "kprintf,1,2,someIdent,someStr,aspect,aspect2", nothrow} {
// LLVM: Function Attrs:
// LLVM-NOT:modular_format
// LLVM-NEXT: define{{.*}}@modular_format({{.*}}) #[[MOD_FORMAT_ATTR:.*]] {
diff --git a/clang/test/CIR/CodeGen/new-delete-deactivation.cpp b/clang/test/CIR/CodeGen/new-delete-deactivation.cpp
index 1375476ebe50b..86431113991d8 100644
--- a/clang/test/CIR/CodeGen/new-delete-deactivation.cpp
+++ b/clang/test/CIR/CodeGen/new-delete-deactivation.cpp
@@ -21,7 +21,7 @@ B makeB();
A *deact_simple() { return new A(makeB()); }
-// CIR-LABEL: cir.func {{.*}} @_Z12deact_simplev() -> !cir.ptr<!rec_A> {
+// CIR-LABEL: cir.func {{.*}} @_Z12deact_simplev() -> !cir.ptr<!rec_A>{{.*}} {
// CIR: %[[RETVAL:.*]] = cir.alloca !cir.ptr<!rec_A>, !cir.ptr<!cir.ptr<!rec_A>>, ["__retval"]
// CIR: %[[NEW_RESULT:.*]] = cir.alloca !cir.ptr<!rec_A>, !cir.ptr<!cir.ptr<!rec_A>>, ["__new_result"]
// CIR: %[[TMP:.*]] = cir.alloca !rec_B, !cir.ptr<!rec_B>, ["ref.tmp0"]
diff --git a/clang/test/CIR/CodeGen/new-delete.cpp b/clang/test/CIR/CodeGen/new-delete.cpp
index b1dfeba0a1a55..fa8ed0dfeb947 100644
--- a/clang/test/CIR/CodeGen/new-delete.cpp
+++ b/clang/test/CIR/CodeGen/new-delete.cpp
@@ -12,7 +12,7 @@ A *a() {
return new A(5);
}
-// CIR: cir.func {{.*}} @_Z1av() -> !cir.ptr<!rec_A> {
+// CIR: cir.func {{.*}} @_Z1av() -> !cir.ptr<!rec_A>{{.*}} {
// CIR: %[[RETVAL:.*]] = cir.alloca !cir.ptr<!rec_A>, !cir.ptr<!cir.ptr<!rec_A>>, ["__retval"]
// CIR: %[[NEW_RESULT:.*]] = cir.alloca !cir.ptr<!rec_A>, !cir.ptr<!cir.ptr<!rec_A>>, ["__new_result"]
// CIR: %[[ALLOC_SIZE:.*]] = cir.const #cir.int<8> : !u64i
@@ -87,7 +87,7 @@ A *b() {
return new A(foo());
}
-// CIR: cir.func {{.*}} @_Z1bv() -> !cir.ptr<!rec_A> {
+// CIR: cir.func {{.*}} @_Z1bv() -> !cir.ptr<!rec_A>{{.*}} {
// CIR: %[[RETVAL:.*]] = cir.alloca !cir.ptr<!rec_A>, !cir.ptr<!cir.ptr<!rec_A>>, ["__retval"]
// CIR: %[[NEW_RESULT:.*]] = cir.alloca !cir.ptr<!rec_A>, !cir.ptr<!cir.ptr<!rec_A>>, ["__new_result"]
// CIR: %[[ALLOC_SIZE:.*]] = cir.const #cir.int<8> : !u64i
@@ -177,7 +177,7 @@ B *c() {
return new B(5);
}
-// CIR: cir.func {{.*}} @_Z1cv() -> !cir.ptr<!rec_B> {
+// CIR: cir.func {{.*}} @_Z1cv() -> !cir.ptr<!rec_B>{{.*}} {
// CIR: %[[RETVAL:.*]] = cir.alloca !cir.ptr<!rec_B>, !cir.ptr<!cir.ptr<!rec_B>>, ["__retval"]
// CIR: %[[NEW_RESULT:.*]] = cir.alloca !cir.ptr<!rec_B>, !cir.ptr<!cir.ptr<!rec_B>>, ["__new_result"]
// CIR: %[[ALLOC_SIZE:.*]] = cir.const #cir.int<4> : !u64i
diff --git a/clang/test/CIR/CodeGen/noreturn.cpp b/clang/test/CIR/CodeGen/noreturn.cpp
index f55628ba08e03..2ae59e3adfcac 100644
--- a/clang/test/CIR/CodeGen/noreturn.cpp
+++ b/clang/test/CIR/CodeGen/noreturn.cpp
@@ -6,7 +6,7 @@
// RUN: FileCheck --input-file=%t.ll %s -check-prefixes=LLVM,OGCG
extern "C" {
-// CIR: cir.func {{.*}} @bar() -> !s32i attributes {noreturn, nothrow} {
+// CIR: cir.func {{.*}} @bar() -> !s32i attributes {{{.*}}noreturn, nothrow} {
// LLVM: Function Attrs:{{.*}} noreturn
// LLVM-NEXT: define {{.*}} i32 @bar() #[[BAR_FOO_ATTR:.*]] {
__attribute((noreturn))
@@ -17,7 +17,7 @@ int bar() { }
// lowering puts the trap decl at the end, so it isn't here to worry about.
// OGCG: declare void @llvm.trap
-// CIR: cir.func {{.*}} @foo() -> !s32i attributes {noreturn, nothrow} {
+// CIR: cir.func {{.*}} @foo() -> !s32i attributes {{{.*}}noreturn, nothrow} {
// LLVM: Function Attrs:{{.*}} noreturn
// LLVM-NEXT: define {{.*}} i32 @foo() #[[BAR_FOO_ATTR]] {
[[noreturn]]
diff --git a/clang/test/CIR/CodeGen/optsize-func-attr.cpp b/clang/test/CIR/CodeGen/optsize-func-attr.cpp
index 38983452b3cab..599bb8fd99435 100644
--- a/clang/test/CIR/CodeGen/optsize-func-attr.cpp
+++ b/clang/test/CIR/CodeGen/optsize-func-attr.cpp
@@ -1,16 +1,16 @@
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -disable-llvm-passes -Os -fclangir -emit-cir %s -o %t.cir
// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -disable-llvm-passes -Os -fclangir -emit-llvm %s -o %t.ll
-// RUN: FileCheck --input-file=%t.ll %s -check-prefix=LLVM,BOTH
+// RUN: FileCheck --input-file=%t.ll %s -check-prefix=BOTH
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -disable-llvm-passes -Os -emit-llvm %s -o %t.ll
-// RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG,BOTH
+// RUN: FileCheck --input-file=%t.ll %s -check-prefix=BOTH
//
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -disable-llvm-passes -Oz -fclangir -emit-cir %s -o %t.cir
// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR,CIROZ
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -disable-llvm-passes -Oz -fclangir -emit-llvm %s -o %t.ll
-// RUN: FileCheck --input-file=%t.ll %s -check-prefix=LLVM,BOTH,BOTHOZ
+// RUN: FileCheck --input-file=%t.ll %s -check-prefix=BOTH,BOTHOZ
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -disable-llvm-passes -Oz -emit-llvm %s -o %t.ll
-// RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG,OGCGOZ,BOTH,BOTHOZ
+// RUN: FileCheck --input-file=%t.ll %s -check-prefix=BOTH,BOTHOZ
extern "C" {
__attribute__((hot))
@@ -34,14 +34,12 @@ extern "C" {
// CIR: cir.call{{.*}}@normal()
// CIROZ-SAME: minsize
// CIR-SAME: optsize
- // LLVM: call void @normal() #[[NORMAL_ATTR]]
- // OGCG: call void @normal() #[[NORMAL_CALL_ATTR:.*]]
+ // BOTH: call void @normal() #[[NORMAL_CALL_ATTR:.*]]
optnone();
// CIR: cir.call{{.*}}@optnone()
// CIR-NOT: optsize
// CIR-NOT: minsize
- // LLVM: call void @optnone() #[[OPTNONE_ATTR]]
- // OGCG: call void @optnone() #[[OPTNONE_CALL_ATTR:.*]]
+ // BOTH: call void @optnone() #[[OPTNONE_CALL_ATTR:.*]]
// CIR: cir.return
}
@@ -58,17 +56,13 @@ extern "C" {
// attributes for caller, to block the 'NOT'.
// BOTH: attributes
//
-// CIR doesn't have sufficiently
diff erent 'attributes' implemented for the
-// caller and the callee to be
diff erent when doing -O settings (as 'optnone'
-// is the only
diff erence). So the below call attributes are only necessary
-// for classic codegen.
-// OGCG: attributes #[[NORMAL_CALL_ATTR]]
-// OGCGOZ-SAME: minsize
-// OGCG-SAME: optsize
+// BOTH: attributes #[[NORMAL_CALL_ATTR]]
+// BOTHOZ-SAME: minsize
+// BOTH-SAME: optsize
//
-// OGCG: attributes #[[OPTNONE_CALL_ATTR]]
-// OGCG-NOT: optsize
-// OGCG-NOT: minsize
+// BOTH: attributes #[[OPTNONE_CALL_ATTR]]
+// BOTH-NOT: optsize
+// BOTH-NOT: minsize
//
// to block the 'NOT'.
// BOTH: !llvm.
diff --git a/clang/test/CIR/CodeGen/paren-list-agg-init.cpp b/clang/test/CIR/CodeGen/paren-list-agg-init.cpp
index ac822e7e38a94..a5e84669d914d 100644
--- a/clang/test/CIR/CodeGen/paren-list-agg-init.cpp
+++ b/clang/test/CIR/CodeGen/paren-list-agg-init.cpp
@@ -570,7 +570,7 @@ void foo18() {
// LLVM-NEXT: [[F:%.*]] = getelementptr {{.*}}i8, ptr [[G]], i64 4
// LLVM-NEXT: call void @{{.*F.*}}(ptr noundef nonnull align 1 dereferenceable(1) [[F]], i32 noundef 1)
// LLVM: ret void
-// CIR: cir.func no_inline dso_local @_Z5foo19v() attributes {nothrow} {
+// CIR: cir.func no_inline dso_local @_Z5foo19v() attributes {{{.*}}nothrow} {
// CIR: %[[G_ALLOCA:.*]] = cir.alloca ![[STRUCT_G]], !cir.ptr<![[STRUCT_G]]>, ["g", init]
// CIR: %[[GET_A:.*]] = cir.get_member %[[G_ALLOCA]][0] {name = "a"} : !cir.ptr<![[STRUCT_G]]> -> !cir.ptr<!s32i>
// CIR: %[[TWO:.*]] = cir.const #cir.int<2> : !s32i
diff --git a/clang/test/CIR/CodeGen/pointer-to-data-member.cpp b/clang/test/CIR/CodeGen/pointer-to-data-member.cpp
index 3d88b989d23cb..2100987462f8b 100644
--- a/clang/test/CIR/CodeGen/pointer-to-data-member.cpp
+++ b/clang/test/CIR/CodeGen/pointer-to-data-member.cpp
@@ -52,14 +52,14 @@ int Point::*pt_member_nested_region = test1();
// Checks for test1()
-// CIR-BEFORE: cir.func {{.*}} @_Z5test1v() -> !cir.data_member<!s32i in !rec_Point> attributes {nothrow} {
+// CIR-BEFORE: cir.func {{.*}} @_Z5test1v() -> !cir.data_member<!s32i in !rec_Point> attributes {{{.*}}nothrow} {
// CIR-BEFORE: %[[RETVAL:.*]] = cir.alloca !cir.data_member<!s32i in !rec_Point>, !cir.ptr<!cir.data_member<!s32i in !rec_Point>>, ["__retval"]
// CIR-BEFORE: %[[MEMBER:.*]] = cir.const #cir.data_member<1> : !cir.data_member<!s32i in !rec_Point>
// CIR-BEFORE: cir.store %[[MEMBER]], %[[RETVAL]] : !cir.data_member<!s32i in !rec_Point>, !cir.ptr<!cir.data_member<!s32i in !rec_Point>>
// CIR-BEFORE: %[[RET:.*]] = cir.load %[[RETVAL]] : !cir.ptr<!cir.data_member<!s32i in !rec_Point>>, !cir.data_member<!s32i in !rec_Point>
// CIR-BEFORE: cir.return %[[RET]] : !cir.data_member<!s32i in !rec_Point>
-// CIR-AFTER: cir.func {{.*}} @_Z5test1v() -> !s64i attributes {nothrow} {
+// CIR-AFTER: cir.func {{.*}} @_Z5test1v() -> !s64i attributes {{{.*}}nothrow} {
// CIR-AFTER: %[[RETVAL:.*]] = cir.alloca !s64i, !cir.ptr<!s64i>, ["__retval"]
// CIR-AFTER: %[[OFFSET:.*]] = cir.const #cir.int<4> : !s64i
// CIR-AFTER: cir.store %[[OFFSET]], %[[RETVAL]] : !s64i, !cir.ptr<!s64i>
diff --git a/clang/test/CIR/CodeGen/pointer-to-member-func.cpp b/clang/test/CIR/CodeGen/pointer-to-member-func.cpp
index a7bf04519575f..1388e38882617 100644
--- a/clang/test/CIR/CodeGen/pointer-to-member-func.cpp
+++ b/clang/test/CIR/CodeGen/pointer-to-member-func.cpp
@@ -52,7 +52,7 @@ auto make_non_virtual() -> void (Foo::*)(int) {
// CIR-BEFORE: %[[RET:.*]] = cir.load %[[RETVAL]]
// CIR-BEFORE: cir.return %[[RET]] : !cir.method<!cir.func<(!cir.ptr<!rec_Foo>, !s32i)> in !rec_Foo>
-// CIR-AFTER: cir.func {{.*}} @_Z16make_non_virtualv() -> !rec_anon_struct attributes {nothrow} {
+// CIR-AFTER: cir.func {{.*}} @_Z16make_non_virtualv() -> !rec_anon_struct attributes {{{.*}}nothrow} {
// CIR-AFTER: %[[RETVAL:.*]] = cir.alloca !rec_anon_struct, !cir.ptr<!rec_anon_struct>, ["__retval"]
// CIR-AFTER: %[[METHOD_PTR:.*]] = cir.get_global @[[NONVIRT_RET]] : !cir.ptr<!rec_anon_struct>
// CIR-AFTER: cir.copy %[[METHOD_PTR]] to %[[RETVAL]] : !cir.ptr<!rec_anon_struct>
diff --git a/clang/test/CIR/CodeGen/ret-attrs.cpp b/clang/test/CIR/CodeGen/ret-attrs.cpp
index 9ef836791fea5..472b8cf5c20fd 100644
--- a/clang/test/CIR/CodeGen/ret-attrs.cpp
+++ b/clang/test/CIR/CodeGen/ret-attrs.cpp
@@ -16,7 +16,7 @@ struct Struct{};
using MemPtrTy = void (Struct::*)();
MemPtrTy not_noundef_memptr(MemPtrTy t){}
-// CIR: cir.func no_inline dso_local @_Z18not_noundef_memptrM6StructFvvE({{.*}}) -> !rec_anon_struct attributes {nothrow} {
+// CIR: cir.func no_inline dso_local @_Z18not_noundef_memptrM6StructFvvE({{.*}}) -> !rec_anon_struct attributes {{{.*}}nothrow} {
// LLVM: define dso_local { i64, i64 } @_Z18not_noundef_memptrM6StructFvvE({{.*}})
void not_noundef_void(){}
diff --git a/clang/test/CIR/CodeGen/side-effect.cpp b/clang/test/CIR/CodeGen/side-effect.cpp
index 0a47e1493a747..a1e0fbbeb0e2c 100644
--- a/clang/test/CIR/CodeGen/side-effect.cpp
+++ b/clang/test/CIR/CodeGen/side-effect.cpp
@@ -8,7 +8,7 @@ extern "C" {
// FIXME: We should figure out how to better print this on functions in the
// future.
-// CIR: cir.func{{.*}}@pure_func() -> !s32i side_effect(pure) attributes {nothrow} {
+// CIR: cir.func{{.*}}@pure_func() -> !s32i side_effect(pure) attributes {{{.*}}nothrow} {
// LLVM: Function Attrs: {{.*}}nounwind{{.*}}willreturn{{.*}}memory(read)
// LLVM: define{{.*}} @pure_func() #{{.*}} {
// OGCG: Function Attrs: {{.*}}nounwind{{.*}}willreturn{{.*}}memory(read)
@@ -16,7 +16,7 @@ extern "C" {
__attribute__((pure))
int pure_func() { return 2;}
-// CIR: cir.func{{.*}}@const_func() -> !s32i side_effect(const) attributes {nothrow} {
+// CIR: cir.func{{.*}}@const_func() -> !s32i side_effect(const) attributes {{{.*}}nothrow} {
// LLVM: Function Attrs: {{.*}}nounwind{{.*}}willreturn{{.*}}memory(none)
// LLVM: define{{.*}} @const_func() #{{.*}} {
// OGCG: Function Attrs: {{.*}}nounwind{{.*}}willreturn{{.*}}memory(none)
diff --git a/clang/test/CIR/CodeGen/ternary.cpp b/clang/test/CIR/CodeGen/ternary.cpp
index 51b9ab7c7db40..c3b8d86af2cd0 100644
--- a/clang/test/CIR/CodeGen/ternary.cpp
+++ b/clang/test/CIR/CodeGen/ternary.cpp
@@ -10,7 +10,7 @@ int x(int y) {
}
// CIR-LABEL: cir.func{{.*}} @_Z1xi(
-// CIR-SAME: %[[ARG0:.*]]: !s32i {{.*}}) -> (!s32i{{.*}}) attributes {nothrow} {
+// CIR-SAME: %[[ARG0:.*]]: !s32i {{.*}}) -> (!s32i{{.*}}) attributes {{{.*}}nothrow} {
// CIR: [[Y:%.+]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["y", init]
// CIR: [[RETVAL:%.+]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["__retval"]
// CIR: cir.store %[[ARG0]], [[Y]] : !s32i, !cir.ptr<!s32i>
@@ -52,7 +52,7 @@ int foo(int a, int b) {
}
// CIR-LABEL: cir.func{{.*}} @_Z3fooii(
-// CIR-SAME: %[[ARG0:.*]]: !s32i {{.*}}, %[[ARG1:.*]]: !s32i {{.*}}) -> (!s32i{{.*}}) attributes {nothrow} {
+// CIR-SAME: %[[ARG0:.*]]: !s32i {{.*}}, %[[ARG1:.*]]: !s32i {{.*}}) -> (!s32i{{.*}}) attributes {{{.*}}nothrow} {
// CIR: [[A:%.+]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["a", init]
// CIR: [[B:%.+]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["b", init]
// CIR: [[RETVAL:%.+]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["__retval"]
diff --git a/clang/test/CIR/CodeGen/try-catch.cpp b/clang/test/CIR/CodeGen/try-catch.cpp
index fd8b9294c9f92..7f50c7ca6ec28 100644
--- a/clang/test/CIR/CodeGen/try-catch.cpp
+++ b/clang/test/CIR/CodeGen/try-catch.cpp
@@ -899,7 +899,7 @@ void cleanup_inside_try_body() {
}
}
-// CIR: cir.func {{.*}} @_Z23cleanup_inside_try_bodyv(){{.*}} personality(@__gxx_personality_v0) {
+// CIR: cir.func {{.*}} @_Z23cleanup_inside_try_bodyv(){{.*}} personality(@__gxx_personality_v0){{.*}} {
// CIR: cir.scope {
// CIR: %[[S:.*]] = cir.alloca !rec_S, !cir.ptr<!rec_S>, ["s"]
// CIR: cir.try {
@@ -1009,7 +1009,7 @@ void call_function_inside_try_catch_with_aggregate_exception_type() {
}
-// CIR: cir.func {{.*}} @_Z60call_function_inside_try_catch_with_aggregate_exception_typev(){{.*}} personality(@__gxx_personality_v0) {
+// CIR: cir.func {{.*}} @_Z60call_function_inside_try_catch_with_aggregate_exception_typev(){{.*}} personality(@__gxx_personality_v0){{.*}} {
// CIR: cir.scope {
// CIR: %[[E_ADDR:.*]] = cir.alloca !rec_CustomError, !cir.ptr<!rec_CustomError>, ["e"]
// CIR: cir.try {
@@ -1129,7 +1129,7 @@ void call_function_inside_try_catch_with_ref_ptr_of_record_exception_type() {
}
}
-// CIR: cir.func {{.*}} @_Z68call_function_inside_try_catch_with_ref_ptr_of_record_exception_typev(){{.*}} personality(@__gxx_personality_v0) {
+// CIR: cir.func {{.*}} @_Z68call_function_inside_try_catch_with_ref_ptr_of_record_exception_typev(){{.*}} personality(@__gxx_personality_v0){{.*}} {
// CIR: %[[E_ADDR:.*]] = cir.alloca !cir.ptr<!cir.ptr<!rec_Record>>, !cir.ptr<!cir.ptr<!cir.ptr<!rec_Record>>>, ["ref_ptr", const]
// CIR: %[[EXN_BYREF_TMP:.*]] = cir.alloca !cir.ptr<!rec_Record>, !cir.ptr<!cir.ptr<!rec_Record>>, ["exn.byref.tmp"]
// CIR: cir.try {
@@ -1251,7 +1251,7 @@ void call_function_inside_try_catch_with_exception_member_ptr_type() {
}
}
-// CIR: cir.func {{.*}} @_Z61call_function_inside_try_catch_with_exception_member_ptr_typev(){{.*}} personality(@__gxx_personality_v0) {
+// CIR: cir.func {{.*}} @_Z61call_function_inside_try_catch_with_exception_member_ptr_typev(){{.*}} personality(@__gxx_personality_v0){{.*}} {
// CIR: cir.scope {
// CIR: %[[E_ADDR:.*]] = cir.alloca !s64i, !cir.ptr<!s64i>, ["memberPtr"]
// CIR: cir.try {
diff --git a/clang/test/CIR/CodeGen/tune-cpu.c b/clang/test/CIR/CodeGen/tune-cpu.c
new file mode 100644
index 0000000000000..d51a6b7ac1abc
--- /dev/null
+++ b/clang/test/CIR/CodeGen/tune-cpu.c
@@ -0,0 +1,13 @@
+// RUN: %clang_cc1 -triple i686-linux-gnu -target-cpu i686 -tune-cpu nehalem -fclangir -emit-cir %s -o - | FileCheck %s -check-prefix=CIR
+// RUN: %clang_cc1 -triple i686-linux-gnu -target-cpu i686 -tune-cpu nehalem -fclangir -emit-llvm %s -o - | FileCheck %s -check-prefix=LLVM
+// RUN: %clang_cc1 -triple i686-linux-gnu -target-cpu i686 -tune-cpu nehalem -emit-llvm %s -o - | FileCheck %s -check-prefix=LLVM
+
+int baz(int a) { return 4; }
+
+// CIR: cir.func{{.*}} @baz
+// CIR-SAME: "cir.target-cpu" = "i686"
+// CIR-SAME: "cir.target-features" = "+cmov,+cx8,+x87"
+// CIR-SAME: "cir.tune-cpu" = "nehalem"
+
+// LLVM: baz{{.*}} #[[ATTR:[0-9]+]]
+// LLVM: attributes #[[ATTR]] = {{.*}}"target-cpu"="i686" "target-features"="+cmov,+cx8,+x87" "tune-cpu"="nehalem"
diff --git a/clang/test/CIR/CodeGenCUDA/kernel-stub-name.cu b/clang/test/CIR/CodeGenCUDA/kernel-stub-name.cu
index 42c8f10430b1f..045b836bdbe47 100644
--- a/clang/test/CIR/CodeGenCUDA/kernel-stub-name.cu
+++ b/clang/test/CIR/CodeGenCUDA/kernel-stub-name.cu
@@ -6,17 +6,17 @@
#include "Inputs/cuda.h"
-// CHECK: cir.func {{.*}} @[[CSTUB:__device_stub__ckernel]]() attributes {cu.kernel_name = #cir.cu.kernel_name<"ckernel">{{.*}}}
+// CHECK: cir.func {{.*}} @[[CSTUB:__device_stub__ckernel]]() attributes {{{.*}}cu.kernel_name = #cir.cu.kernel_name<"ckernel">{{.*}}}
// CHECK: cir.return
// CHECK-NEXT: }
extern "C" __global__ void ckernel() {}
-// CHECK: cir.func {{.*}} @_ZN2ns23__device_stub__nskernelEv() attributes {cu.kernel_name = #cir.cu.kernel_name<"_ZN2ns8nskernelEv">{{.*}}}
+// CHECK: cir.func {{.*}} @_ZN2ns23__device_stub__nskernelEv() attributes {{{.*}}cu.kernel_name = #cir.cu.kernel_name<"_ZN2ns8nskernelEv">{{.*}}}
namespace ns {
__global__ void nskernel() {}
} // namespace ns
-// CHECK: cir.func {{.*}} @_Z25__device_stub__kernelfuncIiEvv() attributes {cu.kernel_name = #cir.cu.kernel_name<"_Z10kernelfuncIiEvv">{{.*}}}
+// CHECK: cir.func {{.*}} @_Z25__device_stub__kernelfuncIiEvv() attributes {{{.*}}cu.kernel_name = #cir.cu.kernel_name<"_Z10kernelfuncIiEvv">{{.*}}}
template <class T>
__global__ void kernelfunc() {}
template __global__ void kernelfunc<int>();
diff --git a/clang/test/CIR/CodeGenHIP/simple.cpp b/clang/test/CIR/CodeGenHIP/simple.cpp
index b3df34aed6afb..f4d2645c7803a 100644
--- a/clang/test/CIR/CodeGenHIP/simple.cpp
+++ b/clang/test/CIR/CodeGenHIP/simple.cpp
@@ -42,7 +42,7 @@ __global__ void global_fn(int a) {}
// CIR-DEVICE: cir.func {{.*}}{{.*}} @_Z9global_fni
// OGCG-DEVICE: define protected amdgpu_kernel void @_Z9global_fni
-// CIR-HOST: @_Z24__device_stub__global_fni{{.*}}attributes {cu.kernel_name = #cir.cu.kernel_name<"_Z9global_fni">{{.*}}}
+// CIR-HOST: @_Z24__device_stub__global_fni{{.*}}attributes {{{.*}}cu.kernel_name = #cir.cu.kernel_name<"_Z9global_fni">{{.*}}}
// CIR-HOST: %[[#CIRKernelArgs:]] = cir.alloca {{.*}}"kernel_args"
// CIR-HOST: %[[#Decayed:]] = cir.cast array_to_ptrdecay %[[#CIRKernelArgs]]
// CIR-HOST: cir.call @__hipPopCallConfiguration
diff --git a/clang/test/CIR/CodeGenHIP/target-features.hip b/clang/test/CIR/CodeGenHIP/target-features.hip
new file mode 100644
index 0000000000000..8d414edcd8e2c
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/target-features.hip
@@ -0,0 +1,65 @@
+#include "../CodeGenCUDA/Inputs/cuda.h"
+
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
+// RUN: -fcuda-is-device -target-cpu gfx900 -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR %s --input-file=%t.cir
+
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
+// RUN: -fcuda-is-device -target-cpu gfx900 -emit-llvm %s -o %t.cir.ll
+// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.cir.ll
+
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip \
+// RUN: -fcuda-is-device -target-cpu gfx900 -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.ll
+
+// Verify that with extra target-features
diff ering from the CPU baseline,
+// only the delta (
diff ering features) is emitted on cir.target-features.
+
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
+// RUN: -fcuda-is-device -target-cpu gfx900 -target-feature +xnack \
+// RUN: -emit-cir %s -o %t-delta.cir
+// RUN: FileCheck --check-prefix=CIR-DELTA %s --input-file=%t-delta.cir
+
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
+// RUN: -fcuda-is-device -target-cpu gfx900 -target-feature +xnack \
+// RUN: -emit-llvm %s -o %t-delta.cir.ll
+// RUN: FileCheck --check-prefix=LLVM-DELTA %s --input-file=%t-delta.cir.ll
+
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip \
+// RUN: -fcuda-is-device -target-cpu gfx900 -target-feature +xnack \
+// RUN: -emit-llvm %s -o %t-delta.ll
+// RUN: FileCheck --check-prefix=LLVM-DELTA %s --input-file=%t-delta.ll
+
+// AMDGPU with gfx900: target-cpu is set, target-features is absent.
+
+// CIR: cir.func{{.*}} @_Z6kernelv()
+// CIR-SAME: "cir.target-cpu" = "gfx900"
+// CIR-NOT: cir.target-features
+__global__ void kernel() {}
+
+// CIR: cir.func{{.*}} @_Z9device_fnv()
+// CIR-SAME: "cir.target-cpu" = "gfx900"
+// CIR-NOT: cir.target-features
+__device__ void device_fn() {}
+
+// LLVM: define{{.*}} void @_Z6kernelv(){{.*}} #[[K_ATTR:[0-9]+]]
+// LLVM: define{{.*}} void @_Z9device_fnv(){{.*}} #[[D_ATTR:[0-9]+]]
+// LLVM-DAG: attributes #[[K_ATTR]] = {{.*}}"target-cpu"="gfx900"
+// LLVM-DAG: attributes #[[D_ATTR]] = {{.*}}"target-cpu"="gfx900"
+
+// AMDGPU with gfx900 + an extra +xnack feature: only the delta is emitted.
+
+// CIR-DELTA: cir.func{{.*}} @_Z6kernelv()
+// CIR-DELTA-SAME: "cir.target-cpu" = "gfx900"
+// CIR-DELTA-SAME: "cir.target-features" = "+xnack"
+
+// CIR-DELTA: cir.func{{.*}} @_Z9device_fnv()
+// CIR-DELTA-SAME: "cir.target-cpu" = "gfx900"
+// CIR-DELTA-SAME: "cir.target-features" = "+xnack"
+
+// LLVM-DELTA: define{{.*}} void @_Z6kernelv(){{.*}} #[[K_ATTR_D:[0-9]+]]
+// LLVM-DELTA: define{{.*}} void @_Z9device_fnv(){{.*}} #[[D_ATTR_D:[0-9]+]]
+// LLVM-DELTA-DAG: attributes #[[K_ATTR_D]] = {{.*}}"target-cpu"="gfx900"{{.*}}"target-features"="+xnack"
+// LLVM-DELTA-DAG: attributes #[[D_ATTR_D]] = {{.*}}"target-cpu"="gfx900"{{.*}}"target-features"="+xnack"
More information about the cfe-commits
mailing list