[AMDGPU] Disable bool range metadata to workaround backend issue

[yaxunl]

Summary

Currently there is a middle-end or backend issue
https://github.com/llvm/llvm-project/issues/58176
which causes values loaded from bool pointer incorrect when
bool range metadata is emitted. Temporarily
disable bool range metadata until the backend issue
is fixed.

[yaxunl]

fix comments in test

[tra]

Is there more info about the issue? What does AMDGPU currently emit for the test case?

AFAICT from running it on CE (https://godbolt.org/z/ccq3vnbrM) llvm optimizes it to essentially *y = *x and generates a 1-byte load+store for both NVPTX and AMDGPU.

[yaxunl]

In D135269#3837394, @tra wrote:

Is there more info about the issue? What does AMDGPU currently emit for the test case?

AFAICT from running it on CE (https://godbolt.org/z/ccq3vnbrM) llvm optimizes it to essentially *y = *x and generates a 1-byte load+store for both NVPTX and AMDGPU.

The issue happens to more complicated test cases which I cannot reduce right now.

Basically 8018d6be3459780e81a5da128a9915eb27909902 caused regressions in some PyTorch tests. Investigation shows the propagation of range metadata for bool type triggered some optimizations which caused some bool values to be loaded incorrectly. I will continue investigating the issue. However, I need a workaround for now.

[tra]

clang/lib/CodeGen/CGExpr.cpp

1792–1794

It would be great to open a github issue, if we don't have one yet, and reference it here, so we can tell later what exactly it is we're working around here and know for sure when/whether we can undo the change.

clang/test/CodeGenCUDA/bool-range.cu

15

Ditto, a bug reference would help.

[yaxunl]

clang/lib/CodeGen/CGExpr.cpp

1792–1794

opened github issue:

https://github.com/llvm/llvm-project/issues/58176

will update the comments

[yaxunl]

update comments with issue link

[nikic]

Checking back here, have you made any progress on reducing the issue?

cc @arsenm for awareness

[yaxunl]

In D135269#3863470, @nikic wrote:

Checking back here, have you made any progress on reducing the issue?

cc @arsenm for awareness

No. I am busy with other work and have not got time to get back on it.

[nikic]

Checking back here again on whether there is any progress on finding the root cause of the issue. If no progress is expected in the near future, I'd ask for this patch to be reverted.

[yaxunl]

In D135269#3981561, @nikic wrote:

Checking back here again on whether there is any progress on finding the root cause of the issue. If no progress is expected in the near future, I'd ask for this patch to be reverted.

@jrbyrnes is working on the root cause of this issue. Any updates? Thanks.

[jrbyrnes]

In D135269#3981856, @yaxunl wrote:

In D135269#3981561, @nikic wrote:

Checking back here again on whether there is any progress on finding the root cause of the issue. If no progress is expected in the near future, I'd ask for this patch to be reverted.

@jrbyrnes is working on the root cause of this issue. Any updates? Thanks.

Thanks for the ping. I would also like to see this reverted as it enables some optimizations. I do not have a definitive answer at the moment (w.r.t reverting this), but hope to provide one soon

As for now, the issue we are seeing from (https://github.com/llvm/llvm-project/commit/8018d6be3459780e81a5da128a9915eb27909902) seems most likely to be a source code issue (first document of issue https://github.com/pytorch/pytorch/issues/54789 . upstream PyTorch currently skips problematic test https://github.com/pytorch/pytorch/blob/b738da8c8e4d9142ad38a1bd8c35d0bfef4b5e3c/torch/testing/_internal/common_methods_invocations.py#L14891) . I will provide a better update soon.

[yaxunl]

In D135269#3982103, @jrbyrnes wrote:

In D135269#3981856, @yaxunl wrote:

In D135269#3981561, @nikic wrote:

Checking back here again on whether there is any progress on finding the root cause of the issue. If no progress is expected in the near future, I'd ask for this patch to be reverted.

@jrbyrnes is working on the root cause of this issue. Any updates? Thanks.

Thanks for the ping. I would also like to see this reverted as it enables some optimizations. I do not have a definitive answer at the moment (w.r.t reverting this), but hope to provide one soon

As for now, the issue we are seeing from (https://github.com/llvm/llvm-project/commit/8018d6be3459780e81a5da128a9915eb27909902) seems most likely to be a source code issue (first document of issue https://github.com/pytorch/pytorch/issues/54789 . upstream PyTorch currently skips problematic test https://github.com/pytorch/pytorch/blob/b738da8c8e4d9142ad38a1bd8c35d0bfef4b5e3c/torch/testing/_internal/common_methods_invocations.py#L14891) . I will provide a better update soon.

I will revert this patch. Thanks.

clang/lib/CodeGen/CGExpr.cpp

if (isNontemporal) {
     Load->setMetadata(CGM.getModule().getMDKindID("nontemporal"), Node);
   }
 
   CGM.DecorateInstructionWithTBAA(Load, TBAAInfo);
 
   if (EmitScalarRangeCheck(Load, Ty, Loc)) {
     // In order to prevent the optimizer from throwing away the check, don't
     // attach range metadata to the load.
-  } else if (CGM.getCodeGenOpts().OptimizationLevel > 0)
+    // TODO: Enable range metadata for AMDGCN after backend issue is fixed.
+  } else if (CGM.getCodeGenOpts().OptimizationLevel > 0 &&
+             !CGM.getTriple().isAMDGCN())
     if (llvm::MDNode *RangeInfo = getRangeForLoadFromType(Ty))
       Load->setMetadata(llvm::LLVMContext::MD_range, RangeInfo);
 
   return EmitFromMemory(Load, Ty);
 }
 
 llvm::Value *CodeGenFunction::EmitToMemory(llvm::Value *Value, QualType Ty) {
   // Bool has a different representation in memory than in registers.

clang/test/CodeGenCUDA/bool-range.cu

+// RUN: %clang_cc1 -emit-llvm %s -O3 -o - -fcuda-is-device \
+// RUN:   -triple nvptx64-unknown-unknown | FileCheck %s -check-prefixes=NV
+// RUN: %clang_cc1 -emit-llvm %s -O3 -o - -fcuda-is-device \
+// RUN:   -triple amdgcn-amd-amdhsa | FileCheck %s -check-prefixes=AMD
+
+#include "Inputs/cuda.h"
+
+// Make sure bool loaded from memory is truncated and
+// range metadata is not emitted.
+
+// NV:  %[[LD:[0-9]+]] = load i8, ptr %x,{{.*}} !range ![[MD:[0-9]+]]
+// NV:  store i8 %[[LD]], ptr %y
+// NV: ![[MD]] = !{i8 0, i8 2}
+
+// TODO: Re-enable range metadata after backend issue is fixed.
+
+// AMD:  %[[LD:[0-9]+]] = load i8, ptr addrspace(1) %x.global
+// AMD-NOT: !range
+// AMD:  %[[AND:[0-9]+]] = and i8 %[[LD]], 1
+// AMD:  store i8 %[[AND]], ptr addrspace(1) %y.global
+__global__ void test1(bool *x, bool *y) {
+  *y = *x != false;
+}