[wgsl] Add num_subgroups builtin
This is just the WGSL frontend support. The builtin is not hooked up
in any backend yet.
Bug: 454652043
Change-Id: Ia5ad1e1a5e37185ae5195201d9b45a1473210c62
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/268576
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Commit-Queue: James Price <jrprice@google.com>
diff --git a/src/tint/cmd/bench/enums_core_bench.cc b/src/tint/cmd/bench/enums_core_bench.cc
index c0532be..cbf34b7 100644
--- a/src/tint/cmd/bench/enums_core_bench.cc
+++ b/src/tint/cmd/bench/enums_core_bench.cc
@@ -1074,76 +1074,83 @@
"ltcal_invocation_index",
"local_in1ocation_index",
"localAAinvocation_inDDex",
- "Om_woP2grups",
- "num_workgoups",
- "num_workgro77ps",
+ "Pum_2uboOps",
+ "numsubgroups",
+ "num_77ubgroups",
+ "num_subgroups",
+ "num_smmUbAAro1s",
+ "nAAm_Xubgroufs",
+ "nuv_VVuvgrous",
+ "num_workgrppups",
+ "numTTworkgroups",
+ "num_workgjjok0ps",
"num_workgroups",
- "n1m_UUorkAAroups",
- "Xum_workgrfupAA",
- "VVum_wrkgrouvs",
- "posippion",
- "posTTtion",
- "kosjj0tion",
+ "um_workgroups",
+ "nu_yyorkgroups",
+ "nuzz_wokgrouYYs",
+ "positWion",
+ "0o7tiUUn",
+ "UUosipi",
"position",
- "posiion",
- "posiyyon",
- "poYizzion",
- "priWitiiie_index",
- "6rimitv7_indUUx",
- "prUUitive_inpe",
+ "hTstion",
+ "posi5on",
+ "itio1",
+ "priIitimme_i44dex",
+ "priYRYYDie_index",
+ "primjjtive_index",
"primitive_index",
- "hrimitiveiTTex",
- "primitie_inde5",
- "primtveind1x",
- "sam44mme_iIdex",
- "saYDpRe_indYY",
- "sampjje_index",
+ "prmitive_index",
+ "QQrimtide_index",
+ "primtive_in5Hex",
+ "saxpkFF_index",
+ "sapreindex",
+ "sampleKKi_Nex",
"sample_index",
- "sample_idex",
- "smple_iQQdex",
- "55amHle_idex",
- "saxFple_mask",
- "sapD_mrsk",
- "sNmple__KKsk",
+ "sample_aTdex",
+ "samplevvindex",
+ "amplejjindex",
+ "samplnnOOmas",
+ "sggOOpleOmsk",
+ "sawwQpe_mask",
"sample_mask",
- "saxalT_mask",
- "sample_vvask",
- "sjjmple_mak",
- "subgrnnOOp_i",
- "sggOOgroOpid",
- "suwwQgoup_id",
+ "sampl_masJ",
+ "sample_masD",
+ "OOample_msk",
+ "subgroup_d",
+ "subro0p_id",
+ "lbgroup_d",
"subgroup_id",
- "subgrup_iJ",
- "subgroup_iD",
- "OOubgroupid",
- "subgroup_nvocation_id",
- "subro00p_invocation_i",
- "subgroup_vqcation_ld",
+ "subgroup_iu",
+ "su33group_id",
+ "subgrouNNid",
+ "subgroup_inocaton_Md",
+ "subg22oup_Rnocationbbid",
+ "subgroup_ivocatiiin_id",
"subgroup_invocation_id",
- "subgroup_invocation_iu",
- "subgroup_invo33ation_id",
- "subgrouNNinvocation_id",
- "subrMup_sze",
- "ubgr22up_sRzbb",
- "subgiiup_size",
+ "ubgroup_invocation_id",
+ "subgup_inCoocution_id",
+ "subHHoup_inUoaJion_id",
+ "subghrou_yize",
+ "WWugroup_BBze",
+ "EEbccroup_sCEze",
"subgroup_size",
- "subgrup_size",
- "soobgop_sCue",
- "subgrHHu_izJ",
- "vrtyg_ihdggx",
- "vrex_WWndeBB",
- "vertex_inECEcc",
+ "subgreep_siGyy",
+ "uUUroup44size",
+ "subgroup_sie",
+ "verteN_index",
+ "MMertexssindex",
+ "vertexDindCx",
"vertex_index",
- "veGtex_ieedyy",
- "ve44teUindex",
- "vertex_inde",
- "workgNoup_id",
- "MMorkgrssup_id",
- "workgrDup_Cd",
+ "ggTYtex_index",
+ "vertjjx_index",
+ "vrtDDx_inde",
+ "99MMkEEroup_id",
+ "wNrkgrou_id",
+ "55orkgroup_id",
"workgroup_id",
- "ggTYkgroup_id",
- "workjjroup_id",
- "wrkDDroup_i",
+ "Qorkgroup_id",
+ "workgronnp_i",
+ "workFFroWWid",
};
for (auto _ : state) {
for (auto* str : kStrings) {
@@ -1156,41 +1163,41 @@
BENCHMARK(BuiltinValueParser);
void InterpolationSamplingParser(::benchmark::State& state) {
const char* kStrings[] = {
- "cEEMM9er",
- "cNter",
- "55enter",
+ "ce66tew",
+ "cvFter",
+ "enWWr",
"center",
- "Qenter",
- "nnnter",
- "cenWFr",
- "centro66w",
- "cvntroFd",
- "cWWntrd",
+ "cntr",
+ "5tK",
+ "cen4er",
+ "cnPtrggd",
+ "UentrSS",
+ "centro2V",
"centroid",
- "entrod",
- "5Ktri",
- "c4ntroid",
- "gtnnPr",
- "eSShe",
- "2itVer",
+ "cpntroid",
+ "X9VVn9r2id",
+ "ntrnd",
+ "ser",
+ "8BBte",
+ "AyACCher",
"either",
- "epther",
- "eiXV9992",
- "nte",
- "Bs",
- "fiBB",
- "yCCAAAt",
+ "eihBv",
+ "A4QQter",
+ "ZZtheT4",
+ "Xirlt",
+ "fiddt",
+ "fis",
"first",
- "ivit",
- "A4rsQQ",
- "T4Zist",
- "Xamlle",
- "ddamle",
- "sadd",
+ "dis1",
+ "GGvs",
+ "jjst",
+ "samwle",
+ "Xamnle",
+ "ja9ple",
"sample",
- "si1le",
- "vmpGG",
- "jjmpl",
+ "sape",
+ "sam2WWe",
+ "samle",
};
for (auto _ : state) {
for (auto* str : kStrings) {
@@ -1203,27 +1210,27 @@
BENCHMARK(InterpolationSamplingParser);
void InterpolationTypeParser(::benchmark::State& state) {
const char* kStrings[] = {
- "fwat",
- "flXn",
- "fljt",
+ "flV",
+ "fnzal",
+ "Ylt",
"flat",
- "fz",
- "flWW2",
- "fat",
- "lneVr",
- "lnz4ear",
- "li__Y",
+ "HMkat",
+ "fzzDI",
+ "dla",
+ "lineh",
+ "lFMMlaN",
+ "lina",
"linear",
- "linHkaM",
- "lInzzaD",
- "qnedr",
- "persehtive",
- "perlMMNctiFe",
- "perspetiv",
+ "9Anx",
+ "VVnnar",
+ "iKer",
+ "perBBCecive",
+ "perpoctive",
+ "perspeccwtEE8e",
"perspective",
- "xersA9tiv",
- "VVnrspetive",
- "perspecKv",
+ "erspectivbb",
+ "6erspective",
+ "erspectgg",
};
for (auto _ : state) {
for (auto* str : kStrings) {
@@ -1236,286 +1243,286 @@
BENCHMARK(InterpolationTypeParser);
void TexelFormatParser(::benchmark::State& state) {
const char* kStrings[] = {
- "bgra8nCBBm",
- "bgrauoorm",
- "bgra8ccnE88rw",
+ "braIun11Em",
+ "bgra8unrm",
+ "pHgra8uorm",
"bgra8unorm",
- "bgr8bbnorm",
- "b6ra8unorm",
- "gKaggnrm",
- "r16fE1at",
- "r16floa",
- "rHpfloat",
+ "bgqqaa8uGorm",
+ "bgra8unm",
+ "bgra8unrxxxx",
+ "r1zzfloaSS",
+ "r16foat",
+ "1o4oat",
"r16float",
- "qqaa6flGat",
- "r6loat",
- "xxxx6_loat",
- "r16SSzznt",
- "r16int",
- "16ont",
+ "r16fla",
+ "rKflat",
+ "rzz6flct",
+ "1gg88ant",
+ "r1jjsint",
+ "riA6sin1",
"r16sint",
- "r16nt",
- "rK6st",
- "rczsint",
- "r88an5ggm",
- "r16snsjrm",
- "116snoriA",
+ "rlVVs66nt",
+ "6sic",
+ "r16GGnht",
+ "jiiyysnor",
+ "rv_QQsnttm",
+ "r6soLLm",
"r16snorm",
- "rVV6sl66rm",
- "6norm",
- "tt1Ghmnom",
- "iiyy6int",
- "QQt1_vuin",
- "rLuint",
+ "r16Wno7m",
+ "6svvoFr",
+ "ryysnzzrm",
+ "16uio",
+ "r16uI_c",
+ "rdd6int",
"r16uint",
- "r1uWcct",
- "r16uFt",
- "rzz6yyin",
- "runorm",
- "I16Rco_m",
- "r16uddrm",
+ "lM6ui2",
+ "r1Ouint",
+ "rKK6muiTT",
+ "rRR6no",
+ "4uno00",
+ "r16nqor",
"r16unorm",
- "M1l22or",
- "O16unorm",
- "KK6TTunmrm",
- "rRR2lo",
- "4flo00",
- "r32nqoa",
+ "r16unorBB",
+ "r16u2WBarm",
+ "r16ujjorm",
+ "rt6HHlot",
+ "3JJflat",
+ "r3ooSlttat",
"r32float",
- "r32floaBB",
- "r32f2WBaat",
- "r32fjjoat",
- "66tinHH",
- "r3JJst",
- "rttoosiSt",
- "r32sint",
- "2r_SSnt",
- "rWW2sint",
- "r32sink",
+ "rSS2far_",
+ "r32flWWat",
+ "r32flokt",
"r3bbitt",
- "r32u1n",
- "rf2uiN",
- "r32uint",
+ "r32s1n",
+ "rf2siN",
+ "r32sint",
"rii2iVt",
"r32winU",
- "r3Funt",
- "448s8LLnt",
- "r8sinZZ",
- "rMMsint",
+ "r3Fsnt",
+ "r38LLuin44",
+ "r32uiZZt",
+ "MM32uint",
+ "r32uint",
+ "r3uint",
+ "rSSMMu6At",
+ "GGgxuinTT",
+ "rIsint",
+ "l8sint",
+ "r8OOn",
"r8sint",
- "8sint",
- "SSAM6int",
- "GGxTTint",
- "r8snorI",
- "r8slorm",
- "rOOsnm",
+ "r8s5n",
+ "r8zRint",
+ "r8siPLL",
+ "r8snhhm",
+ "rsnorm",
+ "8snjm",
"r8snorm",
- "8aKo5m",
- "Rf__zzrm",
- "r8LLnorP",
- "ruihht",
- "ruint",
- "ruit",
+ "r8BB0",
+ "xxL8nrm",
+ "rIIsnorMM",
+ "rXFFig",
+ "rzssunp",
+ "rZZrrinF",
"r8uint",
- "08BB",
- "xxL8un",
- "r8uiMMIt",
- "r8XnFFm",
- "r8unsszz",
- "r8rrnZZFm",
+ "r8jir00",
+ "rwwuirr",
+ "r8iZQ",
+ "rn__nrh",
+ "r8worm",
+ "r3unorm",
"r8unorm",
- "rrrjor00",
- "p8wrrorm",
- "QZuorm",
- "rg__hb1nfloat",
- "rg11w10ufloa",
- "rg11b103float",
+ "runoiXm",
+ "r83m",
+ "rOvv8orm",
+ "rg1Pbvju5loat",
+ "rgi1b1Nfloat",
+ "rg21b10ufloa",
"rg11b10ufloat",
- "rg11b0XXfloai",
- "rg31b10loa",
- "rg11c10u8vvoOO",
- "rg16vlPaj",
- "ugi6plot",
- "g16f2oat",
+ "rg11b1ufDJqat",
+ "r11b10mQloat",
+ "rg11bJ0uflat",
+ "rg16fWWat",
+ "wQg16gglwat",
+ "rg1zfloat",
"rg16float",
- "rgD6fJqat",
- "rg1Qfmoa",
- "rg16fJat",
- "rg1sWWnt",
- "rQQ1wswgnt",
- "rg16sinz",
+ "rWW6flzzat",
+ "rg6flo",
+ "rKUU6fla99q",
+ "RRgEE6snt",
+ "r1sint",
+ "rg16sttt",
"rg16sint",
- "rgWWsinzz",
- "rg1it",
- "rK996qqUUn",
- "EEg16sRRom",
- "rg16nor",
- "rg1tsnorm",
+ "rg1MMin",
+ "raaKKTBBint",
+ "l5516sIIn",
+ "rg16sooBBw",
+ "PPg1__snorm",
+ "rg16RRnorm",
"rg16snorm",
- "rgqnoMMm",
- "rgT6aanBBrKK",
- "rl16sII5rm",
- "rg16wBBoo",
- "rg1__PPint",
- "rRR16uint",
+ "rghSom",
+ "Q1srm",
+ "r6sn8HHrm",
+ "rg1YYuDnt",
+ "rg1Ouint",
+ "Jg8IIuinZ",
"rg16uint",
- "rgSut",
- "riQt",
- "g6uHH8nt",
- "rD16unYYrm",
- "rg16unoOm",
- "r8II6ZJorm",
+ "ogY6uint",
+ "r0X6uint",
+ "rg6KKSncn",
+ "GGg1633nr88",
+ "rXX133unoqqm",
+ "rz16no4m",
"rg16unorm",
- "rgoYunorm",
- "rg6unor0X",
- "rcKn6unoSm",
- "GGg3233la88",
- "rXX333floqqt",
- "rz32lo4t",
+ "rWttWKnhhrm",
+ "rg1rxorm",
+ "rg16CQerm",
+ "g3rYloat",
+ "rg32xmloat",
+ "r3float",
"rg32float",
- "rWttWKlhhat",
- "rg3rxoat",
- "rg32CQeat",
- "rg32sYr",
- "xxgm2sinS",
- "rg3rsit",
+ "rg3244ylyyat",
+ "r3wwfloat",
+ "r32fBBaii",
+ "rgdd2ssint",
+ "rgMsinee",
+ "Tg32sint",
"rg32sint",
- "rgyy2sit",
- "r82i8t",
- "rge2sint",
- "hkkg3uint",
- "rg2uOOBnt",
- "g32ui9t",
+ "r82s00l",
+ "6mmg1Lsint",
+ "Gg32IYYnt",
+ "rYY2uLVVt",
+ "m62s32uinI",
+ "r32uinjj",
"rg32uint",
- "rgg2uDnt",
- "rgna2in1",
- "6gIuinS",
- "rgsinYY",
- "Lg8YVVin",
- "rg8ms2I66",
+ "ppg2u7",
+ "Ng32ui6t",
+ "rg32uinsgg",
+ "g8sint",
+ "JgOOsqnt",
+ "rg00sinMM",
"rg8sint",
- "jjg8snt",
- "pp7nt",
- "rgN6int",
- "rg8snorsgg",
- "rgsnorm",
- "rg8sJoOOm",
+ "rMM8smmnt",
+ "rH8sqqt",
+ "gGsin88",
+ "gRRswworm",
+ "rgqsnorm",
+ "gUsom",
"rg8snorm",
- "MMg8snor00",
- "rg8snmmMMm",
- "r8snqHrm",
- "gGuin88",
- "wwgRRuit",
- "rg8uinq",
+ "g8sBo5MQ",
+ "33g8snorm",
+ "Og8sno33m",
+ "w8unt",
+ "rgAkkM_t",
+ "rg8int",
"rg8uint",
- "8uiU",
- "BQ85nMM",
- "rg8u33nt",
- "Og8uno33m",
- "rg8uom",
- "kkgAunb_M",
+ "dg8i",
+ "rXXI55wwnt",
+ "g8uint",
+ "g8dorm",
+ "dJJunom",
+ "rg8_orm",
"rg8unorm",
- "rg8unom",
- "rduom",
- "rg55wXXnorm",
- "rb10a2uint",
- "rgb10a2ut",
- "rgb1Jadint",
+ "rgnorL",
+ "99g8unrm",
+ "xU8unorm",
+ "9gb10a2uit",
+ "rvb10a2uint",
+ "LLgb10a22uint",
"rgb10a2uint",
- "rb_1a2uint",
- "rgLL0aint",
- "rg10G2uin99",
- "xxgb102UUnom",
- "rgb90a2unom",
- "rgb10a2uvorm",
+ "rg1bf1uint",
+ "rgb1AAa2uin",
+ "CU10a2ui",
+ "rgb1zza2unwnm",
+ "MMwVVb10a2unonm",
+ "rgb10aNNXorm",
"rgb10a2unorm",
- "rgb10LL22unorm",
- "rgb11a2ufom",
- "rgb0aAAungrm",
- "CUa16flo",
- "rgbw1zzflont",
- "rgna16flowVVMM",
+ "rgb10a2un77Wee",
+ "vvgb133a2PPRorm",
+ "rg2210a2unorjj",
+ "rgbaAA6flot",
+ "kkba1NfRRoat",
+ "rQQba16float",
"rgba16float",
- "rgNNaX6flat",
- "rgbee77flWWat",
- "rgb3316PPloavv",
- "22gba16sjjnt",
- "cgAA16sint",
- "Ngbakk6RRin",
+ "rgb16fuuoat",
+ "uuba16flat",
+ "rga1floaZ",
+ "rgHa16snt",
+ "rgdSSS16sint",
+ "rgba16sTTnt",
"rgba16sint",
- "rgba16sinQQ",
- "rgbuu6sint",
- "rba6suunt",
- "rga1snorZ",
- "rgba1sHorm",
- "rgba1dsnoSSSm",
+ "rmb16ctnt",
+ "rgba16shhnt",
+ "rgba16siG",
+ "HHgbaww6snAArbb",
+ "rgba1BBsnorm",
+ "rgbM2hsnogm",
"rgba16snorm",
- "rgba16snTTrm",
- "rba16cmort",
- "rgba16snohhm",
- "rgba16uiG",
- "rHbbwAAa16uint",
- "rgba16uBBnt",
+ "rgba1snorWW",
+ "rgbqq16snorm",
+ "rgbar6snUUm",
+ "rgba6uiEt",
+ "rrVVba16uRnt",
+ "rgbH31GGOint",
"rgba16uint",
- "hga226Mint",
- "rgba16unWW",
- "rqqba16uint",
- "rgbar6unUUm",
- "rgba16Eorm",
- "rgRa16unoVVrr",
+ "rgba1ui55",
+ "rgBa16uiSSt",
+ "rga16uazz",
+ "rga16noAA66",
+ "rgbw6_norm",
+ "rgba6unorm",
"rgba16unorm",
- "rg33OGG1Hunom",
- "55g16Dnorm",
- "rSSbB16unoQm",
- "zzgba32lat",
- "rga32loAA66",
- "rgbw2_loat",
+ "rgbr16norm",
+ "rgba16unRRrG",
+ "KK1ba1LLunoxm",
+ "rgba32NWWlFat",
+ "rgba3GfoUt",
+ "r6bEEFF2flat",
"rgba32float",
- "rgba2float",
- "rgbr32loat",
- "rgba32flRRaG",
- "KxUbLL31sint",
- "rgba32NNWWFt",
- "gbfGUsint",
- "rgba32sint",
- "rgba3FFE6nt",
- "rgba39yyJTTnt",
- "zgba32sit",
- "rb32uinI",
+ "99TJJbyy3float",
+ "rzb32float",
+ "gba3Iloat",
"gb32pt",
- "ttgnn32uint",
+ "ttgnn32sint",
+ "rgba3isiKt",
+ "rgba32sint",
+ "reeba32sint",
+ "rgbg_2si11t",
+ "rgb32sint",
+ "rgba3ont",
+ "rpbX3ui",
+ "rgYY6ss2uin",
"rgba32uint",
- "rgba3iuiKt",
- "reeba32uint",
- "rgbg_2ui11t",
- "rgba8snt",
- "rgbaino",
- "rpX8si",
+ "r77ba3YuInt",
+ "rgbaSSeeuinG",
+ "rgbQluint",
+ "gba8int",
+ "rgba__sint",
+ "rDDba8sin",
"rgba8sint",
- "g6ss8sYYnt",
- "rIIba8s7Yt",
- "ngbGSSeeint",
- "rgbQlnorm",
- "ba8snorm",
- "rgba8sn__rm",
- "rgba8snorm",
- "rgba8soDDm",
- "rgb8snrm",
- "rgb5l55nor1",
+ "rga8int",
+ "551balsint",
"rgbc4int",
- "gba8uint",
- "gba6Ou11nt",
- "rgba8uint",
- "rsa8unt",
- "rgbaXuint",
- "rgbw844int",
+ "gba8snorm",
+ "rgOa811no6m",
+ "rgasnosm",
+ "rgba8snorm",
+ "rgba8snXrm",
+ "rg44w8snorm",
"rgba8YYnorm",
- "gba8unorm",
- "rTa8unyrm",
+ "rgba8int",
+ "gTTa8ynt",
+ "rga8vvbbt",
+ "rgba8uint",
+ "rgbakuint",
+ "g99a5uit",
+ "lRRbauinjj",
+ "YYgb8u77orm",
+ "rbaLLJrm",
+ "rgnnunorm",
"rgba8unorm",
- "bbba8unovv",
- "rgba8unokm",
- "Vb58unor9",
+ "rooba8uord",
+ "rba8Enobm",
+ "rba8unorm",
};
for (auto _ : state) {
for (auto* str : kStrings) {
diff --git a/src/tint/cmd/fuzz/wgsl/dictionary.txt b/src/tint/cmd/fuzz/wgsl/dictionary.txt
index 8940160..d450174 100644
--- a/src/tint/cmd/fuzz/wgsl/dictionary.txt
+++ b/src/tint/cmd/fuzz/wgsl/dictionary.txt
@@ -320,6 +320,7 @@
"must_use"
"normalize"
"num_levels"
+"num_subgroups"
"num_workgroups"
"off"
"original_value"
diff --git a/src/tint/lang/core/core.def b/src/tint/lang/core/core.def
index 49076dd..14c414e 100644
--- a/src/tint/lang/core/core.def
+++ b/src/tint/lang/core/core.def
@@ -65,6 +65,7 @@
subgroup_id
subgroup_invocation_id
subgroup_size
+ num_subgroups
barycentric_coord
primitive_index
@internal __point_size
diff --git a/src/tint/lang/core/enums.cc b/src/tint/lang/core/enums.cc
index fe2f146..3d3e8a0 100644
--- a/src/tint/lang/core/enums.cc
+++ b/src/tint/lang/core/enums.cc
@@ -997,6 +997,9 @@
if (str == "local_invocation_index") {
return BuiltinValue::kLocalInvocationIndex;
}
+ if (str == "num_subgroups") {
+ return BuiltinValue::kNumSubgroups;
+ }
if (str == "num_workgroups") {
return BuiltinValue::kNumWorkgroups;
}
@@ -1053,6 +1056,8 @@
return "local_invocation_id";
case BuiltinValue::kLocalInvocationIndex:
return "local_invocation_index";
+ case BuiltinValue::kNumSubgroups:
+ return "num_subgroups";
case BuiltinValue::kNumWorkgroups:
return "num_workgroups";
case BuiltinValue::kPosition:
diff --git a/src/tint/lang/core/enums.h b/src/tint/lang/core/enums.h
index 181cb7b..3f38a8c 100644
--- a/src/tint/lang/core/enums.h
+++ b/src/tint/lang/core/enums.h
@@ -583,6 +583,7 @@
kInstanceIndex,
kLocalInvocationId,
kLocalInvocationIndex,
+ kNumSubgroups,
kNumWorkgroups,
kPosition,
kPrimitiveIndex,
@@ -622,6 +623,7 @@
"instance_index",
"local_invocation_id",
"local_invocation_index",
+ "num_subgroups",
"num_workgroups",
"position",
"primitive_index",
diff --git a/src/tint/lang/core/enums_test.cc b/src/tint/lang/core/enums_test.cc
index 03eb456..df6ee7a 100644
--- a/src/tint/lang/core/enums_test.cc
+++ b/src/tint/lang/core/enums_test.cc
@@ -959,6 +959,7 @@
{"instance_index", BuiltinValue::kInstanceIndex},
{"local_invocation_id", BuiltinValue::kLocalInvocationId},
{"local_invocation_index", BuiltinValue::kLocalInvocationIndex},
+ {"num_subgroups", BuiltinValue::kNumSubgroups},
{"num_workgroups", BuiltinValue::kNumWorkgroups},
{"position", BuiltinValue::kPosition},
{"primitive_index", BuiltinValue::kPrimitiveIndex},
@@ -996,36 +997,39 @@
{"locaZZ_invocation_index", BuiltinValue::kUndefined},
{"local_2nvocationOinIIPUx", BuiltinValue::kUndefined},
{"local_invocatiZZn_index", BuiltinValue::kUndefined},
- {"num_workgroupnn", BuiltinValue::kUndefined},
- {"nu22_ZZHrkgkkoups", BuiltinValue::kUndefined},
- {"num_workgrops", BuiltinValue::kUndefined},
- {"osiiRn", BuiltinValue::kUndefined},
- {"posqqt99on", BuiltinValue::kUndefined},
- {"posit77on", BuiltinValue::kUndefined},
- {"prmitive3inex", BuiltinValue::kUndefined},
- {"primitiveindeccu", BuiltinValue::kUndefined},
- {"priRRitive_i1ex", BuiltinValue::kUndefined},
- {"sllple_inJJex", BuiltinValue::kUndefined},
- {"sImplMM_ix", BuiltinValue::kUndefined},
- {"s66mpleinTex", BuiltinValue::kUndefined},
- {"sQQmpleJmask", BuiltinValue::kUndefined},
- {"suumpemask", BuiltinValue::kUndefined},
- {"sampl_mak", BuiltinValue::kUndefined},
- {"ygqoup_i33", BuiltinValue::kUndefined},
- {"subgrouxx_id", BuiltinValue::kUndefined},
- {"subgrrN_d", BuiltinValue::kUndefined},
- {"su99group_nvocation_id", BuiltinValue::kUndefined},
- {"subgoup_invocation_id", BuiltinValue::kUndefined},
- {"subgloup_inHocaton_id", BuiltinValue::kUndefined},
- {"sug_oup_sie", BuiltinValue::kUndefined},
- {"subgroup_sze", BuiltinValue::kUndefined},
- {"jbgroup_size", BuiltinValue::kUndefined},
- {"EEetttmmx_index", BuiltinValue::kUndefined},
- {"verte_ndex", BuiltinValue::kUndefined},
- {"vertex_irrdx", BuiltinValue::kUndefined},
- {"workgrxupid", BuiltinValue::kUndefined},
- {"zzorkgou_id", BuiltinValue::kUndefined},
- {"workgroup_ed", BuiltinValue::kUndefined},
+ {"num_subgroupnn", BuiltinValue::kUndefined},
+ {"kkZ22m_subgrHups", BuiltinValue::kUndefined},
+ {"num_subgrous", BuiltinValue::kUndefined},
+ {"nu_workrouRs", BuiltinValue::kUndefined},
+ {"num_w99rkqqroups", BuiltinValue::kUndefined},
+ {"num_workgro77ps", BuiltinValue::kUndefined},
+ {"posi3o", BuiltinValue::kUndefined},
+ {"posiuuccn", BuiltinValue::kUndefined},
+ {"osiRR1on", BuiltinValue::kUndefined},
+ {"primitive_indeJJ", BuiltinValue::kUndefined},
+ {"rimitie_MMIex", BuiltinValue::kUndefined},
+ {"p66Tmitve7index", BuiltinValue::kUndefined},
+ {"QQampJe_index", BuiltinValue::kUndefined},
+ {"sAuue_index", BuiltinValue::kUndefined},
+ {"ampe_index", BuiltinValue::kUndefined},
+ {"ypqe_mas33", BuiltinValue::kUndefined},
+ {"sample_xxask", BuiltinValue::kUndefined},
+ {"samprrNak", BuiltinValue::kUndefined},
+ {"su99roup_id", BuiltinValue::kUndefined},
+ {"subgoup_id", BuiltinValue::kUndefined},
+ {"slbgrup_Hd", BuiltinValue::kUndefined},
+ {"subgrp_invocation__d", BuiltinValue::kUndefined},
+ {"subgroup_invocaion_id", BuiltinValue::kUndefined},
+ {"sujgroup_nvocation_id", BuiltinValue::kUndefined},
+ {"suttgrommp_sizEE", BuiltinValue::kUndefined},
+ {"ubgoup_size", BuiltinValue::kUndefined},
+ {"subgrourr_siz", BuiltinValue::kUndefined},
+ {"vertexxinex", BuiltinValue::kUndefined},
+ {"zzerte_idex", BuiltinValue::kUndefined},
+ {"vertex_ndex", BuiltinValue::kUndefined},
+ {"worppgroup_id", BuiltinValue::kUndefined},
+ {"wWrTTgrup_iuu", BuiltinValue::kUndefined},
+ {"workgrut0_iTT", BuiltinValue::kUndefined},
};
using BuiltinValueParseTest = testing::TestWithParam<BuiltinValueCase>;
@@ -1084,60 +1088,60 @@
};
static constexpr AttributeCase kInvalidAttributeCases[] = {
- {"uigZp", Attribute::kUndefined},
- {"00ui7TT", Attribute::kUndefined},
- {"vvJJ", Attribute::kUndefined},
- {"biQding", Attribute::kUndefined},
- {"bCdRng", Attribute::kUndefined},
+ {"vvn", Attribute::kUndefined},
+ {"aliQn", Attribute::kUndefined},
+ {"al3R", Attribute::kUndefined},
{"iCi", Attribute::kUndefined},
- {"blnPPp_srr", Attribute::kUndefined},
- {"xx8DDuen_src", Attribute::kUndefined},
+ {"bidiPPpr", Attribute::kUndefined},
+ {"buDDndixxg", Attribute::kUndefined},
{"lldqqendYYsrc", Attribute::kUndefined},
- {"uiFFti__", Attribute::kUndefined},
- {"rrGGNNtin", Attribute::kUndefined},
+ {"bFF__nd_sr", Attribute::kUndefined},
+ {"rrleNNGd_sr", Attribute::kUndefined},
{"buiMlin", Attribute::kUndefined},
- {"lo", Attribute::kUndefined},
- {"xllor", Attribute::kUndefined},
+ {"utin", Attribute::kUndefined},
+ {"blilqix", Attribute::kUndefined},
{"olor", Attribute::kUndefined},
- {"camp44e", Attribute::kUndefined},
- {"WWGmpute", Attribute::kUndefined},
+ {"ao44r", Attribute::kUndefined},
+ {"coloWW", Attribute::kUndefined},
{"cjjmpue", Attribute::kUndefined},
- {"djjagnostic", Attribute::kUndefined},
- {"diagnoj1c", Attribute::kUndefined},
+ {"cojjpute", Attribute::kUndefined},
+ {"co1jte", Attribute::kUndefined},
{"diagnotic", Attribute::kUndefined},
- {"99ragment", Attribute::kUndefined},
- {"fVVagmeny", Attribute::kUndefined},
+ {"diag99ostic", Attribute::kUndefined},
+ {"dyagnosVVic", Attribute::kUndefined},
{"frxmeZZt", Attribute::kUndefined},
- {"g33vvp", Attribute::kUndefined},
- {"gros9", Attribute::kUndefined},
+ {"fravvme33", Attribute::kUndefined},
+ {"fsa9ment", Attribute::kUndefined},
{"grFu", Attribute::kUndefined},
- {"ue", Attribute::kUndefined},
- {"Z", Attribute::kUndefined},
- {"i", Attribute::kUndefined},
- {"inpMt_attFchmnt_index", Attribute::kUndefined},
- {"inputWWatZZachment66index", Attribute::kUndefined},
+ {"ueoup", Attribute::kUndefined},
+ {"rIKK", Attribute::kUndefined},
+ {"J", Attribute::kUndefined},
+ {"CSSXd", Attribute::kUndefined},
+ {"JWWZZ", Attribute::kUndefined},
{"inpt_attacme5t_indmdx", Attribute::kUndefined},
- {"BBnterpBlaUe", Attribute::kUndefined},
- {"inter0olatJ11", Attribute::kUndefined},
+ {"BBnput_atacBBmeUt_index", Attribute::kUndefined},
+ {"input0attaJJhme11t_index", Attribute::kUndefined},
{"intfrpottate", Attribute::kUndefined},
- {"inXaittn", Attribute::kUndefined},
- {"inLwriant", Attribute::kUndefined},
+ {"intrttolte", Attribute::kUndefined},
+ {"interpLlwte", Attribute::kUndefined},
{"in1ariant", Attribute::kUndefined},
- {"lowwation", Attribute::kUndefined},
- {"latien", Attribute::kUndefined},
+ {"invarianww", Attribute::kUndefined},
+ {"ineeiat", Attribute::kUndefined},
{"loction", Attribute::kUndefined},
- {"muNNt_use", Attribute::kUndefined},
- {"mustUlRRs__", Attribute::kUndefined},
+ {"loNNation", Attribute::kUndefined},
+ {"locaUlRRo__", Attribute::kUndefined},
{"mHst_use", Attribute::kUndefined},
- {"srCCe", Attribute::kUndefined},
- {"ize", Attribute::kUndefined},
+ {"muCCr_use", Attribute::kUndefined},
+ {"ust_use", Attribute::kUndefined},
{"dzIp", Attribute::kUndefined},
- {"vetex", Attribute::kUndefined},
- {"LNtx", Attribute::kUndefined},
+ {"ize", Attribute::kUndefined},
+ {"LN", Attribute::kUndefined},
{"r", Attribute::kUndefined},
- {"wxxrkgqqoup_GizRR", Attribute::kUndefined},
- {"workSroup_siGGe", Attribute::kUndefined},
+ {"vxxGteqqR", Attribute::kUndefined},
+ {"GGerteS", Attribute::kUndefined},
{"oqkccr8up_size", Attribute::kUndefined},
+ {"workgroup_sze", Attribute::kUndefined},
+ {"woppkgroup_sie", Attribute::kUndefined},
};
using AttributeParseTest = testing::TestWithParam<AttributeCase>;
diff --git a/src/tint/lang/core/ir/binary/decode.cc b/src/tint/lang/core/ir/binary/decode.cc
index 7d7e118..7de9940 100644
--- a/src/tint/lang/core/ir/binary/decode.cc
+++ b/src/tint/lang/core/ir/binary/decode.cc
@@ -1714,6 +1714,8 @@
return core::BuiltinValue::kSubgroupInvocationId;
case pb::BuiltinValue::subgroup_size:
return core::BuiltinValue::kSubgroupSize;
+ case pb::BuiltinValue::num_subgroups:
+ return core::BuiltinValue::kNumSubgroups;
case pb::BuiltinValue::vertex_index:
return core::BuiltinValue::kVertexIndex;
case pb::BuiltinValue::workgroup_id:
diff --git a/src/tint/lang/core/ir/binary/encode.cc b/src/tint/lang/core/ir/binary/encode.cc
index 6f6c2e9..88bdbde 100644
--- a/src/tint/lang/core/ir/binary/encode.cc
+++ b/src/tint/lang/core/ir/binary/encode.cc
@@ -1080,6 +1080,8 @@
return pb::BuiltinValue::subgroup_invocation_id;
case core::BuiltinValue::kSubgroupSize:
return pb::BuiltinValue::subgroup_size;
+ case core::BuiltinValue::kNumSubgroups:
+ return pb::BuiltinValue::num_subgroups;
case core::BuiltinValue::kVertexIndex:
return pb::BuiltinValue::vertex_index;
case core::BuiltinValue::kWorkgroupId:
diff --git a/src/tint/lang/core/ir/validator.cc b/src/tint/lang/core/ir/validator.cc
index d6d956b..6ea8d9f 100644
--- a/src/tint/lang/core/ir/validator.cc
+++ b/src/tint/lang/core/ir/validator.cc
@@ -413,6 +413,14 @@
/* type_error */ "local_invocation_index must be an u32",
};
+constexpr BuiltinChecker kNumSubgroupsChecker{
+ /* name */ "num_subgroups",
+ /* stages */ EnumSet<Function::PipelineStage>(Function::PipelineStage::kCompute),
+ /* direction */ IODirection::kInput,
+ /* type_check */ [](const core::type::Type* ty) -> bool { return ty->Is<core::type::U32>(); },
+ /* type_error */ "num_subgroups must be an u32",
+};
+
constexpr BuiltinChecker kNumWorkgroupsChecker{
/* name */ "num_workgroups",
/* stages */ EnumSet<Function::PipelineStage>(Function::PipelineStage::kCompute),
@@ -519,6 +527,8 @@
return kLocalInvocationIdChecker;
case BuiltinValue::kLocalInvocationIndex:
return kLocalInvocationIndexChecker;
+ case BuiltinValue::kNumSubgroups:
+ return kNumSubgroupsChecker;
case BuiltinValue::kNumWorkgroups:
return kNumWorkgroupsChecker;
case BuiltinValue::kSampleIndex:
diff --git a/src/tint/lang/core/ir/validator_builtin_test.cc b/src/tint/lang/core/ir/validator_builtin_test.cc
index 5fbf43a..9804d4b 100644
--- a/src/tint/lang/core/ir/validator_builtin_test.cc
+++ b/src/tint/lang/core/ir/validator_builtin_test.cc
@@ -972,6 +972,53 @@
)")) << res.Failure();
}
+TEST_F(IR_ValidatorTest, Builtin_NumSubgroups_WrongStage) {
+ auto* f = VertexEntryPoint();
+ AddBuiltinParam(f, "num", BuiltinValue::kNumSubgroups, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_THAT(res.Failure().reason,
+ testing::HasSubstr(
+ R"(:1:19 error: num_subgroups must be used in a compute shader entry point
+%f = @vertex func(%num:u32 [@num_subgroups]):vec4<f32> [@position] {
+ ^^^^^^^^
+)")) << res.Failure();
+}
+
+TEST_F(IR_ValidatorTest, Builtin_NumSubgroups_WrongIODirection) {
+ auto* f = ComputeEntryPoint();
+ AddBuiltinReturn(f, "num", BuiltinValue::kNumSubgroups, ty.u32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_THAT(
+ res.Failure().reason,
+ testing::HasSubstr(R"(:1:1 error: num_subgroups must be an input of a shader entry point
+%f = @compute @workgroup_size(1u, 1u, 1u) func():u32 [@num_subgroups] {
+^^
+)")) << res.Failure();
+}
+
+TEST_F(IR_ValidatorTest, Builtin_NumSubgroups_WrongType) {
+ auto* f = ComputeEntryPoint();
+ AddBuiltinParam(f, "num", BuiltinValue::kNumSubgroups, ty.i32());
+
+ b.Append(f->Block(), [&] { b.Unreachable(); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_THAT(res.Failure().reason,
+ testing::HasSubstr(R"(:1:48 error: num_subgroups must be an u32
+%f = @compute @workgroup_size(1u, 1u, 1u) func(%num:i32 [@num_subgroups]):void {
+ ^^^^^^^^
+)")) << res.Failure();
+}
+
TEST_F(IR_ValidatorTest, Bitcast_MissingArg) {
auto* f = b.Function("f", ty.void_());
b.Append(f->Block(), [&] {
diff --git a/src/tint/lang/spirv/writer/printer/printer.cc b/src/tint/lang/spirv/writer/printer/printer.cc
index ad1ea9e..feec193 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -393,6 +393,8 @@
case core::BuiltinValue::kSubgroupSize:
module_.PushCapability(SpvCapabilityGroupNonUniform);
return SpvBuiltInSubgroupSize;
+ case core::BuiltinValue::kNumSubgroups:
+ TINT_UNIMPLEMENTED();
case core::BuiltinValue::kVertexIndex:
return SpvBuiltInVertexIndex;
case core::BuiltinValue::kWorkgroupId:
diff --git a/src/tint/lang/wgsl/reader/parser/variable_attribute_test.cc b/src/tint/lang/wgsl/reader/parser/variable_attribute_test.cc
index 36cdb0c..044ea55 100644
--- a/src/tint/lang/wgsl/reader/parser/variable_attribute_test.cc
+++ b/src/tint/lang/wgsl/reader/parser/variable_attribute_test.cc
@@ -307,7 +307,7 @@
EXPECT_EQ(attr.value, nullptr);
EXPECT_TRUE(p->has_error());
EXPECT_EQ(p->error(), R"(1:9: expected builtin value name
-Possible values: 'barycentric_coord', 'clip_distances', 'frag_depth', 'front_facing', 'global_invocation_id', 'instance_index', 'local_invocation_id', 'local_invocation_index', 'num_workgroups', 'position', 'primitive_index', 'sample_index', 'sample_mask', 'subgroup_id', 'subgroup_invocation_id', 'subgroup_size', 'vertex_index', 'workgroup_id')");
+Possible values: 'barycentric_coord', 'clip_distances', 'frag_depth', 'front_facing', 'global_invocation_id', 'instance_index', 'local_invocation_id', 'local_invocation_index', 'num_subgroups', 'num_workgroups', 'position', 'primitive_index', 'sample_index', 'sample_mask', 'subgroup_id', 'subgroup_invocation_id', 'subgroup_size', 'vertex_index', 'workgroup_id')");
}
TEST_F(WGSLParserTest, Attribute_Builtin_MisspelledValue) {
@@ -319,7 +319,7 @@
EXPECT_TRUE(p->has_error());
EXPECT_EQ(p->error(), R"(1:9: expected builtin value name
Did you mean 'position'?
-Possible values: 'barycentric_coord', 'clip_distances', 'frag_depth', 'front_facing', 'global_invocation_id', 'instance_index', 'local_invocation_id', 'local_invocation_index', 'num_workgroups', 'position', 'primitive_index', 'sample_index', 'sample_mask', 'subgroup_id', 'subgroup_invocation_id', 'subgroup_size', 'vertex_index', 'workgroup_id')");
+Possible values: 'barycentric_coord', 'clip_distances', 'frag_depth', 'front_facing', 'global_invocation_id', 'instance_index', 'local_invocation_id', 'local_invocation_index', 'num_subgroups', 'num_workgroups', 'position', 'primitive_index', 'sample_index', 'sample_mask', 'subgroup_id', 'subgroup_invocation_id', 'subgroup_size', 'vertex_index', 'workgroup_id')");
}
TEST_F(WGSLParserTest, Attribute_Interpolate_Flat) {
diff --git a/src/tint/lang/wgsl/resolver/subgroups_extension_test.cc b/src/tint/lang/wgsl/resolver/subgroups_extension_test.cc
index 17de108..16ead1d 100644
--- a/src/tint/lang/wgsl/resolver/subgroups_extension_test.cc
+++ b/src/tint/lang/wgsl/resolver/subgroups_extension_test.cc
@@ -274,6 +274,84 @@
"1:2 error: '@builtin(subgroup_id)' is only valid as a compute shader input");
}
+TEST_F(ResolverSubgroupsExtensionTest, UseNumSubgroupsAttribWithoutExtensionError) {
+ Structure("Inputs",
+ Vector{
+ Member("a", ty.u32(), Vector{Builtin(core::BuiltinValue::kNumSubgroups)}),
+ });
+
+ EXPECT_FALSE(r()->Resolve());
+ EXPECT_EQ(r()->error(),
+ "error: use of '@builtin(num_subgroups)' attribute requires enabling extension "
+ "'subgroups'");
+}
+
+TEST_F(ResolverSubgroupsExtensionTest, UseNumSubgroupsAttribWithoutLanguageExtension) {
+ Enable(wgsl::Extension::kSubgroups);
+ Structure("Inputs",
+ Vector{
+ Member("a", ty.u32(), Vector{Builtin(core::BuiltinValue::kNumSubgroups)}),
+ });
+
+ wgsl::AllowedFeatures allowed_features{};
+ allowed_features.extensions.insert(wgsl::Extension::kSubgroups);
+ Resolver resolver{this, allowed_features};
+ EXPECT_FALSE(resolver.Resolve());
+ EXPECT_EQ(resolver.error(),
+ "error: use of '@builtin(num_subgroups)' attribute requires the 'subgroup_id' "
+ "language feature");
+}
+
+TEST_F(ResolverSubgroupsExtensionTest, UseNumSubgroupsAttribWithExtension) {
+ Enable(wgsl::Extension::kSubgroups);
+ Structure("Inputs",
+ Vector{
+ Member("a", ty.u32(), Vector{Builtin(core::BuiltinValue::kNumSubgroups)}),
+ });
+
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+}
+
+TEST_F(ResolverSubgroupsExtensionTest, NumSubgroupsI32Error) {
+ Enable(wgsl::Extension::kSubgroups);
+ Structure("Inputs",
+ Vector{
+ Member("a", ty.i32(), Vector{Builtin(core::BuiltinValue::kNumSubgroups)}),
+ });
+
+ EXPECT_FALSE(r()->Resolve());
+ EXPECT_EQ(r()->error(), "error: store type of '@builtin(num_subgroups)' must be 'u32'");
+}
+
+TEST_F(ResolverSubgroupsExtensionTest, NumSubgroupsFragmentShader) {
+ Enable(wgsl::Extension::kSubgroups);
+ Func("main",
+ Vector{Param("size", ty.u32(), Vector{Builtin(core::BuiltinValue::kNumSubgroups)})},
+ ty.void_(), Empty, Vector{Stage(ast::PipelineStage::kFragment)});
+
+ EXPECT_FALSE(r()->Resolve());
+ EXPECT_EQ(r()->error(),
+ "error: '@builtin(num_subgroups)' is only valid as a compute shader input");
+}
+
+TEST_F(ResolverSubgroupsExtensionTest, NumSubgroupsComputeShaderOutput) {
+ Enable(wgsl::Extension::kSubgroups);
+
+ Func("main", tint::Empty, ty.u32(),
+ Vector{
+ Return(Call<u32>()),
+ },
+ Vector{
+ Stage(ast::PipelineStage::kCompute),
+ WorkgroupSize(1_i),
+ },
+ Vector{Builtin(Source{{1, 2}}, core::BuiltinValue::kNumSubgroups)});
+
+ EXPECT_FALSE(r()->Resolve());
+ EXPECT_EQ(r()->error(),
+ "1:2 error: '@builtin(num_subgroups)' is only valid as a compute shader input");
+}
+
// Using the subgroup_uniformity diagnostic rule without subgroups enabled should succeed.
TEST_F(ResolverSubgroupsExtensionTest, UseSubgroupUniformityRuleWithoutExtensionError) {
DiagnosticDirective(wgsl::DiagnosticSeverity::kOff, "subgroup_uniformity");
diff --git a/src/tint/lang/wgsl/resolver/validator.cc b/src/tint/lang/wgsl/resolver/validator.cc
index b574b29..e6a3809 100644
--- a/src/tint/lang/wgsl/resolver/validator.cc
+++ b/src/tint/lang/wgsl/resolver/validator.cc
@@ -1125,6 +1125,34 @@
return false;
}
break;
+ case core::BuiltinValue::kNumSubgroups:
+ if (!enabled_extensions_.Contains(wgsl::Extension::kSubgroups)) {
+ AddError(attr->source)
+ << "use of " << style::Attribute("@builtin")
+ << style::Code("(", style::Enum(builtin), ")")
+ << " attribute requires enabling extension " << style::Code("subgroups");
+ return false;
+ }
+ // TODO(crbug.com/454654105): Remove this check.
+ if (!allowed_features_.features.contains(wgsl::LanguageFeature::kSubgroupId)) {
+ AddError(attr->source)
+ << "use of " << style::Attribute("@builtin")
+ << style::Code("(", style::Enum(builtin), ")") << " attribute requires the "
+ << style::Code("subgroup_id") << " language feature";
+ return false;
+ }
+ if (!type->Is<core::type::U32>()) {
+ err_builtin_type("u32");
+ return false;
+ }
+ if (stage != ast::PipelineStage::kNone &&
+ !(stage == ast::PipelineStage::kCompute && is_input)) {
+ AddError(attr->source)
+ << style::Attribute("@builtin") << style::Code("(", style::Enum(builtin), ")")
+ << " is only valid as a compute shader input";
+ return false;
+ }
+ break;
case core::BuiltinValue::kSubgroupInvocationId:
case core::BuiltinValue::kSubgroupSize:
if (!enabled_extensions_.Contains(wgsl::Extension::kSubgroups)) {
diff --git a/src/tint/utils/protos/ir/ir.proto b/src/tint/utils/protos/ir/ir.proto
index 0f4c772..fe63070 100644
--- a/src/tint/utils/protos/ir/ir.proto
+++ b/src/tint/utils/protos/ir/ir.proto
@@ -584,6 +584,7 @@
subgroup_id = 17;
primitive_index = 18;
barycentric_coord = 19;
+ num_subgroups = 20;
}
enum BuiltinFn {
