[clang] [llvm] AMDGPU: Add missing flat-global-insts to gfx13 frontend feature map (PR #213121)
via cfe-commits
cfe-commits at lists.llvm.org
Thu Jul 30 12:51:25 PDT 2026
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
More information about the cfe-commits
mailing list