Bug 199895 – [WHLSL] MTLComputePipelineState creation fails with "Compiler encountered an internal error." message

• Status: RESOLVED
• Resolution: MOVED
• Component: WebGPU

[Justin Fan]

The following WHLSL code fails to compile:

  [numthreads(1, 1, 1)]
  compute void sort_0(device uint[] numbers : register(u0),
                      float3 localID        : SV_GroupThreadID) 
  {
      device uint[1] data;

      uint localIndex = uint(localID.x);

      data[localIndex] = numbers[0];
  }

It makes it all the way to MTLComputePipelineState creation; the Metal creation call then fails with the error "Compiler encountered an internal error."

[Justin Fan]

The MSL generated output is below:

#include <metal_stdlib>
#include <metal_atomic>
#include <metal_math>
#include <metal_relational>
#include <metal_compute>
#include <metal_texture>

using namespace metal;

struct type90;

typedef float2 type49;
typedef type49 type48;
typedef array<half, 9> type63;
typedef type63 type62;
typedef array<float, 6> type75;
typedef type75 type74;
typedef short4 type35;
typedef type35 type34;
typedef bool3 type3;
typedef type3 type2;
typedef int3 type39;
typedef type39 type38;
typedef half2 type43;
typedef type43 type42;
typedef array<float, 12> type83;
typedef type83 type82;
typedef half4 type47;
typedef type47 type46;
typedef float4 type53;
typedef type53 type52;
typedef uchar3 type9;
typedef type9 type8;
typedef uchar2 type7;
typedef type7 type6;
typedef array<half, 8> type59;
typedef type59 type58;
typedef array<float, 9> type81;
typedef type81 type80;
typedef bool4 type5;
typedef type5 type4;
typedef char2 type25;
typedef type25 type24;
typedef array<half, 12> type69;
typedef type69 type68;
typedef uint4 type23;
typedef type23 type22;
typedef array<float, 12> type87;
typedef type87 type86;
typedef array<half, 16> type71;
typedef type71 type70;
typedef int2 type37;
typedef type37 type36;
typedef array<float, 6> type79;
typedef type79 type78;
typedef ushort2 type13;
typedef type13 type12;
typedef bool2 type1;
typedef type1 type0;
typedef half3 type45;
typedef type45 type44;
typedef float3 type51;
typedef type51 type50;
typedef type50 type91;
typedef type34 type92;
typedef type46 type93;
typedef ushort4 type17;
typedef type17 type16;
typedef type16 type94;
typedef type6 type95;
typedef type8 type96;
typedef type86 type97;
typedef type62 type98;
typedef array<half, 4> type55;
typedef type55 type54;
typedef type54 type99;
typedef char4 type29;
typedef type29 type28;
typedef type28 type100;
typedef type52 type101;
typedef type4 type102;
typedef uchar4 type11;
typedef type11 type10;
typedef type10 type103;
typedef type44 type104;
typedef type38 type105;
typedef uint3 type21;
typedef type21 type20;
typedef type20 type106;
typedef short2 type31;
typedef type31 type30;
typedef type30 type107;
typedef char3 type27;
typedef type27 type26;
typedef type26 type108;
typedef type0 type109;
typedef type24 type110;
typedef type36 type111;
typedef short3 type33;
typedef type33 type32;
typedef type32 type112;
typedef int4 type41;
typedef type41 type40;
typedef type40 type113;
typedef ushort3 type15;
typedef type15 type14;
typedef type14 type114;
typedef uint2 type19;
typedef type19 type18;
typedef type18 type115;
typedef type78 type116;
typedef type12 type117;
typedef type48 type118;
typedef type22 type119;
typedef type42 type120;
typedef type2 type121;
typedef type82 type122;
typedef type80 type123;
typedef type74 type124;
typedef array<float, 8> type85;
typedef type85 type84;
typedef type84 type125;
typedef type68 type126;
typedef array<half, 8> type67;
typedef type67 type66;
typedef type66 type127;
typedef type70 type128;
typedef array<float, 4> type73;
typedef type73 type72;
typedef type72 type129;
typedef array<half, 6> type57;
typedef type57 type56;
typedef type56 type130;
typedef array<float, 16> type89;
typedef type89 type88;
typedef type88 type131;
typedef array<float, 8> type77;
typedef type77 type76;
typedef type76 type132;
typedef array<half, 12> type65;
typedef type65 type64;
typedef type64 type133;
typedef array<half, 6> type61;
typedef type61 type60;
typedef type60 type134;
typedef type58 type135;
typedef uint32_t type136;
typedef array<type136, 1> type137;
struct type90 {
    type91 structureElement0;
    type92 structureElement1;
    type93 structureElement2;
    type94 structureElement3;
    type95 structureElement4;
    type96 structureElement5;
    type97 structureElement6;
    type98 structureElement7;
    type99 structureElement8;
    type100 structureElement9;
    type101 structureElement10;
    type102 structureElement11;
    type94 structureElement12;
    type94 structureElement13;
    type103 structureElement14;
    type104 structureElement15;
    type94 structureElement16;
    type105 structureElement17;
    type106 structureElement18;
    type107 structureElement19;
    type108 structureElement20;
    type95 structureElement21;
    type93 structureElement22;
    type106 structureElement23;
    type109 structureElement24;
    type92 structureElement25;
    type110 structureElement26;
    type111 structureElement27;
    type101 structureElement28;
    type102 structureElement29;
    type104 structureElement30;
    type109 structureElement31;
    type112 structureElement32;
    type113 structureElement33;
    type106 structureElement34;
    type105 structureElement35;
    type114 structureElement36;
    type107 structureElement37;
    type102 structureElement38;
    type100 structureElement39;
    type102 structureElement40;
    type96 structureElement41;
    type91 structureElement42;
    type115 structureElement43;
    type110 structureElement44;
    type116 structureElement45;
    type94 structureElement46;
    type113 structureElement47;
    type117 structureElement48;
    type118 structureElement49;
    type102 structureElement50;
    type119 structureElement51;
    type119 structureElement52;
    type119 structureElement53;
    type104 structureElement54;
    type120 structureElement55;
    type93 structureElement56;
    type101 structureElement57;
    type91 structureElement58;
    type113 structureElement59;
    type92 structureElement60;
    type101 structureElement61;
    type113 structureElement62;
    type105 structureElement63;
    type119 structureElement64;
    type100 structureElement65;
    type103 structureElement66;
    type91 structureElement67;
    type93 structureElement68;
    type121 structureElement69;
    type100 structureElement70;
    type113 structureElement71;
    type92 structureElement72;
    type96 structureElement73;
    type112 structureElement74;
    type105 structureElement75;
    type92 structureElement76;
    type103 structureElement77;
    type122 structureElement78;
    type119 structureElement79;
    type123 structureElement80;
    type120 structureElement81;
    type103 structureElement82;
    type101 structureElement83;
    type103 structureElement84;
    type101 structureElement85;
    type113 structureElement86;
    type113 structureElement87;
    type119 structureElement88;
    type124 structureElement89;
    type121 structureElement90;
    type125 structureElement91;
    type108 structureElement92;
    type100 structureElement93;
    type102 structureElement94;
    type114 structureElement95;
    type126 structureElement96;
    type108 structureElement97;
    type127 structureElement98;
    type101 structureElement99;
    type93 structureElement100;
    type92 structureElement101;
    type128 structureElement102;
    type118 structureElement103;
    type94 structureElement104;
    type129 structureElement105;
    type93 structureElement106;
    type112 structureElement107;
    type119 structureElement108;
    type130 structureElement109;
    type92 structureElement110;
    type93 structureElement111;
    type103 structureElement112;
    type112 structureElement113;
    type113 structureElement114;
    type103 structureElement115;
    type117 structureElement116;
    type94 structureElement117;
    type92 structureElement118;
    type111 structureElement119;
    type106 structureElement120;
    type131 structureElement121;
    type132 structureElement122;
    type133 structureElement123;
    type103 structureElement124;
    type108 structureElement125;
    type102 structureElement126;
    type100 structureElement127;
    type102 structureElement128;
    type115 structureElement129;
    type101 structureElement130;
    type134 structureElement131;
    type135 structureElement132;
    type137 structureElement133;
    type93 structureElement134;
    type96 structureElement135;
    type100 structureElement136;
    type94 structureElement137;
    type114 structureElement138;
    type100 structureElement139;
    type121 structureElement140;
    type119 structureElement141;
    type114 structureElement142;
    type104 structureElement143;
    type121 structureElement144;
};
typedef thread type91* type178;
typedef thread type92* type170;
typedef thread type93* type176;
typedef thread type94* type161;
typedef thread type95* type156;
typedef thread type96* type157;
typedef thread type97* type196;
typedef thread type98* type184;
typedef thread type99* type180;
typedef thread type100* type167;
typedef thread type101* type179;
typedef thread type102* type155;
typedef thread type103* type158;
typedef thread type104* type175;
typedef thread type105* type172;
typedef thread type106* type163;
typedef thread type107* type168;
typedef thread type108* type166;
typedef thread type109* type153;
typedef thread type110* type165;
typedef thread type111* type171;
typedef thread type112* type169;
typedef thread type113* type173;
typedef thread type114* type160;
typedef thread type115* type162;
typedef thread type116* type192;
typedef thread type117* type159;
typedef thread type118* type177;
typedef thread type119* type164;
typedef thread type120* type174;
typedef thread type121* type154;
typedef thread type122* type194;
typedef thread type123* type193;
typedef thread type124* type190;
typedef thread type125* type195;
typedef thread type126* type187;
typedef thread type127* type186;
typedef thread type128* type188;
typedef thread type129* type189;
typedef thread type130* type181;
typedef thread type131* type197;
typedef thread type132* type191;
typedef thread type133* type185;
typedef thread type134* type183;
typedef thread type135* type182;
struct type139{
    device type136* pointer;
    uint32_t length;
};
typedef thread type136* type143;
struct type144{
    thread type136* pointer;
    uint32_t length;
};
typedef device type136* type145;
typedef void type138;
typedef type90 type140;
typedef thread type140* type141;
typedef float type142;
typedef int32_t type146;
typedef bool type147;
typedef uint8_t type148;
typedef uint16_t type149;
typedef int8_t type150;
typedef int16_t type151;
typedef half type152;
typedef sampler type198;
typedef texture1d<ushort, access::sample> type199;
typedef texture1d<ushort, access::sample> type200;
typedef texture1d<ushort, access::sample> type201;
typedef texture1d<ushort, access::sample> type202;
typedef texture1d<uint, access::sample> type203;
typedef texture1d<uint, access::sample> type204;
typedef texture1d<uint, access::sample> type205;
typedef texture1d<uint, access::sample> type206;
typedef texture1d<short, access::sample> type207;
typedef texture1d<short, access::sample> type208;
typedef texture1d<short, access::sample> type209;
typedef texture1d<short, access::sample> type210;
typedef texture1d<int, access::sample> type211;
typedef texture1d<int, access::sample> type212;
typedef texture1d<int, access::sample> type213;
typedef texture1d<int, access::sample> type214;
typedef texture1d<half, access::sample> type215;
typedef texture1d<half, access::sample> type216;
typedef texture1d<half, access::sample> type217;
typedef texture1d<half, access::sample> type218;
typedef texture1d<float, access::sample> type219;
typedef texture1d<float, access::sample> type220;
typedef texture1d<float, access::sample> type221;
typedef texture1d<float, access::sample> type222;
typedef texture1d_array<ushort, access::sample> type223;
typedef texture1d_array<ushort, access::sample> type224;
typedef texture1d_array<ushort, access::sample> type225;
typedef texture1d_array<ushort, access::sample> type226;
typedef texture1d_array<uint, access::sample> type227;
typedef texture1d_array<uint, access::sample> type228;
typedef texture1d_array<uint, access::sample> type229;
typedef texture1d_array<uint, access::sample> type230;
typedef texture1d_array<short, access::sample> type231;
typedef texture1d_array<short, access::sample> type232;
typedef texture1d_array<short, access::sample> type233;
typedef texture1d_array<short, access::sample> type234;
typedef texture1d_array<int, access::sample> type235;
typedef texture1d_array<int, access::sample> type236;
typedef texture1d_array<int, access::sample> type237;
typedef texture1d_array<int, access::sample> type238;
typedef texture1d_array<half, access::sample> type239;
typedef texture1d_array<half, access::sample> type240;
typedef texture1d_array<half, access::sample> type241;
typedef texture1d_array<half, access::sample> type242;
typedef texture1d_array<float, access::sample> type243;
typedef texture1d_array<float, access::sample> type244;
typedef texture1d_array<float, access::sample> type245;
typedef texture1d_array<float, access::sample> type246;
typedef texture2d<ushort, access::sample> type247;
typedef texture2d<ushort, access::sample> type248;
typedef texture2d<ushort, access::sample> type249;
typedef texture2d<ushort, access::sample> type250;
typedef texture2d<uint, access::sample> type251;
typedef texture2d<uint, access::sample> type252;
typedef texture2d<uint, access::sample> type253;
typedef texture2d<uint, access::sample> type254;
typedef texture2d<short, access::sample> type255;
typedef texture2d<short, access::sample> type256;
typedef texture2d<short, access::sample> type257;
typedef texture2d<short, access::sample> type258;
typedef texture2d<int, access::sample> type259;
typedef texture2d<int, access::sample> type260;
typedef texture2d<int, access::sample> type261;
typedef texture2d<int, access::sample> type262;
typedef texture2d<half, access::sample> type263;
typedef texture2d<half, access::sample> type264;
typedef texture2d<half, access::sample> type265;
typedef texture2d<half, access::sample> type266;
typedef texture2d<float, access::sample> type267;
typedef texture2d<float, access::sample> type268;
typedef texture2d<float, access::sample> type269;
typedef texture2d<float, access::sample> type270;
typedef texture2d_array<ushort, access::sample> type271;
typedef texture2d_array<ushort, access::sample> type272;
typedef texture2d_array<ushort, access::sample> type273;
typedef texture2d_array<ushort, access::sample> type274;
typedef texture2d_array<uint, access::sample> type275;
typedef texture2d_array<uint, access::sample> type276;
typedef texture2d_array<uint, access::sample> type277;
typedef texture2d_array<uint, access::sample> type278;
typedef texture2d_array<short, access::sample> type279;
typedef texture2d_array<short, access::sample> type280;
typedef texture2d_array<short, access::sample> type281;
typedef texture2d_array<short, access::sample> type282;
typedef texture2d_array<int, access::sample> type283;
typedef texture2d_array<int, access::sample> type284;
typedef texture2d_array<int, access::sample> type285;
typedef texture2d_array<int, access::sample> type286;
typedef texture2d_array<half, access::sample> type287;
typedef texture2d_array<half, access::sample> type288;
typedef texture2d_array<half, access::sample> type289;
typedef texture2d_array<half, access::sample> type290;
typedef texture2d_array<float, access::sample> type291;
typedef texture2d_array<float, access::sample> type292;
typedef texture2d_array<float, access::sample> type293;
typedef texture2d_array<float, access::sample> type294;
typedef texture3d<ushort, access::sample> type295;
typedef texture3d<ushort, access::sample> type296;
typedef texture3d<ushort, access::sample> type297;
typedef texture3d<ushort, access::sample> type298;
typedef texture3d<uint, access::sample> type299;
typedef texture3d<uint, access::sample> type300;
typedef texture3d<uint, access::sample> type301;
typedef texture3d<uint, access::sample> type302;
typedef texture3d<short, access::sample> type303;
typedef texture3d<short, access::sample> type304;
typedef texture3d<short, access::sample> type305;
typedef texture3d<short, access::sample> type306;
typedef texture3d<int, access::sample> type307;
typedef texture3d<int, access::sample> type308;
typedef texture3d<int, access::sample> type309;
typedef texture3d<int, access::sample> type310;
typedef texture3d<half, access::sample> type311;
typedef texture3d<half, access::sample> type312;
typedef texture3d<half, access::sample> type313;
typedef texture3d<half, access::sample> type314;
typedef texture3d<float, access::sample> type315;
typedef texture3d<float, access::sample> type316;
typedef texture3d<float, access::sample> type317;
typedef texture3d<float, access::sample> type318;
typedef texturecube<ushort, access::sample> type319;
typedef texturecube<ushort, access::sample> type320;
typedef texturecube<ushort, access::sample> type321;
typedef texturecube<ushort, access::sample> type322;
typedef texturecube<uint, access::sample> type323;
typedef texturecube<uint, access::sample> type324;
typedef texturecube<uint, access::sample> type325;
typedef texturecube<uint, access::sample> type326;
typedef texturecube<short, access::sample> type327;
typedef texturecube<short, access::sample> type328;
typedef texturecube<short, access::sample> type329;
typedef texturecube<short, access::sample> type330;
typedef texturecube<int, access::sample> type331;
typedef texturecube<int, access::sample> type332;
typedef texturecube<int, access::sample> type333;
typedef texturecube<int, access::sample> type334;
typedef texturecube<half, access::sample> type335;
typedef texturecube<half, access::sample> type336;
typedef texturecube<half, access::sample> type337;
typedef texturecube<half, access::sample> type338;
typedef texturecube<float, access::sample> type339;
typedef texturecube<float, access::sample> type340;
typedef texturecube<float, access::sample> type341;
typedef texturecube<float, access::sample> type342;
typedef depth2d<float, access::sample> type343;
typedef depth2d_array<float, access::sample> type344;
typedef depthcube<float, access::sample> type345;
type142 function43(type91);
type136 function110(type142);
type137 function260();
type143 function519(type144, type136);
type145 function520(type139, type136);

