Bug 199424 - [WHLSL] Change whlsl-two-dimensional-array.html to not be flaky on AMD Radeon Pro GPUs
Summary: [WHLSL] Change whlsl-two-dimensional-array.html to not be flaky on AMD Radeon...
Status: RESOLVED FIXED
Alias: None
Product: WebKit
Classification: Unclassified
Component: WebGPU (show other bugs)
Version: WebKit Nightly Build
Hardware: Unspecified Unspecified
: P2 Normal
Assignee: Saam Barati
URL:
Keywords: InRadar
Depends on:
Blocks:
 
Reported: 2019-07-02 15:29 PDT by Saam Barati
Modified: 2019-07-02 16:54 PDT (History)
8 users (show)

See Also:


Attachments
patch (2.41 KB, patch)
2019-07-02 15:43 PDT, Saam Barati
no flags Details | Formatted Diff | Diff

Note You need to log in before you can comment on or make changes to this bug.
Description Saam Barati 2019-07-02 15:29:03 PDT
...
Comment 1 Saam Barati 2019-07-02 15:29:50 PDT
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.
Comment 2 Saam Barati 2019-07-02 15:43:28 PDT
Created attachment 373366 [details]
patch
Comment 3 Myles C. Maxfield 2019-07-02 16:13:58 PDT
Comment on attachment 373366 [details]
patch

🤮
Comment 4 WebKit Commit Bot 2019-07-02 16:53:23 PDT
Comment on attachment 373366 [details]
patch

Clearing flags on attachment: 373366

Committed r247081: <https://trac.webkit.org/changeset/247081>
Comment 5 WebKit Commit Bot 2019-07-02 16:53:24 PDT
All reviewed patches have been landed.  Closing bug.
Comment 6 Radar WebKit Bug Importer 2019-07-02 16:54:15 PDT
<rdar://problem/52556060>