Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
20 commits
Select commit Hold shift + click to select a range
72bfa28
[X86] avx2-builtins.c - fix copy+paste typo in _mm256_cmpeq_epi8 cons…
RKSimon Nov 24, 2025
74f5548
[HLSL][SPIR-V] Implements SV_Position for VS/PS I/O (#168735)
Keenuts Nov 24, 2025
e4cff3c
[mlir] Avoid else after return in ScalableValueBounds (NFC) (#169211)
MacDue Nov 24, 2025
65fd9f1
[Attributor] Support nested conditional branches (#168532)
c-rhodes Nov 24, 2025
999deef
Desugar complex element types for promoted complex division (#168943)
zahiraam Nov 24, 2025
e575539
[milr][memref]: Fold expand_shape + transfer_read (#167679)
FranklandJack Nov 24, 2025
a27842c
[X86][NFC] Add `-show-mc-encoding` to check register misuse (#169264)
phoebewang Nov 24, 2025
d148407
[X86][AVX512] Add pseudos for `AVX512_*_SETALLONES` (#169009)
abhishek-kaushik22 Nov 24, 2025
83765f4
[Utils][update_mc_test_checks] Support generating asm tests from temp…
kosarev Nov 24, 2025
d5927a6
[LLDB] Add unary plus and minus to DIL (#155617)
kuilpd Nov 24, 2025
cd13d9f
[lldb] Add test showing UnwindAssemblyInstEmulation can't handle back…
felipepiovezan Nov 24, 2025
4a567e3
[llvm][utils][lit] Fix imports in ManyTests.py example (#169328)
DavidSpickett Nov 24, 2025
24abb06
[OpenAC][CIR] func-local-declare 'copy' clause lowering (#169115)
erichkeane Nov 24, 2025
ceea07d
[libc++][forward_list] Applied `[[nodiscard]]` (#169019)
H-G-Hristov Nov 24, 2025
456b051
[VPlan] Set ZeroIsPoison=false for FirstActiveLane (#169298)
lukel97 Nov 24, 2025
1580f4b
[AArch64] Update costs for fshl/r and add rotr/l variants. NFC
davemgreen Nov 24, 2025
ad0acf4
AMDGPU/GlobalISel: Combine S16 copy-trunc-readanylane-anyext (#168410)
petar-avramovic Nov 24, 2025
71952df
[OpenMP][SPIRV] Disable exceptions for OpenMP SPIR-V (#169094)
sarnex Nov 24, 2025
d542dce
[OpenACC][CIR] copyin lowering for func-local- declare (#169336)
erichkeane Nov 24, 2025
c08f05e
merge main into amd-staging
ronlieb Nov 24, 2025
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
11 changes: 8 additions & 3 deletions clang/include/clang/Sema/SemaHLSL.h
Original file line number Diff line number Diff line change
Expand Up @@ -250,15 +250,20 @@ class SemaHLSL : public SemaBase {
const RecordType *RT);

void checkSemanticAnnotation(FunctionDecl *EntryPoint, const Decl *Param,
const HLSLAppliedSemanticAttr *SemanticAttr);
const HLSLAppliedSemanticAttr *SemanticAttr,
bool IsInput);

bool determineActiveSemanticOnScalar(FunctionDecl *FD,
DeclaratorDecl *OutputDecl,
DeclaratorDecl *D,
SemanticInfo &ActiveSemantic,
llvm::StringSet<> &ActiveInputSemantics);
llvm::StringSet<> &ActiveSemantics,
bool IsInput);

bool determineActiveSemantic(FunctionDecl *FD, DeclaratorDecl *OutputDecl,
DeclaratorDecl *D, SemanticInfo &ActiveSemantic,
llvm::StringSet<> &ActiveInputSemantics);
llvm::StringSet<> &ActiveSemantics,
bool IsInput);

void processExplicitBindingsOnDecl(VarDecl *D);

Expand Down
62 changes: 54 additions & 8 deletions clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -19,18 +19,64 @@ using namespace clang::CIRGen;

namespace {
struct OpenACCDeclareCleanup final : EHScopeStack::Cleanup {
SourceRange declareRange;
mlir::acc::DeclareEnterOp enterOp;

OpenACCDeclareCleanup(mlir::acc::DeclareEnterOp enterOp) : enterOp(enterOp) {}
OpenACCDeclareCleanup(SourceRange declareRange,
mlir::acc::DeclareEnterOp enterOp)
: declareRange(declareRange), enterOp(enterOp) {}

template <typename OutTy, typename InTy>
void createOutOp(CIRGenFunction &cgf, InTy inOp) {
if constexpr (std::is_same_v<OutTy, mlir::acc::DeleteOp>) {
auto outOp =
OutTy::create(cgf.getBuilder(), inOp.getLoc(), inOp,
inOp.getStructured(), inOp.getImplicit(),
llvm::Twine(inOp.getNameAttr()), inOp.getBounds());
outOp.setDataClause(inOp.getDataClause());
outOp.setModifiers(inOp.getModifiers());
} else {
auto outOp =
OutTy::create(cgf.getBuilder(), inOp.getLoc(), inOp, inOp.getVarPtr(),
inOp.getStructured(), inOp.getImplicit(),
llvm::Twine(inOp.getNameAttr()), inOp.getBounds());
outOp.setDataClause(inOp.getDataClause());
outOp.setModifiers(inOp.getModifiers());
}
}

void emit(CIRGenFunction &cgf) override {
mlir::acc::DeclareExitOp::create(cgf.getBuilder(), enterOp.getLoc(),
enterOp, {});
auto exitOp = mlir::acc::DeclareExitOp::create(
cgf.getBuilder(), enterOp.getLoc(), enterOp, {});

// TODO(OpenACC): Some clauses require that we add info about them to the
// DeclareExitOp. However, we don't have any of those implemented yet, so
// we should add infrastructure here to do that once we have one
// implemented.
// Some data clauses need to be referenced in 'exit', AND need to have an
// operation after the exit. Copy these from the enter operation.
for (mlir::Value val : enterOp.getDataClauseOperands()) {
if (auto copyin = val.getDefiningOp<mlir::acc::CopyinOp>()) {
switch (copyin.getDataClause()) {
default:
cgf.cgm.errorNYI(declareRange,
"OpenACC local declare clause copyin cleanup");
break;
case mlir::acc::DataClause::acc_copy:
createOutOp<mlir::acc::CopyoutOp>(cgf, copyin);
break;
case mlir::acc::DataClause::acc_copyin:
createOutOp<mlir::acc::DeleteOp>(cgf, copyin);
break;
}
} else if (val.getDefiningOp<mlir::acc::DeclareLinkOp>()) {
// Link has no exit clauses, and shouldn't be copied.
continue;
} else if (val.getDefiningOp<mlir::acc::DevicePtrOp>()) {
// DevicePtr has no exit clauses, and shouldn't be copied.
continue;
} else {
cgf.cgm.errorNYI(declareRange, "OpenACC local declare clause cleanup");
continue;
}
exitOp.getDataClauseOperandsMutable().append(val);
}
}
};
} // namespace
Expand All @@ -45,7 +91,7 @@ void CIRGenFunction::emitOpenACCDeclare(const OpenACCDeclareDecl &d) {
d.clauses());

ehStack.pushCleanup<OpenACCDeclareCleanup>(CleanupKind::NormalCleanup,
enterOp);
d.getSourceRange(), enterOp);
}

void CIRGenFunction::emitOpenACCRoutine(const OpenACCRoutineDecl &d) {
Expand Down
20 changes: 14 additions & 6 deletions clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -800,12 +800,16 @@ class OpenACCClauseCIREmitter final
var, mlir::acc::DataClause::acc_copy, clause.getModifierList(),
/*structured=*/true,
/*implicit=*/false);
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
for (const Expr *var : clause.getVarList())
addDataOperand<mlir::acc::CopyinOp>(
var, mlir::acc::DataClause::acc_copy, clause.getModifierList(),
/*structured=*/true,
/*implicit=*/false);
} else if constexpr (isCombinedType<OpTy>) {
applyToComputeOp(clause);
} else {
// TODO: When we've implemented this for everything, switch this to an
// unreachable. declare construct remains.
return clauseNotImplemented(clause);
llvm_unreachable("Unknown construct kind in VisitCopyClause");
}
}

Expand All @@ -822,12 +826,16 @@ class OpenACCClauseCIREmitter final
addDataOperand<mlir::acc::CopyinOp>(
var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(),
/*structured=*/false, /*implicit=*/false);
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
for (const Expr *var : clause.getVarList())
addDataOperand<mlir::acc::CopyinOp>(
var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(),
/*structured=*/true,
/*implicit=*/false);
} else if constexpr (isCombinedType<OpTy>) {
applyToComputeOp(clause);
} else {
// TODO: When we've implemented this for everything, switch this to an
// unreachable. declare construct remains.
return clauseNotImplemented(clause);
llvm_unreachable("Unknown construct kind in VisitCopyInClause");
}
}

Expand Down
4 changes: 2 additions & 2 deletions clang/lib/CodeGen/CGException.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -450,7 +450,7 @@ void CodeGenFunction::EmitCXXThrowExpr(const CXXThrowExpr *E,
// Therefore, we emit a trap which will abort the program, and
// prompt a warning indicating that a trap will be emitted.
const llvm::Triple &T = Target.getTriple();
if (CGM.getLangOpts().OpenMPIsTargetDevice && (T.isNVPTX() || T.isAMDGCN())) {
if (CGM.getLangOpts().OpenMPIsTargetDevice && T.isGPU()) {
EmitTrapCall(llvm::Intrinsic::trap);
return;
}
Expand Down Expand Up @@ -627,7 +627,7 @@ void CodeGenFunction::EmitCXXTryStmt(const CXXTryStmt &S) {
// If we encounter a try statement on in an OpenMP target region offloaded to
// a GPU, we treat it as a basic block.
const bool IsTargetDevice =
(CGM.getLangOpts().OpenMPIsTargetDevice && (T.isNVPTX() || T.isAMDGCN()));
(CGM.getLangOpts().OpenMPIsTargetDevice && T.isGPU());
if (!IsTargetDevice)
EnterCXXTryStmt(S);
EmitStmt(S.getTryBlock());
Expand Down
2 changes: 1 addition & 1 deletion clang/lib/CodeGen/CGExprComplex.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -320,7 +320,7 @@ class ComplexExprEmitter
QualType getPromotionType(FPOptionsOverride Features, QualType Ty,
bool IsComplexDivisor) {
if (auto *CT = Ty->getAs<ComplexType>()) {
QualType ElementType = CT->getElementType();
QualType ElementType = CT->getElementType().getCanonicalType();
bool IsFloatingType = ElementType->isFloatingType();
bool IsComplexRangePromoted = CGF.getLangOpts().getComplexRange() ==
LangOptions::ComplexRangeKind::CX_Promoted;
Expand Down
41 changes: 30 additions & 11 deletions clang/lib/CodeGen/CGHLSLRuntime.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -731,13 +731,22 @@ llvm::Value *CGHLSLRuntime::emitSystemSemanticLoad(
}

if (SemanticName == "SV_POSITION") {
if (CGM.getTriple().getEnvironment() == Triple::EnvironmentType::Pixel)
return createSPIRVBuiltinLoad(B, CGM.getModule(), Type,
Semantic->getAttrName()->getName(),
/* BuiltIn::FragCoord */ 15);
if (CGM.getTriple().getEnvironment() == Triple::EnvironmentType::Pixel) {
if (CGM.getTarget().getTriple().isSPIRV())
return createSPIRVBuiltinLoad(B, CGM.getModule(), Type,
Semantic->getAttrName()->getName(),
/* BuiltIn::FragCoord */ 15);
if (CGM.getTarget().getTriple().isDXIL())
return emitDXILUserSemanticLoad(B, Type, Semantic, Index);
}

if (CGM.getTriple().getEnvironment() == Triple::EnvironmentType::Vertex) {
return emitUserSemanticLoad(B, Type, Decl, Semantic, Index);
}
}

llvm_unreachable("non-handled system semantic. FIXME.");
llvm_unreachable(
"Load hasn't been implemented yet for this system semantic. FIXME");
}

static void createSPIRVBuiltinStore(IRBuilder<> &B, llvm::Module &M,
Expand All @@ -760,12 +769,22 @@ void CGHLSLRuntime::emitSystemSemanticStore(IRBuilder<> &B, llvm::Value *Source,
std::optional<unsigned> Index) {

std::string SemanticName = Semantic->getAttrName()->getName().upper();
if (SemanticName == "SV_POSITION")
createSPIRVBuiltinStore(B, CGM.getModule(), Source,
Semantic->getAttrName()->getName(),
/* BuiltIn::Position */ 0);
else
llvm_unreachable("non-handled system semantic. FIXME.");
if (SemanticName == "SV_POSITION") {
if (CGM.getTarget().getTriple().isDXIL()) {
emitDXILUserSemanticStore(B, Source, Semantic, Index);
return;
}

if (CGM.getTarget().getTriple().isSPIRV()) {
createSPIRVBuiltinStore(B, CGM.getModule(), Source,
Semantic->getAttrName()->getName(),
/* BuiltIn::Position */ 0);
return;
}
}

llvm_unreachable(
"Store hasn't been implemented yet for this system semantic. FIXME");
}

llvm::Value *CGHLSLRuntime::handleScalarSemanticLoad(
Expand Down
3 changes: 1 addition & 2 deletions clang/lib/Frontend/CompilerInvocation.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -4360,8 +4360,7 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args,

// Set the flag to prevent the implementation from emitting device exception
// handling code for those requiring so.
if ((Opts.OpenMPIsTargetDevice && (T.isNVPTX() || T.isAMDGCN())) ||
Opts.OpenCLCPlusPlus) {
if ((Opts.OpenMPIsTargetDevice && T.isGPU()) || Opts.OpenCLCPlusPlus) {

Opts.Exceptions = 0;
Opts.CXXExceptions = 0;
Expand Down
2 changes: 1 addition & 1 deletion clang/lib/Sema/SemaExpr.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -10743,7 +10743,7 @@ static void DetectPrecisionLossInComplexDivision(Sema &S, QualType DivisorTy,
if (!CT)
return;

QualType ElementType = CT->getElementType();
QualType ElementType = CT->getElementType().getCanonicalType();
bool IsComplexRangePromoted = S.getLangOpts().getComplexRange() ==
LangOptions::ComplexRangeKind::CX_Promoted;
if (!ElementType->isFloatingType() || !IsComplexRangePromoted)
Expand Down
35 changes: 21 additions & 14 deletions clang/lib/Sema/SemaHLSL.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -771,9 +771,12 @@ void SemaHLSL::ActOnTopLevelFunction(FunctionDecl *FD) {
}
}

bool SemaHLSL::determineActiveSemanticOnScalar(
FunctionDecl *FD, DeclaratorDecl *OutputDecl, DeclaratorDecl *D,
SemanticInfo &ActiveSemantic, llvm::StringSet<> &UsedSemantics) {
bool SemaHLSL::determineActiveSemanticOnScalar(FunctionDecl *FD,
DeclaratorDecl *OutputDecl,
DeclaratorDecl *D,
SemanticInfo &ActiveSemantic,
llvm::StringSet<> &UsedSemantics,
bool IsInput) {
if (ActiveSemantic.Semantic == nullptr) {
ActiveSemantic.Semantic = D->getAttr<HLSLParsedSemanticAttr>();
if (ActiveSemantic.Semantic)
Expand All @@ -792,7 +795,7 @@ bool SemaHLSL::determineActiveSemanticOnScalar(
if (!A)
return false;

checkSemanticAnnotation(FD, D, A);
checkSemanticAnnotation(FD, D, A, IsInput);
OutputDecl->addAttr(A);

unsigned Location = ActiveSemantic.Index.value_or(0);
Expand Down Expand Up @@ -820,7 +823,8 @@ bool SemaHLSL::determineActiveSemantic(FunctionDecl *FD,
DeclaratorDecl *OutputDecl,
DeclaratorDecl *D,
SemanticInfo &ActiveSemantic,
llvm::StringSet<> &UsedSemantics) {
llvm::StringSet<> &UsedSemantics,
bool IsInput) {
if (ActiveSemantic.Semantic == nullptr) {
ActiveSemantic.Semantic = D->getAttr<HLSLParsedSemanticAttr>();
if (ActiveSemantic.Semantic)
Expand All @@ -833,12 +837,13 @@ bool SemaHLSL::determineActiveSemantic(FunctionDecl *FD,
const RecordType *RT = dyn_cast<RecordType>(T);
if (!RT)
return determineActiveSemanticOnScalar(FD, OutputDecl, D, ActiveSemantic,
UsedSemantics);
UsedSemantics, IsInput);

const RecordDecl *RD = RT->getDecl();
for (FieldDecl *Field : RD->fields()) {
SemanticInfo Info = ActiveSemantic;
if (!determineActiveSemantic(FD, OutputDecl, Field, Info, UsedSemantics)) {
if (!determineActiveSemantic(FD, OutputDecl, Field, Info, UsedSemantics,
IsInput)) {
Diag(Field->getLocation(), diag::note_hlsl_semantic_used_here) << Field;
return false;
}
Expand Down Expand Up @@ -920,7 +925,7 @@ void SemaHLSL::CheckEntryPoint(FunctionDecl *FD) {

// FIXME: Verify output semantics in parameters.
if (!determineActiveSemantic(FD, Param, Param, ActiveSemantic,
ActiveInputSemantics)) {
ActiveInputSemantics, /* IsInput= */ true)) {
Diag(Param->getLocation(), diag::note_previous_decl) << Param;
FD->setInvalidDecl();
}
Expand All @@ -932,12 +937,13 @@ void SemaHLSL::CheckEntryPoint(FunctionDecl *FD) {
if (ActiveSemantic.Semantic)
ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
if (!FD->getReturnType()->isVoidType())
determineActiveSemantic(FD, FD, FD, ActiveSemantic, ActiveOutputSemantics);
determineActiveSemantic(FD, FD, FD, ActiveSemantic, ActiveOutputSemantics,
/* IsInput= */ false);
}

void SemaHLSL::checkSemanticAnnotation(
FunctionDecl *EntryPoint, const Decl *Param,
const HLSLAppliedSemanticAttr *SemanticAttr) {
const HLSLAppliedSemanticAttr *SemanticAttr, bool IsInput) {
auto *ShaderAttr = EntryPoint->getAttr<HLSLShaderAttr>();
assert(ShaderAttr && "Entry point has no shader attribute");
llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
Expand All @@ -961,11 +967,12 @@ void SemaHLSL::checkSemanticAnnotation(
}

if (SemanticName == "SV_POSITION") {
// TODO(#143523): allow use on other shader types & output once the overall
// semantic logic is implemented.
if (ST == llvm::Triple::Pixel)
// SV_Position can be an input or output in vertex shaders,
// but only an input in pixel shaders.
if (ST == llvm::Triple::Vertex || (ST == llvm::Triple::Pixel && IsInput))
return;
DiagnoseAttrStageMismatch(SemanticAttr, ST, {llvm::Triple::Pixel});
DiagnoseAttrStageMismatch(SemanticAttr, ST,
{llvm::Triple::Pixel, llvm::Triple::Vertex});
return;
}

Expand Down
21 changes: 21 additions & 0 deletions clang/test/AST/HLSL/semantic-input-struct-shadow.hlsl
Original file line number Diff line number Diff line change
@@ -0,0 +1,21 @@
// RUN: %clang_cc1 -triple spirv-unknown-vulkan1.3-vertex -finclude-default-header -ast-dump -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple dxil-pc-shadermodel6.8-vertex -finclude-default-header -ast-dump -o - %s | FileCheck %s


// CHECK: CXXRecordDecl {{.*}} referenced struct S definition
// CHECK: FieldDecl {{.*}} field1 'int'
// CHECK-NEXT: HLSLParsedSemanticAttr {{.*}} "A" 0
// CHECK: FieldDecl {{.*}} field2 'int'
// CHECK-NEXT: HLSLParsedSemanticAttr {{.*}} "B" 4

struct S {
int field1 : A;
int field2 : B4;
};

// CHECK: FunctionDecl {{.*}} main 'void (S)'
// CHECK-NEXT: ParmVarDecl {{.*}} s 'S'
// CHECK-NEXT: HLSLParsedSemanticAttr {{.*}} "C" 0
// CHECK-NEXT: HLSLAppliedSemanticAttr {{.*}} "C" 0
// CHECK-NEXT: HLSLAppliedSemanticAttr {{.*}} "C" 1
void main(S s : C) {}
20 changes: 20 additions & 0 deletions clang/test/AST/HLSL/semantic-input-struct.hlsl
Original file line number Diff line number Diff line change
@@ -0,0 +1,20 @@
// RUN: %clang_cc1 -triple spirv-unknown-vulkan1.3-vertex -finclude-default-header -ast-dump -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple dxil-pc-shadermodel6.8-vertex -finclude-default-header -ast-dump -o - %s | FileCheck %s


// CHECK: CXXRecordDecl {{.*}} referenced struct S definition
// CHECK: FieldDecl {{.*}} field1 'int'
// CHECK-NEXT: HLSLParsedSemanticAttr {{.*}} "A" 0
// CHECK: FieldDecl {{.*}} field2 'int'
// CHECK-NEXT: HLSLParsedSemanticAttr {{.*}} "B" 4

struct S {
int field1 : A;
int field2 : B4;
};

// CHECK: FunctionDecl {{.*}} main 'void (S)'
// CHECK-NEXT: ParmVarDecl {{.*}} s 'S'
// CHECK-NEXT: HLSLAppliedSemanticAttr {{.*}} "A" 0
// CHECK-NEXT: HLSLAppliedSemanticAttr {{.*}} "B" 4
void main(S s) {}
9 changes: 9 additions & 0 deletions clang/test/AST/HLSL/semantic-input.hlsl
Original file line number Diff line number Diff line change
@@ -0,0 +1,9 @@
// RUN: %clang_cc1 -triple spirv-unknown-vulkan1.3-vertex -finclude-default-header -ast-dump -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple dxil-pc-shadermodel6.8-vertex -finclude-default-header -ast-dump -o - %s | FileCheck %s

// CHECK: ParmVarDecl {{.*}} a 'float4':'vector<float, 4>'
// CHECK-NEXT: HLSLParsedSemanticAttr {{.*}} "ABC" 0
// CHECK-NEXT: HLSLAppliedSemanticAttr {{.*}} "ABC" 0

void main(float4 a : ABC) {
}
Loading