199424 2019-07-02 15:29:03 -0700 [WHLSL] Change whlsl-two-dimensional-array.html to not be flaky on AMD Radeon Pro GPUs 2019-07-02 16:54:15 -0700 1 1 1 Unclassified WebKit WebGPU WebKit Nightly Build Unspecified Unspecified RESOLVED FIXED InRadar P2 Normal --- 1 saam saam commit-queue dino fpizlo jonlee justin_fan mmaxfield rmorisset webkit-bug-importer oldest_to_newest 1549739 0 saam 2019-07-02 15:29:03 -0700 ... 1549740 1 saam 2019-07-02 15:29:50 -0700 See <rdar://problem/52545976> The metal reduction is: ``` #include <metal_stdlib> #include <metal_atomic> #include <metal_math> #include <metal_relational> #include <metal_compute> #include <metal_texture> using namespace metal; typedef array<float, 10> BaseArray; typedef array<BaseArray, 5> ParentArray; struct WrapperStruct { ParentArray parentArray; uint outerIndex; BaseArray baseArray; uint index; }; struct InputStruct { device int32_t* ptr [[id(0)]]; uint2 structureElement20 [[id(1)]]; }; kernel void computeShader(device int32_t* ptr) { WrapperStruct wrapperStruct = { 0 }; wrapperStruct.parentArray = { 0 }; wrapperStruct.outerIndex = 0; for (; wrapperStruct.outerIndex < wrapperStruct.parentArray.size(); ++wrapperStruct.outerIndex) { wrapperStruct.baseArray = { 0 }; for (unsigned j = 0; j < wrapperStruct.baseArray.size(); ++j) { wrapperStruct.baseArray[j] = wrapperStruct.outerIndex; } wrapperStruct.parentArray[wrapperStruct.outerIndex] = wrapperStruct.baseArray; } wrapperStruct.outerIndex = 0; for (; wrapperStruct.outerIndex < wrapperStruct.parentArray.size(); wrapperStruct.outerIndex++) { wrapperStruct.baseArray = wrapperStruct.parentArray[wrapperStruct.outerIndex]; wrapperStruct.index = 0; for ( ; wrapperStruct.index < wrapperStruct.baseArray.size(); wrapperStruct.index += 1) { if (wrapperStruct.baseArray[wrapperStruct.index] != wrapperStruct.outerIndex) { *ptr = 42; return; } } } *ptr = 1; } ``` *ptr should be 1, but is 42. 1549752 2 373366 saam 2019-07-02 15:43:28 -0700 Created attachment 373366 patch 1549766 3 373366 mmaxfield 2019-07-02 16:13:58 -0700 Comment on attachment 373366 patch 🤮 1549784 4 373366 commit-queue 2019-07-02 16:53:23 -0700 Comment on attachment 373366 patch Clearing flags on attachment: 373366 Committed r247081: <https://trac.webkit.org/changeset/247081> 1549785 5 commit-queue 2019-07-02 16:53:24 -0700 All reviewed patches have been landed. Closing bug. 1549786 6 webkit-bug-importer 2019-07-02 16:54:15 -0700 <rdar://problem/52556060> 373366 2019-07-02 15:43:28 -0700 2019-07-02 16:53:23 -0700 patch a-backup.diff text/plain 2464 saam SW5kZXg6IExheW91dFRlc3RzL0NoYW5nZUxvZwo9PT09PT09PT09PT09PT09PT09PT09PT09PT09 PT09PT09PT09PT09PT09PT09PT09PT09PT09PT09PT09PT09PT09Ci0tLSBMYXlvdXRUZXN0cy9D aGFuZ2VMb2cJKHJldmlzaW9uIDI0NzA3MykKKysrIExheW91dFRlc3RzL0NoYW5nZUxvZwkod29y a2luZyBjb3B5KQpAQCAtMSwzICsxLDEyIEBACisyMDE5LTA3LTAyICBTYWFtIEJhcmF0aSAgPHNi YXJhdGlAYXBwbGUuY29tPgorCisgICAgICAgIFtXSExTTF0gQ2hhbmdlIHdobHNsLXR3by1kaW1l bnNpb25hbC1hcnJheS5odG1sIHRvIG5vdCBiZSBmbGFreSBvbiBBTUQgUmFkZW9uIFBybyBHUFVz CisgICAgICAgIGh0dHBzOi8vYnVncy53ZWJraXQub3JnL3Nob3dfYnVnLmNnaT9pZD0xOTk0MjQK KworICAgICAgICBSZXZpZXdlZCBieSBOT0JPRFkgKE9PUFMhKS4KKworICAgICAgICAqIHdlYmdw dS93aGxzbC10d28tZGltZW5zaW9uYWwtYXJyYXkuaHRtbDoKKwogMjAxOS0wNy0wMiAgQW5kcmVz IEdvbnphbGV6ICA8YW5kcmVzZ18yMkBhcHBsZS5jb20+CiAKICAgICAgICAgRW5oYW5jZSBzdXBw b3J0IG9mIGFyaWEtaGFzcG9wdXAgcGVyIEFSSUEgMS4xIHNwZWNpZmljYXRpb24uCkluZGV4OiBM YXlvdXRUZXN0cy93ZWJncHUvd2hsc2wtdHdvLWRpbWVuc2lvbmFsLWFycmF5Lmh0bWwKPT09PT09 PT09PT09PT09PT09PT09PT09PT09PT09PT09PT09PT09PT09PT09PT09PT09PT09PT09PT09PT09 PT09PQotLS0gTGF5b3V0VGVzdHMvd2ViZ3B1L3dobHNsLXR3by1kaW1lbnNpb25hbC1hcnJheS5o dG1sCShyZXZpc2lvbiAyNDcwMzQpCisrKyBMYXlvdXRUZXN0cy93ZWJncHUvd2hsc2wtdHdvLWRp bWVuc2lvbmFsLWFycmF5Lmh0bWwJKHdvcmtpbmcgY29weSkKQEAgLTgsMjAgKzgsMjAgQEAKIDxz Y3JpcHQ+CiBjb25zdCBzaGFkZXJTb3VyY2UgPSBgCiB2b2lkIGZpbGwodGhyZWFkIGZsb2F0WzEw XVtdIGFycmF5LCBmbG9hdFsxMF0gdmFsdWUpIHsKLSAgICBmb3IgKHVpbnQgaSA9IDA7IGkgPCBh cnJheS5sZW5ndGg7IGkrKykgeworICAgIGZvciAodWludCBpID0gMDsgaSA8IGFycmF5Lmxlbmd0 aDsgaSA9IGkgKyAxKSB7CiAgICAgICAgIGFycmF5W2ldID0gdmFsdWU7CiAgICAgfQogfQogCiB2 b2lkIGZpbGwodGhyZWFkIGZsb2F0W10gYXJyYXksIGZsb2F0IHZhbHVlKSB7Ci0gICAgZm9yICh1 aW50IGkgPSAwOyBpIDwgYXJyYXkubGVuZ3RoOyBpKyspIHsKKyAgICBmb3IgKHVpbnQgaSA9IDA7 IGkgPCBhcnJheS5sZW5ndGg7IGkgPSBpICsgMSkgewogICAgICAgICBhcnJheVtpXSA9IHZhbHVl OwogICAgIH0KIH0KIAogYm9vbCBjb250YWlucyh0aHJlYWQgZmxvYXRbMTBdW10gYXJyYXksIGZs b2F0IHZhbHVlKSB7Ci0gICAgZm9yICh1aW50IGkgPSAwOyBpIDwgYXJyYXkubGVuZ3RoOyBpKysp IHsKLSAgICAgICAgZm9yICh1aW50IGogPSAwOyBqIDwgYXJyYXlbal0ubGVuZ3RoOyBqKyspIHsK KyAgICBmb3IgKHVpbnQgaSA9IDA7IGkgPCBhcnJheS5sZW5ndGg7IGkgPSBpICsgMSkgeworICAg ICAgICBmb3IgKHVpbnQgaiA9IDA7IGogPCBhcnJheVtqXS5sZW5ndGg7IGogPSBqICsgMSkgewog ICAgICAgICAgICAgaWYgKGFycmF5W2ldW2pdICE9IHZhbHVlKQogICAgICAgICAgICAgICAgIHJl dHVybiBmYWxzZTsKICAgICAgICAgfQpAQCAtMzAsNyArMzAsNyBAQCBib29sIGNvbnRhaW5zKHRo cmVhZCBmbG9hdFsxMF1bXSBhcnJheSwgCiB9CiAKIGJvb2wgY29udGFpbnModGhyZWFkIGZsb2F0 W10gYXJyYXksIGZsb2F0IHZhbHVlKSB7Ci0gICAgZm9yICh1aW50IGkgPSAwOyBpIDwgYXJyYXku bGVuZ3RoOyBpKyspIHsKKyAgICBmb3IgKHVpbnQgaSA9IDA7IGkgPCBhcnJheS5sZW5ndGg7IGkg PSBpICsgMSkgewogICAgICAgICBpZiAoYXJyYXlbaV0gIT0gdmFsdWUpCiAgICAgICAgICAgICBy ZXR1cm4gZmFsc2U7CiAgICAgfQpAQCAtNDUsMTMgKzQ1LDEzIEBAIGNvbXB1dGUgdm9pZCBjb21w dXRlU2hhZGVyKGRldmljZSBpbnRbXSAKICAgICBpZiAoIWNvbnRhaW5zKEBhcnJheSwgMCkpCiAg ICAgICAgIHJldHVybjsKIAotICAgIGZvciAodWludCBpID0gMDsgaSA8IGFycmF5Lmxlbmd0aDsg KytpKSB7CisgICAgZm9yICh1aW50IGkgPSAwOyBpIDwgYXJyYXkubGVuZ3RoOyBpID0gaSArIDEp IHsKICAgICAgICAgZmxvYXRbMTBdIHZhbHVlOwogICAgICAgICBmaWxsKEB2YWx1ZSwgZmxvYXQo aSkpOwogICAgICAgICBhcnJheVtpXSA9IHZhbHVlOwogICAgIH0KIAotICAgIGZvciAodWludCBp ID0gMDsgaSA8IGFycmF5Lmxlbmd0aDsgKytpKSB7CisgICAgZm9yICh1aW50IGkgPSAwOyBpIDwg YXJyYXkubGVuZ3RoOyBpID0gaSArIDEpIHsKICAgICAgICAgZmxvYXRbMTBdIHZhbHVlID0gYXJy YXlbaV07CiAgICAgICAgIGlmICghY29udGFpbnMoQHZhbHVlLCBmbG9hdChpKSkpCiAgICAgICAg ICAgICByZXR1cm47Cg==