llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: Matt Arsenault (arsenm) <details> <summary>Changes</summary> fillAMDGCNFeatureMap omitted flat-global-insts for gfx1310/gfx13-generic, so clang wrongly rejected __builtin_amdgcn_av_load_b128 / __builtin_amdgcn_av_store_b128 on those targets even though the backend enables the feature. Add it to the gfx13 case. Co-authored-by: Claude (Claude-Opus-4.8) --- Patch is 250.60 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/213121.diff 3 Files Affected: - (modified) clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-store.cl (+1) - (modified) llvm/lib/TargetParser/AMDGPUTargetParser.cpp (+1) - (modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.av.load.b128.ll (+3291) ``````````diff diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-store.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-store.cl index cd34500a94a24..7f8651c837bfe 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-store.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-store.cl @@ -5,6 +5,7 @@ // RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu11.00-unknown-unknown -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu12.00-unknown-unknown -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu12.50-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu13.10-unknown-unknown -emit-llvm -o - %s | FileCheck %s typedef __attribute__((__vector_size__(4 * sizeof(unsigned int)))) unsigned int v4u32; typedef v4u32 __global *global_ptr_to_v4u32; diff --git a/llvm/lib/TargetParser/AMDGPUTargetParser.cpp b/llvm/lib/TargetParser/AMDGPUTargetParser.cpp index 6e16bc8dc9507..c55e81bbdf6ff 100644 --- a/llvm/lib/TargetParser/AMDGPUTargetParser.cpp +++ b/llvm/lib/TargetParser/AMDGPUTargetParser.cpp @@ -500,6 +500,7 @@ static void fillAMDGCNFeatureMap(StringRef GPU, const Triple &T, Features["gfx12-insts"] = true; Features["gfx1250-insts"] = true; Features["gfx13-insts"] = true; + Features["flat-global-insts"] = true; Features["bitop3-insts"] = true; Features["prng-inst"] = true; Features["tanh-insts"] = true; diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.av.load.b128.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.av.load.b128.ll index fad60eb81d28e..2495fa8a43b29 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.av.load.b128.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.av.load.b128.ll @@ -5,12 +5,14 @@ ; RUN: llc -global-isel=0 -mtriple=amdgpu10.12 < %s | FileCheck -check-prefix=GFX1012-SDAG %s ; RUN: llc -global-isel=0 -mtriple=amdgpu11.00 < %s | FileCheck -check-prefix=GFX1100-SDAG %s ; RUN: llc -global-isel=0 -mtriple=amdgpu12.50 < %s | FileCheck -check-prefix=GFX1250-SDAG %s +; RUN: llc -global-isel=0 -mtriple=amdgpu13.10 < %s | FileCheck -check-prefix=GFX1310-SDAG %s ; RUN: llc -global-isel=1 -mtriple=amdgpu9.06 < %s | FileCheck -check-prefix=GFX906-ISEL %s ; RUN: llc -global-isel=1 -mtriple=amdgpu9.42 < %s | FileCheck -check-prefix=GFX942-ISEL %s ; RUN: llc -global-isel=1 -mtriple=amdgpu10.12 < %s | FileCheck -check-prefix=GFX1012-ISEL %s ; RUN: llc -global-isel=1 -mtriple=amdgpu11.00 < %s | FileCheck -check-prefix=GFX1100-ISEL %s ; RUN: llc -global-isel=1 -mtriple=amdgpu12.50 < %s | FileCheck -check-prefix=GFX1250-ISEL %s +; RUN: llc -global-isel=1 -mtriple=amdgpu13.10 < %s | FileCheck -check-prefix=GFX1310-ISEL %s ;;============================================================================== ;; A few basic test cases @@ -54,6 +56,17 @@ define <4 x i32> @av_global_load_b128_0_00(ptr addrspace(1) %addr) { ; GFX1250-SDAG-NEXT: s_wait_loadcnt 0x0 ; GFX1250-SDAG-NEXT: s_set_pc_i64 s[30:31] ; +; GFX1310-SDAG-LABEL: av_global_load_b128_0_00: +; GFX1310-SDAG: ; %bb.0: ; %entry +; GFX1310-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_expcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_samplecnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX1310-SDAG-NEXT: global_load_b128 v[0:3], v[0:1], off +; GFX1310-SDAG-NEXT: s_wait_loadcnt 0x0 +; GFX1310-SDAG-NEXT: s_set_pc_i64 s[30:31] +; ; GFX906-ISEL-LABEL: av_global_load_b128_0_00: ; GFX906-ISEL: ; %bb.0: ; %entry ; GFX906-ISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) @@ -89,6 +102,17 @@ define <4 x i32> @av_global_load_b128_0_00(ptr addrspace(1) %addr) { ; GFX1250-ISEL-NEXT: global_load_b128 v[0:3], v[0:1], off ; GFX1250-ISEL-NEXT: s_wait_loadcnt 0x0 ; GFX1250-ISEL-NEXT: s_set_pc_i64 s[30:31] +; +; GFX1310-ISEL-LABEL: av_global_load_b128_0_00: +; GFX1310-ISEL: ; %bb.0: ; %entry +; GFX1310-ISEL-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_expcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_samplecnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_kmcnt 0x0 +; GFX1310-ISEL-NEXT: global_load_b128 v[0:3], v[0:1], off +; GFX1310-ISEL-NEXT: s_wait_loadcnt 0x0 +; GFX1310-ISEL-NEXT: s_set_pc_i64 s[30:31] entry: %data = call <4 x i32> @llvm.amdgcn.av.load.b128.p1(ptr addrspace(1) %addr, metadata !0) ret <4 x i32> %data @@ -132,6 +156,17 @@ define <4 x i32> @av_global_load_b128_0_01(ptr addrspace(1) %addr) { ; GFX1250-SDAG-NEXT: s_wait_loadcnt 0x0 ; GFX1250-SDAG-NEXT: s_set_pc_i64 s[30:31] ; +; GFX1310-SDAG-LABEL: av_global_load_b128_0_01: +; GFX1310-SDAG: ; %bb.0: ; %entry +; GFX1310-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_expcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_samplecnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX1310-SDAG-NEXT: global_load_b128 v[0:3], v[0:1], off scope:SCOPE_SE +; GFX1310-SDAG-NEXT: s_wait_loadcnt 0x0 +; GFX1310-SDAG-NEXT: s_set_pc_i64 s[30:31] +; ; GFX906-ISEL-LABEL: av_global_load_b128_0_01: ; GFX906-ISEL: ; %bb.0: ; %entry ; GFX906-ISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) @@ -167,6 +202,17 @@ define <4 x i32> @av_global_load_b128_0_01(ptr addrspace(1) %addr) { ; GFX1250-ISEL-NEXT: global_load_b128 v[0:3], v[0:1], off ; GFX1250-ISEL-NEXT: s_wait_loadcnt 0x0 ; GFX1250-ISEL-NEXT: s_set_pc_i64 s[30:31] +; +; GFX1310-ISEL-LABEL: av_global_load_b128_0_01: +; GFX1310-ISEL: ; %bb.0: ; %entry +; GFX1310-ISEL-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_expcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_samplecnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_kmcnt 0x0 +; GFX1310-ISEL-NEXT: global_load_b128 v[0:3], v[0:1], off scope:SCOPE_SE +; GFX1310-ISEL-NEXT: s_wait_loadcnt 0x0 +; GFX1310-ISEL-NEXT: s_set_pc_i64 s[30:31] entry: %data = call <4 x i32> @llvm.amdgcn.av.load.b128.p1(ptr addrspace(1) %addr, metadata !1) ret <4 x i32> %data @@ -210,6 +256,17 @@ define <4 x i32> @av_global_load_b128_0_10(ptr addrspace(1) %addr) { ; GFX1250-SDAG-NEXT: s_wait_loadcnt 0x0 ; GFX1250-SDAG-NEXT: s_set_pc_i64 s[30:31] ; +; GFX1310-SDAG-LABEL: av_global_load_b128_0_10: +; GFX1310-SDAG: ; %bb.0: ; %entry +; GFX1310-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_expcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_samplecnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX1310-SDAG-NEXT: global_load_b128 v[0:3], v[0:1], off scope:SCOPE_DEV +; GFX1310-SDAG-NEXT: s_wait_loadcnt 0x0 +; GFX1310-SDAG-NEXT: s_set_pc_i64 s[30:31] +; ; GFX906-ISEL-LABEL: av_global_load_b128_0_10: ; GFX906-ISEL: ; %bb.0: ; %entry ; GFX906-ISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) @@ -245,6 +302,17 @@ define <4 x i32> @av_global_load_b128_0_10(ptr addrspace(1) %addr) { ; GFX1250-ISEL-NEXT: global_load_b128 v[0:3], v[0:1], off scope:SCOPE_DEV ; GFX1250-ISEL-NEXT: s_wait_loadcnt 0x0 ; GFX1250-ISEL-NEXT: s_set_pc_i64 s[30:31] +; +; GFX1310-ISEL-LABEL: av_global_load_b128_0_10: +; GFX1310-ISEL: ; %bb.0: ; %entry +; GFX1310-ISEL-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_expcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_samplecnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_kmcnt 0x0 +; GFX1310-ISEL-NEXT: global_load_b128 v[0:3], v[0:1], off scope:SCOPE_DEV +; GFX1310-ISEL-NEXT: s_wait_loadcnt 0x0 +; GFX1310-ISEL-NEXT: s_set_pc_i64 s[30:31] entry: %data = call <4 x i32> @llvm.amdgcn.av.load.b128.p1(ptr addrspace(1) %addr, metadata !2) ret <4 x i32> %data @@ -288,6 +356,17 @@ define <4 x i32> @av_global_load_b128_0_11(ptr addrspace(1) %addr) { ; GFX1250-SDAG-NEXT: s_wait_loadcnt 0x0 ; GFX1250-SDAG-NEXT: s_set_pc_i64 s[30:31] ; +; GFX1310-SDAG-LABEL: av_global_load_b128_0_11: +; GFX1310-SDAG: ; %bb.0: ; %entry +; GFX1310-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_expcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_samplecnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX1310-SDAG-NEXT: global_load_b128 v[0:3], v[0:1], off scope:SCOPE_SYS +; GFX1310-SDAG-NEXT: s_wait_loadcnt 0x0 +; GFX1310-SDAG-NEXT: s_set_pc_i64 s[30:31] +; ; GFX906-ISEL-LABEL: av_global_load_b128_0_11: ; GFX906-ISEL: ; %bb.0: ; %entry ; GFX906-ISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) @@ -323,6 +402,17 @@ define <4 x i32> @av_global_load_b128_0_11(ptr addrspace(1) %addr) { ; GFX1250-ISEL-NEXT: global_load_b128 v[0:3], v[0:1], off scope:SCOPE_SYS ; GFX1250-ISEL-NEXT: s_wait_loadcnt 0x0 ; GFX1250-ISEL-NEXT: s_set_pc_i64 s[30:31] +; +; GFX1310-ISEL-LABEL: av_global_load_b128_0_11: +; GFX1310-ISEL: ; %bb.0: ; %entry +; GFX1310-ISEL-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_expcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_samplecnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_kmcnt 0x0 +; GFX1310-ISEL-NEXT: global_load_b128 v[0:3], v[0:1], off scope:SCOPE_SYS +; GFX1310-ISEL-NEXT: s_wait_loadcnt 0x0 +; GFX1310-ISEL-NEXT: s_set_pc_i64 s[30:31] entry: %data = call <4 x i32> @llvm.amdgcn.av.load.b128.p1(ptr addrspace(1) %addr, metadata !3) ret <4 x i32> %data @@ -371,6 +461,18 @@ define <4 x i32> @av_global_load_b128_saddr_0_00(ptr addrspace(1) inreg %addr) { ; GFX1250-SDAG-NEXT: s_wait_loadcnt 0x0 ; GFX1250-SDAG-NEXT: s_set_pc_i64 s[30:31] ; +; GFX1310-SDAG-LABEL: av_global_load_b128_saddr_0_00: +; GFX1310-SDAG: ; %bb.0: ; %entry +; GFX1310-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_expcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_samplecnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX1310-SDAG-NEXT: v_mov_b32_e32 v0, 0 +; GFX1310-SDAG-NEXT: global_load_b128 v[0:3], v0, s[0:1] +; GFX1310-SDAG-NEXT: s_wait_loadcnt 0x0 +; GFX1310-SDAG-NEXT: s_set_pc_i64 s[30:31] +; ; GFX906-ISEL-LABEL: av_global_load_b128_saddr_0_00: ; GFX906-ISEL: ; %bb.0: ; %entry ; GFX906-ISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) @@ -411,6 +513,18 @@ define <4 x i32> @av_global_load_b128_saddr_0_00(ptr addrspace(1) inreg %addr) { ; GFX1250-ISEL-NEXT: global_load_b128 v[0:3], v0, s[0:1] ; GFX1250-ISEL-NEXT: s_wait_loadcnt 0x0 ; GFX1250-ISEL-NEXT: s_set_pc_i64 s[30:31] +; +; GFX1310-ISEL-LABEL: av_global_load_b128_saddr_0_00: +; GFX1310-ISEL: ; %bb.0: ; %entry +; GFX1310-ISEL-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_expcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_samplecnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_kmcnt 0x0 +; GFX1310-ISEL-NEXT: v_mov_b32_e32 v0, 0 +; GFX1310-ISEL-NEXT: global_load_b128 v[0:3], v0, s[0:1] +; GFX1310-ISEL-NEXT: s_wait_loadcnt 0x0 +; GFX1310-ISEL-NEXT: s_set_pc_i64 s[30:31] entry: %data = call <4 x i32> @llvm.amdgcn.av.load.b128.p1(ptr addrspace(1) %addr, metadata !0) ret <4 x i32> %data @@ -459,6 +573,18 @@ define <4 x i32> @av_global_load_b128_saddr_0_01(ptr addrspace(1) inreg %addr) { ; GFX1250-SDAG-NEXT: s_wait_loadcnt 0x0 ; GFX1250-SDAG-NEXT: s_set_pc_i64 s[30:31] ; +; GFX1310-SDAG-LABEL: av_global_load_b128_saddr_0_01: +; GFX1310-SDAG: ; %bb.0: ; %entry +; GFX1310-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_expcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_samplecnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX1310-SDAG-NEXT: v_mov_b32_e32 v0, 0 +; GFX1310-SDAG-NEXT: global_load_b128 v[0:3], v0, s[0:1] scope:SCOPE_SE +; GFX1310-SDAG-NEXT: s_wait_loadcnt 0x0 +; GFX1310-SDAG-NEXT: s_set_pc_i64 s[30:31] +; ; GFX906-ISEL-LABEL: av_global_load_b128_saddr_0_01: ; GFX906-ISEL: ; %bb.0: ; %entry ; GFX906-ISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) @@ -499,6 +625,18 @@ define <4 x i32> @av_global_load_b128_saddr_0_01(ptr addrspace(1) inreg %addr) { ; GFX1250-ISEL-NEXT: global_load_b128 v[0:3], v0, s[0:1] ; GFX1250-ISEL-NEXT: s_wait_loadcnt 0x0 ; GFX1250-ISEL-NEXT: s_set_pc_i64 s[30:31] +; +; GFX1310-ISEL-LABEL: av_global_load_b128_saddr_0_01: +; GFX1310-ISEL: ; %bb.0: ; %entry +; GFX1310-ISEL-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_expcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_samplecnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_kmcnt 0x0 +; GFX1310-ISEL-NEXT: v_mov_b32_e32 v0, 0 +; GFX1310-ISEL-NEXT: global_load_b128 v[0:3], v0, s[0:1] scope:SCOPE_SE +; GFX1310-ISEL-NEXT: s_wait_loadcnt 0x0 +; GFX1310-ISEL-NEXT: s_set_pc_i64 s[30:31] entry: %data = call <4 x i32> @llvm.amdgcn.av.load.b128.p1(ptr addrspace(1) %addr, metadata !1) ret <4 x i32> %data @@ -547,6 +685,18 @@ define <4 x i32> @av_global_load_b128_saddr_0_02(ptr addrspace(1) inreg %addr) { ; GFX1250-SDAG-NEXT: s_wait_loadcnt 0x0 ; GFX1250-SDAG-NEXT: s_set_pc_i64 s[30:31] ; +; GFX1310-SDAG-LABEL: av_global_load_b128_saddr_0_02: +; GFX1310-SDAG: ; %bb.0: ; %entry +; GFX1310-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_expcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_samplecnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX1310-SDAG-NEXT: v_mov_b32_e32 v0, 0 +; GFX1310-SDAG-NEXT: global_load_b128 v[0:3], v0, s[0:1] scope:SCOPE_DEV +; GFX1310-SDAG-NEXT: s_wait_loadcnt 0x0 +; GFX1310-SDAG-NEXT: s_set_pc_i64 s[30:31] +; ; GFX906-ISEL-LABEL: av_global_load_b128_saddr_0_02: ; GFX906-ISEL: ; %bb.0: ; %entry ; GFX906-ISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) @@ -587,6 +737,18 @@ define <4 x i32> @av_global_load_b128_saddr_0_02(ptr addrspace(1) inreg %addr) { ; GFX1250-ISEL-NEXT: global_load_b128 v[0:3], v0, s[0:1] scope:SCOPE_DEV ; GFX1250-ISEL-NEXT: s_wait_loadcnt 0x0 ; GFX1250-ISEL-NEXT: s_set_pc_i64 s[30:31] +; +; GFX1310-ISEL-LABEL: av_global_load_b128_saddr_0_02: +; GFX1310-ISEL: ; %bb.0: ; %entry +; GFX1310-ISEL-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_expcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_samplecnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_kmcnt 0x0 +; GFX1310-ISEL-NEXT: v_mov_b32_e32 v0, 0 +; GFX1310-ISEL-NEXT: global_load_b128 v[0:3], v0, s[0:1] scope:SCOPE_DEV +; GFX1310-ISEL-NEXT: s_wait_loadcnt 0x0 +; GFX1310-ISEL-NEXT: s_set_pc_i64 s[30:31] entry: %data = call <4 x i32> @llvm.amdgcn.av.load.b128.p1(ptr addrspace(1) %addr, metadata !2) ret <4 x i32> %data @@ -635,6 +797,18 @@ define <4 x i32> @av_global_load_b128_saddr_0_03(ptr addrspace(1) inreg %addr) { ; GFX1250-SDAG-NEXT: s_wait_loadcnt 0x0 ; GFX1250-SDAG-NEXT: s_set_pc_i64 s[30:31] ; +; GFX1310-SDAG-LABEL: av_global_load_b128_saddr_0_03: +; GFX1310-SDAG: ; %bb.0: ; %entry +; GFX1310-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_expcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_samplecnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX1310-SDAG-NEXT: v_mov_b32_e32 v0, 0 +; GFX1310-SDAG-NEXT: global_load_b128 v[0:3], v0, s[0:1] scope:SCOPE_SYS +; GFX1310-SDAG-NEXT: s_wait_loadcnt 0x0 +; GFX1310-SDAG-NEXT: s_set_pc_i64 s[30:31] +; ; GFX906-ISEL-LABEL: av_global_load_b128_saddr_0_03: ; GFX906-ISEL: ; %bb.0: ; %entry ; GFX906-ISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) @@ -675,6 +849,18 @@ define <4 x i32> @av_global_load_b128_saddr_0_03(ptr addrspace(1) inreg %addr) { ; GFX1250-ISEL-NEXT: global_load_b128 v[0:3], v0, s[0:1] scope:SCOPE_SYS ; GFX1250-ISEL-NEXT: s_wait_loadcnt 0x0 ; GFX1250-ISEL-NEXT: s_set_pc_i64 s[30:31] +; +; GFX1310-ISEL-LABEL: av_global_load_b128_saddr_0_03: +; GFX1310-ISEL: ; %bb.0: ; %entry +; GFX1310-ISEL-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_expcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_samplecnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_kmcnt 0x0 +; GFX1310-ISEL-NEXT: v_mov_b32_e32 v0, 0 +; GFX1310-ISEL-NEXT: global_load_b128 v[0:3], v0, s[0:1] scope:SCOPE_SYS +; GFX1310-ISEL-NEXT: s_wait_loadcnt 0x0 +; GFX1310-ISEL-NEXT: s_set_pc_i64 s[30:31] entry: %data = call <4 x i32> @llvm.amdgcn.av.load.b128.p1(ptr addrspace(1) %addr, metadata !3) ret <4 x i32> %data @@ -725,6 +911,17 @@ define <4 x float> @global_load_i8_offset_0(ptr addrspace(1) %sbase) { ; GFX1250-SDAG-NEXT: s_wait_loadcnt 0x0 ; GFX1250-SDAG-NEXT: s_set_pc_i64 s[30:31] ; +; GFX1310-SDAG-LABEL: global_load_i8_offset_0: +; GFX1310-SDAG: ; %bb.0: +; GFX1310-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_expcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_samplecnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX1310-SDAG-NEXT: global_load_b128 v[0:3], v[0:1], off +; GFX1310-SDAG-NEXT: s_wait_loadcnt 0x0 +; GFX1310-SDAG-NEXT: s_set_pc_i64 s[30:31] +; ; GFX906-ISEL-LABEL: global_load_i8_offset_0: ; GFX906-ISEL: ; %bb.0: ; GFX906-ISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) @@ -760,6 +957,17 @@ define <4 x float> @global_load_i8_offset_0(ptr addrspace(1) %sbase) { ; GFX1250-ISEL-NEXT: global_load_b128 v[0:3], v[0:1], off ; GFX1250-ISEL-NEXT: s_wait_loadcnt 0x0 ; GFX1250-ISEL-NEXT: s_set_pc_i64 s[30:31] +; +; GFX1310-ISEL-LABEL: global_load_i8_offset_0: +; GFX1310-ISEL: ; %bb.0: +; GFX1310-ISEL-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_expcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_samplecnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_kmcnt 0x0 +; GFX1310-ISEL-NEXT: global_load_b128 v[0:3], v[0:1], off +; GFX1310-ISEL-NEXT: s_wait_loadcnt 0x0 +; GFX1310-ISEL-NEXT: s_set_pc_i64 s[30:31] %load = call <4 x i32> @llvm.amdgcn.av.load.b128.p1(ptr addrspace(1) %sbase, metadata !0) %cast.load = bitcast <4 x i32> %load to <4 x float> ret <4 x float> %cast.load @@ -805,6 +1013,17 @@ define <4 x float> @global_load_i8_offset_4095(ptr addrspace(1) %sbase) { ; GFX1250-SDAG-NEXT: s_wait_loadcnt 0x0 ; GFX1250-SDAG-NEXT: s_set_pc_i64 s[30:31] ; +; GFX1310-SDAG-LABEL: global_load_i8_offset_4095: +; GFX1310-SDAG: ; %bb.0: +; GFX1310-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_expcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_samplecnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX1310-SDAG-NEXT: global_load_b128 v[0:3], v[0:1], off offset:4095 scope:SCOPE_SE +; GFX1310-SDAG-NEXT: s_wait_loadcnt 0x0 +; GFX1310-SDAG-NEXT: s_set_pc_i64 s[30:31] +; ; GFX906-ISEL-LABEL: global_load_i8_offset_4095: ; GFX906-ISEL: ; %bb.0: ; GFX906-ISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) @@ -842,6 +1061,17 @@ define <4 x float> @global_load_i8_offset_4095(ptr addrspace(1) %sbase) { ; GFX1250-ISEL-NEXT: global_load_b128 v[0:3], v[0:1], off offset:4095 ; GFX1250-ISEL-NEXT: s_wait_loadcnt 0x0 ; GFX1250-ISEL-NEXT: s_set_pc_i64 s[30:31] +; +; GFX1310-ISEL-LABEL: global_load_i8_offset_4095: +; GFX1310-ISEL: ; %bb.0: +; GFX1310-ISEL-NEXT: s_wait_loadcnt_dscnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_expcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_samplecnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_bvhcnt 0x0 +; GFX1310-ISEL-NEXT: s_wait_kmcnt 0x0 +; GFX1310-ISEL-NEXT: global_load_b128 v[0:3], v[0:1], off offset:4095 scope:SCOPE_SE +; GFX1310-ISEL-NEXT: s_wait_loadcnt 0x0 +; GFX1310-ISEL-NEXT: s_set_pc_i64 s[30:31] %gep0 = getelementptr inbounds i8, ptr addrspace(1) %sbase, i64 4095 %load = call <4 x i32> @llvm.amdgcn.av.load.b128.p1(ptr addrspace(1) %gep0, metadata !1) %cast.load = bitcast <4 x i32> %load to <4 x float> @@ -896,6 +1126,17 @@ define <4 x float> @global_load_i8_offset_4096(ptr addrspace(1) %sbase) { ; GFX1250-SDAG-NEXT: s_wait_loadcnt 0x0 ; GFX1250-SDAG-NEXT: s_set_pc_i64 s[30:31] ; +; GFX1310-SDAG-LABEL: global_load_i8_offset_4096: +; GFX1310-SDAG: ; %... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/213121 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
