Skip to content

Conversation

@Keenuts
Copy link
Contributor

@Keenuts Keenuts commented Sep 16, 2025

This PR is an incremental improvement regarding semantics I/O in HLSL. This PR allows
system semantics to be used on struct type in addition to parameters (state today).
This PR doesn't consider implicit indexing increment that happens when placing a semantic on an aggregate/array as implemented system semantics don't allow such use yet.

The next step will be to enable user semantics, which will bring the need to properly determine semantic indices depending on context.
This PR diverge from the initial wg-hlsl proposal as all diagnostics are done in Sema (initial proposal suggested running diags in codegen).

This is not yet a solid semantic implementation, but increases the test coverage and improves the status from where we are now.

@Keenuts Keenuts requested a review from s-perron September 16, 2025 11:42
@llvmbot llvmbot added clang Clang issues not falling into any other category clang:frontend Language frontend issues, e.g. anything involving "Sema" clang:codegen IR generation bugs: mangling, exceptions, etc. HLSL HLSL Language Support labels Sep 16, 2025
@llvmbot
Copy link
Member

llvmbot commented Sep 16, 2025

@llvm/pr-subscribers-clang-codegen

@llvm/pr-subscribers-clang

Author: Nathan Gauër (Keenuts)

Changes

This PR is an incremental improvement regarding semantics I/O in HLSL. This PR allows
system semantics to be used on struct type in addition to parameters (state today).
This PR doesn't consider implicit indexing increment that happens when placing a semantic on an aggregate/array as implemented system semantics don't allow such use yet.

The next step will be to enable user semantics, which will bring the need to properly determine semantic indices depending on context.
This PR diverge from the initial wg-hlsl proposal as all diagnostics are done in Sema (initial proposal suggested running diags in codegen).

This is not yet a solid semantic implementation, but increases the test coverage and improves the status from where we are now.


Patch is 25.93 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/159047.diff

12 Files Affected:

  • (modified) clang/include/clang/AST/Attr.h (+2)
  • (modified) clang/include/clang/Basic/DiagnosticFrontendKinds.td (-4)
  • (modified) clang/include/clang/Basic/DiagnosticSemaKinds.td (+1)
  • (modified) clang/include/clang/Sema/SemaHLSL.h (+16-5)
  • (modified) clang/lib/CodeGen/CGHLSLRuntime.cpp (+60-28)
  • (modified) clang/lib/CodeGen/CGHLSLRuntime.h (+7-10)
  • (modified) clang/lib/Sema/SemaHLSL.cpp (+74-16)
  • (added) clang/test/CodeGenHLSL/semantics/semantic-struct-1.hlsl (+23)
  • (added) clang/test/CodeGenHLSL/semantics/semantic-struct-2.hlsl (+25)
  • (added) clang/test/CodeGenHLSL/semantics/semantic-struct-nested-inherit.hlsl (+30)
  • (added) clang/test/CodeGenHLSL/semantics/semantic-struct-nested-shadow.hlsl (+30)
  • (added) clang/test/CodeGenHLSL/semantics/semantic-struct-nested.hlsl (+30)
