[AMDGPU] Numerical error related to load/store vectorization

[qedawkins]

The following input example: https://gist.github.com/qedawkins/c620832f96a5c504295f9694cc8956e2

Which is approximately compiled from

// int8_t a[144]; // [2 x 8 x 3 x 3]
// int8_t b[144]; // [2 x 8 x 3 x 3]
// int c[2];
__global__ void helloworld(int8_t* a, int8_t* b, int* c)
{
  int id = hipThreadIdx_x
  if (id < 2) {
    for (int c = 0; c < 8; ++c) {
      for (int h = 0; h < 3; ++h) {
        for (int w = 0; w < 3; ++w) {
          int x = a[id * 72 + c * 9 + h * 3 + w];
          int y = b[id * 72 + c * 9 + h * 3 + w];
          c[id] += x * y;
        }
      }
    }
  }
}

When compiled with

llc -O3 module_test_dispatch_0_rocm_hsaco_fb.optimized.ll -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1100

Produces incorrect numerics. For inputs

a = np.ones((2, 8, 3, 3), dtype=np.int8)
b = np.broadcast_to(np.array([1, 2, 1], dtype=np.int8), (2, 8, 3, 3))

In other words only the values of b vary along the inner most loop, this gives c = [88, 88] when the correct values should be (1 + 2 + 1) * 3 * 8 = [96, 96]

Compiling with

llc -O3 module_test_dispatch_0_rocm_hsaco_fb.optimized.ll -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1100 -amdgpu-load-store-vectorizer=false

(turning off the load/store vectorizer) or just with -O1 instead of -O3 gives the correct values. I am not sure exactly how the generated code is miscompiling and don't know whether load/store vectorization is the root cause or if disabling it is just masking the problem. Further information about the downstream issue here: iree-org/iree#18798

cc @MaheshRavishankar

[llvmbot]

@llvm/issue-subscribers-backend-amdgpu

Author: Quinn Dawkins (qedawkins)

The following input example: https://gist.github.com/qedawkins/c620832f96a5c504295f9694cc8956e2

Which is approximately compiled from

// int8_t a[144]; // [2 x 8 x 3 x 3]
// int8_t b[144]; // [2 x 8 x 3 x 3]
// int c[2];
__global__ void helloworld(int8_t* a, int8_t* b, int* c)
{
  int id = hipThreadIdx_x
  if (id < 2) {
    for (int c = 0; c < 8; ++c) {
      for (int h = 0; h < 3; ++h) {
        for (int w = 0; w < 3; ++w) {
          int x = a[id * 72 + c * 9 + h * 3 + w];
          int y = b[id * 72 + c * 9 + h * 3 + w];
          c[id] += x * y;
        }
      }
    }
  }
}

When compiled with

llc -O3 module_test_dispatch_0_rocm_hsaco_fb.optimized.ll -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1100

Produces incorrect numerics. For inputs

a = np.ones((2, 8, 3, 3), dtype=np.int8)
b = np.broadcast_to(np.array([1, 2, 1], dtype=np.int8), (2, 8, 3, 3))

In other words only the values of b vary along the inner most loop, this gives c = [88, 88] when the correct values should be (1 + 2 + 1) * 3 * 8 = [96, 96]

Compiling with

llc -O3 module_test_dispatch_0_rocm_hsaco_fb.optimized.ll -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1100 -amdgpu-load-store-vectorizer=false

(turning off the load/store vectorizer) or just with -O1 instead of -O3 gives the correct values. I am not sure exactly how the generated code is miscompiling and don't know whether load/store vectorization is the root cause or if disabling it is just masking the problem. Further information about the downstream issue here: iree-org/iree#18798

cc @MaheshRavishankar

[arsenm]

No idea what's wrong but there are some weird things with this IR. Something is quite suboptimal with whatever produced it.

First it has a sequence of assumes at the start of the function, asserting that the pointers are aligned. The assumes are implemented using ptrtoint (a better way would be in terms of llvm.ptrmask). But the assumes are entirely unnecessary, they are redundant with the align attribute on the pointer parameters. Assumes are in general harmful, so these should only be added when they definitely do something.