template <typename T>
inline void memsetZero(thread T& value)
{
    thread char* ptr = static_cast<thread char*>(static_cast<thread void*>(&value));
    for (size_t i = 0; i < sizeof(T); ++i)
        ptr[i] = 0;
}
type142 function43(type91 v) {
    return v.x;
}
type136 function110(type142 x) {
    return static_cast<type136>(x);
}
type137 function260() {
    type137 x;
    memsetZero(x);
    return x;
}
type143 function519(type144 v, type136 n) {
    if (n < v.length) return &(v.pointer[n]);
    return nullptr;
}
type145 function520(type139 v, type136 n) {
    if (n < v.length) return &(v.pointer[n]);
    return nullptr;
}
struct type346 {
    device type136* structureElement145 [[id(0)]];
    uint2 structureElement146 [[id(1)]];
};


kernel void function521(device type346& variable0 [[buffer(0)]], uint3 variable2 [[thread_position_in_grid]], uint3 variable3 [[thread_position_in_threadgroup]], uint3 variable4 [[threadgroup_position_in_grid]]) {
type139 variable5;
type91 variable6;
type91 variable7;
type91 variable8;
size_t variable1 = variable0.structureElement146.y;
variable1 = variable1 << 32;
variable1 = variable1 | variable0.structureElement146.x;
variable1 = variable1 / sizeof(type136);
if (variable1 > 0xFFFFFFFF) variable1 = 0xFFFFFFFF;
variable5 = { variable0.structureElement145, static_cast<uint32_t>(variable1) };
variable6 = type51(variable2);
variable7 = type51(variable3);
variable8 = type51(variable4);
{
type140 variable9;
thread type140* variable11 = &variable9;
type141 variable12 = variable11;
type141 variable10 = variable12;
type137 variable14 = function260();
type137 variable13 = variable14;
thread type141* variable17 = &variable10;
thread type137* variable16 = &variable10->structureElement133;
type137 variable15 = *variable16;
thread type137* variable18 = &variable13;
if (variable16) *variable16 = variable13;
thread type91* variable20 = &variable7;
type142 variable21 = function43(variable7);
type136 variable22 = function110(variable21);
type136 variable19 = variable22;
thread type141* variable25 = &variable10;
thread type137* variable24 = &variable10->structureElement133;
type137 variable23 = *variable24;
type144 variable26 = { variable24->data(), 1 };
thread type136* variable27 = &variable19;
type143 variable28 = function519(variable26, variable19);
type143 variable30 = variable28;
type136 variable29;
if (variable30) variable29 = *variable28;
else memsetZero(variable29);
thread type139* variable31 = &variable5;
type146 variable32 = static_cast<type146>(0);
type145 variable33 = function520(variable5, variable32);
type145 variable35 = variable33;
type136 variable34;
if (variable35) variable34 = *variable33;
else memsetZero(variable34);
if (variable30) *variable30 = variable34;
}
}

