[Webkit-unassigned] [Bug 199895] [WHLSL] MTLComputePipelineState creation fails with "Compiler encountered an internal error." message
bugzilla-daemon at webkit.org
bugzilla-daemon at webkit.org
Wed Jul 17 19:57:09 PDT 2019
https://bugs.webkit.org/show_bug.cgi?id=199895
--- Comment #1 from Justin Fan <justin_fan at apple.com> ---
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;
}
}
--
You are receiving this mail because:
You are the assignee for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.webkit.org/pipermail/webkit-unassigned/attachments/20190718/3c3b539f/attachment-0001.html>
More information about the webkit-unassigned
mailing list