diff --git a/clang/include/clang/AST/Attr.h b/clang/include/clang/AST/Attr.h
index fe388b9fa045e..00d4a0035fccf 100644
--- a/clang/include/clang/AST/Attr.h
+++ b/clang/include/clang/AST/Attr.h
@@ -259,6 +259,8 @@ class HLSLSemanticAttr : public HLSLAnnotationAttr {
 
   unsigned getSemanticIndex() const { return SemanticIndex; }
 
+  bool isSemanticIndexExplicit() const { return SemanticExplicitIndex; }
+
   // Implement isa/cast/dyncast/etc.
   static bool classof(const Attr *A) {
     return A->getKind() >= attr::FirstHLSLSemanticAttr &&
diff --git a/clang/include/clang/Basic/DiagnosticFrontendKinds.td b/clang/include/clang/Basic/DiagnosticFrontendKinds.td
index 2fd2ae434d7c5..15447558cf952 100644
--- a/clang/include/clang/Basic/DiagnosticFrontendKinds.td
+++ b/clang/include/clang/Basic/DiagnosticFrontendKinds.td
@@ -400,10 +400,6 @@ def warn_hlsl_langstd_minimal :
           "recommend using %1 instead">,
   InGroup<HLSLDXCCompat>;
 
-def err_hlsl_semantic_missing : Error<"semantic annotations must be present "
-                                      "for all input and outputs of an entry "
-                                      "function or patch constant function">;
-
 // ClangIR frontend errors
 def err_cir_to_cir_transform_failed : Error<
     "CIR-to-CIR transformation failed">, DefaultFatal;
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 757404a3f5eac..6950fdd994b3d 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -13123,6 +13123,7 @@ def err_hlsl_duplicate_parameter_modifier : Error<"duplicate parameter modifier
 def err_hlsl_missing_semantic_annotation : Error<
   "semantic annotations must be present for all parameters of an entry "
   "function or patch constant function">;
+def note_hlsl_semantic_used_here : Note<"%0 used here">;
 def err_hlsl_unknown_semantic : Error<"unknown HLSL semantic %0">;
 def err_hlsl_semantic_output_not_supported
     : Error<"semantic %0 does not support output">;
diff --git a/clang/include/clang/Sema/SemaHLSL.h b/clang/include/clang/Sema/SemaHLSL.h
index b5ddca0fe2ca5..f0a940353515c 100644
--- a/clang/include/clang/Sema/SemaHLSL.h
+++ b/clang/include/clang/Sema/SemaHLSL.h
@@ -130,9 +130,6 @@ class SemaHLSL : public SemaBase {
   bool ActOnUninitializedVarDecl(VarDecl *D);
   void ActOnEndOfTranslationUnit(TranslationUnitDecl *TU);
   void CheckEntryPoint(FunctionDecl *FD);
-  bool isSemanticValid(FunctionDecl *FD, DeclaratorDecl *D);
-  void CheckSemanticAnnotation(FunctionDecl *EntryPoint, const Decl *Param,
-                               const HLSLAnnotationAttr *AnnotationAttr);
   void DiagnoseAttrStageMismatch(
       const Attr *A, llvm::Triple::EnvironmentType Stage,
       std::initializer_list<llvm::Triple::EnvironmentType> AllowedStages);
@@ -177,9 +174,9 @@ class SemaHLSL : public SemaBase {
   bool handleResourceTypeAttr(QualType T, const ParsedAttr &AL);
 
   template <typename T>
-  T *createSemanticAttr(const ParsedAttr &AL,
+  T *createSemanticAttr(const AttributeCommonInfo &ACI,
                         std::optional<unsigned> Location) {
-    T *Attr = ::new (getASTContext()) T(getASTContext(), AL);
+    T *Attr = ::new (getASTContext()) T(getASTContext(), ACI);
     if (Attr->isSemanticIndexable())
       Attr->setSemanticIndex(Location ? *Location : 0);
     else if (Location.has_value()) {
@@ -246,10 +243,24 @@ class SemaHLSL : public SemaBase {
 
   IdentifierInfo *RootSigOverrideIdent = nullptr;
 
+  struct SemanticInfo {
+    HLSLSemanticAttr *Semantic;
+    std::optional<uint32_t> Index;
+  };
+
 private:
   void collectResourceBindingsOnVarDecl(VarDecl *D);
   void collectResourceBindingsOnUserRecordDecl(const VarDecl *VD,
                                                const RecordType *RT);
+
+  void checkSemanticAnnotation(FunctionDecl *EntryPoint, const Decl *Param,
+                               const HLSLSemanticAttr *SemanticAttr);
+  HLSLSemanticAttr *createSemantic(const SemanticInfo &Semantic);
+  bool isSemanticOnScalarValid(FunctionDecl *FD, DeclaratorDecl *D,
+                               SemanticInfo &ActiveSemantic);
+  bool isSemanticValid(FunctionDecl *FD, DeclaratorDecl *D,
+                       SemanticInfo &ActiveSemantic);
+
   void processExplicitBindingsOnDecl(VarDecl *D);
 
   void diagnoseAvailabilityViolations(TranslationUnitDecl *TU);
diff --git a/clang/lib/CodeGen/CGHLSLRuntime.cpp b/clang/lib/CodeGen/CGHLSLRuntime.cpp
index afee1198e0988..c68d85ac837f3 100644
--- a/clang/lib/CodeGen/CGHLSLRuntime.cpp
+++ b/clang/lib/CodeGen/CGHLSLRuntime.cpp
@@ -566,17 +566,16 @@ static llvm::Value *createSPIRVBuiltinLoad(IRBuilder<> &B, llvm::Module &M,
   return B.CreateLoad(Ty, GV);
 }
 
-llvm::Value *
-CGHLSLRuntime::emitSystemSemanticLoad(IRBuilder<> &B, llvm::Type *Type,
-                                      const clang::DeclaratorDecl *Decl,
-                                      SemanticInfo &ActiveSemantic) {
-  if (isa<HLSLSV_GroupIndexAttr>(ActiveSemantic.Semantic)) {
+llvm::Value *CGHLSLRuntime::emitSystemSemanticLoad(
+    IRBuilder<> &B, llvm::Type *Type, const clang::DeclaratorDecl *Decl,
+    Attr *Semantic, std::optional<unsigned> Index) {
+  if (isa<HLSLSV_GroupIndexAttr>(Semantic)) {
     llvm::Function *GroupIndex =
         CGM.getIntrinsic(getFlattenedThreadIdInGroupIntrinsic());
     return B.CreateCall(FunctionCallee(GroupIndex));
   }
 
-  if (isa<HLSLSV_DispatchThreadIDAttr>(ActiveSemantic.Semantic)) {
+  if (isa<HLSLSV_DispatchThreadIDAttr>(Semantic)) {
     llvm::Intrinsic::ID IntrinID = getThreadIdIntrinsic();
     llvm::Function *ThreadIDIntrinsic =
         llvm::Intrinsic::isOverloaded(IntrinID)
@@ -585,7 +584,7 @@ CGHLSLRuntime::emitSystemSemanticLoad(IRBuilder<> &B, llvm::Type *Type,
     return buildVectorInput(B, ThreadIDIntrinsic, Type);
   }
 
-  if (isa<HLSLSV_GroupThreadIDAttr>(ActiveSemantic.Semantic)) {
+  if (isa<HLSLSV_GroupThreadIDAttr>(Semantic)) {
     llvm::Intrinsic::ID IntrinID = getGroupThreadIdIntrinsic();
     llvm::Function *GroupThreadIDIntrinsic =
         llvm::Intrinsic::isOverloaded(IntrinID)
@@ -594,7 +593,7 @@ CGHLSLRuntime::emitSystemSemanticLoad(IRBuilder<> &B, llvm::Type *Type,
     return buildVectorInput(B, GroupThreadIDIntrinsic, Type);
   }
 
-  if (isa<HLSLSV_GroupIDAttr>(ActiveSemantic.Semantic)) {
+  if (isa<HLSLSV_GroupIDAttr>(Semantic)) {
     llvm::Intrinsic::ID IntrinID = getGroupIdIntrinsic();
     llvm::Function *GroupIDIntrinsic =
         llvm::Intrinsic::isOverloaded(IntrinID)
@@ -603,8 +602,7 @@ CGHLSLRuntime::emitSystemSemanticLoad(IRBuilder<> &B, llvm::Type *Type,
     return buildVectorInput(B, GroupIDIntrinsic, Type);
   }
 
-  if (HLSLSV_PositionAttr *S =
-          dyn_cast<HLSLSV_PositionAttr>(ActiveSemantic.Semantic)) {
+  if (HLSLSV_PositionAttr *S = dyn_cast<HLSLSV_PositionAttr>(Semantic)) {
     if (CGM.getTriple().getEnvironment() == Triple::EnvironmentType::Pixel)
       return createSPIRVBuiltinLoad(B, CGM.getModule(), Type,
                                     S->getAttrName()->getName(),
@@ -616,28 +614,45 @@ CGHLSLRuntime::emitSystemSemanticLoad(IRBuilder<> &B, llvm::Type *Type,
 
 llvm::Value *
 CGHLSLRuntime::handleScalarSemanticLoad(IRBuilder<> &B, llvm::Type *Type,
-                                        const clang::DeclaratorDecl *Decl,
-                                        SemanticInfo &ActiveSemantic) {
-
-  if (!ActiveSemantic.Semantic) {
-    ActiveSemantic.Semantic = Decl->getAttr<HLSLSemanticAttr>();
-    if (!ActiveSemantic.Semantic) {
-      CGM.getDiags().Report(Decl->getInnerLocStart(),
-                            diag::err_hlsl_semantic_missing);
-      return nullptr;
-    }
-    ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
+                                        const clang::DeclaratorDecl *Decl) {
+  HLSLSemanticAttr *Semantic = Decl->getAttr<HLSLSemanticAttr>();
+  // Sema either attached a semantic to each field/param, or raised an error.
+  assert(Semantic);
+
+  std::optional<unsigned> Index = std::nullopt;
+  if (Semantic->isSemanticIndexExplicit())
+    Index = Semantic->getSemanticIndex();
+  return emitSystemSemanticLoad(B, Type, Decl, Semantic, Index);
+}
+
+llvm::Value *
+CGHLSLRuntime::handleStructSemanticLoad(IRBuilder<> &B, llvm::Type *Type,
+                                        const clang::DeclaratorDecl *Decl) {
+  const llvm::StructType *ST = cast<StructType>(Type);
+  const clang::RecordDecl *RD = Decl->getType()->getAsRecordDecl();
+
+  assert(std::distance(RD->field_begin(), RD->field_end()) ==
+         ST->getNumElements());
+
+  llvm::Value *Aggregate = llvm::PoisonValue::get(Type);
+  auto FieldDecl = RD->field_begin();
+  for (unsigned I = 0; I < ST->getNumElements(); ++I) {
+    llvm::Value *ChildValue =
+        handleSemanticLoad(B, ST->getElementType(I), *FieldDecl);
+    assert(ChildValue);
+    Aggregate = B.CreateInsertValue(Aggregate, ChildValue, I);
+    ++FieldDecl;
   }
 
-  return emitSystemSemanticLoad(B, Type, Decl, ActiveSemantic);
+  return Aggregate;
 }
 
 llvm::Value *
 CGHLSLRuntime::handleSemanticLoad(IRBuilder<> &B, llvm::Type *Type,
-                                  const clang::DeclaratorDecl *Decl,
-                                  SemanticInfo &ActiveSemantic) {
-  assert(!Type->isStructTy());
-  return handleScalarSemanticLoad(B, Type, Decl, ActiveSemantic);
+                                  const clang::DeclaratorDecl *Decl) {
+  if (Type->isStructTy())
+    return handleStructSemanticLoad(B, Type, Decl);
+  return handleScalarSemanticLoad(B, Type, Decl);
 }
 
 void CGHLSLRuntime::emitEntryFunction(const FunctionDecl *FD,
@@ -684,8 +699,25 @@ void CGHLSLRuntime::emitEntryFunction(const FunctionDecl *FD,
     }
 
     const ParmVarDecl *PD = FD->getParamDecl(Param.getArgNo() - SRetOffset);
-    SemanticInfo ActiveSemantic = {nullptr, 0};
-    Args.push_back(handleSemanticLoad(B, Param.getType(), PD, ActiveSemantic));
+    llvm::Value *SemanticValue = nullptr;
+    if ([[maybe_unused]] HLSLParamModifierAttr *MA =
+            PD->getAttr<HLSLParamModifierAttr>()) {
+      llvm_unreachable("Not handled yet");
+    } else {
+      llvm::Type *ParamType =
+          Param.hasByValAttr() ? Param.getParamByValType() : Param.getType();
+      SemanticValue = handleSemanticLoad(B, ParamType, PD);
+      if (!SemanticValue)
+        return;
+      if (Param.hasByValAttr()) {
+        llvm::Value *Var = B.CreateAlloca(Param.getParamByValType());
+        B.CreateStore(SemanticValue, Var);
+        SemanticValue = Var;
+      }
+    }
+
+    assert(SemanticValue);
+    Args.push_back(SemanticValue);
   }
 
   CallInst *CI = B.CreateCall(FunctionCallee(Fn), Args, OB);
diff --git a/clang/lib/CodeGen/CGHLSLRuntime.h b/clang/lib/CodeGen/CGHLSLRuntime.h
index 370f3d5c5d30d..039b881d2c9ee 100644
--- a/clang/lib/CodeGen/CGHLSLRuntime.h
+++ b/clang/lib/CodeGen/CGHLSLRuntime.h
@@ -144,22 +144,19 @@ class CGHLSLRuntime {
                             llvm::Type *Type,
                             SmallVectorImpl<llvm::Value *> &Inputs);
 
-  struct SemanticInfo {
-    clang::HLSLSemanticAttr *Semantic;
-    uint32_t Index;
-  };
-
   llvm::Value *emitSystemSemanticLoad(llvm::IRBuilder<> &B, llvm::Type *Type,
                                       const clang::DeclaratorDecl *Decl,
-                                      SemanticInfo &ActiveSemantic);
+                                      Attr *Semantic,
+                                      std::optional<unsigned> Index);
 
   llvm::Value *handleScalarSemanticLoad(llvm::IRBuilder<> &B, llvm::Type *Type,
-                                        const clang::DeclaratorDecl *Decl,
-                                        SemanticInfo &ActiveSemantic);
+                                        const clang::DeclaratorDecl *Decl);
+
+  llvm::Value *handleStructSemanticLoad(llvm::IRBuilder<> &B, llvm::Type *Type,
+                                        const clang::DeclaratorDecl *Decl);
 
   llvm::Value *handleSemanticLoad(llvm::IRBuilder<> &B, llvm::Type *Type,
-                                  const clang::DeclaratorDecl *Decl,
-                                  SemanticInfo &ActiveSemantic);
+                                  const clang::DeclaratorDecl *Decl);
 
 public:
   CGHLSLRuntime(CodeGenModule &CGM) : CGM(CGM) {}
diff --git a/clang/lib/Sema/SemaHLSL.cpp b/clang/lib/Sema/SemaHLSL.cpp
index 0af38472b0fec..ee4088d528c98 100644
--- a/clang/lib/Sema/SemaHLSL.cpp
+++ b/clang/lib/Sema/SemaHLSL.cpp
@@ -769,23 +769,78 @@ void SemaHLSL::ActOnTopLevelFunction(FunctionDecl *FD) {
   }
 }
 
-bool SemaHLSL::isSemanticValid(FunctionDecl *FD, DeclaratorDecl *D) {
-  const auto *AnnotationAttr = D->getAttr<HLSLAnnotationAttr>();
-  if (AnnotationAttr) {
-    CheckSemanticAnnotation(FD, D, AnnotationAttr);
-    return true;
+HLSLSemanticAttr *SemaHLSL::createSemantic(const SemanticInfo &Info) {
+  std::string SemanticName = Info.Semantic->getAttrName()->getName().upper();
+
+  if (SemanticName == "SV_DISPATCHTHREADID") {
+    return createSemanticAttr<HLSLSV_DispatchThreadIDAttr>(*Info.Semantic,
+                                                           Info.Index);
+  } else if (SemanticName == "SV_GROUPINDEX") {
+    return createSemanticAttr<HLSLSV_GroupIndexAttr>(*Info.Semantic,
+                                                     Info.Index);
+  } else if (SemanticName == "SV_GROUPTHREADID") {
+    return createSemanticAttr<HLSLSV_GroupThreadIDAttr>(*Info.Semantic,
+                                                        Info.Index);
+  } else if (SemanticName == "SV_GROUPID") {
+    return createSemanticAttr<HLSLSV_GroupIDAttr>(*Info.Semantic, Info.Index);
+  } else if (SemanticName == "SV_POSITION") {
+    return createSemanticAttr<HLSLSV_PositionAttr>(*Info.Semantic, Info.Index);
+  } else
+    Diag(Info.Semantic->getLoc(), diag::err_hlsl_unknown_semantic)
+        << *Info.Semantic;
+
+  return nullptr;
+}
+
+bool SemaHLSL::isSemanticOnScalarValid(FunctionDecl *FD, DeclaratorDecl *D,
+                                       SemanticInfo &ActiveSemantic) {
+  if (ActiveSemantic.Semantic == nullptr) {
+    ActiveSemantic.Semantic = D->getAttr<HLSLSemanticAttr>();
+    if (ActiveSemantic.Semantic &&
+        ActiveSemantic.Semantic->isSemanticIndexExplicit())
+      ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
+  }
+
+  if (!ActiveSemantic.Semantic) {
+    Diag(D->getLocation(), diag::err_hlsl_missing_semantic_annotation);
+    return false;
+  }
+
+  auto *A = createSemantic(ActiveSemantic);
+  if (!A)
+    return false;
+
+  checkSemanticAnnotation(FD, D, A);
+  D->dropAttrs<HLSLSemanticAttr>();
+  D->addAttr(A);
+  return true;
+}
+
+bool SemaHLSL::isSemanticValid(FunctionDecl *FD, DeclaratorDecl *D,
+                               SemanticInfo &ActiveSemantic) {
+  if (ActiveSemantic.Semantic == nullptr) {
+    ActiveSemantic.Semantic = D->getAttr<HLSLSemanticAttr>();
+    if (ActiveSemantic.Semantic &&
+        ActiveSemantic.Semantic->isSemanticIndexExplicit())
+      ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
   }
 
   const Type *T = D->getType()->getUnqualifiedDesugaredType();
   const RecordType *RT = dyn_cast<RecordType>(T);
   if (!RT)
-    return false;
+    return isSemanticOnScalarValid(FD, D, ActiveSemantic);
 
   const RecordDecl *RD = RT->getOriginalDecl();
   for (FieldDecl *Field : RD->fields()) {
-    if (!isSemanticValid(FD, Field))
+    SemanticInfo Info = ActiveSemantic;
+    if (!isSemanticValid(FD, Field, Info)) {
+      Diag(Field->getLocation(), diag::note_hlsl_semantic_used_here) << Field;
       return false;
+    }
+    if (ActiveSemantic.Semantic)
+      ActiveSemantic = Info;
   }
+
   return true;
 }
 
@@ -852,8 +907,11 @@ void SemaHLSL::CheckEntryPoint(FunctionDecl *FD) {
   }
 
   for (ParmVarDecl *Param : FD->parameters()) {
-    if (!isSemanticValid(FD, Param)) {
-      Diag(FD->getLocation(), diag::err_hlsl_missing_semantic_annotation);
+    SemanticInfo ActiveSemantic;
+    ActiveSemantic.Semantic = nullptr;
+    ActiveSemantic.Index = std::nullopt;
+
+    if (!isSemanticValid(FD, Param, ActiveSemantic)) {
       Diag(Param->getLocation(), diag::note_previous_decl) << Param;
       FD->setInvalidDecl();
     }
@@ -861,31 +919,31 @@ void SemaHLSL::CheckEntryPoint(FunctionDecl *FD) {
   // FIXME: Verify return type semantic annotation.
 }
 
-void SemaHLSL::CheckSemanticAnnotation(
-    FunctionDecl *EntryPoint, const Decl *Param,
-    const HLSLAnnotationAttr *AnnotationAttr) {
+void SemaHLSL::checkSemanticAnnotation(FunctionDecl *EntryPoint,
+                                       const Decl *Param,
+                                       const HLSLSemanticAttr *SemanticAttr) {
   auto *ShaderAttr = EntryPoint->getAttr<HLSLShaderAttr>();
   assert(ShaderAttr && "Entry point has no shader attribute");
   llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
 
-  switch (AnnotationAttr->getKind()) {
+  switch (SemanticAttr->getKind()) {
   case attr::HLSLSV_DispatchThreadID:
   case attr::HLSLSV_GroupIndex:
   case attr::HLSLSV_GroupThreadID:
   case attr::HLSLSV_GroupID:
     if (ST == llvm::Triple::Compute)
       return;
-    DiagnoseAttrStageMismatch(AnnotationAttr, ST, {llvm::Triple::Compute});
+    DiagnoseAttrStageMismatch(SemanticAttr, ST, {llvm::Triple::Compute});
     break;
   case attr::HLSLSV_Position:
     // TODO(#143523): allow use on other shader types & output once the overall
     // semantic logic is implemented.
     if (ST == llvm::Triple::Pixel)
       return;
-    DiagnoseAttrStageMismatch(AnnotationAttr, ST, {llvm::Triple::Pixel});
+    DiagnoseAttrStageMismatch(SemanticAttr, ST, {llvm::Triple::Pixel});
     break;
   default:
-    llvm_unreachable("Unknown HLSLAnnotationAttr");
+    llvm_unreachable("Unknown SemanticAttr");
   }
 }
 
diff --git a/clang/test/CodeGenHLSL/semantics/semantic-struct-1.hlsl b/clang/test/CodeGenHLSL/semantics/semantic-struct-1.hlsl
new file mode 100644
index 0000000000000..ddd0baed41f37
--- /dev/null
+++ b/clang/test/CodeGenHLSL/semantics/semantic-struct-1.hlsl
@@ -0,0 +1,23 @@
+// RUN: %clang_cc1 -triple dxil-pc-shadermodel6.3-library -x hlsl -emit-llvm -finclude-default-header -disable-llvm-passes -o - %s | FileCheck %s --check-prefixes=CHECK,CHECK-DXIL -DTARGET=dx
+// RUN: %clang_cc1 -triple spirv-linux-vulkan-library -x hlsl -emit-llvm -finclude-default-header -disable-llvm-passes -o - %s | FileCheck %s --check-prefixes=CHECK,CHECK-SPIRV -DTARGET=spv
+
+
+struct Input {
+  uint Idx : SV_DispatchThreadID;
+
+};
+
+// Make sure SV_DispatchThreadID translated into dx.thread.id.
+
+// CHECK:       define void @foo()
+// CHECK-DXIL:  %[[#ID:]] = call i32 @llvm.[[TARGET]].thread.id(i32 0)
+// CHECK-SPIRV: %[[#ID:]] = call i32 @llvm.[[TARGET]].thread.id.i32(i32 0)
+// CHECK:      %[[#TMP:]] = insertvalue %struct.Input poison, i32 %[[#ID]], 0
+// CHECK:      %[[#VAR:]] = alloca %struct.Input, align 8
+// CHECK:                   store %struct.Input %[[#TMP]], ptr %[[#VAR]], align 4
+// CHECK-DXIL:              call void @{{.*}}foo{{.*}}(ptr %[[#VAR]])
+// CHECK-SPIRV:             call spir_func void @{{.*}}foo{{.*}}(ptr %[[#VAR]])
+[shader("compute")]
+[numthreads(8,8,1)]
+void foo(Input input) {}
+
diff --git a/clang/test/CodeGenHLSL/semantics/semantic-struct-2.hlsl b/clang/test/CodeGenHLSL/semantics/semantic-struct-2.hlsl
new file mode 100644
index 0000000000000..0d9c91e746454
--- /dev/null
+++ b/clang/test/CodeGenHLSL/semantics/semantic-struct-2.hlsl
@@ -0,0 +1,25 @@
+// RUN: %clang_cc1 -triple dxil-pc-shadermodel6.3-library -x hlsl -emit-llvm -finclude-default-header -disable-llvm-passes -o - %s | FileCheck %s --check-prefixes=CHECK,CHECK-DXIL -DTARGET=dx
+// RUN: %clang_cc1 -triple spirv-linux-vulkan-library -x hlsl -emit-llvm -finclude-default-header -disable-llvm-passes -o - %s | FileCheck %s --check-prefixes=CHECK,CHECK-SPIRV -DTARGET=spv
+
+
+struct Input {
+  uint Idx : SV_DispatchThreadID;
+  uint Gid : SV_GroupID;
+};
+
+// Make sure SV_DispatchThreadID translated into dx.thread.id.
+
+// CHECK:       define void @foo()
+// CHECK-DXIL:  %[[#ID:]] = call i32 @llvm.[[TARGET]].thread.id(i32 0)
+// CHECK-SPIRV: %[[#ID:]] = call i32 @llvm.[[TARGET]].thread.id.i32(i32 0)
+// CHECK:     %[[#TMP1:]] = insertvalue %struct.Input poison, i32 %[[#ID]], 0
+// CHECK-DXIL: %[[#GID:]] = call i32 @llvm...
[truncated]

Copy link
Contributor

@s-perron s-perron left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

You should have some AST test to make sure the attribute is correctly added to the entry point. I tried looking at an example, and the connection between the attribute on the function parameter does not seem obvious. That might need to be better designed.

EDIT: I was looking at the ast-dump, which was showing an attribute without any parameters. Looking closer, I saw that you add the decl to the attribute. That creates an obvious link. We just need to make sure the ast-dump reflects that.

if (!A)
return false;

checkSemanticAnnotation(FD, D, A);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Does not have to be fixed in this PR, but should this function return false if the check fails?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We could yes, but still emitting a diagnostic will make the compilation fail, so here we add an attribute to something we are going to discard anyway.

@Keenuts
Copy link
Contributor Author

Keenuts commented Sep 19, 2025

EDIT: I was looking at the ast-dump, which was showing an attribute without any parameters. Looking closer, I saw that you add the decl to the attribute. That creates an obvious link. We just need to make sure the ast-dump reflects that.

Modified the code so ast-dump shows this correctly. We already have some AST checks but expanded them to check the newly generated attributes on the entrypoint

@Keenuts Keenuts requested a review from s-perron September 19, 2025 18:29
Copy link
Contributor

@s-perron s-perron left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This LGTM, but there are others will have to approve.

@Keenuts Keenuts requested review from llvm-beanz and tex3d September 22, 2025 13:22
Copy link
Collaborator

@llvm-beanz llvm-beanz left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We should have tests for some cases like:

struct S {
  float4 f0 : SV_Position;
  float4 f1;  // error due to missing semantic here!
};

float4 main(S s) { }

// Semantics on return objects are required but not supported yet (right?)
float4 main2(float4 p : SV_POSITION, float4 f) { } // error for missing semantic on `f`

@Keenuts
Copy link
Contributor Author

Keenuts commented Oct 15, 2025

We should have tests for some cases like:

struct S {
  float4 f0 : SV_Position;
  float4 f1;  // error due to missing semantic here!
};

float4 main(S s) { }

// Semantics on return objects are required but not supported yet (right?)
float4 main2(float4 p : SV_POSITION, float4 f) { } // error for missing semantic on `f`

Yes, I'll add several sema checks now that those are validated in sema.

Copy link
Contributor

@tex3d tex3d left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'm approving this, but think the approach will need structural changes to handle the general struct case.

Specifically:

I think the attribute added to the function decl to describe a loaded/stored element should be a different attribute than the semantic attribute. It's a shader I/O element, aka. signature element, constructed from a path through parameter space with a semantic applied.

Another issue is that the user adds an explicit semantic to the function decl which gets applied to the return value, so adding more copies of semantics for other things on the same decl when traversing parameters and structs in parameters makes it hard to differentiate an explicit semantic for the return value from semantics inferred from interpreting parameters and fields in types.

I don't think pointing at the leaf field in a type is enough. You'd need the full path through the decls that reaches particular instance of the leaf field, since that type could be used multiple times in the shader inputs/outputs/return type, and containing types. Recording a full path of decls from the param to the field is one approach, but I don't think it's necessarily the right one.

I think a different approach for storing interpreted signature elements would be better. We could instead add signature element attributes to each parameter decl or to the function for the return type, in order of traversal to leaf fields. These would form a sort of worklist of items to generate code for loading or storing as you traverse each parameter type. Once one is processed, it's only looking at the next (field) that applies to the same param decl or return type. There's no need to search all elements for each field. You're just matching the next leaf field with the next signature element and asserting that certain captured properties are consistent with those properties picked up during this traversal (like the dimensionality and type).

This would be different from the current approach which traverses the parameters, recurses into fields, then for each field, iterates all semantics on the function decl looking for one with TargetDecl pointing to the current field. While adding the full path could fix the problems with multiple instances, it's more complicated and unnecessary if you take the approach I suggested instead, where we know the next field must match the next element, no searching necessary.

Technically, we don't need to store anything more to the AST than what was explicitly declared in HLSL, as long as we reconstruct the same elements that were constructed and diagnosed in Sema in exactly the same way. We could do this if we had shared code for traversal and element construction from the AST.

Each parameter or return type, based on shader kind and parameter modifier/type/location slots into one "Signature Point" (SigPoint for short) that identifies a grouping of shader I/O parameters handled in a particular way (more than just input/output, it accounts for things like separating per-invocation system values from sets of input vertices (GS/HS/DS), or patch constant outputs (HS) or inputs (DS), and separating per-primitive mesh shader outputs from per-vertex outputs.

Whether stored back to the AST, or temporarily constructed for Sema diagnostics and again for CodeGen, signature elements need some properties for the DirectX target at least:

  • shape (rows, cols, HLSL component type)
  • system value index expansion (one index per row)
    • Full type traversal from decl with active semantic is necessary to assign these, when this decl is (optionally an array of) a struct type, since other fields will consume indices from the same set in traversal order.
  • which signature this belongs to (and SigPoint), if any
  • system value interpretation (user or specific system value)
  • interpolation mode
  • packing location if element type (by interpretation) is packed into a signature
    Diagnostics will need to keep track of some things that CodeGen doesn't need:
  • per element: decl location with applicable user-specified semantic attribute for semantics inherited by contained fields
  • per sigpoint: map of previously defined elements using semantic name/index key for diagnosing duplicate/overlapping semantics

This commit adds the support for semantics annotations on structs, but
only for inputs. Due to the current semantics implemented, we cannot
test much more than nesting/shadowing.

Once user semantics are implemented, we'll be able to test arrays in
structs and more complex cases.

As-is, this commit has one weakness vs DXC: semantics type validation is
not looking at the inner-most type, but the outermost type:

```hlsl
struct Inner {
  uint tid;
};

Inner inner : SV_GroupID
```

This sample would fail today because `SV_GroupID` require the type to be
an integer. This works in DXC as the inner type is a integer.
Because GroupIndex is not correctly validated, I uses this semantic to
test the inheritance/shadowing. But this will need to be fixed in a
later commit.

Requires llvm#152537
The previous solution had a major drawback: if a stuct was used
by multiple entrypoints, we had conflicting attribute.

This commit moves the attribute to the function declaration:
 - each field with an active semantic will have a related attribute
   attached to the corresponding entrypoint.
   This means the semantic list is per-entrypoint.
@Keenuts
Copy link
Contributor Author

Keenuts commented Oct 23, 2025

I'm approving this, but think the approach will need structural changes to handle the general struct case.

Yes, especially with the struct reuse case. As per the last meeting discussion, merging this as-is as it is an improvement over the status quo, but it does require significant additional work before we can call this done.

As for the rest, I'll move this discussion to the wg-hlsl, but some early thought:

Specifically:

I think the attribute added to the function decl to describe a loaded/stored element should be a different attribute than the semantic attribute. It's a shader I/O element, aka. signature element, constructed from a path through parameter space with a semantic applied.

I see the benefit of having an early attribute which only carries the text info with no additional meaning, and then having a post-sema attribute which carries the more complex semantic like "input/output" and other bits.

What I don't agree with is that the notion of signature element or packing should be part of this.

I don't think pointing at the leaf field in a type is enough.

That's correct. We either need the full path to disambiguate, or another representation like the flat you are suggesting. I'll move this discussion to a wg-hlsl proposal.

Technically, we don't need to store anything more to the AST than what was explicitly declared in HLSL [...]

Each parameter or return type, based on shader kind and parameter modifier/type/location slots into one "Signature Point" (SigPoint for short) [...]

Whether stored back to the AST, or temporarily constructed for Sema diagnostics and again for CodeGen, signature elements need some properties for the DirectX target at least:

  • shape (rows, cols, HLSL component type)
  • system value index expansion (one index per row)

IMO the HLSL frontend should not care about DXIL nor SPIR-V for those. We should emit enough so codegen knows how to emit something sensible, but the signature packing should be a DXIL backend specificity.

What HLSL requires is that semantic index are not overlapping, and that some system semantics are to be used only in specific stages and this kind of high-level constraints. This is IMO what the AST should capture. Hence a per-field/per-index semantic+index.

The fact the the input signature cannot be emitted because they are no registers lefts or DXIL specific behaviors should IMO be a backend issue.
Maybe it's simpler to build those checks in Sema, but it should I think not be stored in the AST.

For GS/HS/DS and other specific inputs, I'd think we'd want to add an additional attribute to carry the additional behavior: like if you have the triangle keyword before a parameter with the SV_POSITION semantic. One attribute for the semantic, and one for the triangle modifier. Moving this to the wg-hlsl discussion also.

@Keenuts Keenuts closed this Oct 23, 2025
@Keenuts Keenuts reopened this Oct 23, 2025
@Keenuts
Copy link
Contributor Author

Keenuts commented Oct 23, 2025

As per the last meeting discussion, merging this once the CI passes so we can move a bit forward.
I'll also open a few new items in wg-hlsl to discuss some aspects that have been raised here.

@Keenuts Keenuts merged commit 865cd8e into llvm:main Oct 23, 2025
18 of 21 checks passed
@Keenuts Keenuts deleted the hlsl-semantics-2 branch October 23, 2025 13:49
@llvm-ci
Copy link
Collaborator

llvm-ci commented Oct 24, 2025

LLVM Buildbot has detected a new failure on builder ppc64le-flang-rhel-clang running on ppc64le-flang-rhel-test while building clang at step 5 "build-unified-tree".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/157/builds/41714

Here is the relevant piece of the build log for the reference
Step 5 (build-unified-tree) failure: build (failure) (timed out)
...
1668.605 [61/9/7043] Generating ../../../../include/flang/ieee_features.mod
1668.689 [61/8/7044] Generating ../../../../include/flang/iso_c_binding.mod
1668.806 [61/7/7045] Generating ../../../../include/flang/ieee_arithmetic.mod
1668.808 [61/6/7046] Generating ../../../../include/flang/__cuda_device.mod
1668.934 [60/6/7047] Generating ../../../../include/flang/ieee_exceptions.mod
1669.407 [60/5/7048] Generating ../../../../include/flang/iso_fortran_env.mod
1669.423 [60/4/7049] Generating ../../../../include/flang/mma.mod
1670.106 [60/3/7050] Generating ../../../../include/flang/__ppc_intrinsics.mod
1670.116 [60/2/7051] Generating ../../../../include/flang/cudadevice.mod
1671.958 [59/2/7052] Generating ../../../../include/flang/cooperative_groups.mod
command timed out: 1200 seconds without output running [b'ninja'], attempting to kill
process killed by signal 9
program finished with exit code -1
elapsedTime=2872.983032

dvbuka pushed a commit to dvbuka/llvm-project that referenced this pull request Oct 27, 2025
This PR is an incremental improvement regarding semantics I/O in HLSL.
This PR allows
system semantics to be used on struct type in addition to parameters
(state today).
This PR doesn't consider implicit indexing increment that happens when
placing a semantic on an aggregate/array as implemented system semantics
don't allow such use yet.

The next step will be to enable user semantics, which will bring the
need to properly determine semantic indices depending on context.
This PR diverge from the initial wg-hlsl proposal as all diagnostics are
done in Sema (initial proposal suggested running diags in codegen).

This is not yet a solid semantic implementation, but increases the test
coverage and improves the status from where we are now.
Lukacma pushed a commit to Lukacma/llvm-project that referenced this pull request Oct 29, 2025
This PR is an incremental improvement regarding semantics I/O in HLSL.
This PR allows
system semantics to be used on struct type in addition to parameters
(state today).
This PR doesn't consider implicit indexing increment that happens when
placing a semantic on an aggregate/array as implemented system semantics
don't allow such use yet.

The next step will be to enable user semantics, which will bring the
need to properly determine semantic indices depending on context.
This PR diverge from the initial wg-hlsl proposal as all diagnostics are
done in Sema (initial proposal suggested running diags in codegen).

This is not yet a solid semantic implementation, but increases the test
coverage and improves the status from where we are now.
aokblast pushed a commit to aokblast/llvm-project that referenced this pull request Oct 30, 2025
This PR is an incremental improvement regarding semantics I/O in HLSL.
This PR allows
system semantics to be used on struct type in addition to parameters
(state today).
This PR doesn't consider implicit indexing increment that happens when
placing a semantic on an aggregate/array as implemented system semantics
don't allow such use yet.

The next step will be to enable user semantics, which will bring the
need to properly determine semantic indices depending on context.
This PR diverge from the initial wg-hlsl proposal as all diagnostics are
done in Sema (initial proposal suggested running diags in codegen).

This is not yet a solid semantic implementation, but increases the test
coverage and improves the status from where we are now.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

clang:codegen IR generation bugs: mangling, exceptions, etc. clang:frontend Language frontend issues, e.g. anything involving "Sema" clang Clang issues not falling into any other category HLSL HLSL Language Support

Projects

None yet

Development

Successfully merging this pull request may close these issues.

6 participants