[Justin Fan]

Problem seems to occur whenever writing to 'data' with a calculated index. Using a literal constant index works fine.

[Justin Fan]

(In reply to Justin Fan from comment #2)
> Problem seems to occur whenever writing to 'data' with a calculated index.
> Using a literal constant index works fine.

Nevermind. In a more complex version of the example shader adding the line "sharedData[0] = 0" also causes this issue.

[Justin Fan]

Even simpler reduction compatible with the test harness:

    void test(device uint[] buffer) {
        device uint[34] cache;
        cache[buffer[0]] = buffer[0];
    }

[Justin Fan]

[numthreads(${threadsPerThreadgroup}, 1, 1)]
compute void horizontal(device uint[] origBuffer : register(u${originalBufferBindingNum}),
                        device uint[] outputBuffer : register(u${outputBufferBindingNum}),
                        float3 groupThreadID : SV_GroupThreadID,
                        float3 dispatchThreadID : SV_DispatchThreadID)
{
    uint localIndex = uint(groupThreadID.x);
    uint2 globalIndex = uint2(uint(dispatchThreadID.x), uint(dispatchThreadID.y));

    threadgroup uint[${cacheSize}] gCache;

    if (localIndex < ${blurRadius}) {
        uint x = uint(max(int(localIndex) - ${blurRadius}, 0));
        gCache[localIndex] = origBuffer[globalIndex.y * ${image.width} + x];
    }

    if (localIndex >= ${nMinusBlurRadius}) {
        uint x = min(globalIndex.x + ${blurRadius}, uint(${image.width} - 1));
        gCache[localIndex] = origBuffer[globalIndex.y * ${image.width} + x];
    }

}

[Myles C. Maxfield]

This is a Metal compiler bug. I filed <rdar://problem/53291128>