[AMDGPU] Add ptr.s.buffer.load intrinsic, use it from Clang#209243
Open
krzysz00 wants to merge 2 commits into
Open
[AMDGPU] Add ptr.s.buffer.load intrinsic, use it from Clang#209243krzysz00 wants to merge 2 commits into
krzysz00 wants to merge 2 commits into
Conversation
This commit adds a version of the existing s_buffer_load intrinsic that more accurately models the memory semantics of the s_buffer_load instruction, namely that it is, in fact, a memory load. To preserve the existing behavior that the "nomem" s.buffer.load intrinsic was using, Clang and MLIR add !invariant.load metadata when constructing the intrinsic (matching documented requirements on scalarazable buffer loads) and a late codegen pass adds the metadata just to be safe. Tests that were "about" s.buffer.load have been copied to create versions that use the new intrinsic, as was done for the other *.ptr.buffer.* operations. Other tests have been upgraded to use the new intrinsic. This has mainly resulted in minor instruction ordering changes in prologues, if any change at all. However, CodeGen/AMDGPU/dagcombine-fma-fmad.ll has seen a v_fma => v_mad pattern fail to match in one case, I don't know if this is a real regression. While I was here, the s_buffer_load => buffer_load fallback has been updated to preserve cachepolity. AI disclosure: Primarily AI-written code, but I have at least looked at and tried to find the worst of the silliness. Co-Authored-By: Codex <codex@openai.com>
krzysz00
requested review from
cmc-rep,
jayfoad,
macurtis-amd,
nhaehnle,
slinder1 and
zGoldthorpe
llvmorg-github-actions
Bot
added
backend:AMDGPU
clang:frontend
|
@llvm/pr-subscribers-mlir-llvm @llvm/pr-subscribers-clang-codegen Author: Krzysztof Drewniak (krzysz00) ChangesThis commit adds a version of the existing s_buffer_load intrinsic that more accurately models the memory semantics of the s_buffer_load instruction, namely that it is, in fact, a memory load. To preserve the existing behavior that the "nomem" s.buffer.load intrinsic was using, Clang and MLIR add !invariant.load metadata when constructing the intrinsic (matching documented requirements on scalarazable buffer loads) and a late codegen pass adds the metadata just to be safe. Tests that were "about" s.buffer.load have been copied to create versions that use the new intrinsic, as was done for the other .ptr.buffer. operations. Other tests have been upgraded to use the new intrinsic. This has mainly resulted in minor instruction ordering changes in prologues, if any change at all. However, CodeGen/AMDGPU/dagcombine-fma-fmad.ll has seen a v_fma => v_mad pattern fail to match in one case, I don't know if this is a real regression. While I was here, the s_buffer_load => buffer_load fallback has been updated to preserve cachepolity. AI disclosure: Primarily AI-written code, but I have at least looked at and tried to find the worst of the silliness. Patch is 783.77 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/209243.diff 51 Files Affected:
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
index c78c8ec291eef..0b5215381c72a 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
@@ -718,7 +718,7 @@ value gives an undefined result.
def DocCatSBufferLoad : DocumentationCategory<"S-Buffer Load Builtins"> {
let Content = [{
-These builtins lower to ``llvm.amdgcn.s.buffer.load`` and issue an
+These builtins lower to ``llvm.amdgcn.ptr.s.buffer.load`` and issue an
``s_buffer_load`` when the byte offset is uniform across the wavefront.
When the offset is divergent, the backend may lower to a ``buffer_load``.
@@ -735,8 +735,7 @@ When the offset is divergent, the backend may lower to a ``buffer_load``.
These builtins take the buffer resource as ``__amdgpu_buffer_rsrc_t``, which
can be created with ``__builtin_amdgcn_make_buffer_rsrc``. The backend
-lowers the resource to the 4-dword SGPR descriptor required by
-``llvm.amdgcn.s.buffer.load``.
+lowers the resource to the 4-dword SGPR descriptor required by the instruction.
The return type selects the load width. Separate builtins are provided for
each supported scalar and vector element type.
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index 66e4e688e33ee..119cdc7b8ab19 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -34,17 +34,16 @@ namespace {
static Value *emitAMDGPUSBufferLoadBuiltin(CodeGenFunction &CGF,
const CallExpr *E) {
llvm::Type *RetTy = CGF.ConvertType(E->getType());
- Function *F = CGF.CGM.getIntrinsic(Intrinsic::amdgcn_s_buffer_load, RetTy);
+ Function *F =
+ CGF.CGM.getIntrinsic(Intrinsic::amdgcn_ptr_s_buffer_load, RetTy);
Value *RsrcPtr = CGF.EmitScalarExpr(E->getArg(0));
- llvm::Type *I128Ty = llvm::IntegerType::get(CGF.getLLVMContext(), 128);
- llvm::Type *RsrcVecTy =
- llvm::FixedVectorType::get(CGF.Builder.getInt32Ty(), 4);
- Value *RsrcInt = CGF.Builder.CreatePtrToInt(RsrcPtr, I128Ty);
- Value *Rsrc = CGF.Builder.CreateBitCast(RsrcInt, RsrcVecTy);
-
- return CGF.Builder.CreateCall(F, {Rsrc, CGF.EmitScalarExpr(E->getArg(1)),
- CGF.EmitScalarExpr(E->getArg(2))});
+ CallInst *Call =
+ CGF.Builder.CreateCall(F, {RsrcPtr, CGF.EmitScalarExpr(E->getArg(1)),
+ CGF.EmitScalarExpr(E->getArg(2))});
+ Call->setMetadata(llvm::LLVMContext::MD_invariant_load,
+ llvm::MDNode::get(CGF.Builder.getContext(), {}));
+ return Call;
}
// Has second type mangled argument.
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-s-buffer-load.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-s-buffer-load.cl
index 0c640e5020267..570becbf1bf88 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-s-buffer-load.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-s-buffer-load.cl
@@ -23,9 +23,7 @@ typedef half v4f16 __attribute__((ext_vector_type(4)));
// CHECK-LABEL: @test_amdgcn_s_buffer_load_i32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.s.buffer.load.i32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.ptr.s.buffer.load.i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8:![0-9]+]]
// CHECK-NEXT: ret i32 [[TMP0]]
//
int test_amdgcn_s_buffer_load_i32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -34,9 +32,7 @@ int test_amdgcn_s_buffer_load_i32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_i32_non_const_offset(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.s.buffer.load.i32(<4 x i32> [[RSRC_VEC:%.*]], i32 [[OFFSET:%.*]], i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.ptr.s.buffer.load.i32(ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret i32 [[TMP0]]
//
int test_amdgcn_s_buffer_load_i32_non_const_offset(__amdgpu_buffer_rsrc_t rsrc, int offset) {
@@ -45,9 +41,7 @@ int test_amdgcn_s_buffer_load_i32_non_const_offset(__amdgpu_buffer_rsrc_t rsrc,
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v2i32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.s.buffer.load.v2i32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.ptr.s.buffer.load.v2i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <2 x i32> [[TMP0]]
//
v2i32 test_amdgcn_s_buffer_load_v2i32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -56,9 +50,7 @@ v2i32 test_amdgcn_s_buffer_load_v2i32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v3i32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x i32> @llvm.amdgcn.s.buffer.load.v3i32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x i32> @llvm.amdgcn.ptr.s.buffer.load.v3i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <3 x i32> [[TMP0]]
//
v3i32 test_amdgcn_s_buffer_load_v3i32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -67,9 +59,7 @@ v3i32 test_amdgcn_s_buffer_load_v3i32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v4i32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.s.buffer.load.v4i32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.ptr.s.buffer.load.v4i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <4 x i32> [[TMP0]]
//
v4i32 test_amdgcn_s_buffer_load_v4i32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -78,9 +68,7 @@ v4i32 test_amdgcn_s_buffer_load_v4i32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v8i32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <8 x i32> @llvm.amdgcn.s.buffer.load.v8i32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <8 x i32> @llvm.amdgcn.ptr.s.buffer.load.v8i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <8 x i32> [[TMP0]]
//
v8i32 test_amdgcn_s_buffer_load_v8i32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -89,9 +77,7 @@ v8i32 test_amdgcn_s_buffer_load_v8i32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v16i32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <16 x i32> @llvm.amdgcn.s.buffer.load.v16i32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <16 x i32> @llvm.amdgcn.ptr.s.buffer.load.v16i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <16 x i32> [[TMP0]]
//
v16i32 test_amdgcn_s_buffer_load_v16i32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -100,9 +86,7 @@ v16i32 test_amdgcn_s_buffer_load_v16i32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_f32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call float @llvm.amdgcn.s.buffer.load.f32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call float @llvm.amdgcn.ptr.s.buffer.load.f32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret float [[TMP0]]
//
float test_amdgcn_s_buffer_load_f32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -111,9 +95,7 @@ float test_amdgcn_s_buffer_load_f32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v2f32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x float> @llvm.amdgcn.s.buffer.load.v2f32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x float> @llvm.amdgcn.ptr.s.buffer.load.v2f32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <2 x float> [[TMP0]]
//
v2f32 test_amdgcn_s_buffer_load_v2f32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -122,9 +104,7 @@ v2f32 test_amdgcn_s_buffer_load_v2f32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v3f32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x float> @llvm.amdgcn.s.buffer.load.v3f32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x float> @llvm.amdgcn.ptr.s.buffer.load.v3f32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <3 x float> [[TMP0]]
//
v3f32 test_amdgcn_s_buffer_load_v3f32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -133,9 +113,7 @@ v3f32 test_amdgcn_s_buffer_load_v3f32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v4f32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x float> @llvm.amdgcn.s.buffer.load.v4f32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x float> @llvm.amdgcn.ptr.s.buffer.load.v4f32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <4 x float> [[TMP0]]
//
v4f32 test_amdgcn_s_buffer_load_v4f32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -144,9 +122,7 @@ v4f32 test_amdgcn_s_buffer_load_v4f32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v8f32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <8 x float> @llvm.amdgcn.s.buffer.load.v8f32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <8 x float> @llvm.amdgcn.ptr.s.buffer.load.v8f32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <8 x float> [[TMP0]]
//
v8f32 test_amdgcn_s_buffer_load_v8f32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -155,9 +131,7 @@ v8f32 test_amdgcn_s_buffer_load_v8f32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v16f32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <16 x float> @llvm.amdgcn.s.buffer.load.v16f32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <16 x float> @llvm.amdgcn.ptr.s.buffer.load.v16f32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <16 x float> [[TMP0]]
//
v16f32 test_amdgcn_s_buffer_load_v16f32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -166,9 +140,7 @@ v16f32 test_amdgcn_s_buffer_load_v16f32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_i8(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call i8 @llvm.amdgcn.s.buffer.load.i8(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call i8 @llvm.amdgcn.ptr.s.buffer.load.i8(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret i8 [[TMP0]]
//
char test_amdgcn_s_buffer_load_i8(__amdgpu_buffer_rsrc_t rsrc) {
@@ -177,9 +149,7 @@ char test_amdgcn_s_buffer_load_i8(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_u8(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call i8 @llvm.amdgcn.s.buffer.load.i8(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call i8 @llvm.amdgcn.ptr.s.buffer.load.i8(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret i8 [[TMP0]]
//
unsigned char test_amdgcn_s_buffer_load_u8(__amdgpu_buffer_rsrc_t rsrc) {
@@ -188,9 +158,7 @@ unsigned char test_amdgcn_s_buffer_load_u8(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_i16(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call i16 @llvm.amdgcn.s.buffer.load.i16(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call i16 @llvm.amdgcn.ptr.s.buffer.load.i16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret i16 [[TMP0]]
//
short test_amdgcn_s_buffer_load_i16(__amdgpu_buffer_rsrc_t rsrc) {
@@ -199,9 +167,7 @@ short test_amdgcn_s_buffer_load_i16(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_u16(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call i16 @llvm.amdgcn.s.buffer.load.i16(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call i16 @llvm.amdgcn.ptr.s.buffer.load.i16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret i16 [[TMP0]]
//
unsigned short test_amdgcn_s_buffer_load_u16(__amdgpu_buffer_rsrc_t rsrc) {
@@ -210,9 +176,7 @@ unsigned short test_amdgcn_s_buffer_load_u16(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_f16(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call half @llvm.amdgcn.s.buffer.load.f16(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call half @llvm.amdgcn.ptr.s.buffer.load.f16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret half [[TMP0]]
//
half test_amdgcn_s_buffer_load_f16(__amdgpu_buffer_rsrc_t rsrc) {
@@ -221,9 +185,7 @@ half test_amdgcn_s_buffer_load_f16(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v2f16(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x half> @llvm.amdgcn.s.buffer.load.v2f16(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x half> @llvm.amdgcn.ptr.s.buffer.load.v2f16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <2 x half> [[TMP0]]
//
v2f16 test_amdgcn_s_buffer_load_v2f16(__amdgpu_buffer_rsrc_t rsrc) {
@@ -232,9 +194,7 @@ v2f16 test_amdgcn_s_buffer_load_v2f16(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v3f16(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x half> @llvm.amdgcn.s.buffer.load.v3f16(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x half> @llvm.amdgcn.ptr.s.buffer.load.v3f16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <3 x half> [[TMP0]]
//
v3f16 test_amdgcn_s_buffer_load_v3f16(__amdgpu_buffer_rsrc_t rsrc) {
@@ -243,9 +203,7 @@ v3f16 test_amdgcn_s_buffer_load_v3f16(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v4f16(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x half> @llvm.amdgcn.s.buffer.load.v4f16(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x half> @llvm.amdgcn.ptr.s.buffer.load.v4f16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <4 x half> [[TMP0]]
//
v4f16 test_amdgcn_s_buffer_load_v4f16(__amdgpu_buffer_rsrc_t rsrc) {
@@ -254,9 +212,7 @@ v4f16 test_amdgcn_s_buffer_load_v4f16(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v2i8(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x i8> @llvm.amdgcn.s.buffer.load.v2i8(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x i8> @llvm.amdgcn.ptr.s.buffer.load.v2i8(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <2 x i8> [[TMP0]]
//
v2i8 test_amdgcn_s_buffer_load_v2i8(__amdgpu_buffer_rsrc_t rsrc) {
@@ -265,9 +221,7 @@ v2i8 test_amdgcn_s_buffer_load_v2i8(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v3i8(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x i8> @llvm.amdgcn.s.buffer.load.v3i8(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x i8> @llvm.amdgcn.ptr.s.buffer.load.v3i8(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <3 x i8> [[TMP0]]
//
v3i8 test_amdgcn_s_buffer_load_v3i8(__amdgpu_buffer_rsrc_t rsrc) {
@@ -276,9 +230,7 @@ v3i8 test_amdgcn_s_buffer_load_v3i8(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v4i8(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x i8> @llvm.amdgcn.s.buffer.load.v4i8(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x i8> @llvm.amdgcn.ptr.s.buffer.load.v4i8(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
...
[truncated]
|
|
@llvm/pr-subscribers-llvm-globalisel Author: Krzysztof Drewniak (krzysz00) ChangesThis commit adds a version of the existing s_buffer_load intrinsic that more accurately models the memory semantics of the s_buffer_load instruction, namely that it is, in fact, a memory load. To preserve the existing behavior that the "nomem" s.buffer.load intrinsic was using, Clang and MLIR add !invariant.load metadata when constructing the intrinsic (matching documented requirements on scalarazable buffer loads) and a late codegen pass adds the metadata just to be safe. Tests that were "about" s.buffer.load have been copied to create versions that use the new intrinsic, as was done for the other .ptr.buffer. operations. Other tests have been upgraded to use the new intrinsic. This has mainly resulted in minor instruction ordering changes in prologues, if any change at all. However, CodeGen/AMDGPU/dagcombine-fma-fmad.ll has seen a v_fma => v_mad pattern fail to match in one case, I don't know if this is a real regression. While I was here, the s_buffer_load => buffer_load fallback has been updated to preserve cachepolity. AI disclosure: Primarily AI-written code, but I have at least looked at and tried to find the worst of the silliness. Patch is 783.77 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/209243.diff 51 Files Affected:
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
index c78c8ec291eef..0b5215381c72a 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
@@ -718,7 +718,7 @@ value gives an undefined result.
def DocCatSBufferLoad : DocumentationCategory<"S-Buffer Load Builtins"> {
let Content = [{
-These builtins lower to ``llvm.amdgcn.s.buffer.load`` and issue an
+These builtins lower to ``llvm.amdgcn.ptr.s.buffer.load`` and issue an
``s_buffer_load`` when the byte offset is uniform across the wavefront.
When the offset is divergent, the backend may lower to a ``buffer_load``.
@@ -735,8 +735,7 @@ When the offset is divergent, the backend may lower to a ``buffer_load``.
These builtins take the buffer resource as ``__amdgpu_buffer_rsrc_t``, which
can be created with ``__builtin_amdgcn_make_buffer_rsrc``. The backend
-lowers the resource to the 4-dword SGPR descriptor required by
-``llvm.amdgcn.s.buffer.load``.
+lowers the resource to the 4-dword SGPR descriptor required by the instruction.
The return type selects the load width. Separate builtins are provided for
each supported scalar and vector element type.
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index 66e4e688e33ee..119cdc7b8ab19 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -34,17 +34,16 @@ namespace {
static Value *emitAMDGPUSBufferLoadBuiltin(CodeGenFunction &CGF,
const CallExpr *E) {
llvm::Type *RetTy = CGF.ConvertType(E->getType());
- Function *F = CGF.CGM.getIntrinsic(Intrinsic::amdgcn_s_buffer_load, RetTy);
+ Function *F =
+ CGF.CGM.getIntrinsic(Intrinsic::amdgcn_ptr_s_buffer_load, RetTy);
Value *RsrcPtr = CGF.EmitScalarExpr(E->getArg(0));
- llvm::Type *I128Ty = llvm::IntegerType::get(CGF.getLLVMContext(), 128);
- llvm::Type *RsrcVecTy =
- llvm::FixedVectorType::get(CGF.Builder.getInt32Ty(), 4);
- Value *RsrcInt = CGF.Builder.CreatePtrToInt(RsrcPtr, I128Ty);
- Value *Rsrc = CGF.Builder.CreateBitCast(RsrcInt, RsrcVecTy);
-
- return CGF.Builder.CreateCall(F, {Rsrc, CGF.EmitScalarExpr(E->getArg(1)),
- CGF.EmitScalarExpr(E->getArg(2))});
+ CallInst *Call =
+ CGF.Builder.CreateCall(F, {RsrcPtr, CGF.EmitScalarExpr(E->getArg(1)),
+ CGF.EmitScalarExpr(E->getArg(2))});
+ Call->setMetadata(llvm::LLVMContext::MD_invariant_load,
+ llvm::MDNode::get(CGF.Builder.getContext(), {}));
+ return Call;
}
// Has second type mangled argument.
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-s-buffer-load.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-s-buffer-load.cl
index 0c640e5020267..570becbf1bf88 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-s-buffer-load.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-s-buffer-load.cl
@@ -23,9 +23,7 @@ typedef half v4f16 __attribute__((ext_vector_type(4)));
// CHECK-LABEL: @test_amdgcn_s_buffer_load_i32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.s.buffer.load.i32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.ptr.s.buffer.load.i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8:![0-9]+]]
// CHECK-NEXT: ret i32 [[TMP0]]
//
int test_amdgcn_s_buffer_load_i32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -34,9 +32,7 @@ int test_amdgcn_s_buffer_load_i32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_i32_non_const_offset(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.s.buffer.load.i32(<4 x i32> [[RSRC_VEC:%.*]], i32 [[OFFSET:%.*]], i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.ptr.s.buffer.load.i32(ptr addrspace(8) [[RSRC:%.*]], i32 [[OFFSET:%.*]], i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret i32 [[TMP0]]
//
int test_amdgcn_s_buffer_load_i32_non_const_offset(__amdgpu_buffer_rsrc_t rsrc, int offset) {
@@ -45,9 +41,7 @@ int test_amdgcn_s_buffer_load_i32_non_const_offset(__amdgpu_buffer_rsrc_t rsrc,
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v2i32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.s.buffer.load.v2i32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x i32> @llvm.amdgcn.ptr.s.buffer.load.v2i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <2 x i32> [[TMP0]]
//
v2i32 test_amdgcn_s_buffer_load_v2i32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -56,9 +50,7 @@ v2i32 test_amdgcn_s_buffer_load_v2i32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v3i32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x i32> @llvm.amdgcn.s.buffer.load.v3i32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x i32> @llvm.amdgcn.ptr.s.buffer.load.v3i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <3 x i32> [[TMP0]]
//
v3i32 test_amdgcn_s_buffer_load_v3i32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -67,9 +59,7 @@ v3i32 test_amdgcn_s_buffer_load_v3i32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v4i32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.s.buffer.load.v4i32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.ptr.s.buffer.load.v4i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <4 x i32> [[TMP0]]
//
v4i32 test_amdgcn_s_buffer_load_v4i32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -78,9 +68,7 @@ v4i32 test_amdgcn_s_buffer_load_v4i32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v8i32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <8 x i32> @llvm.amdgcn.s.buffer.load.v8i32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <8 x i32> @llvm.amdgcn.ptr.s.buffer.load.v8i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <8 x i32> [[TMP0]]
//
v8i32 test_amdgcn_s_buffer_load_v8i32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -89,9 +77,7 @@ v8i32 test_amdgcn_s_buffer_load_v8i32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v16i32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <16 x i32> @llvm.amdgcn.s.buffer.load.v16i32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <16 x i32> @llvm.amdgcn.ptr.s.buffer.load.v16i32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <16 x i32> [[TMP0]]
//
v16i32 test_amdgcn_s_buffer_load_v16i32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -100,9 +86,7 @@ v16i32 test_amdgcn_s_buffer_load_v16i32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_f32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call float @llvm.amdgcn.s.buffer.load.f32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call float @llvm.amdgcn.ptr.s.buffer.load.f32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret float [[TMP0]]
//
float test_amdgcn_s_buffer_load_f32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -111,9 +95,7 @@ float test_amdgcn_s_buffer_load_f32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v2f32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x float> @llvm.amdgcn.s.buffer.load.v2f32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x float> @llvm.amdgcn.ptr.s.buffer.load.v2f32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <2 x float> [[TMP0]]
//
v2f32 test_amdgcn_s_buffer_load_v2f32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -122,9 +104,7 @@ v2f32 test_amdgcn_s_buffer_load_v2f32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v3f32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x float> @llvm.amdgcn.s.buffer.load.v3f32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x float> @llvm.amdgcn.ptr.s.buffer.load.v3f32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <3 x float> [[TMP0]]
//
v3f32 test_amdgcn_s_buffer_load_v3f32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -133,9 +113,7 @@ v3f32 test_amdgcn_s_buffer_load_v3f32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v4f32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x float> @llvm.amdgcn.s.buffer.load.v4f32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x float> @llvm.amdgcn.ptr.s.buffer.load.v4f32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <4 x float> [[TMP0]]
//
v4f32 test_amdgcn_s_buffer_load_v4f32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -144,9 +122,7 @@ v4f32 test_amdgcn_s_buffer_load_v4f32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v8f32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <8 x float> @llvm.amdgcn.s.buffer.load.v8f32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <8 x float> @llvm.amdgcn.ptr.s.buffer.load.v8f32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <8 x float> [[TMP0]]
//
v8f32 test_amdgcn_s_buffer_load_v8f32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -155,9 +131,7 @@ v8f32 test_amdgcn_s_buffer_load_v8f32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v16f32(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <16 x float> @llvm.amdgcn.s.buffer.load.v16f32(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <16 x float> @llvm.amdgcn.ptr.s.buffer.load.v16f32(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <16 x float> [[TMP0]]
//
v16f32 test_amdgcn_s_buffer_load_v16f32(__amdgpu_buffer_rsrc_t rsrc) {
@@ -166,9 +140,7 @@ v16f32 test_amdgcn_s_buffer_load_v16f32(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_i8(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call i8 @llvm.amdgcn.s.buffer.load.i8(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call i8 @llvm.amdgcn.ptr.s.buffer.load.i8(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret i8 [[TMP0]]
//
char test_amdgcn_s_buffer_load_i8(__amdgpu_buffer_rsrc_t rsrc) {
@@ -177,9 +149,7 @@ char test_amdgcn_s_buffer_load_i8(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_u8(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call i8 @llvm.amdgcn.s.buffer.load.i8(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call i8 @llvm.amdgcn.ptr.s.buffer.load.i8(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret i8 [[TMP0]]
//
unsigned char test_amdgcn_s_buffer_load_u8(__amdgpu_buffer_rsrc_t rsrc) {
@@ -188,9 +158,7 @@ unsigned char test_amdgcn_s_buffer_load_u8(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_i16(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call i16 @llvm.amdgcn.s.buffer.load.i16(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call i16 @llvm.amdgcn.ptr.s.buffer.load.i16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret i16 [[TMP0]]
//
short test_amdgcn_s_buffer_load_i16(__amdgpu_buffer_rsrc_t rsrc) {
@@ -199,9 +167,7 @@ short test_amdgcn_s_buffer_load_i16(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_u16(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call i16 @llvm.amdgcn.s.buffer.load.i16(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call i16 @llvm.amdgcn.ptr.s.buffer.load.i16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret i16 [[TMP0]]
//
unsigned short test_amdgcn_s_buffer_load_u16(__amdgpu_buffer_rsrc_t rsrc) {
@@ -210,9 +176,7 @@ unsigned short test_amdgcn_s_buffer_load_u16(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_f16(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call half @llvm.amdgcn.s.buffer.load.f16(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call half @llvm.amdgcn.ptr.s.buffer.load.f16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret half [[TMP0]]
//
half test_amdgcn_s_buffer_load_f16(__amdgpu_buffer_rsrc_t rsrc) {
@@ -221,9 +185,7 @@ half test_amdgcn_s_buffer_load_f16(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v2f16(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x half> @llvm.amdgcn.s.buffer.load.v2f16(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x half> @llvm.amdgcn.ptr.s.buffer.load.v2f16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <2 x half> [[TMP0]]
//
v2f16 test_amdgcn_s_buffer_load_v2f16(__amdgpu_buffer_rsrc_t rsrc) {
@@ -232,9 +194,7 @@ v2f16 test_amdgcn_s_buffer_load_v2f16(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v3f16(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x half> @llvm.amdgcn.s.buffer.load.v3f16(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x half> @llvm.amdgcn.ptr.s.buffer.load.v3f16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <3 x half> [[TMP0]]
//
v3f16 test_amdgcn_s_buffer_load_v3f16(__amdgpu_buffer_rsrc_t rsrc) {
@@ -243,9 +203,7 @@ v3f16 test_amdgcn_s_buffer_load_v3f16(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v4f16(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x half> @llvm.amdgcn.s.buffer.load.v4f16(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x half> @llvm.amdgcn.ptr.s.buffer.load.v4f16(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <4 x half> [[TMP0]]
//
v4f16 test_amdgcn_s_buffer_load_v4f16(__amdgpu_buffer_rsrc_t rsrc) {
@@ -254,9 +212,7 @@ v4f16 test_amdgcn_s_buffer_load_v4f16(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v2i8(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x i8> @llvm.amdgcn.s.buffer.load.v2i8(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x i8> @llvm.amdgcn.ptr.s.buffer.load.v2i8(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <2 x i8> [[TMP0]]
//
v2i8 test_amdgcn_s_buffer_load_v2i8(__amdgpu_buffer_rsrc_t rsrc) {
@@ -265,9 +221,7 @@ v2i8 test_amdgcn_s_buffer_load_v2i8(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v3i8(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x i8> @llvm.amdgcn.s.buffer.load.v3i8(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <3 x i8> @llvm.amdgcn.ptr.s.buffer.load.v3i8(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
// CHECK-NEXT: ret <3 x i8> [[TMP0]]
//
v3i8 test_amdgcn_s_buffer_load_v3i8(__amdgpu_buffer_rsrc_t rsrc) {
@@ -276,9 +230,7 @@ v3i8 test_amdgcn_s_buffer_load_v3i8(__amdgpu_buffer_rsrc_t rsrc) {
// CHECK-LABEL: @test_amdgcn_s_buffer_load_v4i8(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[RSRC_INT:%.*]] = ptrtoint ptr addrspace(8) [[RSRC:%.*]] to i128
-// CHECK-NEXT: [[RSRC_VEC:%.*]] = bitcast i128 [[RSRC_INT]] to <4 x i32>
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x i8> @llvm.amdgcn.s.buffer.load.v4i8(<4 x i32> [[RSRC_VEC:%.*]], i32 0, i32 0)
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <4 x i8> @llvm.amdgcn.ptr.s.buffer.load.v4i8(ptr addrspace(8) [[RSRC:%.*]], i32 0, i32 0), !invariant.load [[META8]]
...
[truncated]
|
You can test this locally with the following command:git diff -U0 --pickaxe-regex -S '([^a-zA-Z0-9#_-]undef([^a-zA-Z0-9_-]|$)|UndefValue::get)' 'HEAD~1' HEAD llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.ptr.s.buffer.load.ll llvm/test/CodeGen/AMDGPU/GlobalISel/regbanklegalize-amdgcn.ptr.s.buffer.load.ll llvm/test/CodeGen/AMDGPU/GlobalISel/regbanklegalize-amdgcn.ptr.s.buffer.load.subdword.ll llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-amdgcn.ptr.s.buffer.load.ll llvm/test/CodeGen/AMDGPU/gfx12_scalar_subword_ptr_s_buffer_loads.ll llvm/test/CodeGen/AMDGPU/llvm.amdgcn.ptr.s.buffer.load-gfx12.ll llvm/test/CodeGen/AMDGPU/llvm.amdgcn.ptr.s.buffer.load-mmo.ll llvm/test/CodeGen/AMDGPU/llvm.amdgcn.ptr.s.buffer.load.ll llvm/test/CodeGen/AMDGPU/ptr-s-buffer-load-mmo-offsets.ll clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp llvm/lib/Target/AMDGPU/AMDGPUAsanInstrumentation.cpp llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp llvm/lib/Target/AMDGPU/AMDGPULowerIntrinsics.cpp llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp llvm/lib/Target/AMDGPU/SIISelLowering.cpp llvm/lib/Target/AMDGPU/SIISelLowering.h llvm/test/CodeGen/AMDGPU/GlobalISel/is-safe-to-sink-bug.ll llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.set.inactive.ll llvm/test/CodeGen/AMDGPU/amdgpu-uniform-intrinsic-wwm-single-lane.ll llvm/test/CodeGen/AMDGPU/bug-deadlanes.ll llvm/test/CodeGen/AMDGPU/bug-vopc-commute.ll llvm/test/CodeGen/AMDGPU/dagcombine-fma-fmad.ll llvm/test/CodeGen/AMDGPU/fneg-modifier-casting.ll llvm/test/CodeGen/AMDGPU/group-image-instructions.ll llvm/test/CodeGen/AMDGPU/insert_waitcnt_for_precise_memory.ll llvm/test/CodeGen/AMDGPU/llvm.amdgcn.set.inactive.ll llvm/test/CodeGen/AMDGPU/scalar-float-sop2.ll llvm/test/CodeGen/AMDGPU/scheduler-subrange-crash.ll llvm/test/CodeGen/AMDGPU/set_kill_i1_for_floation_point_comparison.ll llvm/test/CodeGen/AMDGPU/sgpr-copy.ll llvm/test/CodeGen/AMDGPU/si-scheduler.ll llvm/test/CodeGen/AMDGPU/si-sgpr-spill.ll llvm/test/CodeGen/AMDGPU/si-spill-cf.ll llvm/test/CodeGen/AMDGPU/splitkit-getsubrangeformask.ll llvm/test/CodeGen/AMDGPU/vgpr-spill-emergency-stack-slot.ll llvm/test/CodeGen/AMDGPU/wqm.ll llvm/test/CodeGen/AMDGPU/wwm-reserved-spill.ll llvm/test/CodeGen/AMDGPU/wwm-reserved.ll llvm/test/Transforms/EarlyCSE/AMDGPU/intrinsics.ll llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-demanded-vector-elts-inseltpoison.ll llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-demanded-vector-elts.llThe following files introduce new uses of undef:
Undef is now deprecated and should only be used in the rare cases where no replacement is possible. For example, a load of uninitialized memory yields In tests, avoid using For example, this is considered a bad practice: define void @fn() {
...
br i1 undef, ...
}Please use the following instead: define void @fn(i1 %cond) {
...
br i1 %cond, ...
}Please refer to the Undefined Behavior Manual for more information. |
Contributor
Author
|
False positive undef detector |
shiltian
reviewed
Jul 13, 2026
| AMDGPURsrcIntrinsic<0>; | ||
|
|
||
| // Generate an s_buffer_load instruction, falling back to buffer_load if the offset | ||
| // is not uniform. This operates like the s.buffer.load intrinsic, but uses a buffer |
Contributor
There was a problem hiding this comment.
Why does int_amdgcn_s_buffer_load do the other way around, aka generate buffer_load by default but optimize to s_buffer_load?
Contributor
There was a problem hiding this comment.
the scalar cache isn't coherent with vector, it's difficult to migrate that direction
Contributor
Author
There was a problem hiding this comment.
I think s_buffer_load is documented weird, given that it's got those "no memory" semantics and all.
Both of these are an s_buffer_load when all the arguments are uniform and a buffer_load otherwise, I just wrote the docs on this one as "you're explicitly requesting a uniform load".
Contributor
Author
There was a problem hiding this comment.
the scalar cache isn't coherent with vector, it's difficult to migrate that direction
Which is why (though this was underdocumented for the "old form") you must only s_buffer_load invariant things/constants, therefore preventing the coherence issue from coming up in any realistic workload
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
This commit adds a version of the existing s_buffer_load intrinsic that more accurately models the memory semantics of the s_buffer_load instruction, namely that it is, in fact, a memory load.
To preserve the existing behavior that the "nomem" s.buffer.load intrinsic was using, Clang and MLIR add !invariant.load metadata when constructing the intrinsic (matching documented requirements on scalarazable buffer loads) and a late codegen pass adds the metadata just to be safe.
Tests that were "about" s.buffer.load have been copied to create versions that use the new intrinsic, as was done for the other .ptr.buffer. operations.
Other tests have been upgraded to use the new intrinsic. This has mainly resulted in minor instruction ordering changes in prologues, if any change at all. However, CodeGen/AMDGPU/dagcombine-fma-fmad.ll has seen a v_fma => v_mad pattern fail to match in one case, I don't know if this is a real regression.
While I was here, the s_buffer_load => buffer_load fallback has been updated to preserve cachepolity.
AI disclosure: Primarily AI-written code, but I have at least looked at and tried to find the worst of the silliness.