[clang] [llvm] [Clang][SYCL] Define Intel IGCA preprocessor macro definitions for SYCL during device compilation (PR #227877)
Ian Li via cfe-commits
cfe-commits at lists.llvm.org
Sat Oct 24 12:36:27 PDT 2026
https://github.com/ianayl updated https://github.com/llvm/llvm-project/pull/227877
>From 2c070fcf7fdc8965e9949ffd6ea4ca9a2c4a49ca Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Mon, 28 Sep 2026 17:35:39 -0700
Subject: [PATCH 01/30] WIP addition of IGCA to OffloadArch
---
clang/include/clang/Basic/OffloadArch.h | 14 ++++-
clang/lib/Basic/OffloadArch.cpp | 17 ++++++
clang/unittests/Basic/OffloadArchTest.cpp | 2 +
.../llvm/TargetParser/IntelGPUTargetParser.h | 57 ++++++++++++++++++
.../lib/TargetParser/IntelGPUTargetParser.cpp | 60 +++++++++++++++++++
5 files changed, 149 insertions(+), 1 deletion(-)
diff --git a/clang/include/clang/Basic/OffloadArch.h b/clang/include/clang/Basic/OffloadArch.h
index ba8040ebe5dc2c..16ed08bcb1d00a 100644
--- a/clang/include/clang/Basic/OffloadArch.h
+++ b/clang/include/clang/Basic/OffloadArch.h
@@ -22,6 +22,9 @@ enum GPUKind : uint8_t;
namespace AMDGPU {
enum GPUKind : uint8_t;
}
+namespace IntelGPU {
+struct IGCATarget;
+}
} // namespace llvm
namespace clang {
@@ -39,6 +42,7 @@ class OffloadArch {
AMDGCNSPIRV, // The 'amdgcnspirv' pseudo target.
IntelCPU, // Kind is an IntelArch.
IntelGPU, // Kind is an IntelArch.
+ IntelIGCA, // Kind is an Intel IGCA Target
Generic, // The 'generic' processor model.
};
@@ -64,6 +68,7 @@ class OffloadArch {
static OffloadArch getAMDGPU(llvm::AMDGPU::GPUKind K) {
return {TargetArch::AMDGPU, static_cast<uint32_t>(K)};
}
+ static OffloadArch getIntelIGCA(llvm::IntelGPU::IGCATarget T);
static constexpr OffloadArch getIntel(TargetArch V, IntelArch A) {
return {V, static_cast<uint32_t>(A)};
}
@@ -85,18 +90,23 @@ class OffloadArch {
bool isAMDGCNSPIRV() const { return V == TargetArch::AMDGCNSPIRV; }
bool isIntelCPU() const { return V == TargetArch::IntelCPU; }
bool isIntelGPU() const { return V == TargetArch::IntelGPU; }
+ // TODO: Decide whether IntelIGCA counts as Intel here. clang-sycl-linker
+ // uses isIntel() to decide on AOT compilation, and runAOTCompile would then
+ // need an IntelIGCA dispatch.
bool isIntel() const { return isIntelCPU() || isIntelGPU(); }
bool isGeneric() const { return V == TargetArch::Generic; }
bool isUnused() const { return V == TargetArch::Unused; }
bool isUnknown() const { return V == TargetArch::Unknown; }
+ bool isIntelIGCA() const { return V == TargetArch::IntelIGCA; }
- // Only valid when isNVPTX() / isAMDGPU() respectively.
+ // Only valid when isNVPTX() / isAMDGPU() / isIntelIGCA() respectively.
llvm::NVPTX::GPUKind nvptxKind() const {
return static_cast<llvm::NVPTX::GPUKind>(Kind);
}
llvm::AMDGPU::GPUKind amdgpuKind() const {
return static_cast<llvm::AMDGPU::GPUKind>(Kind);
}
+ llvm::IntelGPU::IGCATarget igcaTarget() const;
bool operator==(const OffloadArch &Other) const {
return V == Other.V && Kind == Other.Kind;
@@ -106,6 +116,8 @@ class OffloadArch {
bool operator<(const OffloadArch &Other) const {
return std::tie(V, Kind) < std::tie(Other.V, Other.Kind);
}
+ // TODO test cases for whether or not IntelIGCA functions properly for these
+ // comparisons
};
const char *OffloadArchToString(OffloadArch A);
diff --git a/clang/lib/Basic/OffloadArch.cpp b/clang/lib/Basic/OffloadArch.cpp
index ca8dfab142dbb1..8ec02edfa6649c 100644
--- a/clang/lib/Basic/OffloadArch.cpp
+++ b/clang/lib/Basic/OffloadArch.cpp
@@ -11,6 +11,7 @@
#include "llvm/ADT/StringRef.h"
#include "llvm/TargetParser/AMDGPUTargetParser.h"
#include "llvm/TargetParser/NVPTXTargetParser.h"
+#include "llvm/TargetParser/IntelGPUTargetParser.h"
#include "llvm/TargetParser/Triple.h"
namespace clang {
@@ -23,6 +24,14 @@ OffloadArch OffloadArch::HIPDefault() {
return getAMDGPU(llvm::AMDGPU::parseArchAMDGCN("gfx906"));
}
+OffloadArch OffloadArch::getIntelIGCA(llvm::IntelGPU::IGCATarget T) {
+ return {TargetArch::IntelIGCA, T.pack()};
+}
+
+llvm::IntelGPU::IGCATarget OffloadArch::igcaTarget() const {
+ return llvm::IntelGPU::IGCATarget::unpack(Kind);
+}
+
const char *OffloadArchToString(OffloadArch A) {
switch (A.targetArch()) {
case OffloadArch::TargetArch::Unused:
@@ -39,6 +48,8 @@ const char *OffloadArchToString(OffloadArch A) {
return "graniterapids";
case OffloadArch::TargetArch::IntelGPU:
return "bmg_g21";
+ case OffloadArch::TargetArch::IntelIGCA:
+ return llvm::IntelGPU::getIGCATargetName(A.igcaTarget()).data();
case OffloadArch::TargetArch::Generic:
return "generic";
}
@@ -54,6 +65,8 @@ const char *OffloadArchToVirtualArchString(OffloadArch A) {
return "compute_amdgcn";
case OffloadArch::TargetArch::Unknown:
return "unknown";
+ case OffloadArch::TargetArch::IntelIGCA:
+ return llvm::IntelGPU::getIGCATargetName(A.igcaTarget()).data();
case OffloadArch::TargetArch::Unused:
case OffloadArch::TargetArch::IntelCPU:
case OffloadArch::TargetArch::IntelGPU:
@@ -85,6 +98,8 @@ OffloadArch StringToOffloadArch(llvm::StringRef S) {
return OffloadArch::getNVPTX(NV);
if (llvm::AMDGPU::GPUKind AK = llvm::AMDGPU::parseArchAMDGCN(S))
return OffloadArch::getAMDGPU(AK);
+ if (llvm::IntelGPU::IGCATarget IT = llvm::IntelGPU::parseIGCATarget(S))
+ return OffloadArch::getIntelIGCA(IT);
return OffloadArch::getUnknown();
}
@@ -119,6 +134,8 @@ llvm::Triple OffloadArchToTriple(const llvm::Triple &DefaultToolchainTriple,
return llvm::Triple(llvm::Triple::amdgpu, llvm::Triple::NoSubArch,
llvm::Triple::AMD, llvm::Triple::AMDHSA);
+ // TODO: Handle IntelIGCA, which currently gets an empty triple, so
+ // --offload-arch=igca_* cannot pick an offload toolchain.
return {};
}
diff --git a/clang/unittests/Basic/OffloadArchTest.cpp b/clang/unittests/Basic/OffloadArchTest.cpp
index 013ed75fbd57eb..f46ec566c7f037 100644
--- a/clang/unittests/Basic/OffloadArchTest.cpp
+++ b/clang/unittests/Basic/OffloadArchTest.cpp
@@ -54,6 +54,8 @@ TEST(OffloadArchTest, Unknown) {
}
// Names must round-trip through parse -> string.
+// TODO: Add IGCA targets (e.g. "igca_20ca"), plus IntelIGCA classification and
+// ==/< ordering tests.
TEST(OffloadArchTest, RoundTrip) {
for (const char *Name :
{"sm_52", "sm_90a", "gfx906", "gfx1201", "gfx12-generic", "amdgcnspirv",
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
index 4ae635e97167ba..b344e99070b485 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
@@ -34,6 +34,48 @@ enum GPUKind : uint8_t {
#include "llvm/TargetParser/IntelGPUTargetParser.def"
};
+// TODO Do I want the IGCA_ prefix here?
+enum class IGCAFeatureSet : uint8_t {
+ IGCA_CORE = 0,
+ IGCA_COMPUTE,
+ IGCA_RENDER
+};
+
+// TODO these need to be tested:
+struct IGCATarget {
+ uint16_t Target = 0;
+ IGCAFeatureSet FeatureSet = IGCAFeatureSet::IGCA_CORE;
+ bool IsExactFeatureSet = false;
+
+ /// Return true if the IGCATarget is a valid, properly initialized target.
+ bool isValid() const { return Target != 0; }
+ explicit operator bool() const { return isValid(); }
+
+ bool isCore() const { return FeatureSet == IGCAFeatureSet::IGCA_CORE; }
+ bool isExact() const { return IsExactFeatureSet; }
+ bool isCompute() const { return FeatureSet == IGCAFeatureSet::IGCA_COMPUTE; }
+ bool isRender() const { return FeatureSet == IGCAFeatureSet::IGCA_RENDER; }
+ bool isComputeExact() const { return isCompute() && isExact(); };
+ bool isRenderExact() const { return isRender() && isExact(); };
+
+ /// Pack an IGCATarget into an uint32_t.
+ ///
+ /// We encode as [31:16] Target, [2:1] FeatureSet, [0] IsExactFeatureSet.
+ /// [15:3] is left as reserved.
+ uint32_t pack() const;
+
+ /// Obtain an IGCATarget from a packed uint32_t.
+ static IGCATarget unpack(uint32_t V);
+
+ friend bool operator==(IGCATarget A, IGCATarget B) {
+ // TODO is this bad?
+ return A.pack() == B.pack();
+ }
+
+ /// Return an invalid IGCA Target that returns false on IGCATarget::isValid().
+ static IGCATarget invalid() { return IGCATarget{}; }
+};
+
/// \return the device name that \p GPUIPVersion identifies, as reported by the
/// driver, e.g. "xe-pvc". If the table lists no such device, return an empty
/// string. Only the major and minor versions are looked at, so every revision
@@ -44,6 +86,21 @@ LLVM_ABI StringRef getArchName(uint32_t GPUIPVersion);
/// \return the numeric name of \p GPUIPVersion, e.g. "xe_35.11.0".
LLVM_ABI std::string getNumericArchName(uint32_t GPUIPVersion);
+/// Parse an IGCA Target.
+IGCATarget parseIGCATarget(StringRef TargetStr);
+
+/// Return the IGCA string for a given IGCA Target, i.e. "igca_60ca". Returns
+/// empty string if Target is not a valid target.
+std::string getIGCATargetName(IGCATarget Target);
+// TODO to carry stringref's, I need to create that IGCA Target table
+// If that IGCA Target table gets really verbose, we should just use the
+// strings from that table and use StringRef here instead of std::string.
+// TODO DO STRINGREF'S
+
+// TODO
+// - get igcatarget
+// - translate other targets?
+
} // namespace IntelGPU
} // namespace llvm
diff --git a/llvm/lib/TargetParser/IntelGPUTargetParser.cpp b/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
index 4b4d162ffbfff7..67721b5fda0fda 100644
--- a/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
+++ b/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
@@ -63,3 +63,63 @@ std::string llvm::IntelGPU::getNumericArchName(uint32_t GPUIPVersion) {
return ("xe_" + Twine(Major) + "." + Twine(Minor) + "." + Twine(Revision))
.str();
}
+
+IGCATarget llvm::IntelGPU::parseIGCATarget(StringRef MaybeTarget) {
+ if (!MaybeTarget.consume_front("igca_"))
+ return IGCATarget::invalid();
+
+ uint16_t Target;
+ if (MaybeTarget.consumeInteger(10, Target))
+ return IGCATarget::invalid();
+ // TODO check that Target contains a valid target level
+
+ IGCAFeatureSet FS = IGCAFeatureSet::IGCA_CORE;
+ if (MaybeTarget.consume_front("c"))
+ FS = IGCAFeatureSet::IGCA_COMPUTE;
+ else if (MaybeTarget.consume_front("r"))
+ FS = IGCAFeatureSet::IGCA_RENDER;
+ bool IsExactFS = MaybeTarget.consume_front("a");
+ if (!MaybeTarget.empty())
+ return IGCATarget::invalid();
+
+ return { Target, FS, IsExactFS };
+}
+
+static constexpr uint32_t IGCATargetShift = 16;
+static constexpr uint32_t IGCAFeatureSetShift = 1;
+static constexpr uint32_t IGCAIsExactFSShift = 0;
+static constexpr uint32_t IGCAFeatureSetMask = 0x03;
+static constexpr uint32_t IGCAIsExactFSMask = 0x1;
+
+// An IGCA Target is packed as follows:
+//
+// 31 16 15 3 2 1 0 0
+// +------------------+--------------+-------------+-----------+
+// | Target | Reserved | Feature set | Is Exact? |
+// +------------------+--------------+-------------+-----------+
+// 16 bits 13 bits 2 bits 1 bit
+//
+uint32_t llvm::IntelGPU::IGCATarget::pack() const {
+ // TODO add debug mode asserts here
+ return uint32_t(Target) << IGCATargetShift |
+ (uint32_t(FeatureSet) & IGCAFeatureSetMask) << IGCAFeatureSetShift |
+ (uint32_t(IsExactFeatureSet) & IGCAIsExactFSMask)
+ << IGCAIsExactFSShift;
+}
+
+IGCATarget llvm::IntelGPU::IGCATarget::unpack(uint32_t V) {
+ return {
+ uint16_t(V >> IGCATargetShift),
+ IGCAFeatureSet((V >> IGCAFeatureSetShift) & IGCAFeatureSetMask),
+ bool((V >> IGCAIsExactFSShift) & IGCAIsExactFSMask)
+ };
+}
+
+std::string llvm::IntelGPU::getIGCATargetName(IGCATarget T) {
+ if (!T)
+ return "";
+
+ StringRef FS = T.isCompute() ? "c" : T.isRender() ? "r" : "";
+ StringRef Exact = T.isExact() ? "a" : "";
+ return ("igca_" + Twine(T.Target) + FS + Exact).str();
+}
>From 5f13c9b7c26a1626318a387a8d8b0c4eb203ecaa Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Tue, 29 Sep 2026 12:22:01 -0700
Subject: [PATCH 02/30] Properly define IGCA targets in targetparser.def, and
collapse IGCATarget struct to wrap around uint32_t packed representation
instead
---
clang/include/clang/Basic/OffloadArch.h | 2 +-
.../TargetParser/IntelGPUTargetParser.def | 27 ++++-
.../llvm/TargetParser/IntelGPUTargetParser.h | 103 ++++++++++--------
.../lib/TargetParser/IntelGPUTargetParser.cpp | 80 +++++++-------
.../TargetParser/IntelGPUTargetParserTest.cpp | 89 +++++++++++++++
5 files changed, 218 insertions(+), 83 deletions(-)
diff --git a/clang/include/clang/Basic/OffloadArch.h b/clang/include/clang/Basic/OffloadArch.h
index 16ed08bcb1d00a..bdbbcc2c8b13f2 100644
--- a/clang/include/clang/Basic/OffloadArch.h
+++ b/clang/include/clang/Basic/OffloadArch.h
@@ -23,7 +23,7 @@ namespace AMDGPU {
enum GPUKind : uint8_t;
}
namespace IntelGPU {
-struct IGCATarget;
+class IGCATarget;
}
} // namespace llvm
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def
index bc9ba1d3e438b3..dba4da8b7c0a86 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def
@@ -17,7 +17,8 @@
// GK_<KIND>.
// MAJOR - Major component of the device's GPU IP version.
// MINOR - Minor component of the device's GPU IP version.
-// IGCA_TARGET - Numeric IGCA target.
+// IGCA_TARGET - Numeric IGCA target: Targets should also be defined
+// in the INTEL_IGCA_TARGET table.
// IGCA_FEATURE_SETS - Token naming the IGCA feature sets the device
// implements: Core (none), Compute ("c"),
// Render ("r"), ComputeExact ("ca") or
@@ -37,6 +38,15 @@
// Several devices can share a major and a minor version. The rows are ordered
// so that the name to print for such a group comes first.
//
+// This file also includes a list of valid IGCA targets, which can be used to
+// compile binaries with support different GPU architectures and feature sets:
+//
+// INTEL_IGCA_TARGET(TARGET)
+// TARGET - Numeric IGCA target level, e.g. 60. corresponds to
+// "igca_60". A defined IGCA target means feature sets,
+// i.e. "igca_<TARGET>c", or "igca_<Target>ra" are also
+// accepted.
+//
//===----------------------------------------------------------------------===//
#ifndef INTEL_GPU
@@ -92,5 +102,20 @@ INTEL_GPU_COMPAT("xe-bmg", XE_BMG, 40, Render)
INTEL_GPU_COMPAT("xe-mtl", XE_MTL, 30, Render)
INTEL_GPU_COMPAT("xe-dg2", XE_DG2, 15, RenderExact)
+// IGCA target levels. Only targets defined here are valid.
+#ifndef INTEL_IGCA_TARGET
+#define INTEL_IGCA_TARGET(TARGET)
+#endif
+
+INTEL_IGCA_TARGET(10)
+INTEL_IGCA_TARGET(15)
+INTEL_IGCA_TARGET(20)
+INTEL_IGCA_TARGET(30)
+INTEL_IGCA_TARGET(35)
+INTEL_IGCA_TARGET(40)
+INTEL_IGCA_TARGET(50)
+INTEL_IGCA_TARGET(60)
+
#undef INTEL_GPU
#undef INTEL_GPU_COMPAT
+#undef INTEL_IGCA_TARGET
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
index b344e99070b485..4931937be3af40 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
@@ -22,6 +22,8 @@
#include <string>
namespace llvm {
+template <typename T> class SmallVectorImpl;
+
namespace IntelGPU {
/// Intel GPU architecture names, covering both physical devices and the
@@ -34,46 +36,62 @@ enum GPUKind : uint8_t {
#include "llvm/TargetParser/IntelGPUTargetParser.def"
};
-// TODO Do I want the IGCA_ prefix here?
-enum class IGCAFeatureSet : uint8_t {
+enum IGCAFeatureSet : uint8_t {
IGCA_CORE = 0,
IGCA_COMPUTE,
IGCA_RENDER
};
-// TODO these need to be tested:
-struct IGCATarget {
- uint16_t Target = 0;
- IGCAFeatureSet FeatureSet = IGCAFeatureSet::IGCA_CORE;
- bool IsExactFeatureSet = false;
-
- /// Return true if the IGCATarget is a valid, properly initialized target.
- bool isValid() const { return Target != 0; }
+/// Wrapper around an IGCA target's uint32_t representation, packed using the
+/// following format:
+///
+/// 31 16 15 3 2 1 0 0
+/// +------------------+--------------+-------------+-----------+
+/// | Target | Reserved | Feature set | Is Exact? |
+/// +------------------+--------------+-------------+-----------+
+/// 16 bits 13 bits 2 bits 1 bit
+///
+/// Target == 0 denotes an invalid IGCA Target.
+class IGCATarget {
+ uint32_t V = 0;
+ constexpr explicit IGCATarget(uint32_t V) : V(V) {}
+
+public:
+ static constexpr uint32_t TargetShift = 16;
+ static constexpr uint32_t FeatureSetShift = 1;
+ static constexpr uint32_t IsExactShift = 0;
+ static constexpr uint32_t FeatureSetMask = 0x3;
+ static constexpr uint32_t IsExactMask = 0x1;
+
+ constexpr IGCATarget(uint16_t Target, IGCAFeatureSet FeatureSet, bool IsExact)
+ : V(uint32_t(Target) << TargetShift |
+ (uint32_t(FeatureSet) & FeatureSetMask) << FeatureSetShift |
+ (uint32_t(IsExact) & IsExactMask) << IsExactShift) {}
+
+ /// \return an invalid IGCATarget.
+ static constexpr IGCATarget invalid() { return IGCATarget(0); }
+
+ /// \return the packed uint32_t representation of the IGCA target.
+ constexpr uint32_t pack() const { return V; }
+ /// Wrap a packed uint32_t IGCA target with an IGCATarget class.
+ static constexpr IGCATarget unpack(uint32_t V) { return IGCATarget(V); }
+
+ uint16_t getTarget() const { return uint16_t(V >> TargetShift); }
+ IGCAFeatureSet getFeatureSet() const {
+ return IGCAFeatureSet((V >> FeatureSetShift) & FeatureSetMask);
+ }
+ bool isExact() const { return (V >> IsExactShift) & IsExactMask; }
+ bool isValid() const { return getTarget() != 0; }
explicit operator bool() const { return isValid(); }
- bool isCore() const { return FeatureSet == IGCAFeatureSet::IGCA_CORE; }
- bool isExact() const { return IsExactFeatureSet; }
- bool isCompute() const { return FeatureSet == IGCAFeatureSet::IGCA_COMPUTE; }
- bool isRender() const { return FeatureSet == IGCAFeatureSet::IGCA_RENDER; }
- bool isComputeExact() const { return isCompute() && isExact(); };
- bool isRenderExact() const { return isRender() && isExact(); };
-
- /// Pack an IGCATarget into an uint32_t.
- ///
- /// We encode as [31:16] Target, [2:1] FeatureSet, [0] IsExactFeatureSet.
- /// [15:3] is left as reserved.
- uint32_t pack() const;
-
- /// Obtain an IGCATarget from a packed uint32_t.
- static IGCATarget unpack(uint32_t V);
-
- friend bool operator==(IGCATarget A, IGCATarget B) {
- // TODO is this bad?
- return A.pack() == B.pack();
- }
+ bool isCore() const { return getFeatureSet() == IGCA_CORE; }
+ bool isCompute() const { return getFeatureSet() == IGCA_COMPUTE; }
+ bool isRender() const { return getFeatureSet() == IGCA_RENDER; }
+ bool isComputeExact() const { return isCompute() && isExact(); }
+ bool isRenderExact() const { return isRender() && isExact(); }
- /// Return an invalid IGCA Target that returns false on IGCATarget::isValid().
- static IGCATarget invalid() { return IGCATarget{}; }
+ friend bool operator==(IGCATarget A, IGCATarget B) { return A.V == B.V; }
+ friend bool operator!=(IGCATarget A, IGCATarget B) { return A.V != B.V; }
};
/// \return the device name that \p GPUIPVersion identifies, as reported by the
@@ -86,20 +104,17 @@ LLVM_ABI StringRef getArchName(uint32_t GPUIPVersion);
/// \return the numeric name of \p GPUIPVersion, e.g. "xe_35.11.0".
LLVM_ABI std::string getNumericArchName(uint32_t GPUIPVersion);
-/// Parse an IGCA Target.
-IGCATarget parseIGCATarget(StringRef TargetStr);
+/// Parse an IGCA target string, such as "igca_60ca". \return an invalid
+/// IGCATarget if \p TargetStr is not a known target in
+/// IntelGPUTargetParser.def
+LLVM_ABI IGCATarget parseIGCATarget(StringRef TargetStr);
-/// Return the IGCA string for a given IGCA Target, i.e. "igca_60ca". Returns
-/// empty string if Target is not a valid target.
-std::string getIGCATargetName(IGCATarget Target);
-// TODO to carry stringref's, I need to create that IGCA Target table
-// If that IGCA Target table gets really verbose, we should just use the
-// strings from that table and use StringRef here instead of std::string.
-// TODO DO STRINGREF'S
+/// \return the \p Target as a string, i.e. "igca_60ca". \return an empty string
+/// if \p Target is invalid or not a known target in IntelGPUTargetParser.def.
+LLVM_ABI StringRef getIGCATargetName(IGCATarget Target);
-// TODO
-// - get igcatarget
-// - translate other targets?
+/// Append every legal IGCA target spelling to \p Values.
+LLVM_ABI void fillValidIGCATargetList(SmallVectorImpl<StringRef> &Values);
} // namespace IntelGPU
} // namespace llvm
diff --git a/llvm/lib/TargetParser/IntelGPUTargetParser.cpp b/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
index 67721b5fda0fda..7b861b71e4c6a6 100644
--- a/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
+++ b/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
@@ -64,62 +64,68 @@ std::string llvm::IntelGPU::getNumericArchName(uint32_t GPUIPVersion) {
.str();
}
+/// Check that Level is a known IGCA Target in IntelGPUTargetParser.def.
+bool isKnownIGCATargetLevel(uint16_t Level) {
+ switch (Level) {
+#define INTEL_IGCA_TARGET(TARGET) case TARGET:
+#include "llvm/TargetParser/IntelGPUTargetParser.def"
+ return true;
+ default:
+ return false;
+ }
+}
+
IGCATarget llvm::IntelGPU::parseIGCATarget(StringRef MaybeTarget) {
if (!MaybeTarget.consume_front("igca_"))
return IGCATarget::invalid();
uint16_t Target;
- if (MaybeTarget.consumeInteger(10, Target))
+ if (MaybeTarget.consumeInteger(10, Target) || !isKnownIGCATargetLevel(Target))
return IGCATarget::invalid();
- // TODO check that Target contains a valid target level
IGCAFeatureSet FS = IGCAFeatureSet::IGCA_CORE;
if (MaybeTarget.consume_front("c"))
FS = IGCAFeatureSet::IGCA_COMPUTE;
else if (MaybeTarget.consume_front("r"))
FS = IGCAFeatureSet::IGCA_RENDER;
- bool IsExactFS = MaybeTarget.consume_front("a");
+ // Exact form needs to either be compute or render:
+ bool IsExactFS = (
+ FS != IGCAFeatureSet::IGCA_CORE && MaybeTarget.consume_front("a")
+ );
if (!MaybeTarget.empty())
return IGCATarget::invalid();
return { Target, FS, IsExactFS };
}
-static constexpr uint32_t IGCATargetShift = 16;
-static constexpr uint32_t IGCAFeatureSetShift = 1;
-static constexpr uint32_t IGCAIsExactFSShift = 0;
-static constexpr uint32_t IGCAFeatureSetMask = 0x03;
-static constexpr uint32_t IGCAIsExactFSMask = 0x1;
-
-// An IGCA Target is packed as follows:
-//
-// 31 16 15 3 2 1 0 0
-// +------------------+--------------+-------------+-----------+
-// | Target | Reserved | Feature set | Is Exact? |
-// +------------------+--------------+-------------+-----------+
-// 16 bits 13 bits 2 bits 1 bit
-//
-uint32_t llvm::IntelGPU::IGCATarget::pack() const {
- // TODO add debug mode asserts here
- return uint32_t(Target) << IGCATargetShift |
- (uint32_t(FeatureSet) & IGCAFeatureSetMask) << IGCAFeatureSetShift |
- (uint32_t(IsExactFeatureSet) & IGCAIsExactFSMask)
- << IGCAIsExactFSShift;
-}
-
-IGCATarget llvm::IntelGPU::IGCATarget::unpack(uint32_t V) {
- return {
- uint16_t(V >> IGCATargetShift),
- IGCAFeatureSet((V >> IGCAFeatureSetShift) & IGCAFeatureSetMask),
- bool((V >> IGCAIsExactFSShift) & IGCAIsExactFSMask)
- };
+StringRef llvm::IntelGPU::getIGCATargetName(IGCATarget T) {
+ switch (T.pack()) {
+#define INTEL_IGCA_TARGET(TARGET) \
+ case IGCATarget(TARGET, IGCAFeatureSet::IGCA_CORE, false).pack(): \
+ return "igca_" #TARGET; \
+ case IGCATarget(TARGET, IGCAFeatureSet::IGCA_COMPUTE, false).pack(): \
+ return "igca_" #TARGET "c"; \
+ case IGCATarget(TARGET, IGCAFeatureSet::IGCA_COMPUTE, true).pack(): \
+ return "igca_" #TARGET "ca"; \
+ case IGCATarget(TARGET, IGCAFeatureSet::IGCA_RENDER, false).pack(): \
+ return "igca_" #TARGET "r"; \
+ case IGCATarget(TARGET, IGCAFeatureSet::IGCA_RENDER, true).pack(): \
+ return "igca_" #TARGET "ra";
+#include "llvm/TargetParser/IntelGPUTargetParser.def"
+ default:
+ return "";
+ }
}
-std::string llvm::IntelGPU::getIGCATargetName(IGCATarget T) {
- if (!T)
- return "";
+#define INTEL_IGCA_TARGET_FEATURESETS(TARGET) \
+ "igca_" #TARGET, "igca_" #TARGET "c", "igca_" #TARGET "ca", \
+ "igca_" #TARGET "r", "igca_" #TARGET "ra"
- StringRef FS = T.isCompute() ? "c" : T.isRender() ? "r" : "";
- StringRef Exact = T.isExact() ? "a" : "";
- return ("igca_" + Twine(T.Target) + FS + Exact).str();
+void llvm::IntelGPU::fillValidIGCATargetList(
+ SmallVectorImpl<StringRef> &Values) {
+#define INTEL_IGCA_TARGET(TARGET) \
+ Values.append({INTEL_IGCA_TARGET_FEATURESETS(TARGET)});
+#include "llvm/TargetParser/IntelGPUTargetParser.def"
}
+
+#undef INTEL_IGCA_TARGET_FEATURESETS
diff --git a/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp b/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
index ab8d456769d0da..93a31558e7e03e 100644
--- a/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
+++ b/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
@@ -78,4 +78,93 @@ TEST(IntelGPUTargetParserTest, NumericArchName) {
"xe_12.60.7");
}
+TEST(IntelGPUTargetParserTest, IGCATargetBehavior) {
+ IntelGPU::IGCATarget T(60, IntelGPU::IGCA_COMPUTE, true);
+ EXPECT_TRUE(T.isValid());
+ EXPECT_EQ(T.getTarget(), 60);
+ EXPECT_EQ(T.getFeatureSet(), IntelGPU::IGCA_COMPUTE);
+ EXPECT_TRUE(T.isExact());
+ EXPECT_TRUE(T.isComputeExact());
+ EXPECT_FALSE(T.isRender());
+ EXPECT_EQ(IntelGPU::IGCATarget::unpack(T.pack()), T);
+ EXPECT_NE(T, IntelGPU::IGCATarget(60, IntelGPU::IGCA_COMPUTE, false));
+ // Invalid behavior:
+ EXPECT_FALSE(IntelGPU::IGCATarget::invalid());
+ EXPECT_EQ(IntelGPU::IGCATarget::invalid().pack(), 0u);
+}
+
+TEST(IntelGPUTargetParserTest, ParseIGCATarget) {
+ using IntelGPU::IGCATarget;
+ EXPECT_EQ(IntelGPU::parseIGCATarget("igca_10"),
+ IGCATarget(10, IntelGPU::IGCA_CORE, false));
+ EXPECT_EQ(IntelGPU::parseIGCATarget("igca_20c"),
+ IGCATarget(20, IntelGPU::IGCA_COMPUTE, false));
+ EXPECT_EQ(IntelGPU::parseIGCATarget("igca_20ca"),
+ IGCATarget(20, IntelGPU::IGCA_COMPUTE, true));
+ EXPECT_EQ(IntelGPU::parseIGCATarget("igca_15r"),
+ IGCATarget(15, IntelGPU::IGCA_RENDER, false));
+ EXPECT_EQ(IntelGPU::parseIGCATarget("igca_15ra"),
+ IGCATarget(15, IntelGPU::IGCA_RENDER, true));
+ // Only levels in the table are valid targets.
+ EXPECT_FALSE(IntelGPU::parseIGCATarget("igca_42"));
+ EXPECT_FALSE(IntelGPU::parseIGCATarget("igca_0"));
+ // Cannot have exact without naming feature set:
+ EXPECT_FALSE(IntelGPU::parseIGCATarget("igca_10a"));
+ // Malformed spellings:
+ EXPECT_FALSE(IntelGPU::parseIGCATarget("igca_"));
+ EXPECT_FALSE(IntelGPU::parseIGCATarget("igca_10x"));
+ EXPECT_FALSE(IntelGPU::parseIGCATarget("igca_10cr"));
+ EXPECT_FALSE(IntelGPU::parseIGCATarget("igca_10caa"));
+ EXPECT_FALSE(IntelGPU::parseIGCATarget("igca10"));
+ EXPECT_FALSE(IntelGPU::parseIGCATarget("IGCA_10"));
+ EXPECT_FALSE(IntelGPU::parseIGCATarget(""));
+}
+
+TEST(IntelGPUTargetParserTest, IGCATargetName) {
+ using IntelGPU::IGCATarget;
+ EXPECT_EQ(IntelGPU::getIGCATargetName(
+ IGCATarget(60, IntelGPU::IGCA_CORE, false)),
+ "igca_60");
+ EXPECT_EQ(IntelGPU::getIGCATargetName(
+ IGCATarget(60, IntelGPU::IGCA_COMPUTE, true)),
+ "igca_60ca");
+ EXPECT_EQ(IntelGPU::getIGCATargetName(
+ IGCATarget(15, IntelGPU::IGCA_RENDER, false)),
+ "igca_15r");
+ EXPECT_EQ(IntelGPU::getIGCATargetName(IGCATarget::invalid()), "");
+ // A level that is not in the table has no spelling, and neither does an
+ // exact core target.
+ EXPECT_EQ(IntelGPU::getIGCATargetName(
+ IGCATarget(11, IntelGPU::IGCA_CORE, false)),
+ "");
+ EXPECT_EQ(IntelGPU::getIGCATargetName(
+ IGCATarget(10, IntelGPU::IGCA_CORE, true)),
+ "");
+}
+
+TEST(IntelGPUTargetParserTest, EveryIGCASpellingRoundTrips) {
+ SmallVector<StringRef> Names;
+ IntelGPU::fillValidIGCATargetList(Names);
+ EXPECT_FALSE(Names.empty());
+ for (StringRef Name : Names) {
+ IntelGPU::IGCATarget T = IntelGPU::parseIGCATarget(Name);
+ EXPECT_TRUE(T.isValid()) << Name;
+ EXPECT_EQ(IntelGPU::getIGCATargetName(T), Name);
+ }
+}
+
+TEST(IntelGPUTargetParserTest, EveryDeviceHasValidIGCATargetLevel) {
+ // Every device's defined IGCA target should have a corresponding entry in
+ // INTEL_IGCA_TARGET:
+#define INTEL_IGCA_TARGET_CHECK(NAME, IGCA_TARGET) \
+ EXPECT_TRUE(IntelGPU::parseIGCATarget("igca_" #IGCA_TARGET).isValid()) \
+ << NAME;
+#define INTEL_GPU(NAME, KIND, MAJOR, MINOR, IGCA_TARGET, IGCA_FEATURE_SETS) \
+ INTEL_IGCA_TARGET_CHECK(NAME, IGCA_TARGET)
+#define INTEL_GPU_COMPAT(NAME, KIND, IGCA_TARGET, IGCA_FEATURE_SETS) \
+ INTEL_IGCA_TARGET_CHECK(NAME, IGCA_TARGET)
+#include "llvm/TargetParser/IntelGPUTargetParser.def"
+#undef INTEL_IGCA_TARGET_CHECK
+}
+
} // namespace
>From 9f66c948d302026f4f2ac8b940933bfe351ffcdf Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Tue, 29 Sep 2026 13:24:24 -0700
Subject: [PATCH 03/30] clang-format
---
clang/lib/Basic/OffloadArch.cpp | 2 +-
.../llvm/TargetParser/IntelGPUTargetParser.h | 8 ++---
.../lib/TargetParser/IntelGPUTargetParser.cpp | 9 +++---
.../TargetParser/IntelGPUTargetParserTest.cpp | 30 +++++++++----------
4 files changed, 22 insertions(+), 27 deletions(-)
diff --git a/clang/lib/Basic/OffloadArch.cpp b/clang/lib/Basic/OffloadArch.cpp
index 8ec02edfa6649c..c3004bd3333cb0 100644
--- a/clang/lib/Basic/OffloadArch.cpp
+++ b/clang/lib/Basic/OffloadArch.cpp
@@ -10,8 +10,8 @@
#include "llvm/ADT/SmallVector.h"
#include "llvm/ADT/StringRef.h"
#include "llvm/TargetParser/AMDGPUTargetParser.h"
-#include "llvm/TargetParser/NVPTXTargetParser.h"
#include "llvm/TargetParser/IntelGPUTargetParser.h"
+#include "llvm/TargetParser/NVPTXTargetParser.h"
#include "llvm/TargetParser/Triple.h"
namespace clang {
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
index 4931937be3af40..339becdd102225 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
@@ -36,11 +36,7 @@ enum GPUKind : uint8_t {
#include "llvm/TargetParser/IntelGPUTargetParser.def"
};
-enum IGCAFeatureSet : uint8_t {
- IGCA_CORE = 0,
- IGCA_COMPUTE,
- IGCA_RENDER
-};
+enum IGCAFeatureSet : uint8_t { IGCA_CORE = 0, IGCA_COMPUTE, IGCA_RENDER };
/// Wrapper around an IGCA target's uint32_t representation, packed using the
/// following format:
@@ -105,7 +101,7 @@ LLVM_ABI StringRef getArchName(uint32_t GPUIPVersion);
LLVM_ABI std::string getNumericArchName(uint32_t GPUIPVersion);
/// Parse an IGCA target string, such as "igca_60ca". \return an invalid
-/// IGCATarget if \p TargetStr is not a known target in
+/// IGCATarget if \p TargetStr is not a known target in
/// IntelGPUTargetParser.def
LLVM_ABI IGCATarget parseIGCATarget(StringRef TargetStr);
diff --git a/llvm/lib/TargetParser/IntelGPUTargetParser.cpp b/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
index 7b861b71e4c6a6..0e1163dff9aec8 100644
--- a/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
+++ b/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
@@ -82,20 +82,19 @@ IGCATarget llvm::IntelGPU::parseIGCATarget(StringRef MaybeTarget) {
uint16_t Target;
if (MaybeTarget.consumeInteger(10, Target) || !isKnownIGCATargetLevel(Target))
return IGCATarget::invalid();
-
+
IGCAFeatureSet FS = IGCAFeatureSet::IGCA_CORE;
if (MaybeTarget.consume_front("c"))
FS = IGCAFeatureSet::IGCA_COMPUTE;
else if (MaybeTarget.consume_front("r"))
FS = IGCAFeatureSet::IGCA_RENDER;
// Exact form needs to either be compute or render:
- bool IsExactFS = (
- FS != IGCAFeatureSet::IGCA_CORE && MaybeTarget.consume_front("a")
- );
+ bool IsExactFS =
+ (FS != IGCAFeatureSet::IGCA_CORE && MaybeTarget.consume_front("a"));
if (!MaybeTarget.empty())
return IGCATarget::invalid();
- return { Target, FS, IsExactFS };
+ return {Target, FS, IsExactFS};
}
StringRef llvm::IntelGPU::getIGCATargetName(IGCATarget T) {
diff --git a/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp b/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
index 93a31558e7e03e..f1c55912305ece 100644
--- a/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
+++ b/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
@@ -122,24 +122,24 @@ TEST(IntelGPUTargetParserTest, ParseIGCATarget) {
TEST(IntelGPUTargetParserTest, IGCATargetName) {
using IntelGPU::IGCATarget;
- EXPECT_EQ(IntelGPU::getIGCATargetName(
- IGCATarget(60, IntelGPU::IGCA_CORE, false)),
- "igca_60");
- EXPECT_EQ(IntelGPU::getIGCATargetName(
- IGCATarget(60, IntelGPU::IGCA_COMPUTE, true)),
- "igca_60ca");
- EXPECT_EQ(IntelGPU::getIGCATargetName(
- IGCATarget(15, IntelGPU::IGCA_RENDER, false)),
- "igca_15r");
+ EXPECT_EQ(
+ IntelGPU::getIGCATargetName(IGCATarget(60, IntelGPU::IGCA_CORE, false)),
+ "igca_60");
+ EXPECT_EQ(
+ IntelGPU::getIGCATargetName(IGCATarget(60, IntelGPU::IGCA_COMPUTE, true)),
+ "igca_60ca");
+ EXPECT_EQ(
+ IntelGPU::getIGCATargetName(IGCATarget(15, IntelGPU::IGCA_RENDER, false)),
+ "igca_15r");
EXPECT_EQ(IntelGPU::getIGCATargetName(IGCATarget::invalid()), "");
// A level that is not in the table has no spelling, and neither does an
// exact core target.
- EXPECT_EQ(IntelGPU::getIGCATargetName(
- IGCATarget(11, IntelGPU::IGCA_CORE, false)),
- "");
- EXPECT_EQ(IntelGPU::getIGCATargetName(
- IGCATarget(10, IntelGPU::IGCA_CORE, true)),
- "");
+ EXPECT_EQ(
+ IntelGPU::getIGCATargetName(IGCATarget(11, IntelGPU::IGCA_CORE, false)),
+ "");
+ EXPECT_EQ(
+ IntelGPU::getIGCATargetName(IGCATarget(10, IntelGPU::IGCA_CORE, true)),
+ "");
}
TEST(IntelGPUTargetParserTest, EveryIGCASpellingRoundTrips) {
>From 2f89fa423ffc15e3ca033c574d92c92845e3e980 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 1 Oct 2026 15:58:22 -0700
Subject: [PATCH 04/30] Enforce that only certain IGCA target feature sets can
exist, simplify code and logic
---
.../TargetParser/IntelGPUTargetParser.def | 53 +++++--
.../llvm/TargetParser/IntelGPUTargetParser.h | 62 +++++---
.../lib/TargetParser/IntelGPUTargetParser.cpp | 63 ++------
.../TargetParser/IntelGPUTargetParserTest.cpp | 139 +++++++++++++-----
4 files changed, 199 insertions(+), 118 deletions(-)
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def
index dba4da8b7c0a86..7ad6b201c3081d 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def
@@ -41,11 +41,15 @@
// This file also includes a list of valid IGCA targets, which can be used to
// compile binaries with support different GPU architectures and feature sets:
//
-// INTEL_IGCA_TARGET(TARGET)
-// TARGET - Numeric IGCA target level, e.g. 60. corresponds to
-// "igca_60". A defined IGCA target means feature sets,
-// i.e. "igca_<TARGET>c", or "igca_<Target>ra" are also
-// accepted.
+// INTEL_IGCA_TARGET(NAME, TARGET, FEATURE_SET)
+// One row per valid IGCA target spelling. Only the combinations listed
+// are valid.
+// NAME - Name of the IGCA target, e.g. "igca_15ra". Name must
+// agree with the other columns.
+// TARGET - Numeric IGCA target level, e.g. 60.
+// FEATURE_SET - Feature set of the current target, i.e. "Compute" for
+// "c" suffix, "Render" for "r", "ComputeExact" for
+// "ca", "Core" for "", etc.
//
//===----------------------------------------------------------------------===//
@@ -102,19 +106,38 @@ INTEL_GPU_COMPAT("xe-bmg", XE_BMG, 40, Render)
INTEL_GPU_COMPAT("xe-mtl", XE_MTL, 30, Render)
INTEL_GPU_COMPAT("xe-dg2", XE_DG2, 15, RenderExact)
-// IGCA target levels. Only targets defined here are valid.
+// IGCA targets. Only targets listed here are valid.
#ifndef INTEL_IGCA_TARGET
-#define INTEL_IGCA_TARGET(TARGET)
+#define INTEL_IGCA_TARGET(NAME, TARGET, FEATURE_SET)
#endif
-INTEL_IGCA_TARGET(10)
-INTEL_IGCA_TARGET(15)
-INTEL_IGCA_TARGET(20)
-INTEL_IGCA_TARGET(30)
-INTEL_IGCA_TARGET(35)
-INTEL_IGCA_TARGET(40)
-INTEL_IGCA_TARGET(50)
-INTEL_IGCA_TARGET(60)
+INTEL_IGCA_TARGET("igca_10", 10, Core)
+INTEL_IGCA_TARGET("igca_10r", 10, Render)
+
+INTEL_IGCA_TARGET("igca_15", 15, Core)
+INTEL_IGCA_TARGET("igca_15r", 15, Render)
+INTEL_IGCA_TARGET("igca_15ra", 15, RenderExact)
+
+INTEL_IGCA_TARGET("igca_20", 20, Core)
+INTEL_IGCA_TARGET("igca_20c", 20, Compute)
+INTEL_IGCA_TARGET("igca_20ca", 20, ComputeExact)
+
+INTEL_IGCA_TARGET("igca_30", 30, Core)
+INTEL_IGCA_TARGET("igca_30r", 30, Render)
+
+INTEL_IGCA_TARGET("igca_35", 35, Core)
+INTEL_IGCA_TARGET("igca_35r", 35, Render)
+INTEL_IGCA_TARGET("igca_35ra", 35, RenderExact)
+
+INTEL_IGCA_TARGET("igca_40", 40, Core)
+INTEL_IGCA_TARGET("igca_40r", 40, Render)
+
+INTEL_IGCA_TARGET("igca_50", 50, Core)
+INTEL_IGCA_TARGET("igca_50r", 50, Render)
+
+INTEL_IGCA_TARGET("igca_60", 60, Core)
+INTEL_IGCA_TARGET("igca_60c", 60, Compute)
+INTEL_IGCA_TARGET("igca_60r", 60, Render)
#undef INTEL_GPU
#undef INTEL_GPU_COMPAT
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
index 339becdd102225..7d8e1cf1ab4f8b 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
@@ -16,6 +16,7 @@
#ifndef LLVM_TARGETPARSER_INTELGPUTARGETPARSER_H
#define LLVM_TARGETPARSER_INTELGPUTARGETPARSER_H
+#include "llvm/ADT/STLForwardCompat.h"
#include "llvm/ADT/StringRef.h"
#include "llvm/Support/Compiler.h"
#include <cstdint>
@@ -36,7 +37,16 @@ enum GPUKind : uint8_t {
#include "llvm/TargetParser/IntelGPUTargetParser.def"
};
-enum IGCAFeatureSet : uint8_t { IGCA_CORE = 0, IGCA_COMPUTE, IGCA_RENDER };
+/// Represents an IGCA feature set, including whether or not it is exact. The
+/// values are bitfields where the last bit represents exact. See IGCA Target
+/// representation below for more info.
+enum class IGCAFeatureSet : uint8_t {
+ Core = 0b000,
+ Compute = 0b010,
+ ComputeExact = 0b011,
+ Render = 0b100,
+ RenderExact = 0b101,
+};
/// Wrapper around an IGCA target's uint32_t representation, packed using the
/// following format:
@@ -54,15 +64,12 @@ class IGCATarget {
public:
static constexpr uint32_t TargetShift = 16;
- static constexpr uint32_t FeatureSetShift = 1;
- static constexpr uint32_t IsExactShift = 0;
- static constexpr uint32_t FeatureSetMask = 0x3;
+ static constexpr uint32_t FeatureSetMask = 0x7;
static constexpr uint32_t IsExactMask = 0x1;
- constexpr IGCATarget(uint16_t Target, IGCAFeatureSet FeatureSet, bool IsExact)
+ constexpr IGCATarget(uint16_t Target, IGCAFeatureSet FeatureSet)
: V(uint32_t(Target) << TargetShift |
- (uint32_t(FeatureSet) & FeatureSetMask) << FeatureSetShift |
- (uint32_t(IsExact) & IsExactMask) << IsExactShift) {}
+ (llvm::to_underlying(FeatureSet) & FeatureSetMask)) {}
/// \return an invalid IGCATarget.
static constexpr IGCATarget invalid() { return IGCATarget(0); }
@@ -73,18 +80,39 @@ class IGCATarget {
static constexpr IGCATarget unpack(uint32_t V) { return IGCATarget(V); }
uint16_t getTarget() const { return uint16_t(V >> TargetShift); }
+ /// \return the feature set, including whether it is exact.
+ ///
+ /// To check if a target has a feature set regardless of exactness, use
+ /// hasCompute() or hasRender() instead of raw comparison between
+ /// IGCAFeatureSet enums.
IGCAFeatureSet getFeatureSet() const {
- return IGCAFeatureSet((V >> FeatureSetShift) & FeatureSetMask);
+ return IGCAFeatureSet(V & FeatureSetMask);
}
- bool isExact() const { return (V >> IsExactShift) & IsExactMask; }
+ bool isExact() const { return V & IsExactMask; }
bool isValid() const { return getTarget() != 0; }
explicit operator bool() const { return isValid(); }
- bool isCore() const { return getFeatureSet() == IGCA_CORE; }
- bool isCompute() const { return getFeatureSet() == IGCA_COMPUTE; }
- bool isRender() const { return getFeatureSet() == IGCA_RENDER; }
- bool isComputeExact() const { return isCompute() && isExact(); }
- bool isRenderExact() const { return isRender() && isExact(); }
+ bool isCore() const { return getFeatureSet() == IGCAFeatureSet::Core; }
+ /// \return true if the target has a non-exact Compute feature set.
+ bool isCompute() const { return getFeatureSet() == IGCAFeatureSet::Compute; }
+ /// \return true if the target has a non-exact Render feature set.
+ bool isRender() const { return getFeatureSet() == IGCAFeatureSet::Render; }
+ bool isComputeExact() const {
+ return getFeatureSet() == IGCAFeatureSet::ComputeExact;
+ }
+ bool isRenderExact() const {
+ return getFeatureSet() == IGCAFeatureSet::RenderExact;
+ }
+ /// \return true if target has either Compute or ComputeExact feature set.
+ bool hasCompute() const {
+ return (llvm::to_underlying(getFeatureSet()) & ~IsExactMask) ==
+ llvm::to_underlying(IGCAFeatureSet::Compute);
+ }
+ /// \return true if target has either Render or RenderExact feature set.
+ bool hasRender() const {
+ return (llvm::to_underlying(getFeatureSet()) & ~IsExactMask) ==
+ llvm::to_underlying(IGCAFeatureSet::Render);
+ }
friend bool operator==(IGCATarget A, IGCATarget B) { return A.V == B.V; }
friend bool operator!=(IGCATarget A, IGCATarget B) { return A.V != B.V; }
@@ -100,16 +128,16 @@ LLVM_ABI StringRef getArchName(uint32_t GPUIPVersion);
/// \return the numeric name of \p GPUIPVersion, e.g. "xe_35.11.0".
LLVM_ABI std::string getNumericArchName(uint32_t GPUIPVersion);
-/// Parse an IGCA target string, such as "igca_60ca". \return an invalid
+/// Parse an IGCA target string, such as "igca_20ca". \return an invalid
/// IGCATarget if \p TargetStr is not a known target in
/// IntelGPUTargetParser.def
LLVM_ABI IGCATarget parseIGCATarget(StringRef TargetStr);
-/// \return the \p Target as a string, i.e. "igca_60ca". \return an empty string
+/// \return the \p Target as a string, i.e. "igca_20ca". \return an empty string
/// if \p Target is invalid or not a known target in IntelGPUTargetParser.def.
LLVM_ABI StringRef getIGCATargetName(IGCATarget Target);
-/// Append every legal IGCA target spelling to \p Values.
+/// Append every valid IGCA target spelling to \p Values.
LLVM_ABI void fillValidIGCATargetList(SmallVectorImpl<StringRef> &Values);
} // namespace IntelGPU
diff --git a/llvm/lib/TargetParser/IntelGPUTargetParser.cpp b/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
index 0e1163dff9aec8..20335d30dc392b 100644
--- a/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
+++ b/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
@@ -11,6 +11,7 @@
//===----------------------------------------------------------------------===//
#include "llvm/TargetParser/IntelGPUTargetParser.h"
+#include "llvm/ADT/StringSwitch.h"
#include "llvm/ADT/Twine.h"
#include <cassert>
@@ -64,67 +65,31 @@ std::string llvm::IntelGPU::getNumericArchName(uint32_t GPUIPVersion) {
.str();
}
-/// Check that Level is a known IGCA Target in IntelGPUTargetParser.def.
-bool isKnownIGCATargetLevel(uint16_t Level) {
- switch (Level) {
-#define INTEL_IGCA_TARGET(TARGET) case TARGET:
-#include "llvm/TargetParser/IntelGPUTargetParser.def"
- return true;
- default:
- return false;
- }
-}
-
IGCATarget llvm::IntelGPU::parseIGCATarget(StringRef MaybeTarget) {
- if (!MaybeTarget.consume_front("igca_"))
- return IGCATarget::invalid();
-
- uint16_t Target;
- if (MaybeTarget.consumeInteger(10, Target) || !isKnownIGCATargetLevel(Target))
- return IGCATarget::invalid();
-
- IGCAFeatureSet FS = IGCAFeatureSet::IGCA_CORE;
- if (MaybeTarget.consume_front("c"))
- FS = IGCAFeatureSet::IGCA_COMPUTE;
- else if (MaybeTarget.consume_front("r"))
- FS = IGCAFeatureSet::IGCA_RENDER;
- // Exact form needs to either be compute or render:
- bool IsExactFS =
- (FS != IGCAFeatureSet::IGCA_CORE && MaybeTarget.consume_front("a"));
- if (!MaybeTarget.empty())
- return IGCATarget::invalid();
-
- return {Target, FS, IsExactFS};
+ return StringSwitch<IGCATarget>(MaybeTarget)
+#define INTEL_IGCA_TARGET(NAME, TARGET, FEATURE_SET) \
+ .Case(NAME, IGCATarget(TARGET, IGCAFeatureSet::FEATURE_SET))
+#include "llvm/TargetParser/IntelGPUTargetParser.def"
+ .Default(IGCATarget::invalid());
}
StringRef llvm::IntelGPU::getIGCATargetName(IGCATarget T) {
switch (T.pack()) {
-#define INTEL_IGCA_TARGET(TARGET) \
- case IGCATarget(TARGET, IGCAFeatureSet::IGCA_CORE, false).pack(): \
- return "igca_" #TARGET; \
- case IGCATarget(TARGET, IGCAFeatureSet::IGCA_COMPUTE, false).pack(): \
- return "igca_" #TARGET "c"; \
- case IGCATarget(TARGET, IGCAFeatureSet::IGCA_COMPUTE, true).pack(): \
- return "igca_" #TARGET "ca"; \
- case IGCATarget(TARGET, IGCAFeatureSet::IGCA_RENDER, false).pack(): \
- return "igca_" #TARGET "r"; \
- case IGCATarget(TARGET, IGCAFeatureSet::IGCA_RENDER, true).pack(): \
- return "igca_" #TARGET "ra";
+#define INTEL_IGCA_TARGET(NAME, TARGET, FEATURE_SET) \
+ case IGCATarget(TARGET, IGCAFeatureSet::FEATURE_SET).pack(): \
+ return NAME;
#include "llvm/TargetParser/IntelGPUTargetParser.def"
default:
return "";
}
}
-#define INTEL_IGCA_TARGET_FEATURESETS(TARGET) \
- "igca_" #TARGET, "igca_" #TARGET "c", "igca_" #TARGET "ca", \
- "igca_" #TARGET "r", "igca_" #TARGET "ra"
-
+// TODO: Ensure -fsycl --offload-arch provides a list of valid IGCA
+// architectures, similar to how compiling for an nvptx triple returns a list
+// of valid GPU architectures. The user trying to input an invalid IGCA target
+// without being told what IGCA targets actually exist might get confusing.
void llvm::IntelGPU::fillValidIGCATargetList(
SmallVectorImpl<StringRef> &Values) {
-#define INTEL_IGCA_TARGET(TARGET) \
- Values.append({INTEL_IGCA_TARGET_FEATURESETS(TARGET)});
+#define INTEL_IGCA_TARGET(NAME, TARGET, FEATURE_SET) Values.push_back(NAME);
#include "llvm/TargetParser/IntelGPUTargetParser.def"
}
-
-#undef INTEL_IGCA_TARGET_FEATURESETS
diff --git a/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp b/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
index f1c55912305ece..b9685a61c4271d 100644
--- a/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
+++ b/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
@@ -79,33 +79,68 @@ TEST(IntelGPUTargetParserTest, NumericArchName) {
}
TEST(IntelGPUTargetParserTest, IGCATargetBehavior) {
- IntelGPU::IGCATarget T(60, IntelGPU::IGCA_COMPUTE, true);
- EXPECT_TRUE(T.isValid());
- EXPECT_EQ(T.getTarget(), 60);
- EXPECT_EQ(T.getFeatureSet(), IntelGPU::IGCA_COMPUTE);
- EXPECT_TRUE(T.isExact());
- EXPECT_TRUE(T.isComputeExact());
- EXPECT_FALSE(T.isRender());
- EXPECT_EQ(IntelGPU::IGCATarget::unpack(T.pack()), T);
- EXPECT_NE(T, IntelGPU::IGCATarget(60, IntelGPU::IGCA_COMPUTE, false));
+ using IntelGPU::IGCAFeatureSet;
+ using IntelGPU::IGCATarget;
+ IGCATarget C{60, IGCAFeatureSet::ComputeExact};
+ EXPECT_TRUE(C.isValid());
+ EXPECT_EQ(C.getTarget(), 60);
+ EXPECT_EQ(C.getFeatureSet(), IGCAFeatureSet::ComputeExact);
+ EXPECT_TRUE(C.isExact());
+ // An exact feature set is still of its class:
+ EXPECT_TRUE(C.hasCompute());
+ EXPECT_TRUE(C.isComputeExact());
+ EXPECT_FALSE(C.isCore());
+ EXPECT_FALSE(C.isCompute());
+ EXPECT_FALSE(C.isRender());
+ EXPECT_FALSE(C.hasRender());
+ EXPECT_FALSE(C.isRenderExact());
+ EXPECT_EQ(IGCATarget::unpack(C.pack()), C);
+ EXPECT_NE(C, IGCATarget(60, IGCAFeatureSet::Compute));
+ // A feature set that is not exact:
+ IGCATarget R{60, IGCAFeatureSet::Render};
+ EXPECT_TRUE(R.isRender());
+ EXPECT_FALSE(R.isRenderExact());
+ EXPECT_FALSE(R.isExact());
+ EXPECT_FALSE(R.isCompute());
+ EXPECT_TRUE(R.hasRender());
+ EXPECT_FALSE(R.hasCompute());
+ EXPECT_EQ(IGCATarget::unpack(R.pack()), R);
+ IGCATarget Core{60, IGCAFeatureSet::Core};
+ EXPECT_TRUE(Core.isCore());
// Invalid behavior:
- EXPECT_FALSE(IntelGPU::IGCATarget::invalid());
- EXPECT_EQ(IntelGPU::IGCATarget::invalid().pack(), 0u);
+ EXPECT_FALSE(IGCATarget::invalid());
+ EXPECT_EQ(IGCATarget::invalid().pack(), 0u);
}
TEST(IntelGPUTargetParserTest, ParseIGCATarget) {
+ using IntelGPU::IGCAFeatureSet;
using IntelGPU::IGCATarget;
EXPECT_EQ(IntelGPU::parseIGCATarget("igca_10"),
- IGCATarget(10, IntelGPU::IGCA_CORE, false));
+ IGCATarget(10, IGCAFeatureSet::Core));
EXPECT_EQ(IntelGPU::parseIGCATarget("igca_20c"),
- IGCATarget(20, IntelGPU::IGCA_COMPUTE, false));
+ IGCATarget(20, IGCAFeatureSet::Compute));
EXPECT_EQ(IntelGPU::parseIGCATarget("igca_20ca"),
- IGCATarget(20, IntelGPU::IGCA_COMPUTE, true));
+ IGCATarget(20, IGCAFeatureSet::ComputeExact));
EXPECT_EQ(IntelGPU::parseIGCATarget("igca_15r"),
- IGCATarget(15, IntelGPU::IGCA_RENDER, false));
+ IGCATarget(15, IGCAFeatureSet::Render));
EXPECT_EQ(IntelGPU::parseIGCATarget("igca_15ra"),
- IGCATarget(15, IntelGPU::IGCA_RENDER, true));
- // Only levels in the table are valid targets.
+ IGCATarget(15, IGCAFeatureSet::RenderExact));
+ EXPECT_EQ(IntelGPU::parseIGCATarget("igca_35ra"),
+ IGCATarget(35, IGCAFeatureSet::RenderExact));
+ EXPECT_EQ(IntelGPU::parseIGCATarget("igca_60c"),
+ IGCATarget(60, IGCAFeatureSet::Compute));
+ EXPECT_EQ(IntelGPU::parseIGCATarget("igca_60r"),
+ IGCATarget(60, IGCAFeatureSet::Render));
+ // Not all target levels support all feature sets:
+ EXPECT_FALSE(IntelGPU::parseIGCATarget("igca_10c"));
+ EXPECT_FALSE(IntelGPU::parseIGCATarget("igca_20r"));
+ EXPECT_FALSE(IntelGPU::parseIGCATarget("igca_20ra"));
+ // Only some target levels have an exact feature set:
+ EXPECT_FALSE(IntelGPU::parseIGCATarget("igca_10ra"));
+ EXPECT_FALSE(IntelGPU::parseIGCATarget("igca_30ra"));
+ EXPECT_FALSE(IntelGPU::parseIGCATarget("igca_60ca"));
+ EXPECT_FALSE(IntelGPU::parseIGCATarget("igca_60ra"));
+ // Only target levels in the table are valid targets:
EXPECT_FALSE(IntelGPU::parseIGCATarget("igca_42"));
EXPECT_FALSE(IntelGPU::parseIGCATarget("igca_0"));
// Cannot have exact without naming feature set:
@@ -121,25 +156,25 @@ TEST(IntelGPUTargetParserTest, ParseIGCATarget) {
}
TEST(IntelGPUTargetParserTest, IGCATargetName) {
+ using IntelGPU::IGCAFeatureSet;
using IntelGPU::IGCATarget;
+ EXPECT_EQ(IntelGPU::getIGCATargetName(IGCATarget(60, IGCAFeatureSet::Core)),
+ "igca_60");
EXPECT_EQ(
- IntelGPU::getIGCATargetName(IGCATarget(60, IntelGPU::IGCA_CORE, false)),
- "igca_60");
- EXPECT_EQ(
- IntelGPU::getIGCATargetName(IGCATarget(60, IntelGPU::IGCA_COMPUTE, true)),
- "igca_60ca");
- EXPECT_EQ(
- IntelGPU::getIGCATargetName(IGCATarget(15, IntelGPU::IGCA_RENDER, false)),
- "igca_15r");
+ IntelGPU::getIGCATargetName(IGCATarget(20, IGCAFeatureSet::ComputeExact)),
+ "igca_20ca");
+ EXPECT_EQ(IntelGPU::getIGCATargetName(IGCATarget(15, IGCAFeatureSet::Render)),
+ "igca_15r");
+ // A target that does not exist / is not valid should produce "" to signify
+ // invalid target.
EXPECT_EQ(IntelGPU::getIGCATargetName(IGCATarget::invalid()), "");
- // A level that is not in the table has no spelling, and neither does an
- // exact core target.
+ EXPECT_EQ(IntelGPU::getIGCATargetName(IGCATarget(11, IGCAFeatureSet::Core)),
+ "");
EXPECT_EQ(
- IntelGPU::getIGCATargetName(IGCATarget(11, IntelGPU::IGCA_CORE, false)),
+ IntelGPU::getIGCATargetName(IGCATarget(60, IGCAFeatureSet::ComputeExact)),
"");
EXPECT_EQ(
- IntelGPU::getIGCATargetName(IGCATarget(10, IntelGPU::IGCA_CORE, true)),
- "");
+ IntelGPU::getIGCATargetName(IGCATarget(10, IGCAFeatureSet::Compute)), "");
}
TEST(IntelGPUTargetParserTest, EveryIGCASpellingRoundTrips) {
@@ -153,16 +188,46 @@ TEST(IntelGPUTargetParserTest, EveryIGCASpellingRoundTrips) {
}
}
-TEST(IntelGPUTargetParserTest, EveryDeviceHasValidIGCATargetLevel) {
- // Every device's defined IGCA target should have a corresponding entry in
- // INTEL_IGCA_TARGET:
-#define INTEL_IGCA_TARGET_CHECK(NAME, IGCA_TARGET) \
- EXPECT_TRUE(IntelGPU::parseIGCATarget("igca_" #IGCA_TARGET).isValid()) \
+// Spell an IGCA target from its fields, independently of the NAME column.
+std::string expectedIGCASpelling(uint16_t Target,
+ IntelGPU::IGCAFeatureSet FeatureSet) {
+ std::string Name = "igca_" + std::to_string(Target);
+ switch (FeatureSet) {
+ case IntelGPU::IGCAFeatureSet::Core:
+ return Name;
+ case IntelGPU::IGCAFeatureSet::Compute:
+ return Name + "c";
+ case IntelGPU::IGCAFeatureSet::ComputeExact:
+ return Name + "ca";
+ case IntelGPU::IGCAFeatureSet::Render:
+ return Name + "r";
+ case IntelGPU::IGCAFeatureSet::RenderExact:
+ return Name + "ra";
+ }
+ return "";
+}
+
+TEST(IntelGPUTargetParserTest, IGCASpellingsMatchFields) {
+ // Every INTEL_IGCA_TARGET's NAME must agree with information in the rest of
+ // its row.
+#define INTEL_IGCA_TARGET(NAME, TARGET, FEATURE_SET) \
+ EXPECT_EQ(NAME, expectedIGCASpelling( \
+ TARGET, IntelGPU::IGCAFeatureSet::FEATURE_SET));
+#include "llvm/TargetParser/IntelGPUTargetParser.def"
+}
+
+TEST(IntelGPUTargetParserTest, EveryDeviceHasValidIGCATarget) {
+ // Every device's IGCA target and feature sets should have a corresponding
+ // entry in INTEL_IGCA_TARGET:
+#define INTEL_IGCA_TARGET_CHECK(NAME, IGCA_TARGET, IGCA_FEATURE_SETS) \
+ EXPECT_NE(IntelGPU::getIGCATargetName(IntelGPU::IGCATarget( \
+ IGCA_TARGET, IntelGPU::IGCAFeatureSet::IGCA_FEATURE_SETS)), \
+ "") \
<< NAME;
#define INTEL_GPU(NAME, KIND, MAJOR, MINOR, IGCA_TARGET, IGCA_FEATURE_SETS) \
- INTEL_IGCA_TARGET_CHECK(NAME, IGCA_TARGET)
+ INTEL_IGCA_TARGET_CHECK(NAME, IGCA_TARGET, IGCA_FEATURE_SETS)
#define INTEL_GPU_COMPAT(NAME, KIND, IGCA_TARGET, IGCA_FEATURE_SETS) \
- INTEL_IGCA_TARGET_CHECK(NAME, IGCA_TARGET)
+ INTEL_IGCA_TARGET_CHECK(NAME, IGCA_TARGET, IGCA_FEATURE_SETS)
#include "llvm/TargetParser/IntelGPUTargetParser.def"
#undef INTEL_IGCA_TARGET_CHECK
}
>From 88050039b3e3cd31c802211b58c8d9ff8f65546b Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 22 Oct 2026 13:56:30 -0700
Subject: [PATCH 05/30] Reject bad igca targets when fed into --offload-arch,
add tests
---
clang/lib/Driver/Driver.cpp | 5 +++++
clang/test/Driver/sycl-offload-arch-igca.cpp | 19 +++++++++++++++++++
2 files changed, 24 insertions(+)
create mode 100644 clang/test/Driver/sycl-offload-arch-igca.cpp
diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp
index a22a81a8ba6e1c..741ddefad97e49 100644
--- a/clang/lib/Driver/Driver.cpp
+++ b/clang/lib/Driver/Driver.cpp
@@ -3928,6 +3928,11 @@ static StringRef getCanonicalArchString(Compilation &C,
return StringRef();
}
}
+ } else if (Triple.isSPIRV() && ArchStr.starts_with("igca_") &&
+ !Arch.isIntelIGCA()) {
+ C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
+ << "IGCA" << ArchStr;
+ return StringRef();
}
if (Arch.isNVPTX())
diff --git a/clang/test/Driver/sycl-offload-arch-igca.cpp b/clang/test/Driver/sycl-offload-arch-igca.cpp
new file mode 100644
index 00000000000000..ecb1705bcc36d7
--- /dev/null
+++ b/clang/test/Driver/sycl-offload-arch-igca.cpp
@@ -0,0 +1,19 @@
+/// Tests for Intel IGCA targets in --offload-arch.
+
+// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fsycl \
+// RUN: --offload-targets=spirv64-unknown-unknown --offload-arch=igca_20ca \
+// RUN: -c %s 2>&1 | FileCheck %s
+// CHECK: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spirv64-unknown-unknown,arch=igca_20ca,kind=sycl"
+
+/// Reject IGCA targets not listed in IntelGPUTargetParser.def:
+// RUN: not %clang -### --target=x86_64-unknown-linux-gnu -fsycl \
+// RUN: --offload-targets=spirv64-unknown-unknown --offload-arch=igca_42a \
+// RUN: -c %s 2>&1 | FileCheck -check-prefix=CHK-INVALID %s
+// CHK-INVALID: error: unsupported IGCA gpu architecture: igca_42a
+// CHK-INVALID-NOT: llvm-offload-binary
+
+/// TODO: Enable this when OffloadArchToTriple gives IGCA targets a triple.
+// TODO-RUN %clang -### --target=x86_64-unknown-linux-gnu -fsycl \
+// TODO-RUN --offload-arch=igca_20ca -c %s 2>&1 | FileCheck -check-prefix=CHK-INFER %s
+// CHK-INFER-NOT: error:
+// CHK-INFER: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spirv64{{[^,]*}},arch=igca_20ca,kind=sycl"
>From 51e0883c2d86f6b71fd1e210a8c3ac8c908dd276 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 22 Oct 2026 15:10:39 -0700
Subject: [PATCH 06/30] Address TODOs
---
clang/include/clang/Basic/OffloadArch.h | 7 +---
.../clang-sycl-linker/ClangSYCLLinker.cpp | 4 ++
clang/unittests/Basic/OffloadArchTest.cpp | 41 +++++++++++++++++--
3 files changed, 43 insertions(+), 9 deletions(-)
diff --git a/clang/include/clang/Basic/OffloadArch.h b/clang/include/clang/Basic/OffloadArch.h
index bdbbcc2c8b13f2..cc3de55ca4bf92 100644
--- a/clang/include/clang/Basic/OffloadArch.h
+++ b/clang/include/clang/Basic/OffloadArch.h
@@ -90,10 +90,7 @@ class OffloadArch {
bool isAMDGCNSPIRV() const { return V == TargetArch::AMDGCNSPIRV; }
bool isIntelCPU() const { return V == TargetArch::IntelCPU; }
bool isIntelGPU() const { return V == TargetArch::IntelGPU; }
- // TODO: Decide whether IntelIGCA counts as Intel here. clang-sycl-linker
- // uses isIntel() to decide on AOT compilation, and runAOTCompile would then
- // need an IntelIGCA dispatch.
- bool isIntel() const { return isIntelCPU() || isIntelGPU(); }
+ bool isIntel() const { return isIntelCPU() || isIntelGPU() || isIntelIGCA(); }
bool isGeneric() const { return V == TargetArch::Generic; }
bool isUnused() const { return V == TargetArch::Unused; }
bool isUnknown() const { return V == TargetArch::Unknown; }
@@ -116,8 +113,6 @@ class OffloadArch {
bool operator<(const OffloadArch &Other) const {
return std::tie(V, Kind) < std::tie(Other.V, Other.Kind);
}
- // TODO test cases for whether or not IntelIGCA functions properly for these
- // comparisons
};
const char *OffloadArchToString(OffloadArch A);
diff --git a/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp b/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp
index 53ef6375039ac3..ea5ce961677d19 100644
--- a/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp
+++ b/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp
@@ -748,6 +748,10 @@ static Error runAOTCompile(StringRef InputFile, StringRef OutputFile,
return runAOTCompileIntelGPU(InputFile, OutputFile, Args);
if (OA.isIntelCPU())
return runAOTCompileIntelCPU(InputFile, OutputFile, Args);
+ // FIXME(IGCA): Implement IGCA behavior.
+ if (OA.isIntelIGCA())
+ return createStringError("AOT compilation for IGCA target '" + Arch +
+ "' is not yet implemented");
llvm_unreachable("runAOTCompile dispatched on unsupported arch");
}
diff --git a/clang/unittests/Basic/OffloadArchTest.cpp b/clang/unittests/Basic/OffloadArchTest.cpp
index f46ec566c7f037..319d510ec88562 100644
--- a/clang/unittests/Basic/OffloadArchTest.cpp
+++ b/clang/unittests/Basic/OffloadArchTest.cpp
@@ -7,6 +7,7 @@
//===----------------------------------------------------------------------===//
#include "clang/Basic/OffloadArch.h"
+#include "llvm/TargetParser/IntelGPUTargetParser.h"
#include "llvm/TargetParser/Triple.h"
#include "gtest/gtest.h"
@@ -41,6 +42,15 @@ TEST(OffloadArchTest, TargetArchClassification) {
EXPECT_FALSE(IntelGPU.isIntelCPU());
EXPECT_TRUE(IntelGPU.isIntelGPU());
+ OffloadArch IGCA = parse("igca_20ca");
+ EXPECT_TRUE(IGCA.isIntelIGCA());
+ EXPECT_FALSE(IGCA.isIntelGPU());
+ EXPECT_FALSE(IGCA.isIntelCPU());
+ EXPECT_TRUE(IGCA.isIntel());
+ EXPECT_FALSE(IGCA.isNVPTX());
+ EXPECT_FALSE(IGCA.isAMDGPU());
+ EXPECT_FALSE(IGCA.isAMDGCNSPIRV());
+
OffloadArch Generic = parse("generic");
EXPECT_FALSE(Generic.isNVPTX());
EXPECT_FALSE(Generic.isAMDGPU());
@@ -54,12 +64,11 @@ TEST(OffloadArchTest, Unknown) {
}
// Names must round-trip through parse -> string.
-// TODO: Add IGCA targets (e.g. "igca_20ca"), plus IntelIGCA classification and
-// ==/< ordering tests.
TEST(OffloadArchTest, RoundTrip) {
for (const char *Name :
{"sm_52", "sm_90a", "gfx906", "gfx1201", "gfx12-generic", "amdgcnspirv",
- "graniterapids", "bmg_g21", "generic"}) {
+ "graniterapids", "bmg_g21", "igca_10", "igca_20c", "igca_20ca",
+ "igca_15r", "igca_15ra", "generic"}) {
OffloadArch A = parse(Name);
EXPECT_FALSE(A.isUnknown()) << Name;
EXPECT_STREQ(OffloadArchToString(A), Name);
@@ -94,3 +103,29 @@ TEST(OffloadArchTest, AMDGPUSubArchRoundTrip) {
EXPECT_EQ(getOffloadArchSubArch(OffloadArch::getAMDGCNSPIRV()),
llvm::Triple::NoSubArch);
}
+
+// Test Intel IGCA target equal and comparsion operators:
+TEST(OffloadArchTest, IntelIGCAComparison) {
+ // Listed in ascending order: by target level, then by feature set.
+ const char *Ordered[] = {"igca_10", "igca_15r", "igca_20c", "igca_20ca",
+ "igca_35"};
+ for (size_t I = 0; I < std::size(Ordered); ++I) {
+ OffloadArch A = parse(Ordered[I]);
+ ASSERT_TRUE(A.isIntelIGCA()) << Ordered[I];
+ EXPECT_EQ(A, parse(Ordered[I])) << Ordered[I];
+ EXPECT_FALSE(A < A) << Ordered[I];
+ if (I + 1 < std::size(Ordered)) {
+ OffloadArch Next = parse(Ordered[I + 1]);
+ EXPECT_NE(A, Next) << Ordered[I] << " vs " << Ordered[I + 1];
+ EXPECT_TRUE(A < Next) << Ordered[I] << " vs " << Ordered[I + 1];
+ EXPECT_FALSE(Next < A) << Ordered[I] << " vs " << Ordered[I + 1];
+ }
+ }
+
+ // The same 'Kind' under another TargetArch is still a different architecture.
+ OffloadArch IGCA = parse("igca_20ca");
+ OffloadArch SameKind = OffloadArch::getIntel(
+ OffloadArch::TargetArch::IntelGPU,
+ static_cast<OffloadArch::IntelArch>(IGCA.igcaTarget().pack()));
+ EXPECT_NE(IGCA, SameKind);
+}
>From 0811ea6ff830f23c459006d6ab945b019214cd31 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 22 Oct 2026 15:13:04 -0700
Subject: [PATCH 07/30] add nice label for which TODOs are related to IGCA
---
clang/lib/Basic/OffloadArch.cpp | 2 +-
clang/test/Driver/sycl-offload-arch-igca.cpp | 2 +-
2 files changed, 2 insertions(+), 2 deletions(-)
diff --git a/clang/lib/Basic/OffloadArch.cpp b/clang/lib/Basic/OffloadArch.cpp
index c3004bd3333cb0..00ba58cc8550b0 100644
--- a/clang/lib/Basic/OffloadArch.cpp
+++ b/clang/lib/Basic/OffloadArch.cpp
@@ -134,7 +134,7 @@ llvm::Triple OffloadArchToTriple(const llvm::Triple &DefaultToolchainTriple,
return llvm::Triple(llvm::Triple::amdgpu, llvm::Triple::NoSubArch,
llvm::Triple::AMD, llvm::Triple::AMDHSA);
- // TODO: Handle IntelIGCA, which currently gets an empty triple, so
+ // FIXME(IGCA): Handle IntelIGCA, which currently gets an empty triple, so
// --offload-arch=igca_* cannot pick an offload toolchain.
return {};
}
diff --git a/clang/test/Driver/sycl-offload-arch-igca.cpp b/clang/test/Driver/sycl-offload-arch-igca.cpp
index ecb1705bcc36d7..5e2c45f180a913 100644
--- a/clang/test/Driver/sycl-offload-arch-igca.cpp
+++ b/clang/test/Driver/sycl-offload-arch-igca.cpp
@@ -12,7 +12,7 @@
// CHK-INVALID: error: unsupported IGCA gpu architecture: igca_42a
// CHK-INVALID-NOT: llvm-offload-binary
-/// TODO: Enable this when OffloadArchToTriple gives IGCA targets a triple.
+/// FIXME(IGCA): Enable this when OffloadArchToTriple gives IGCA targets a triple.
// TODO-RUN %clang -### --target=x86_64-unknown-linux-gnu -fsycl \
// TODO-RUN --offload-arch=igca_20ca -c %s 2>&1 | FileCheck -check-prefix=CHK-INFER %s
// CHK-INFER-NOT: error:
>From 12eb2960b1fdb969da5f4bf8c9ddf3fb740cd46c Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 22 Oct 2026 15:41:43 -0700
Subject: [PATCH 08/30] mark TODO as fixme(igca)
---
llvm/lib/TargetParser/IntelGPUTargetParser.cpp | 7 +++----
1 file changed, 3 insertions(+), 4 deletions(-)
diff --git a/llvm/lib/TargetParser/IntelGPUTargetParser.cpp b/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
index 20335d30dc392b..c6c3fe2a9acab1 100644
--- a/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
+++ b/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
@@ -84,10 +84,9 @@ StringRef llvm::IntelGPU::getIGCATargetName(IGCATarget T) {
}
}
-// TODO: Ensure -fsycl --offload-arch provides a list of valid IGCA
-// architectures, similar to how compiling for an nvptx triple returns a list
-// of valid GPU architectures. The user trying to input an invalid IGCA target
-// without being told what IGCA targets actually exist might get confusing.
+// FIXME(IGCA): Make -fsycl --offload-arch provides a list of valid IGCA
+// targets upon a bad IGCA target, similar to how compiling for an nvptx triple
+// returns a list of valid GPU architectures.
void llvm::IntelGPU::fillValidIGCATargetList(
SmallVectorImpl<StringRef> &Values) {
#define INTEL_IGCA_TARGET(NAME, TARGET, FEATURE_SET) Values.push_back(NAME);
>From 4c0f0bf4341df205ae2c02b3307105dd5e3a7e65 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 22 Oct 2026 15:53:40 -0700
Subject: [PATCH 09/30] Fix + adjust comments to match LLVM comment standards
---
clang/include/clang/Basic/OffloadArch.h | 2 +-
.../tools/clang-sycl-linker/ClangSYCLLinker.cpp | 2 +-
clang/unittests/Basic/OffloadArchTest.cpp | 2 +-
.../llvm/TargetParser/IntelGPUTargetParser.def | 7 ++++---
.../llvm/TargetParser/IntelGPUTargetParser.h | 17 +++++++++--------
5 files changed, 16 insertions(+), 14 deletions(-)
diff --git a/clang/include/clang/Basic/OffloadArch.h b/clang/include/clang/Basic/OffloadArch.h
index cc3de55ca4bf92..437f85f99843f4 100644
--- a/clang/include/clang/Basic/OffloadArch.h
+++ b/clang/include/clang/Basic/OffloadArch.h
@@ -42,7 +42,7 @@ class OffloadArch {
AMDGCNSPIRV, // The 'amdgcnspirv' pseudo target.
IntelCPU, // Kind is an IntelArch.
IntelGPU, // Kind is an IntelArch.
- IntelIGCA, // Kind is an Intel IGCA Target
+ IntelIGCA, // Kind is an llvm::IntelGPU::IGCATarget.
Generic, // The 'generic' processor model.
};
diff --git a/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp b/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp
index ea5ce961677d19..690d6ee07a1026 100644
--- a/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp
+++ b/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp
@@ -748,7 +748,7 @@ static Error runAOTCompile(StringRef InputFile, StringRef OutputFile,
return runAOTCompileIntelGPU(InputFile, OutputFile, Args);
if (OA.isIntelCPU())
return runAOTCompileIntelCPU(InputFile, OutputFile, Args);
- // FIXME(IGCA): Implement IGCA behavior.
+ // FIXME(IGCA): Implement AOT compilation for IGCA targets.
if (OA.isIntelIGCA())
return createStringError("AOT compilation for IGCA target '" + Arch +
"' is not yet implemented");
diff --git a/clang/unittests/Basic/OffloadArchTest.cpp b/clang/unittests/Basic/OffloadArchTest.cpp
index 319d510ec88562..ad2bc3a2ae00f5 100644
--- a/clang/unittests/Basic/OffloadArchTest.cpp
+++ b/clang/unittests/Basic/OffloadArchTest.cpp
@@ -104,7 +104,7 @@ TEST(OffloadArchTest, AMDGPUSubArchRoundTrip) {
llvm::Triple::NoSubArch);
}
-// Test Intel IGCA target equal and comparsion operators:
+// Test equality and ordering of Intel IGCA targets.
TEST(OffloadArchTest, IntelIGCAComparison) {
// Listed in ascending order: by target level, then by feature set.
const char *Ordered[] = {"igca_10", "igca_15r", "igca_20c", "igca_20ca",
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def
index 7ad6b201c3081d..de276875351bec 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def
@@ -17,8 +17,8 @@
// GK_<KIND>.
// MAJOR - Major component of the device's GPU IP version.
// MINOR - Minor component of the device's GPU IP version.
-// IGCA_TARGET - Numeric IGCA target: Targets should also be defined
-// in the INTEL_IGCA_TARGET table.
+// IGCA_TARGET - Numeric IGCA target: Every target should also be
+// listed in the INTEL_IGCA_TARGET table.
// IGCA_FEATURE_SETS - Token naming the IGCA feature sets the device
// implements: Core (none), Compute ("c"),
// Render ("r"), ComputeExact ("ca") or
@@ -39,7 +39,8 @@
// so that the name to print for such a group comes first.
//
// This file also includes a list of valid IGCA targets, which can be used to
-// compile binaries with support different GPU architectures and feature sets:
+// compile binaries with support for different GPU architectures and feature
+// sets:
//
// INTEL_IGCA_TARGET(NAME, TARGET, FEATURE_SET)
// One row per valid IGCA target spelling. Only the combinations listed
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
index 7d8e1cf1ab4f8b..40b440431172f1 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
@@ -38,8 +38,8 @@ enum GPUKind : uint8_t {
};
/// Represents an IGCA feature set, including whether or not it is exact. The
-/// values are bitfields where the last bit represents exact. See IGCA Target
-/// representation below for more info.
+/// values are bitfields where the lowest bit represents exact. See IGCATarget
+/// for the packed representation.
enum class IGCAFeatureSet : uint8_t {
Core = 0b000,
Compute = 0b010,
@@ -71,7 +71,7 @@ class IGCATarget {
: V(uint32_t(Target) << TargetShift |
(llvm::to_underlying(FeatureSet) & FeatureSetMask)) {}
- /// \return an invalid IGCATarget.
+ /// \return an invalid IGCATarget (whose packed representation is 0).
static constexpr IGCATarget invalid() { return IGCATarget(0); }
/// \return the packed uint32_t representation of the IGCA target.
@@ -128,13 +128,14 @@ LLVM_ABI StringRef getArchName(uint32_t GPUIPVersion);
/// \return the numeric name of \p GPUIPVersion, e.g. "xe_35.11.0".
LLVM_ABI std::string getNumericArchName(uint32_t GPUIPVersion);
-/// Parse an IGCA target string, such as "igca_20ca". \return an invalid
-/// IGCATarget if \p TargetStr is not a known target in
-/// IntelGPUTargetParser.def
+/// Parse an IGCA target string such as "igca_20ca".
+///
+/// \return an invalid IGCATarget if \p TargetStr is not a known target in
+/// IntelGPUTargetParser.def.
LLVM_ABI IGCATarget parseIGCATarget(StringRef TargetStr);
-/// \return the \p Target as a string, i.e. "igca_20ca". \return an empty string
-/// if \p Target is invalid or not a known target in IntelGPUTargetParser.def.
+/// \return the \p Target as a string, i.e. "igca_20ca", or an empty string if
+/// \p Target is invalid or not a known target in IntelGPUTargetParser.def.
LLVM_ABI StringRef getIGCATargetName(IGCATarget Target);
/// Append every valid IGCA target spelling to \p Values.
>From b782fbe517b0cfd08aaac2646a79da53f3a14b64 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 22 Oct 2026 16:32:15 -0700
Subject: [PATCH 10/30] Unify invalid igcatarget value
---
.../llvm/TargetParser/IntelGPUTargetParser.h | 26 ++++++++++++-------
.../TargetParser/IntelGPUTargetParserTest.cpp | 6 +++++
2 files changed, 22 insertions(+), 10 deletions(-)
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
index 40b440431172f1..7b1cd67d94d038 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
@@ -51,16 +51,22 @@ enum class IGCAFeatureSet : uint8_t {
/// Wrapper around an IGCA target's uint32_t representation, packed using the
/// following format:
///
-/// 31 16 15 3 2 1 0 0
-/// +------------------+--------------+-------------+-----------+
-/// | Target | Reserved | Feature set | Is Exact? |
-/// +------------------+--------------+-------------+-----------+
-/// 16 bits 13 bits 2 bits 1 bit
+/// 31 16 15 3 2 0
+/// +------------------+------------------------+---------------+
+/// | Target | Reserved | Feature set |
+/// +------------------+------------------------+---------------+
+/// 16 bits 13 bits 3 bits
///
-/// Target == 0 denotes an invalid IGCA Target.
+/// Note that bit 0 is a flag that represents if the feature set is "exact",
+/// and Target == 0 denotes an invalid IGCA Target.
class IGCATarget {
uint32_t V = 0;
- constexpr explicit IGCATarget(uint32_t V) : V(V) {}
+ constexpr explicit IGCATarget(uint32_t Packed) {
+ // Target == 0 is invalid, so V remains 0, and any feature set bits are
+ // ignored. This keeps V == 0 as the only canonical invalid representation.
+ if (uint16_t Target = Packed >> TargetShift; Target != 0)
+ V = Packed;
+ }
public:
static constexpr uint32_t TargetShift = 16;
@@ -68,8 +74,8 @@ class IGCATarget {
static constexpr uint32_t IsExactMask = 0x1;
constexpr IGCATarget(uint16_t Target, IGCAFeatureSet FeatureSet)
- : V(uint32_t(Target) << TargetShift |
- (llvm::to_underlying(FeatureSet) & FeatureSetMask)) {}
+ : IGCATarget(uint32_t(Target) << TargetShift |
+ (llvm::to_underlying(FeatureSet) & FeatureSetMask)) {}
/// \return an invalid IGCATarget (whose packed representation is 0).
static constexpr IGCATarget invalid() { return IGCATarget(0); }
@@ -115,7 +121,7 @@ class IGCATarget {
}
friend bool operator==(IGCATarget A, IGCATarget B) { return A.V == B.V; }
- friend bool operator!=(IGCATarget A, IGCATarget B) { return A.V != B.V; }
+ friend bool operator!=(IGCATarget A, IGCATarget B) { return !(A == B); }
};
/// \return the device name that \p GPUIPVersion identifies, as reported by the
diff --git a/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp b/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
index b9685a61c4271d..786394236fd92a 100644
--- a/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
+++ b/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
@@ -110,6 +110,12 @@ TEST(IntelGPUTargetParserTest, IGCATargetBehavior) {
// Invalid behavior:
EXPECT_FALSE(IGCATarget::invalid());
EXPECT_EQ(IGCATarget::invalid().pack(), 0u);
+ EXPECT_EQ(IGCATarget(0, IGCAFeatureSet::Compute), IGCATarget::invalid());
+ // unpack(2) is invalid since it results in Target = 0. It should thus equate
+ // to IGCATarget::invalid().
+ EXPECT_EQ(IGCATarget::unpack(2), IGCATarget::invalid());
+ EXPECT_FALSE(IGCATarget::unpack(2).isValid());
+ EXPECT_EQ(IGCATarget(0, IGCAFeatureSet::RenderExact), IGCATarget::invalid());
}
TEST(IntelGPUTargetParserTest, ParseIGCATarget) {
>From 48aa3071adad1c5e65aa6fd7edcaded17f043cc5 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 22 Oct 2026 17:24:00 -0700
Subject: [PATCH 11/30] Make isvalid actually validate IGCA target validity
---
.../llvm/TargetParser/IntelGPUTargetParser.h | 32 ++++++++++++++-----
.../TargetParser/IntelGPUTargetParserTest.cpp | 29 ++++++++++-------
2 files changed, 42 insertions(+), 19 deletions(-)
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
index 7b1cd67d94d038..f4d44407d0daf7 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
@@ -57,14 +57,30 @@ enum class IGCAFeatureSet : uint8_t {
/// +------------------+------------------------+---------------+
/// 16 bits 13 bits 3 bits
///
-/// Note that bit 0 is a flag that represents if the feature set is "exact",
-/// and Target == 0 denotes an invalid IGCA Target.
+/// Note that bit 0 is a flag that represents if the feature set is "exact".
+/// Only targets listed in INTEL_IGCA_TARGET are valid; any other value is
+/// stored as invalid(), whose packed representation is 0.
class IGCATarget {
uint32_t V = 0;
+
+ /// \return true if \p Packed contains a target listed in INTEL_IGCA_TARGET.
+ static constexpr bool isKnownIGCATargetLevel(uint32_t Packed) {
+ switch (Packed) {
+#define INTEL_IGCA_TARGET(NAME, TARGET, FEATURE_SET) \
+ case uint32_t(TARGET) << TargetShift | \
+ llvm::to_underlying(IGCAFeatureSet::FEATURE_SET):
+#include "llvm/TargetParser/IntelGPUTargetParser.def"
+ return true;
+ default:
+ return false;
+ }
+ }
+
constexpr explicit IGCATarget(uint32_t Packed) {
- // Target == 0 is invalid, so V remains 0, and any feature set bits are
- // ignored. This keeps V == 0 as the only canonical invalid representation.
- if (uint16_t Target = Packed >> TargetShift; Target != 0)
+ // Targets not listed in INTEL_IGCA_TARGET (including Target == 0) are
+ // invalid, so V remains 0. This keeps V == 0 as the only canonical invalid
+ // representation.
+ if (isKnownIGCATargetLevel(Packed))
V = Packed;
}
@@ -85,7 +101,7 @@ class IGCATarget {
/// Wrap a packed uint32_t IGCA target with an IGCATarget class.
static constexpr IGCATarget unpack(uint32_t V) { return IGCATarget(V); }
- uint16_t getTarget() const { return uint16_t(V >> TargetShift); }
+ constexpr uint16_t getTarget() const { return uint16_t(V >> TargetShift); }
/// \return the feature set, including whether it is exact.
///
/// To check if a target has a feature set regardless of exactness, use
@@ -95,8 +111,8 @@ class IGCATarget {
return IGCAFeatureSet(V & FeatureSetMask);
}
bool isExact() const { return V & IsExactMask; }
- bool isValid() const { return getTarget() != 0; }
- explicit operator bool() const { return isValid(); }
+ constexpr bool isValid() const { return getTarget() != 0; }
+ constexpr explicit operator bool() const { return isValid(); }
bool isCore() const { return getFeatureSet() == IGCAFeatureSet::Core; }
/// \return true if the target has a non-exact Compute feature set.
diff --git a/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp b/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
index 786394236fd92a..8c288d1cbf5704 100644
--- a/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
+++ b/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
@@ -81,9 +81,9 @@ TEST(IntelGPUTargetParserTest, NumericArchName) {
TEST(IntelGPUTargetParserTest, IGCATargetBehavior) {
using IntelGPU::IGCAFeatureSet;
using IntelGPU::IGCATarget;
- IGCATarget C{60, IGCAFeatureSet::ComputeExact};
+ IGCATarget C{20, IGCAFeatureSet::ComputeExact};
EXPECT_TRUE(C.isValid());
- EXPECT_EQ(C.getTarget(), 60);
+ EXPECT_EQ(C.getTarget(), 20);
EXPECT_EQ(C.getFeatureSet(), IGCAFeatureSet::ComputeExact);
EXPECT_TRUE(C.isExact());
// An exact feature set is still of its class:
@@ -95,7 +95,7 @@ TEST(IntelGPUTargetParserTest, IGCATargetBehavior) {
EXPECT_FALSE(C.hasRender());
EXPECT_FALSE(C.isRenderExact());
EXPECT_EQ(IGCATarget::unpack(C.pack()), C);
- EXPECT_NE(C, IGCATarget(60, IGCAFeatureSet::Compute));
+ EXPECT_NE(C, IGCATarget(20, IGCAFeatureSet::Compute));
// A feature set that is not exact:
IGCATarget R{60, IGCAFeatureSet::Render};
EXPECT_TRUE(R.isRender());
@@ -116,6 +116,15 @@ TEST(IntelGPUTargetParserTest, IGCATargetBehavior) {
EXPECT_EQ(IGCATarget::unpack(2), IGCATarget::invalid());
EXPECT_FALSE(IGCATarget::unpack(2).isValid());
EXPECT_EQ(IGCATarget(0, IGCAFeatureSet::RenderExact), IGCATarget::invalid());
+ // Targets not listed in INTEL_IGCA_TARGET are invalid:
+ EXPECT_FALSE(IGCATarget(42, IGCAFeatureSet::Core).isValid());
+ EXPECT_EQ(IGCATarget(42, IGCAFeatureSet::Core), IGCATarget::invalid());
+ EXPECT_EQ(IGCATarget(60, IGCAFeatureSet::ComputeExact),
+ IGCATarget::invalid());
+ // If the reserved bits are set, something probably went wrong:
+ EXPECT_EQ(
+ IGCATarget::unpack(IGCATarget(20, IGCAFeatureSet::Compute).pack() | 0x8),
+ IGCATarget::invalid());
}
TEST(IntelGPUTargetParserTest, ParseIGCATarget) {
@@ -222,20 +231,18 @@ TEST(IntelGPUTargetParserTest, IGCASpellingsMatchFields) {
#include "llvm/TargetParser/IntelGPUTargetParser.def"
}
-TEST(IntelGPUTargetParserTest, EveryDeviceHasValidIGCATarget) {
- // Every device's IGCA target and feature sets should have a corresponding
- // entry in INTEL_IGCA_TARGET:
+// Every device's IGCA target and feature sets should have a corresponding
+// entry in INTEL_IGCA_TARGET:
#define INTEL_IGCA_TARGET_CHECK(NAME, IGCA_TARGET, IGCA_FEATURE_SETS) \
- EXPECT_NE(IntelGPU::getIGCATargetName(IntelGPU::IGCATarget( \
- IGCA_TARGET, IntelGPU::IGCAFeatureSet::IGCA_FEATURE_SETS)), \
- "") \
- << NAME;
+ static_assert(IntelGPU::IGCATarget( \
+ IGCA_TARGET, IntelGPU::IGCAFeatureSet::IGCA_FEATURE_SETS) \
+ .isValid(), \
+ NAME " uses an IGCA target not listed in INTEL_IGCA_TARGET");
#define INTEL_GPU(NAME, KIND, MAJOR, MINOR, IGCA_TARGET, IGCA_FEATURE_SETS) \
INTEL_IGCA_TARGET_CHECK(NAME, IGCA_TARGET, IGCA_FEATURE_SETS)
#define INTEL_GPU_COMPAT(NAME, KIND, IGCA_TARGET, IGCA_FEATURE_SETS) \
INTEL_IGCA_TARGET_CHECK(NAME, IGCA_TARGET, IGCA_FEATURE_SETS)
#include "llvm/TargetParser/IntelGPUTargetParser.def"
#undef INTEL_IGCA_TARGET_CHECK
-}
} // namespace
>From dcaf8d2144ed30e2d2837396da5a710c2ed3eb4b Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 22 Oct 2026 18:06:21 -0700
Subject: [PATCH 12/30] address reviewer comments
---
clang/lib/Basic/OffloadArch.cpp | 3 ++
.../llvm/TargetParser/IntelGPUTargetParser.h | 38 ++++++++++++-------
.../TargetParser/IntelGPUTargetParserTest.cpp | 2 +
3 files changed, 30 insertions(+), 13 deletions(-)
diff --git a/clang/lib/Basic/OffloadArch.cpp b/clang/lib/Basic/OffloadArch.cpp
index 00ba58cc8550b0..91094c7c6ae811 100644
--- a/clang/lib/Basic/OffloadArch.cpp
+++ b/clang/lib/Basic/OffloadArch.cpp
@@ -13,6 +13,7 @@
#include "llvm/TargetParser/IntelGPUTargetParser.h"
#include "llvm/TargetParser/NVPTXTargetParser.h"
#include "llvm/TargetParser/Triple.h"
+#include <cassert>
namespace clang {
@@ -25,10 +26,12 @@ OffloadArch OffloadArch::HIPDefault() {
}
OffloadArch OffloadArch::getIntelIGCA(llvm::IntelGPU::IGCATarget T) {
+ assert(T.isValid() && "T is not a valid IGCA target");
return {TargetArch::IntelIGCA, T.pack()};
}
llvm::IntelGPU::IGCATarget OffloadArch::igcaTarget() const {
+ assert(isIntelIGCA() && "cannot get IGCA target from a non-IGCA OffloadArch");
return llvm::IntelGPU::IGCATarget::unpack(Kind);
}
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
index f4d44407d0daf7..a1026d2da195d6 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
@@ -64,7 +64,7 @@ class IGCATarget {
uint32_t V = 0;
/// \return true if \p Packed contains a target listed in INTEL_IGCA_TARGET.
- static constexpr bool isKnownIGCATargetLevel(uint32_t Packed) {
+ static constexpr bool isKnownIGCATarget(uint32_t Packed) {
switch (Packed) {
#define INTEL_IGCA_TARGET(NAME, TARGET, FEATURE_SET) \
case uint32_t(TARGET) << TargetShift | \
@@ -80,7 +80,7 @@ class IGCATarget {
// Targets not listed in INTEL_IGCA_TARGET (including Target == 0) are
// invalid, so V remains 0. This keeps V == 0 as the only canonical invalid
// representation.
- if (isKnownIGCATargetLevel(Packed))
+ if (isKnownIGCATarget(Packed))
V = Packed;
}
@@ -107,37 +107,49 @@ class IGCATarget {
/// To check if a target has a feature set regardless of exactness, use
/// hasCompute() or hasRender() instead of raw comparison between
/// IGCAFeatureSet enums.
- IGCAFeatureSet getFeatureSet() const {
+ constexpr IGCAFeatureSet getFeatureSet() const {
return IGCAFeatureSet(V & FeatureSetMask);
}
- bool isExact() const { return V & IsExactMask; }
+ constexpr bool isExact() const { return V & IsExactMask; }
constexpr bool isValid() const { return getTarget() != 0; }
constexpr explicit operator bool() const { return isValid(); }
- bool isCore() const { return getFeatureSet() == IGCAFeatureSet::Core; }
+ /// \return true if the target has the Core feature set.
+ constexpr bool isCore() const {
+ // invalid() has the same feature set bits as Core, so we check explicitly.
+ return isValid() && getFeatureSet() == IGCAFeatureSet::Core;
+ }
/// \return true if the target has a non-exact Compute feature set.
- bool isCompute() const { return getFeatureSet() == IGCAFeatureSet::Compute; }
+ constexpr bool isCompute() const {
+ return getFeatureSet() == IGCAFeatureSet::Compute;
+ }
/// \return true if the target has a non-exact Render feature set.
- bool isRender() const { return getFeatureSet() == IGCAFeatureSet::Render; }
- bool isComputeExact() const {
+ constexpr bool isRender() const {
+ return getFeatureSet() == IGCAFeatureSet::Render;
+ }
+ constexpr bool isComputeExact() const {
return getFeatureSet() == IGCAFeatureSet::ComputeExact;
}
- bool isRenderExact() const {
+ constexpr bool isRenderExact() const {
return getFeatureSet() == IGCAFeatureSet::RenderExact;
}
/// \return true if target has either Compute or ComputeExact feature set.
- bool hasCompute() const {
+ constexpr bool hasCompute() const {
return (llvm::to_underlying(getFeatureSet()) & ~IsExactMask) ==
llvm::to_underlying(IGCAFeatureSet::Compute);
}
/// \return true if target has either Render or RenderExact feature set.
- bool hasRender() const {
+ constexpr bool hasRender() const {
return (llvm::to_underlying(getFeatureSet()) & ~IsExactMask) ==
llvm::to_underlying(IGCAFeatureSet::Render);
}
- friend bool operator==(IGCATarget A, IGCATarget B) { return A.V == B.V; }
- friend bool operator!=(IGCATarget A, IGCATarget B) { return !(A == B); }
+ friend constexpr bool operator==(IGCATarget A, IGCATarget B) {
+ return A.V == B.V;
+ }
+ friend constexpr bool operator!=(IGCATarget A, IGCATarget B) {
+ return !(A == B);
+ }
};
/// \return the device name that \p GPUIPVersion identifies, as reported by the
diff --git a/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp b/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
index 8c288d1cbf5704..9361c6cbfddece 100644
--- a/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
+++ b/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
@@ -109,6 +109,8 @@ TEST(IntelGPUTargetParserTest, IGCATargetBehavior) {
EXPECT_TRUE(Core.isCore());
// Invalid behavior:
EXPECT_FALSE(IGCATarget::invalid());
+ // invalid() has no feature set, even though its bits match Core:
+ EXPECT_FALSE(IGCATarget::invalid().isCore());
EXPECT_EQ(IGCATarget::invalid().pack(), 0u);
EXPECT_EQ(IGCATarget(0, IGCAFeatureSet::Compute), IGCATarget::invalid());
// unpack(2) is invalid since it results in Target = 0. It should thus equate
>From bb17463a132963d37e8d913ce3be7860373351a5 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 22 Oct 2026 18:09:17 -0700
Subject: [PATCH 13/30] remove fillValidIGCATargetList
---
llvm/include/llvm/TargetParser/IntelGPUTargetParser.h | 3 ---
llvm/lib/TargetParser/IntelGPUTargetParser.cpp | 5 -----
.../TargetParser/IntelGPUTargetParserTest.cpp | 11 -----------
3 files changed, 19 deletions(-)
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
index a1026d2da195d6..64d42f100ec947 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
@@ -172,9 +172,6 @@ LLVM_ABI IGCATarget parseIGCATarget(StringRef TargetStr);
/// \p Target is invalid or not a known target in IntelGPUTargetParser.def.
LLVM_ABI StringRef getIGCATargetName(IGCATarget Target);
-/// Append every valid IGCA target spelling to \p Values.
-LLVM_ABI void fillValidIGCATargetList(SmallVectorImpl<StringRef> &Values);
-
} // namespace IntelGPU
} // namespace llvm
diff --git a/llvm/lib/TargetParser/IntelGPUTargetParser.cpp b/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
index c6c3fe2a9acab1..2f2fb032993afb 100644
--- a/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
+++ b/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
@@ -87,8 +87,3 @@ StringRef llvm::IntelGPU::getIGCATargetName(IGCATarget T) {
// FIXME(IGCA): Make -fsycl --offload-arch provides a list of valid IGCA
// targets upon a bad IGCA target, similar to how compiling for an nvptx triple
// returns a list of valid GPU architectures.
-void llvm::IntelGPU::fillValidIGCATargetList(
- SmallVectorImpl<StringRef> &Values) {
-#define INTEL_IGCA_TARGET(NAME, TARGET, FEATURE_SET) Values.push_back(NAME);
-#include "llvm/TargetParser/IntelGPUTargetParser.def"
-}
diff --git a/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp b/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
index 9361c6cbfddece..119b10352224db 100644
--- a/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
+++ b/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
@@ -194,17 +194,6 @@ TEST(IntelGPUTargetParserTest, IGCATargetName) {
IntelGPU::getIGCATargetName(IGCATarget(10, IGCAFeatureSet::Compute)), "");
}
-TEST(IntelGPUTargetParserTest, EveryIGCASpellingRoundTrips) {
- SmallVector<StringRef> Names;
- IntelGPU::fillValidIGCATargetList(Names);
- EXPECT_FALSE(Names.empty());
- for (StringRef Name : Names) {
- IntelGPU::IGCATarget T = IntelGPU::parseIGCATarget(Name);
- EXPECT_TRUE(T.isValid()) << Name;
- EXPECT_EQ(IntelGPU::getIGCATargetName(T), Name);
- }
-}
-
// Spell an IGCA target from its fields, independently of the NAME column.
std::string expectedIGCASpelling(uint16_t Target,
IntelGPU::IGCAFeatureSet FeatureSet) {
>From 5201fa0fb53f67574ac4830a3911c6d5151fd88a Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 22 Oct 2026 21:22:19 -0700
Subject: [PATCH 14/30] address reviewer comments
---
clang/include/clang/Basic/OffloadArch.h | 2 +-
clang/test/Driver/sycl-offload-arch-igca.cpp | 18 +++++++++++++-----
.../llvm/TargetParser/IntelGPUTargetParser.def | 2 +-
.../llvm/TargetParser/IntelGPUTargetParser.h | 2 +-
4 files changed, 16 insertions(+), 8 deletions(-)
diff --git a/clang/include/clang/Basic/OffloadArch.h b/clang/include/clang/Basic/OffloadArch.h
index 437f85f99843f4..af35004f06d6e4 100644
--- a/clang/include/clang/Basic/OffloadArch.h
+++ b/clang/include/clang/Basic/OffloadArch.h
@@ -90,11 +90,11 @@ class OffloadArch {
bool isAMDGCNSPIRV() const { return V == TargetArch::AMDGCNSPIRV; }
bool isIntelCPU() const { return V == TargetArch::IntelCPU; }
bool isIntelGPU() const { return V == TargetArch::IntelGPU; }
+ bool isIntelIGCA() const { return V == TargetArch::IntelIGCA; }
bool isIntel() const { return isIntelCPU() || isIntelGPU() || isIntelIGCA(); }
bool isGeneric() const { return V == TargetArch::Generic; }
bool isUnused() const { return V == TargetArch::Unused; }
bool isUnknown() const { return V == TargetArch::Unknown; }
- bool isIntelIGCA() const { return V == TargetArch::IntelIGCA; }
// Only valid when isNVPTX() / isAMDGPU() / isIntelIGCA() respectively.
llvm::NVPTX::GPUKind nvptxKind() const {
diff --git a/clang/test/Driver/sycl-offload-arch-igca.cpp b/clang/test/Driver/sycl-offload-arch-igca.cpp
index 5e2c45f180a913..714aa4edf0136e 100644
--- a/clang/test/Driver/sycl-offload-arch-igca.cpp
+++ b/clang/test/Driver/sycl-offload-arch-igca.cpp
@@ -12,8 +12,16 @@
// CHK-INVALID: error: unsupported IGCA gpu architecture: igca_42a
// CHK-INVALID-NOT: llvm-offload-binary
-/// FIXME(IGCA): Enable this when OffloadArchToTriple gives IGCA targets a triple.
-// TODO-RUN %clang -### --target=x86_64-unknown-linux-gnu -fsycl \
-// TODO-RUN --offload-arch=igca_20ca -c %s 2>&1 | FileCheck -check-prefix=CHK-INFER %s
-// CHK-INFER-NOT: error:
-// CHK-INFER: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spirv64{{[^,]*}},arch=igca_20ca,kind=sycl"
+/// A valid IGCA target is still rejected on NVPTX and AMDGPU device triples,
+/// with the existing CUDA/HIP diagnostic:
+// RUN: not %clang -### --target=x86_64-unknown-linux-gnu -fsycl \
+// RUN: --offload-targets=nvptx64-nvidia-cuda --offload-arch=igca_20ca \
+// RUN: -c %s 2>&1 | FileCheck -check-prefix=CHK-NVPTX %s
+// CHK-NVPTX: error: unsupported CUDA gpu architecture: igca_20ca
+// CHK-NVPTX-NOT: llvm-offload-binary
+
+// RUN: not %clang -### --target=x86_64-unknown-linux-gnu -fsycl \
+// RUN: --offload-targets=amdgcn-amd-amdhsa --offload-arch=igca_20ca \
+// RUN: -c %s 2>&1 | FileCheck -check-prefix=CHK-AMDGPU %s
+// CHK-AMDGPU: error: unsupported HIP gpu architecture: igca_20ca
+// CHK-AMDGPU-NOT: llvm-offload-binary
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def
index de276875351bec..d29a70a9ce737c 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def
@@ -48,7 +48,7 @@
// NAME - Name of the IGCA target, e.g. "igca_15ra". Name must
// agree with the other columns.
// TARGET - Numeric IGCA target level, e.g. 60.
-// FEATURE_SET - Feature set of the current target, i.e. "Compute" for
+// FEATURE_SET - Feature set of the current target, e.g. "Compute" for
// "c" suffix, "Render" for "r", "ComputeExact" for
// "ca", "Core" for "", etc.
//
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
index 64d42f100ec947..1e1c67e9a8534a 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
@@ -166,7 +166,7 @@ LLVM_ABI std::string getNumericArchName(uint32_t GPUIPVersion);
///
/// \return an invalid IGCATarget if \p TargetStr is not a known target in
/// IntelGPUTargetParser.def.
-LLVM_ABI IGCATarget parseIGCATarget(StringRef TargetStr);
+LLVM_ABI IGCATarget parseIGCATarget(StringRef MaybeTarget);
/// \return the \p Target as a string, i.e. "igca_20ca", or an empty string if
/// \p Target is invalid or not a known target in IntelGPUTargetParser.def.
>From b9d06d3930b1534e97768a1c7e3db5c4b0be7e9e Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 22 Oct 2026 22:14:00 -0700
Subject: [PATCH 15/30] address more reviewer comments
---
clang/lib/Driver/Driver.cpp | 7 +++++--
clang/test/Driver/sycl-offload-arch-igca.cpp | 8 ++++++++
2 files changed, 13 insertions(+), 2 deletions(-)
diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp
index 741ddefad97e49..14c4fa99be1b82 100644
--- a/clang/lib/Driver/Driver.cpp
+++ b/clang/lib/Driver/Driver.cpp
@@ -3928,8 +3928,11 @@ static StringRef getCanonicalArchString(Compilation &C,
return StringRef();
}
}
- } else if (Triple.isSPIRV() && ArchStr.starts_with("igca_") &&
- !Arch.isIntelIGCA()) {
+ } else if (Triple.isSPIRV() && ArchStr.starts_with("igca") &&
+ !StringToOffloadArch(ArchStr).isIntelIGCA()) {
+ // FIXME(IGCA): AMD targets should reject IGCA targets altogether.
+ // FIXME(IGCA): This only checks ArchStr's starting with "igca". Typos such
+ // as IGCA_20ca still fall through this check.
C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
<< "IGCA" << ArchStr;
return StringRef();
diff --git a/clang/test/Driver/sycl-offload-arch-igca.cpp b/clang/test/Driver/sycl-offload-arch-igca.cpp
index 714aa4edf0136e..215fa9023a9790 100644
--- a/clang/test/Driver/sycl-offload-arch-igca.cpp
+++ b/clang/test/Driver/sycl-offload-arch-igca.cpp
@@ -12,6 +12,14 @@
// CHK-INVALID: error: unsupported IGCA gpu architecture: igca_42a
// CHK-INVALID-NOT: llvm-offload-binary
+/// IGCA targets should not accept feature suffixes, so a valid target followed
+/// by a feature suffix should be rejected:
+// RUN: not %clang -### --target=x86_64-unknown-linux-gnu -fsycl \
+// RUN: --offload-targets=spirv64-unknown-unknown --offload-arch=igca_20ca:foo \
+// RUN: -c %s 2>&1 | FileCheck -check-prefix=CHK-FEATURE %s
+// CHK-FEATURE: error: unsupported IGCA gpu architecture: igca_20ca:foo
+// CHK-FEATURE-NOT: llvm-offload-binary
+
/// A valid IGCA target is still rejected on NVPTX and AMDGPU device triples,
/// with the existing CUDA/HIP diagnostic:
// RUN: not %clang -### --target=x86_64-unknown-linux-gnu -fsycl \
>From f7b51ec6c1ae1727b2851f0b9d6ae6da8ee864ea Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 22 Oct 2026 22:20:20 -0700
Subject: [PATCH 16/30] clang-format
---
clang/lib/Driver/Driver.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp
index 14c4fa99be1b82..e79c7a366aaf63 100644
--- a/clang/lib/Driver/Driver.cpp
+++ b/clang/lib/Driver/Driver.cpp
@@ -3932,7 +3932,7 @@ static StringRef getCanonicalArchString(Compilation &C,
!StringToOffloadArch(ArchStr).isIntelIGCA()) {
// FIXME(IGCA): AMD targets should reject IGCA targets altogether.
// FIXME(IGCA): This only checks ArchStr's starting with "igca". Typos such
- // as IGCA_20ca still fall through this check.
+ // as IGCA_20ca still fall through this check.
C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
<< "IGCA" << ArchStr;
return StringRef();
>From 914bae6f0a29e448755843aba706da47618733c8 Mon Sep 17 00:00:00 2001
From: Ian Li <ianayl.work at gmail.com>
Date: Fri, 23 Oct 2026 13:36:03 -0400
Subject: [PATCH 17/30] Update
llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
Co-authored-by: Tom Honermann <tom at honermann.net>
---
llvm/include/llvm/TargetParser/IntelGPUTargetParser.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
index 1e1c67e9a8534a..7df44f5e4031a3 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
@@ -164,7 +164,7 @@ LLVM_ABI std::string getNumericArchName(uint32_t GPUIPVersion);
/// Parse an IGCA target string such as "igca_20ca".
///
-/// \return an invalid IGCATarget if \p TargetStr is not a known target in
+/// \return an invalid IGCATarget if \p MaybeTarget is not a known target in
/// IntelGPUTargetParser.def.
LLVM_ABI IGCATarget parseIGCATarget(StringRef MaybeTarget);
>From a10cd15c17174c7336fbd6aae843bbdb1df2a0d9 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Fri, 23 Oct 2026 13:20:28 -0700
Subject: [PATCH 18/30] Don't accept IGCA for amdhsa, remove comments
---
clang/lib/Basic/OffloadArch.cpp | 2 --
clang/lib/Driver/Driver.cpp | 17 +++++++++++------
clang/test/Driver/sycl-offload-arch-igca.cpp | 6 ++++++
.../tools/clang-sycl-linker/ClangSYCLLinker.cpp | 1 -
4 files changed, 17 insertions(+), 9 deletions(-)
diff --git a/clang/lib/Basic/OffloadArch.cpp b/clang/lib/Basic/OffloadArch.cpp
index 91094c7c6ae811..8428fa34998f97 100644
--- a/clang/lib/Basic/OffloadArch.cpp
+++ b/clang/lib/Basic/OffloadArch.cpp
@@ -137,8 +137,6 @@ llvm::Triple OffloadArchToTriple(const llvm::Triple &DefaultToolchainTriple,
return llvm::Triple(llvm::Triple::amdgpu, llvm::Triple::NoSubArch,
llvm::Triple::AMD, llvm::Triple::AMDHSA);
- // FIXME(IGCA): Handle IntelIGCA, which currently gets an empty triple, so
- // --offload-arch=igca_* cannot pick an offload toolchain.
return {};
}
diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp
index e79c7a366aaf63..2e24328aedd1b5 100644
--- a/clang/lib/Driver/Driver.cpp
+++ b/clang/lib/Driver/Driver.cpp
@@ -3928,14 +3928,19 @@ static StringRef getCanonicalArchString(Compilation &C,
return StringRef();
}
}
- } else if (Triple.isSPIRV() && ArchStr.starts_with("igca") &&
- !StringToOffloadArch(ArchStr).isIntelIGCA()) {
- // FIXME(IGCA): AMD targets should reject IGCA targets altogether.
+ } else if (Triple.isSPIRV() && ArchStr.starts_with("igca")) {
+ if (Triple.getOS() == llvm::Triple::AMDHSA) {
+ C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
+ << "HIP" << ArchStr;
+ return StringRef();
+ }
// FIXME(IGCA): This only checks ArchStr's starting with "igca". Typos such
// as IGCA_20ca still fall through this check.
- C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
- << "IGCA" << ArchStr;
- return StringRef();
+ if (!StringToOffloadArch(ArchStr).isIntelIGCA()) {
+ C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
+ << "IGCA" << ArchStr;
+ return StringRef();
+ }
}
if (Arch.isNVPTX())
diff --git a/clang/test/Driver/sycl-offload-arch-igca.cpp b/clang/test/Driver/sycl-offload-arch-igca.cpp
index 215fa9023a9790..63c8922a97ac68 100644
--- a/clang/test/Driver/sycl-offload-arch-igca.cpp
+++ b/clang/test/Driver/sycl-offload-arch-igca.cpp
@@ -31,5 +31,11 @@
// RUN: not %clang -### --target=x86_64-unknown-linux-gnu -fsycl \
// RUN: --offload-targets=amdgcn-amd-amdhsa --offload-arch=igca_20ca \
// RUN: -c %s 2>&1 | FileCheck -check-prefix=CHK-AMDGPU %s
+// RUN: not %clang -### --target=x86_64-unknown-linux-gnu -fsycl \
+// RUN: --offload-targets=spirv64-amd-amdhsa --offload-arch=igca_20ca \
+// RUN: -c %s 2>&1 | FileCheck -check-prefix=CHK-AMDGPU %s
+// RUN: not %clang -### --target=x86_64-unknown-linux-gnu -fsycl \
+// RUN: --offload-targets=spirv64-unknown-amdhsa --offload-arch=igca_20ca \
+// RUN: -c %s 2>&1 | FileCheck -check-prefix=CHK-AMDGPU %s
// CHK-AMDGPU: error: unsupported HIP gpu architecture: igca_20ca
// CHK-AMDGPU-NOT: llvm-offload-binary
diff --git a/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp b/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp
index 690d6ee07a1026..8b179f4c37e23b 100644
--- a/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp
+++ b/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp
@@ -748,7 +748,6 @@ static Error runAOTCompile(StringRef InputFile, StringRef OutputFile,
return runAOTCompileIntelGPU(InputFile, OutputFile, Args);
if (OA.isIntelCPU())
return runAOTCompileIntelCPU(InputFile, OutputFile, Args);
- // FIXME(IGCA): Implement AOT compilation for IGCA targets.
if (OA.isIntelIGCA())
return createStringError("AOT compilation for IGCA target '" + Arch +
"' is not yet implemented");
>From f0b83c665e6a53dbe3a868af624440509c1fdb60 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Fri, 23 Oct 2026 13:42:41 -0700
Subject: [PATCH 19/30] remove hasCompute/isCompute and create isAnyCompute
---
.../llvm/TargetParser/IntelGPUTargetParser.h | 14 +++-----------
.../TargetParser/IntelGPUTargetParserTest.cpp | 12 ++++--------
2 files changed, 7 insertions(+), 19 deletions(-)
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
index 7df44f5e4031a3..197111ddd09f12 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
@@ -105,7 +105,7 @@ class IGCATarget {
/// \return the feature set, including whether it is exact.
///
/// To check if a target has a feature set regardless of exactness, use
- /// hasCompute() or hasRender() instead of raw comparison between
+ /// isAnyCompute() or isAnyRender() instead of raw comparison between
/// IGCAFeatureSet enums.
constexpr IGCAFeatureSet getFeatureSet() const {
return IGCAFeatureSet(V & FeatureSetMask);
@@ -119,14 +119,6 @@ class IGCATarget {
// invalid() has the same feature set bits as Core, so we check explicitly.
return isValid() && getFeatureSet() == IGCAFeatureSet::Core;
}
- /// \return true if the target has a non-exact Compute feature set.
- constexpr bool isCompute() const {
- return getFeatureSet() == IGCAFeatureSet::Compute;
- }
- /// \return true if the target has a non-exact Render feature set.
- constexpr bool isRender() const {
- return getFeatureSet() == IGCAFeatureSet::Render;
- }
constexpr bool isComputeExact() const {
return getFeatureSet() == IGCAFeatureSet::ComputeExact;
}
@@ -134,12 +126,12 @@ class IGCATarget {
return getFeatureSet() == IGCAFeatureSet::RenderExact;
}
/// \return true if target has either Compute or ComputeExact feature set.
- constexpr bool hasCompute() const {
+ constexpr bool isAnyCompute() const {
return (llvm::to_underlying(getFeatureSet()) & ~IsExactMask) ==
llvm::to_underlying(IGCAFeatureSet::Compute);
}
/// \return true if target has either Render or RenderExact feature set.
- constexpr bool hasRender() const {
+ constexpr bool isAnyRender() const {
return (llvm::to_underlying(getFeatureSet()) & ~IsExactMask) ==
llvm::to_underlying(IGCAFeatureSet::Render);
}
diff --git a/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp b/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
index 119b10352224db..207207095a9f5a 100644
--- a/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
+++ b/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
@@ -87,23 +87,19 @@ TEST(IntelGPUTargetParserTest, IGCATargetBehavior) {
EXPECT_EQ(C.getFeatureSet(), IGCAFeatureSet::ComputeExact);
EXPECT_TRUE(C.isExact());
// An exact feature set is still of its class:
- EXPECT_TRUE(C.hasCompute());
+ EXPECT_TRUE(C.isAnyCompute());
EXPECT_TRUE(C.isComputeExact());
EXPECT_FALSE(C.isCore());
- EXPECT_FALSE(C.isCompute());
- EXPECT_FALSE(C.isRender());
- EXPECT_FALSE(C.hasRender());
+ EXPECT_FALSE(C.isAnyRender());
EXPECT_FALSE(C.isRenderExact());
EXPECT_EQ(IGCATarget::unpack(C.pack()), C);
EXPECT_NE(C, IGCATarget(20, IGCAFeatureSet::Compute));
// A feature set that is not exact:
IGCATarget R{60, IGCAFeatureSet::Render};
- EXPECT_TRUE(R.isRender());
EXPECT_FALSE(R.isRenderExact());
EXPECT_FALSE(R.isExact());
- EXPECT_FALSE(R.isCompute());
- EXPECT_TRUE(R.hasRender());
- EXPECT_FALSE(R.hasCompute());
+ EXPECT_TRUE(R.isAnyRender());
+ EXPECT_FALSE(R.isAnyCompute());
EXPECT_EQ(IGCATarget::unpack(R.pack()), R);
IGCATarget Core{60, IGCAFeatureSet::Core};
EXPECT_TRUE(Core.isCore());
>From 6fc15d3db330196552ecb5c50ed295adb446c7e5 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Fri, 23 Oct 2026 17:05:07 -0700
Subject: [PATCH 20/30] redo logic to derive from the table
---
.../TargetParser/IntelGPUTargetParser.def | 106 ++++++++++--------
.../llvm/TargetParser/IntelGPUTargetParser.h | 11 +-
.../lib/TargetParser/IntelGPUTargetParser.cpp | 11 +-
3 files changed, 66 insertions(+), 62 deletions(-)
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def
index d29a70a9ce737c..a17f39ccfc3ed8 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.def
@@ -17,8 +17,7 @@
// GK_<KIND>.
// MAJOR - Major component of the device's GPU IP version.
// MINOR - Minor component of the device's GPU IP version.
-// IGCA_TARGET - Numeric IGCA target: Every target should also be
-// listed in the INTEL_IGCA_TARGET table.
+// IGCA_TARGET - Numeric IGCA target, e.g. 60.
// IGCA_FEATURE_SETS - Token naming the IGCA feature sets the device
// implements: Core (none), Compute ("c"),
// Render ("r"), ComputeExact ("ca") or
@@ -38,22 +37,65 @@
// Several devices can share a major and a minor version. The rows are ordered
// so that the name to print for such a group comes first.
//
-// This file also includes a list of valid IGCA targets, which can be used to
-// compile binaries with support for different GPU architectures and feature
-// sets:
+// This file also derives an INTEL_IGCA_TARGET table, which lists the valid
+// IGCA targets derived from the INTEL_GPU and INTEL_GPU_COMPAT rows. A row with
+// target N and feature set FS implies:
+// - igca_N, Core feature set, which every device implements;
+// - igca_N<c|r>, non-exact form of FS, unless FS is Core;
+// - igca_N<ca|ra>, exact form of FS, if FS is exact.
+// E.g. (15, RenderExact) implies igca_15, igca_15r and igca_15ra.
+// Targets come out in row order, once per row that implies them, so the table
+// is unsorted and contains duplicates.
//
// INTEL_IGCA_TARGET(NAME, TARGET, FEATURE_SET)
-// One row per valid IGCA target spelling. Only the combinations listed
-// are valid.
-// NAME - Name of the IGCA target, e.g. "igca_15ra". Name must
-// agree with the other columns.
-// TARGET - Numeric IGCA target level, e.g. 60.
-// FEATURE_SET - Feature set of the current target, e.g. "Compute" for
-// "c" suffix, "Render" for "r", "ComputeExact" for
-// "ca", "Core" for "", etc.
+// NAME - Readable string representing the target, e.g. "igca_15ra".
+// TARGET - Numeric IGCA target, e.g. 15.
+// FEATURE_SET - IGCAFeatureSet enumerator, e.g. RenderExact.
//
//===----------------------------------------------------------------------===//
+// Derive INTEL_IGCA_TARGET table from the rows in INTEL_GPU and
+// INTEL_GPU_COMPAT below; see the top of this file.
+//
+// INTEL_IGCA_TARGET(NAME, TARGET, FEATURE_SET)
+//
+// Note that this derivation produces UNSORTED, DUPLICATE entries. Sorting and
+// deduping needs to be done in code.
+#ifdef INTEL_IGCA_TARGET
+
+#if defined(INTEL_GPU) || defined(INTEL_GPU_COMPAT)
+// This block defines INTEL_GPU and INTEL_GPU_COMPAT macros itself, which would
+// overwrite an includer's own definitions of e.g. INTEL_GPU below. To use both
+// INTEL_GPU and INTEL_IGCA_TARGET tables, include this file once for each use.
+#error "INTEL_IGCA_TARGET cannot be used together with INTEL_GPU(_COMPAT)"
+#endif
+
+#define INTEL_IGCA_IMPLIED_Core(TARGET) \
+ INTEL_IGCA_TARGET("igca_" #TARGET , TARGET, Core)
+
+#define INTEL_IGCA_IMPLIED_Compute(TARGET) \
+ INTEL_IGCA_IMPLIED_Core(TARGET) \
+ INTEL_IGCA_TARGET("igca_" #TARGET "c" , TARGET, Compute)
+
+#define INTEL_IGCA_IMPLIED_ComputeExact(TARGET) \
+ INTEL_IGCA_IMPLIED_Compute(TARGET) \
+ INTEL_IGCA_TARGET("igca_" #TARGET "ca", TARGET, ComputeExact)
+
+#define INTEL_IGCA_IMPLIED_Render(TARGET) \
+ INTEL_IGCA_IMPLIED_Core(TARGET) \
+ INTEL_IGCA_TARGET("igca_" #TARGET "r" , TARGET, Render)
+
+#define INTEL_IGCA_IMPLIED_RenderExact(TARGET) \
+ INTEL_IGCA_IMPLIED_Render(TARGET) \
+ INTEL_IGCA_TARGET("igca_" #TARGET "ra", TARGET, RenderExact)
+
+#define INTEL_GPU(NAME, KIND, MAJOR, MINOR, IGCA_TARGET, IGCA_FEATURE_SETS) \
+ INTEL_IGCA_IMPLIED_##IGCA_FEATURE_SETS(IGCA_TARGET)
+#define INTEL_GPU_COMPAT(NAME, KIND, IGCA_TARGET, IGCA_FEATURE_SETS) \
+ INTEL_IGCA_IMPLIED_##IGCA_FEATURE_SETS(IGCA_TARGET)
+#endif
+
+// Intel GPU definitions
#ifndef INTEL_GPU
#define INTEL_GPU(NAME, KIND, MAJOR, MINOR, IGCA_TARGET, IGCA_FEATURE_SETS)
#endif
@@ -107,39 +149,11 @@ INTEL_GPU_COMPAT("xe-bmg", XE_BMG, 40, Render)
INTEL_GPU_COMPAT("xe-mtl", XE_MTL, 30, Render)
INTEL_GPU_COMPAT("xe-dg2", XE_DG2, 15, RenderExact)
-// IGCA targets. Only targets listed here are valid.
-#ifndef INTEL_IGCA_TARGET
-#define INTEL_IGCA_TARGET(NAME, TARGET, FEATURE_SET)
-#endif
-
-INTEL_IGCA_TARGET("igca_10", 10, Core)
-INTEL_IGCA_TARGET("igca_10r", 10, Render)
-
-INTEL_IGCA_TARGET("igca_15", 15, Core)
-INTEL_IGCA_TARGET("igca_15r", 15, Render)
-INTEL_IGCA_TARGET("igca_15ra", 15, RenderExact)
-
-INTEL_IGCA_TARGET("igca_20", 20, Core)
-INTEL_IGCA_TARGET("igca_20c", 20, Compute)
-INTEL_IGCA_TARGET("igca_20ca", 20, ComputeExact)
-
-INTEL_IGCA_TARGET("igca_30", 30, Core)
-INTEL_IGCA_TARGET("igca_30r", 30, Render)
-
-INTEL_IGCA_TARGET("igca_35", 35, Core)
-INTEL_IGCA_TARGET("igca_35r", 35, Render)
-INTEL_IGCA_TARGET("igca_35ra", 35, RenderExact)
-
-INTEL_IGCA_TARGET("igca_40", 40, Core)
-INTEL_IGCA_TARGET("igca_40r", 40, Render)
-
-INTEL_IGCA_TARGET("igca_50", 50, Core)
-INTEL_IGCA_TARGET("igca_50r", 50, Render)
-
-INTEL_IGCA_TARGET("igca_60", 60, Core)
-INTEL_IGCA_TARGET("igca_60c", 60, Compute)
-INTEL_IGCA_TARGET("igca_60r", 60, Render)
-
#undef INTEL_GPU
#undef INTEL_GPU_COMPAT
#undef INTEL_IGCA_TARGET
+#undef INTEL_IGCA_IMPLIED_Core
+#undef INTEL_IGCA_IMPLIED_Compute
+#undef INTEL_IGCA_IMPLIED_ComputeExact
+#undef INTEL_IGCA_IMPLIED_Render
+#undef INTEL_IGCA_IMPLIED_RenderExact
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
index 197111ddd09f12..035abb848e8fd5 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
@@ -65,15 +65,12 @@ class IGCATarget {
/// \return true if \p Packed contains a target listed in INTEL_IGCA_TARGET.
static constexpr bool isKnownIGCATarget(uint32_t Packed) {
- switch (Packed) {
#define INTEL_IGCA_TARGET(NAME, TARGET, FEATURE_SET) \
- case uint32_t(TARGET) << TargetShift | \
- llvm::to_underlying(IGCAFeatureSet::FEATURE_SET):
+ if (Packed == (uint32_t(TARGET) << TargetShift | \
+ llvm::to_underlying(IGCAFeatureSet::FEATURE_SET))) \
+ return true;
#include "llvm/TargetParser/IntelGPUTargetParser.def"
- return true;
- default:
- return false;
- }
+ return false;
}
constexpr explicit IGCATarget(uint32_t Packed) {
diff --git a/llvm/lib/TargetParser/IntelGPUTargetParser.cpp b/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
index 2f2fb032993afb..b9e1517170e615 100644
--- a/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
+++ b/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
@@ -74,16 +74,9 @@ IGCATarget llvm::IntelGPU::parseIGCATarget(StringRef MaybeTarget) {
}
StringRef llvm::IntelGPU::getIGCATargetName(IGCATarget T) {
- switch (T.pack()) {
#define INTEL_IGCA_TARGET(NAME, TARGET, FEATURE_SET) \
- case IGCATarget(TARGET, IGCAFeatureSet::FEATURE_SET).pack(): \
+ if (T == IGCATarget(TARGET, IGCAFeatureSet::FEATURE_SET)) \
return NAME;
#include "llvm/TargetParser/IntelGPUTargetParser.def"
- default:
- return "";
- }
+ return "";
}
-
-// FIXME(IGCA): Make -fsycl --offload-arch provides a list of valid IGCA
-// targets upon a bad IGCA target, similar to how compiling for an nvptx triple
-// returns a list of valid GPU architectures.
>From 2c98b89fcf259700145914584bd2d5a1ce161a49 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 17 Sep 2026 14:03:24 -0700
Subject: [PATCH 21/30] WIP: Introduce OffloadTargetInfo
---
clang/include/clang/Basic/OffloadArch.h | 1 +
clang/include/clang/Basic/TargetInfo.h | 7 ++++
clang/lib/Basic/Targets/OffloadTarget.h | 54 +++++++++++++++++++++++++
clang/lib/Basic/Targets/SPIR.h | 8 +++-
clang/lib/Frontend/InitPreprocessor.cpp | 3 ++
5 files changed, 71 insertions(+), 2 deletions(-)
create mode 100644 clang/lib/Basic/Targets/OffloadTarget.h
diff --git a/clang/include/clang/Basic/OffloadArch.h b/clang/include/clang/Basic/OffloadArch.h
index af35004f06d6e4..8aacd128f38031 100644
--- a/clang/include/clang/Basic/OffloadArch.h
+++ b/clang/include/clang/Basic/OffloadArch.h
@@ -95,6 +95,7 @@ class OffloadArch {
bool isGeneric() const { return V == TargetArch::Generic; }
bool isUnused() const { return V == TargetArch::Unused; }
bool isUnknown() const { return V == TargetArch::Unknown; }
+ bool isUnknownOrUnused() const { return isUnknown() || isUnused(); }
// Only valid when isNVPTX() / isAMDGPU() / isIntelIGCA() respectively.
llvm::NVPTX::GPUKind nvptxKind() const {
diff --git a/clang/include/clang/Basic/TargetInfo.h b/clang/include/clang/Basic/TargetInfo.h
index a6523e8eaaa437..d7816685f08c61 100644
--- a/clang/include/clang/Basic/TargetInfo.h
+++ b/clang/include/clang/Basic/TargetInfo.h
@@ -21,6 +21,7 @@
#include "clang/Basic/CodeGenOptions.h"
#include "clang/Basic/LLVM.h"
#include "clang/Basic/LangOptions.h"
+#include "clang/Basic/OffloadArch.h"
#include "clang/Basic/Specifiers.h"
#include "clang/Basic/TargetCXXABI.h"
#include "clang/Basic/TargetOptions.h"
@@ -1425,6 +1426,12 @@ class TargetInfo : public TransferrableTargetInfo,
return isValidCPUName(Name);
}
+ /// Returns an OffloadArch if targeting an offload device, otherwise return
+ /// Unused.
+ virtual OffloadArch getOffloadArch() const {
+ return OffloadArch::getUnused();
+ }
+
virtual ParsedTargetAttr parseTargetAttr(StringRef Str) const;
/// Determine whether this TargetInfo supports tune in target attribute.
diff --git a/clang/lib/Basic/Targets/OffloadTarget.h b/clang/lib/Basic/Targets/OffloadTarget.h
new file mode 100644
index 00000000000000..afcc2aa0db6e95
--- /dev/null
+++ b/clang/lib/Basic/Targets/OffloadTarget.h
@@ -0,0 +1,54 @@
+//===--- OffloadTarget.h - Declare generic offload target -------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file declares a generic "offload" TargetInfo object.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_CLANG_LIB_BASIC_TARGETS_OFFLOADTARGET_H
+#define LLVM_CLANG_LIB_BASIC_TARGETS_OFFLOADTARGET_H
+
+#include "clang/Basic/TargetInfo.h"
+#include "clang/Basic/TargetOptions.h"
+#include "clang/Basic/OffloadArch.h"
+#include "llvm/Support/Compiler.h"
+#include "llvm/TargetParser/Triple.h"
+
+namespace clang {
+namespace targets {
+
+class LLVM_LIBRARY_VISIBILITY OffloadTargetInfo : public TargetInfo {
+protected:
+ OffloadArch DeviceArch = OffloadArch::getUnused();
+
+public:
+ OffloadTargetInfo(const llvm::Triple &Triple) : TargetInfo(Triple) { }
+
+ // Clang driver emits -target-cpu to indicate offload device architecture for
+ // both CUDA and SYCL. We override setCPU to capture this offload device
+ // architecture information.
+ bool setCPU(StringRef Name) override {
+ DeviceArch = StringToOffloadArch(Name);
+ return !DeviceArch.isUnknownOrUnused();
+ }
+
+ bool isValidCPUName(StringRef Name) const override {
+ return !StringToOffloadArch(Name).isUnknown();
+ }
+
+ void fillValidCPUList(SmallVectorImpl<StringRef> &Values) const override {
+ fillValidOffloadArchList(Values);
+ }
+
+ OffloadArch getOffloadArch() const override { return DeviceArch; }
+};
+
+} // namespace targets
+} // namespace clang
+
+#endif // LLVM_CLANG_LIB_BASIC_TARGETS_OFFLOADTARGET_H
\ No newline at end of file
diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h
index 43a94cb9d31c89..bb2a61e1eb9f16 100644
--- a/clang/lib/Basic/Targets/SPIR.h
+++ b/clang/lib/Basic/Targets/SPIR.h
@@ -14,6 +14,7 @@
#define LLVM_CLANG_LIB_BASIC_TARGETS_SPIR_H
#include "Targets.h"
+#include "OffloadTarget.h"
#include "clang/Basic/AddressSpaces.h"
#include "clang/Basic/Diagnostic.h"
#include "clang/Basic/DiagnosticFrontend.h"
@@ -70,7 +71,7 @@ static constexpr LangASMap SPIRDefIsGenMap = {
};
// Base class for SPIR and SPIR-V target info.
-class LLVM_LIBRARY_VISIBILITY BaseSPIRTargetInfo : public TargetInfo {
+class LLVM_LIBRARY_VISIBILITY BaseSPIRTargetInfo : public OffloadTargetInfo {
std::unique_ptr<TargetInfo> HostTarget;
protected:
@@ -78,7 +79,7 @@ class LLVM_LIBRARY_VISIBILITY BaseSPIRTargetInfo : public TargetInfo {
const TargetInfo *getHostTarget() const { return HostTarget.get(); }
BaseSPIRTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
- : TargetInfo(Triple) {
+ : OffloadTargetInfo(Triple) {
assert((Triple.isSPIR() || Triple.isSPIRV()) &&
"Invalid architecture for SPIR or SPIR-V.");
TLSSupported = false;
@@ -520,6 +521,9 @@ class LLVM_LIBRARY_VISIBILITY SPIRV64AMDGCNTargetInfo final
// __builtin_amdgcn_processor_is
bool isValidCPUName(StringRef Name) const override;
void fillValidCPUList(SmallVectorImpl<StringRef> &Values) const override;
+ // AMD does not rely on -target-cpu to set/track CPU/GPU architecture; setCPU
+ // for AMD should not be used and thus unconditionally return false.
+ bool setCPU(StringRef Name) override { return false; }
};
class LLVM_LIBRARY_VISIBILITY SPIRV64IntelTargetInfo final
diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp
index 784ff9951a25ad..4cdb79c5f10370 100644
--- a/clang/lib/Frontend/InitPreprocessor.cpp
+++ b/clang/lib/Frontend/InitPreprocessor.cpp
@@ -1505,6 +1505,9 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
// SYCL device compiler which doesn't produce host binary.
if (LangOpts.SYCLIsDevice) {
Builder.defineMacro("__SYCL_DEVICE_ONLY__", "1");
+ if (OffloadArch Arch = TI.getOffloadArch(); !Arch.isUnknownOrUnused()) {
+ // TODO some sort of function to process -target-cpu
+ }
}
// OpenCL definitions.
>From 1aa97d6b902f7cfc484ebe46bebec0e0563ca19f Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Wed, 30 Sep 2026 13:04:21 -0700
Subject: [PATCH 22/30] Add macros indicating IGCA level
---
clang/include/clang/Basic/TargetInfo.h | 6 ++----
clang/lib/Basic/OffloadArch.cpp | 1 +
clang/lib/Basic/TargetInfo.cpp | 5 +++++
clang/lib/Frontend/InitPreprocessor.cpp | 16 ++++++++++++++++
4 files changed, 24 insertions(+), 4 deletions(-)
diff --git a/clang/include/clang/Basic/TargetInfo.h b/clang/include/clang/Basic/TargetInfo.h
index d7816685f08c61..6a46b4bebcc25b 100644
--- a/clang/include/clang/Basic/TargetInfo.h
+++ b/clang/include/clang/Basic/TargetInfo.h
@@ -21,7 +21,6 @@
#include "clang/Basic/CodeGenOptions.h"
#include "clang/Basic/LLVM.h"
#include "clang/Basic/LangOptions.h"
-#include "clang/Basic/OffloadArch.h"
#include "clang/Basic/Specifiers.h"
#include "clang/Basic/TargetCXXABI.h"
#include "clang/Basic/TargetOptions.h"
@@ -55,6 +54,7 @@ class DiagnosticsEngine;
class LangOptions;
class CodeGenOptions;
class MacroBuilder;
+class OffloadArch;
/// Contains information gathered from parsing the contents of TargetAttr.
struct ParsedTargetAttr {
@@ -1428,9 +1428,7 @@ class TargetInfo : public TransferrableTargetInfo,
/// Returns an OffloadArch if targeting an offload device, otherwise return
/// Unused.
- virtual OffloadArch getOffloadArch() const {
- return OffloadArch::getUnused();
- }
+ virtual OffloadArch getOffloadArch() const;
virtual ParsedTargetAttr parseTargetAttr(StringRef Str) const;
diff --git a/clang/lib/Basic/OffloadArch.cpp b/clang/lib/Basic/OffloadArch.cpp
index 8428fa34998f97..f5d9ec497ada39 100644
--- a/clang/lib/Basic/OffloadArch.cpp
+++ b/clang/lib/Basic/OffloadArch.cpp
@@ -101,6 +101,7 @@ OffloadArch StringToOffloadArch(llvm::StringRef S) {
return OffloadArch::getNVPTX(NV);
if (llvm::AMDGPU::GPUKind AK = llvm::AMDGPU::parseArchAMDGCN(S))
return OffloadArch::getAMDGPU(AK);
+ // TODO test that the operator bool causes this to fail if invalid
if (llvm::IntelGPU::IGCATarget IT = llvm::IntelGPU::parseIGCATarget(S))
return OffloadArch::getIntelIGCA(IT);
return OffloadArch::getUnknown();
diff --git a/clang/lib/Basic/TargetInfo.cpp b/clang/lib/Basic/TargetInfo.cpp
index 6e390226d2e4b8..43ed97494a3f66 100644
--- a/clang/lib/Basic/TargetInfo.cpp
+++ b/clang/lib/Basic/TargetInfo.cpp
@@ -16,6 +16,7 @@
#include "clang/Basic/Diagnostic.h"
#include "clang/Basic/DiagnosticFrontend.h"
#include "clang/Basic/LangOptions.h"
+#include "clang/Basic/OffloadArch.h"
#include "llvm/ADT/APFloat.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/StringExtras.h"
@@ -601,6 +602,10 @@ bool TargetInfo::initFeatureMap(
return true;
}
+OffloadArch TargetInfo::getOffloadArch() const {
+ return OffloadArch::getUnused();
+}
+
ParsedTargetAttr TargetInfo::parseTargetAttr(StringRef Features) const {
ParsedTargetAttr Ret;
if (Features == "default")
diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp
index 4cdb79c5f10370..994dd60f4ec84b 100644
--- a/clang/lib/Frontend/InitPreprocessor.cpp
+++ b/clang/lib/Frontend/InitPreprocessor.cpp
@@ -14,6 +14,7 @@
#include "clang/Basic/DiagnosticLex.h"
#include "clang/Basic/HLSLRuntime.h"
#include "clang/Basic/MacroBuilder.h"
+#include "clang/Basic/OffloadArch.h"
#include "clang/Basic/SourceManager.h"
#include "clang/Basic/SyncScope.h"
#include "clang/Basic/TargetInfo.h"
@@ -27,6 +28,7 @@
#include "llvm/ADT/APFloat.h"
#include "llvm/IR/DataLayout.h"
#include "llvm/IR/DerivedTypes.h"
+#include "llvm/TargetParser/IntelGPUTargetParser.h"
using namespace clang;
static bool MacroBodyEndsInBackslash(StringRef MacroBody) {
@@ -1507,6 +1509,20 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
Builder.defineMacro("__SYCL_DEVICE_ONLY__", "1");
if (OffloadArch Arch = TI.getOffloadArch(); !Arch.isUnknownOrUnused()) {
// TODO some sort of function to process -target-cpu
+ if (Arch.isIntelIGCA()) {
+ llvm::IntelGPU::IGCATarget IGCA = Arch.igcaTarget();
+
+ Builder.defineMacro("__OFFLOAD_ARCH_IGCA__", Twine(IGCA.getTarget()));
+ if (IGCA.isCompute())
+ Builder.defineMacro("__OFFLOAD_ARCH_IGCA_COMPUTE__");
+ else if (IGCA.isRender())
+ Builder.defineMacro("__OFFLOAD_ARCH_IGCA_RENDER__");
+
+ if (IGCA.isComputeExact())
+ Builder.defineMacro("__OFFLOAD_ARCH_IGCA_COMPUTE_EXACT__");
+ else if (IGCA.isRenderExact())
+ Builder.defineMacro("__OFFLOAD_ARCH_IGCA_RENDER_EXACT__");
+ }
}
}
>From 967e5f5d9c8106090f52d3d5847d07748a8dd29b Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Wed, 30 Sep 2026 13:42:41 -0700
Subject: [PATCH 23/30] Test that macros are actually working
---
clang/test/Preprocessor/sycl-igca-macros.cpp | 44 ++++++++++++++++++++
1 file changed, 44 insertions(+)
create mode 100644 clang/test/Preprocessor/sycl-igca-macros.cpp
diff --git a/clang/test/Preprocessor/sycl-igca-macros.cpp b/clang/test/Preprocessor/sycl-igca-macros.cpp
new file mode 100644
index 00000000000000..72c4a5cca06fa6
--- /dev/null
+++ b/clang/test/Preprocessor/sycl-igca-macros.cpp
@@ -0,0 +1,44 @@
+/// Check that IGCA target macros are predefined correctly during SYCL device
+/// compilation if an IGCA target was provided through -target-cpu.
+
+/// Core feature set:
+// RUN: %clang_cc1 %s -fsycl-is-device -target-cpu igca_20 -E -dM \
+// RUN: | FileCheck %s --check-prefix=CORE --implicit-check-not=__OFFLOAD_ARCH_IGCA
+// CORE: #define __OFFLOAD_ARCH_IGCA__ 20
+
+/// Compute / render feature sets:
+// RUN: %clang_cc1 %s -fsycl-is-device -target-cpu igca_20c -E -dM \
+// RUN: | FileCheck %s --check-prefix=COMPUTE --implicit-check-not=__OFFLOAD_ARCH_IGCA
+// COMPUTE-DAG: #define __OFFLOAD_ARCH_IGCA__ 20
+// COMPUTE-DAG: #define __OFFLOAD_ARCH_IGCA_COMPUTE__ 1
+
+// RUN: %clang_cc1 %s -fsycl-is-device -target-cpu igca_15r -E -dM \
+// RUN: | FileCheck %s --check-prefix=RENDER --implicit-check-not=__OFFLOAD_ARCH_IGCA
+// RENDER-DAG: #define __OFFLOAD_ARCH_IGCA__ 15
+// RENDER-DAG: #define __OFFLOAD_ARCH_IGCA_RENDER__ 1
+
+/// Compute-exact / render-exact:
+// RUN: %clang_cc1 %s -fsycl-is-device -target-cpu igca_20ca -E -dM \
+// RUN: | FileCheck %s --check-prefix=COMPUTE-EXACT --implicit-check-not=__OFFLOAD_ARCH_IGCA
+// COMPUTE-EXACT-DAG: #define __OFFLOAD_ARCH_IGCA__ 20
+// COMPUTE-EXACT-DAG: #define __OFFLOAD_ARCH_IGCA_COMPUTE__ 1
+// COMPUTE-EXACT-DAG: #define __OFFLOAD_ARCH_IGCA_COMPUTE_EXACT__ 1
+
+// RUN: %clang_cc1 %s -fsycl-is-device -target-cpu igca_15ra -E -dM \
+// RUN: | FileCheck %s --check-prefix=RENDER-EXACT --implicit-check-not=__OFFLOAD_ARCH_IGCA
+// RENDER-EXACT-DAG: #define __OFFLOAD_ARCH_IGCA__ 15
+// RENDER-EXACT-DAG: #define __OFFLOAD_ARCH_IGCA_RENDER__ 1
+// RENDER-EXACT-DAG: #define __OFFLOAD_ARCH_IGCA_RENDER_EXACT__ 1
+
+/// No IGCA macros should be defined without an IGCA target:
+// RUN: %clang_cc1 %s -fsycl-is-device -E -dM \
+// RUN: | FileCheck %s --check-prefix=NONE --implicit-check-not=__OFFLOAD_ARCH_IGCA
+/// IGCA macros should only be defined during device compilation:
+// RUN: %clang_cc1 %s -triple spirv64-unknown-unknown -target-cpu igca_20ca -E -dM \
+// RUN: | FileCheck %s --check-prefix=NONE --implicit-check-not=__OFFLOAD_ARCH_IGCA
+// NONE: #define __STDC__ 1
+
+/// Reject unknown IGCA targets:
+// RUN: not %clang_cc1 %s -fsycl-is-device -target-cpu igca_42 -E -dM 2>&1 \
+// RUN: | FileCheck %s --check-prefix=INVALID
+// INVALID: error: unknown target CPU 'igca_42'
>From 530919fbff994970a1286b2e164b3bce706daab4 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Wed, 30 Sep 2026 13:54:16 -0700
Subject: [PATCH 24/30] Remove TODO
---
clang/lib/Basic/OffloadArch.cpp | 1 -
clang/lib/Frontend/InitPreprocessor.cpp | 1 -
2 files changed, 2 deletions(-)
diff --git a/clang/lib/Basic/OffloadArch.cpp b/clang/lib/Basic/OffloadArch.cpp
index f5d9ec497ada39..8428fa34998f97 100644
--- a/clang/lib/Basic/OffloadArch.cpp
+++ b/clang/lib/Basic/OffloadArch.cpp
@@ -101,7 +101,6 @@ OffloadArch StringToOffloadArch(llvm::StringRef S) {
return OffloadArch::getNVPTX(NV);
if (llvm::AMDGPU::GPUKind AK = llvm::AMDGPU::parseArchAMDGCN(S))
return OffloadArch::getAMDGPU(AK);
- // TODO test that the operator bool causes this to fail if invalid
if (llvm::IntelGPU::IGCATarget IT = llvm::IntelGPU::parseIGCATarget(S))
return OffloadArch::getIntelIGCA(IT);
return OffloadArch::getUnknown();
diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp
index 994dd60f4ec84b..5ca59fa9b5b1da 100644
--- a/clang/lib/Frontend/InitPreprocessor.cpp
+++ b/clang/lib/Frontend/InitPreprocessor.cpp
@@ -1508,7 +1508,6 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
if (LangOpts.SYCLIsDevice) {
Builder.defineMacro("__SYCL_DEVICE_ONLY__", "1");
if (OffloadArch Arch = TI.getOffloadArch(); !Arch.isUnknownOrUnused()) {
- // TODO some sort of function to process -target-cpu
if (Arch.isIntelIGCA()) {
llvm::IntelGPU::IGCATarget IGCA = Arch.igcaTarget();
>From c54d0b6331d936cfed1ddaf222b1f7f6f000a6b9 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Wed, 30 Sep 2026 14:40:22 -0700
Subject: [PATCH 25/30] clang-format
---
clang/lib/Basic/Targets/OffloadTarget.h | 4 ++--
clang/lib/Basic/Targets/SPIR.h | 2 +-
2 files changed, 3 insertions(+), 3 deletions(-)
diff --git a/clang/lib/Basic/Targets/OffloadTarget.h b/clang/lib/Basic/Targets/OffloadTarget.h
index afcc2aa0db6e95..a121ac699a79fa 100644
--- a/clang/lib/Basic/Targets/OffloadTarget.h
+++ b/clang/lib/Basic/Targets/OffloadTarget.h
@@ -13,9 +13,9 @@
#ifndef LLVM_CLANG_LIB_BASIC_TARGETS_OFFLOADTARGET_H
#define LLVM_CLANG_LIB_BASIC_TARGETS_OFFLOADTARGET_H
+#include "clang/Basic/OffloadArch.h"
#include "clang/Basic/TargetInfo.h"
#include "clang/Basic/TargetOptions.h"
-#include "clang/Basic/OffloadArch.h"
#include "llvm/Support/Compiler.h"
#include "llvm/TargetParser/Triple.h"
@@ -27,7 +27,7 @@ class LLVM_LIBRARY_VISIBILITY OffloadTargetInfo : public TargetInfo {
OffloadArch DeviceArch = OffloadArch::getUnused();
public:
- OffloadTargetInfo(const llvm::Triple &Triple) : TargetInfo(Triple) { }
+ OffloadTargetInfo(const llvm::Triple &Triple) : TargetInfo(Triple) {}
// Clang driver emits -target-cpu to indicate offload device architecture for
// both CUDA and SYCL. We override setCPU to capture this offload device
diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h
index bb2a61e1eb9f16..feb7aad713fecc 100644
--- a/clang/lib/Basic/Targets/SPIR.h
+++ b/clang/lib/Basic/Targets/SPIR.h
@@ -13,8 +13,8 @@
#ifndef LLVM_CLANG_LIB_BASIC_TARGETS_SPIR_H
#define LLVM_CLANG_LIB_BASIC_TARGETS_SPIR_H
-#include "Targets.h"
#include "OffloadTarget.h"
+#include "Targets.h"
#include "clang/Basic/AddressSpaces.h"
#include "clang/Basic/Diagnostic.h"
#include "clang/Basic/DiagnosticFrontend.h"
>From baf3dbe184c2b2559839fdd174a16237ec45d2fb Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Fri, 23 Oct 2026 18:59:11 -0700
Subject: [PATCH 26/30] fix change in API from base branch
---
clang/lib/Frontend/InitPreprocessor.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp
index 5ca59fa9b5b1da..64d0aca02768a8 100644
--- a/clang/lib/Frontend/InitPreprocessor.cpp
+++ b/clang/lib/Frontend/InitPreprocessor.cpp
@@ -1512,9 +1512,9 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
llvm::IntelGPU::IGCATarget IGCA = Arch.igcaTarget();
Builder.defineMacro("__OFFLOAD_ARCH_IGCA__", Twine(IGCA.getTarget()));
- if (IGCA.isCompute())
+ if (IGCA.isAnyCompute())
Builder.defineMacro("__OFFLOAD_ARCH_IGCA_COMPUTE__");
- else if (IGCA.isRender())
+ else if (IGCA.isAnyRender())
Builder.defineMacro("__OFFLOAD_ARCH_IGCA_RENDER__");
if (IGCA.isComputeExact())
>From ece9f75cc8a2a0bb3c97acda6bcda805e1418172 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Fri, 23 Oct 2026 21:25:56 -0700
Subject: [PATCH 27/30] delete obsolete fillValidCPUList
---
clang/lib/Basic/Targets/OffloadTarget.h | 4 ----
1 file changed, 4 deletions(-)
diff --git a/clang/lib/Basic/Targets/OffloadTarget.h b/clang/lib/Basic/Targets/OffloadTarget.h
index a121ac699a79fa..c1680f4f188c80 100644
--- a/clang/lib/Basic/Targets/OffloadTarget.h
+++ b/clang/lib/Basic/Targets/OffloadTarget.h
@@ -41,10 +41,6 @@ class LLVM_LIBRARY_VISIBILITY OffloadTargetInfo : public TargetInfo {
return !StringToOffloadArch(Name).isUnknown();
}
- void fillValidCPUList(SmallVectorImpl<StringRef> &Values) const override {
- fillValidOffloadArchList(Values);
- }
-
OffloadArch getOffloadArch() const override { return DeviceArch; }
};
>From 79a58a3fbf039ec6592f3a1428f39a36e1f83735 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Sat, 24 Oct 2026 00:07:04 -0700
Subject: [PATCH 28/30] When an obviously wrong IGCA target is provided,
provide a list of valid IGCA targets
---
clang/lib/Driver/Driver.cpp | 5 +++++
clang/test/Driver/sycl-offload-arch-igca.cpp | 2 ++
.../llvm/TargetParser/IntelGPUTargetParser.h | 3 +++
llvm/lib/TargetParser/IntelGPUTargetParser.cpp | 18 ++++++++++++++++++
.../TargetParser/IntelGPUTargetParserTest.cpp | 11 +++++++++++
5 files changed, 39 insertions(+)
diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp
index 1c07618e48e8bd..6d95f07b83a8e6 100644
--- a/clang/lib/Driver/Driver.cpp
+++ b/clang/lib/Driver/Driver.cpp
@@ -110,6 +110,7 @@
#include "llvm/Support/VirtualFileSystem.h"
#include "llvm/Support/raw_ostream.h"
#include "llvm/TargetParser/Host.h"
+#include "llvm/TargetParser/IntelGPUTargetParser.h"
#include "llvm/TargetParser/RISCVISAInfo.h"
#include <cstdlib> // ::getenv
#include <map>
@@ -3960,6 +3961,10 @@ static StringRef getCanonicalArchString(Compilation &C,
if (!StringToOffloadArch(ArchStr).isIntelIGCA()) {
C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
<< "IGCA" << ArchStr;
+ SmallVector<StringRef, 32> ValidList;
+ llvm::IntelGPU::fillValidIGCATargetList(ValidList);
+ C.getDriver().Diag(clang::diag::note_valid_options)
+ << llvm::join(ValidList, ", ");
return StringRef();
}
}
diff --git a/clang/test/Driver/sycl-offload-arch-igca.cpp b/clang/test/Driver/sycl-offload-arch-igca.cpp
index 63c8922a97ac68..4c6c74dd2a838a 100644
--- a/clang/test/Driver/sycl-offload-arch-igca.cpp
+++ b/clang/test/Driver/sycl-offload-arch-igca.cpp
@@ -10,6 +10,7 @@
// RUN: --offload-targets=spirv64-unknown-unknown --offload-arch=igca_42a \
// RUN: -c %s 2>&1 | FileCheck -check-prefix=CHK-INVALID %s
// CHK-INVALID: error: unsupported IGCA gpu architecture: igca_42a
+// CHK-INVALID-NEXT: note: valid target CPU values are: igca_{{.*}}
// CHK-INVALID-NOT: llvm-offload-binary
/// IGCA targets should not accept feature suffixes, so a valid target followed
@@ -18,6 +19,7 @@
// RUN: --offload-targets=spirv64-unknown-unknown --offload-arch=igca_20ca:foo \
// RUN: -c %s 2>&1 | FileCheck -check-prefix=CHK-FEATURE %s
// CHK-FEATURE: error: unsupported IGCA gpu architecture: igca_20ca:foo
+// CHK-FEATURE-NEXT: note: valid target CPU values are: igca_{{.*}}
// CHK-FEATURE-NOT: llvm-offload-binary
/// A valid IGCA target is still rejected on NVPTX and AMDGPU device triples,
diff --git a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
index 035abb848e8fd5..f2ac7a87f50811 100644
--- a/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
+++ b/llvm/include/llvm/TargetParser/IntelGPUTargetParser.h
@@ -161,6 +161,9 @@ LLVM_ABI IGCATarget parseIGCATarget(StringRef MaybeTarget);
/// \p Target is invalid or not a known target in IntelGPUTargetParser.def.
LLVM_ABI StringRef getIGCATargetName(IGCATarget Target);
+/// Append every valid IGCA target spelling to \p Values.
+LLVM_ABI void fillValidIGCATargetList(SmallVectorImpl<StringRef> &Values);
+
} // namespace IntelGPU
} // namespace llvm
diff --git a/llvm/lib/TargetParser/IntelGPUTargetParser.cpp b/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
index b9e1517170e615..b551227bc496e4 100644
--- a/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
+++ b/llvm/lib/TargetParser/IntelGPUTargetParser.cpp
@@ -11,6 +11,8 @@
//===----------------------------------------------------------------------===//
#include "llvm/TargetParser/IntelGPUTargetParser.h"
+#include "llvm/ADT/STLExtras.h"
+#include "llvm/ADT/SmallVector.h"
#include "llvm/ADT/StringSwitch.h"
#include "llvm/ADT/Twine.h"
#include <cassert>
@@ -80,3 +82,19 @@ StringRef llvm::IntelGPU::getIGCATargetName(IGCATarget T) {
#include "llvm/TargetParser/IntelGPUTargetParser.def"
return "";
}
+
+void llvm::IntelGPU::fillValidIGCATargetList(
+ SmallVectorImpl<StringRef> &Values) {
+ // fillValidIGCATargetList is currently only used for diagnostics, similar to
+ // how other fillValid* functions are used. We can thus afford to generate
+ // this list when we actually need it:
+ SmallVector<IGCATarget> Targets;
+#define INTEL_IGCA_TARGET(NAME, TARGET, FEATURE_SET) \
+ Targets.push_back(IGCATarget(TARGET, IGCAFeatureSet::FEATURE_SET));
+#include "llvm/TargetParser/IntelGPUTargetParser.def"
+ llvm::sort(Targets,
+ [](IGCATarget A, IGCATarget B) { return A.pack() < B.pack(); });
+ Targets.erase(llvm::unique(Targets), Targets.end());
+ for (IGCATarget T : Targets)
+ Values.push_back(getIGCATargetName(T));
+}
diff --git a/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp b/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
index 207207095a9f5a..6f9d35b0cc0d89 100644
--- a/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
+++ b/llvm/unittests/TargetParser/IntelGPUTargetParserTest.cpp
@@ -190,6 +190,17 @@ TEST(IntelGPUTargetParserTest, IGCATargetName) {
IntelGPU::getIGCATargetName(IGCATarget(10, IGCAFeatureSet::Compute)), "");
}
+TEST(IntelGPUTargetParserTest, EveryIGCASpellingRoundTrips) {
+ SmallVector<StringRef> Names;
+ IntelGPU::fillValidIGCATargetList(Names);
+ EXPECT_FALSE(Names.empty());
+ for (StringRef Name : Names) {
+ IntelGPU::IGCATarget T = IntelGPU::parseIGCATarget(Name);
+ EXPECT_TRUE(T.isValid()) << Name;
+ EXPECT_EQ(IntelGPU::getIGCATargetName(T), Name);
+ }
+}
+
// Spell an IGCA target from its fields, independently of the NAME column.
std::string expectedIGCASpelling(uint16_t Target,
IntelGPU::IGCAFeatureSet FeatureSet) {
>From e210efe1d78bbab559a3e6360459d5b061197eac Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Sat, 24 Oct 2026 00:31:50 -0700
Subject: [PATCH 29/30] formatting changes
---
clang/include/clang/Basic/TargetInfo.h | 6 ++++--
clang/lib/Basic/Targets/OffloadTarget.h | 8 ++++----
2 files changed, 8 insertions(+), 6 deletions(-)
diff --git a/clang/include/clang/Basic/TargetInfo.h b/clang/include/clang/Basic/TargetInfo.h
index 6a46b4bebcc25b..7fad8de84fc0d7 100644
--- a/clang/include/clang/Basic/TargetInfo.h
+++ b/clang/include/clang/Basic/TargetInfo.h
@@ -1426,8 +1426,10 @@ class TargetInfo : public TransferrableTargetInfo,
return isValidCPUName(Name);
}
- /// Returns an OffloadArch if targeting an offload device, otherwise return
- /// Unused.
+ /// Returns the OffloadArch if targeting an offload device, otherwise returns
+ /// OffloadArch::getUnused().
+ virtual OffloadArch getOffloadArch() const;
+
virtual OffloadArch getOffloadArch() const;
virtual ParsedTargetAttr parseTargetAttr(StringRef Str) const;
diff --git a/clang/lib/Basic/Targets/OffloadTarget.h b/clang/lib/Basic/Targets/OffloadTarget.h
index c1680f4f188c80..def181b5847dd8 100644
--- a/clang/lib/Basic/Targets/OffloadTarget.h
+++ b/clang/lib/Basic/Targets/OffloadTarget.h
@@ -29,10 +29,10 @@ class LLVM_LIBRARY_VISIBILITY OffloadTargetInfo : public TargetInfo {
public:
OffloadTargetInfo(const llvm::Triple &Triple) : TargetInfo(Triple) {}
- // Clang driver emits -target-cpu to indicate offload device architecture for
- // both CUDA and SYCL. We override setCPU to capture this offload device
- // architecture information.
bool setCPU(StringRef Name) override {
+ // Clang driver emits -target-cpu to indicate offload device architecture
+ // for both CUDA and SYCL. We override setCPU to capture this offload device
+ // architecture information.
DeviceArch = StringToOffloadArch(Name);
return !DeviceArch.isUnknownOrUnused();
}
@@ -47,4 +47,4 @@ class LLVM_LIBRARY_VISIBILITY OffloadTargetInfo : public TargetInfo {
} // namespace targets
} // namespace clang
-#endif // LLVM_CLANG_LIB_BASIC_TARGETS_OFFLOADTARGET_H
\ No newline at end of file
+#endif // LLVM_CLANG_LIB_BASIC_TARGETS_OFFLOADTARGET_H
>From 59e305ebf2e4ecafd8bf393822596688368b9336 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Sat, 24 Oct 2026 12:35:56 -0700
Subject: [PATCH 30/30] editing things in vsode's diff browser is a bad idea
and I should've known
---
clang/include/clang/Basic/TargetInfo.h | 2 --
1 file changed, 2 deletions(-)
diff --git a/clang/include/clang/Basic/TargetInfo.h b/clang/include/clang/Basic/TargetInfo.h
index 7fad8de84fc0d7..7215cb8938ffc7 100644
--- a/clang/include/clang/Basic/TargetInfo.h
+++ b/clang/include/clang/Basic/TargetInfo.h
@@ -1430,8 +1430,6 @@ class TargetInfo : public TransferrableTargetInfo,
/// OffloadArch::getUnused().
virtual OffloadArch getOffloadArch() const;
- virtual OffloadArch getOffloadArch() const;
-
virtual ParsedTargetAttr parseTargetAttr(StringRef Str) const;
/// Determine whether this TargetInfo supports tune in target attribute.
More information about the cfe-commits
mailing list