Second, this IR has not run through the middle end optimizer passes. The vectorization seen in llc -O3 produces strange 7 and 9 element vectors, which produce sloppy looking codegen. The vectorization produced from the middle end is much more reasonable looking

[arsenm]

Does it work if you use -global-isel?

[MaheshRavishankar]

Thanks @arsenm ! We can definitely drop the assumes that you think are redundant. The IR should be produced with opt -O3 though.

[arsenm]

Thanks @arsenm ! We can definitely drop the assumes that you think are redundant. The IR should be produced with opt -O3 though.

It is not. The IR radically changes if I run opt/clang -O3 on it

[qedawkins]

Sorry for the late reply, last week was busy for me.

Does it work if you use -global-isel?

Yes, global-isel appears to work. Looking at the output there is none of the vectorization that I was seeing without it

before:

    global_load_i8 v16, v[9:10], off offset:8
    global_load_b64 v[9:10], v[9:10], off 
    global_load_i8 v17, v[11:12], off offset:7
    global_load_i8 v18, v[13:14], off offset:8
    global_load_d16_b16 v15, v[11:12], off offset:5
    global_load_i8 v13, v[13:14], off 
    global_load_b32 v11, v[11:12], off offset:1

afterwards all loads are global_load_i8. I can try and see whether global isel will work for our use cases downstream.

Additionally enabling the SLPVectorizer (llvm::PipelineTuningOptions::SLPVectorization = true) appears to fix the issue as well. I'm guessing that's the primary pass that's missing in the middle end that is happening with clang -O3?

[MaheshRavishankar]

I checked, we are using -O2 AFAICS. So maybe -O2 to -O3 makes a difference.

[arsenm]

Overall I think the vectorization is a red herring. The SLPVectorizer should run at any optimization level, by default.

My suspicion is that this is something wrong with the v_perm_b32 and related bit shuffling optimizations. Globalisel just happens to not implement those, so it's the closest thing we have to an option to disable them

[MaheshRavishankar]

Overall I think the vectorization is a red herring. The SLPVectorizer should run at any optimization level, by default.

We've had issues with SLP vectorizer using up too many registers and spilling/degrading performance (on other backends). We tried turning it on for AMD GPU backend, and it seems ok but it is something we are scared about cause we prefer to vectorize ourselves instead of using SLP vectorizer.

My suspicion is that this is something wrong with the v_perm_b32 and related bit shuffling optimizations. Globalisel just happens to not implement those, so it's the closest thing we have to an option to disable them

That's probably right.

[krzysz00]

Would like to record some progress notes:

https://github.com/krzysz00/iree-bug-18798-reproduction-harness is now a standalone reproduction harness for the bug

While trying to reduce the problem further, I discovered that removing all the llvm.assume stuff didn't affect the correctness of the generated code (it was still correct), but removing the addition of the Y and Z workitem IDs (which, per !reqd_work_group_size are always 0) does make the code work

Also, some debugging session a while back yielded that applying

diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index bbdc006b9afc..dedaeb2fe34b 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -11932,6 +11932,7 @@ calculateSrcByte(const SDValue Op, uint64_t DestByte, uint64_t SrcIndex = 0,
static const std::optional<ByteProvider<SDValue>>
calculateByteProvider(const SDValue &Op, unsigned Index, unsigned Depth,
                       unsigned StartingIndex = 0) {
+                        return std::nullopt;
   // Finding Src tree of RHS of or typically requires at least 1 additional
   // depth
   if (Depth > 6)

would work around the issue

[krzysz00]

Another observation is that load-store-vectorizer isn't treating the two input arguments consistently for reasons that aren't clear to me

[krzysz00]

Further notes from operation "flail around wildly" - disabling (by patch) loop strength reduction causes the load-store-vectorizer equivalence classes to be the length-9 ones that're expected, which causes the bug to go away

So, it seems we have

1. Loop strength reduction doing something that confuses the vectorizer - but only when there're those additional adds
2. SelectionDAG choking on <7 x i8>

Could someone remind me the incantations for dumping out IR right before it gets selection-DAG'd? There might be two tickets here, it feels ilke