[llvm-branch-commits] [clang] [llvm] [AMDGPU] Add intrinsics and builtins for gfx13 exclusive scan (PR #226078)
via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Sep 24 02:00:43 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Mariusz Sikora (mariusz-sikora-at-amd)
<details>
<summary>Changes</summary>
---
Patch is 109.85 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/226078.diff
11 Files Affected:
- (modified) clang/include/clang/Basic/BuiltinsAMDGPU.td (+21)
- (modified) clang/include/clang/Basic/BuiltinsAMDGPUDocs.td (+20)
- (modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl (+16)
- (modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl (+260)
- (added) clang/test/SemaOpenCL/builtins-amdgcn-gfx13-exclusive-scan-types.cl (+42)
- (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+67)
- (modified) llvm/lib/Target/AMDGPU/AMDGPU.td (+1-1)
- (modified) llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp (+21)
- (modified) llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp (+13)
- (modified) llvm/lib/Target/AMDGPU/VOP3Instructions.td (+20-18)
- (added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.exclusive.scan.ll (+1594)
``````````diff
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index 4fd604390d3ce..ef789e1aa0028 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -928,6 +928,27 @@ def __builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64 : AMDGPUBuiltin<"_ExtVector
def __builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64 : AMDGPUBuiltin<"_ExtVector<4, float>(int, _ExtVector<2, int>, _ExtVector<4, float>, int)", [Const], "swmmac-gfx1200-insts,wavefrontsize64">;
def __builtin_amdgcn_prng_b32 : AMDGPUBuiltin<"unsigned int(unsigned int)", [Const], "prng-inst">;
+
+class AMDGPUExclusiveScanBuiltin<string prototype, list<string> argNames>
+ : AMDGPUBuiltin<prototype, [Const], "exclusive-scan-insts"> {
+ let Documentation = [DocExclusiveScan];
+ let ArgNames = argNames;
+}
+
+def __builtin_amdgcn_exclusive_scan_sum_i32 : AMDGPUExclusiveScanBuiltin<"int(int, int, bool)", ["src0", "src1", "clamp"]>;
+def __builtin_amdgcn_exclusive_scan_sum_u32 : AMDGPUExclusiveScanBuiltin<"unsigned int(unsigned int, unsigned int, bool)", ["src0", "src1", "clamp"]>;
+def __builtin_amdgcn_exclusive_scan_xor_b32 : AMDGPUExclusiveScanBuiltin<"int(int, int)", ["src0", "src1"]>;
+def __builtin_amdgcn_exclusive_scan_or_b32 : AMDGPUExclusiveScanBuiltin<"int(int, int)", ["src0", "src1"]>;
+def __builtin_amdgcn_exclusive_scan_and_b32 : AMDGPUExclusiveScanBuiltin<"int(int, int)", ["src0", "src1"]>;
+def __builtin_amdgcn_exclusive_scan_min_i16 : AMDGPUExclusiveScanBuiltin<"short(short, int)", ["src0", "src1"]>;
+def __builtin_amdgcn_exclusive_scan_min_u16 : AMDGPUExclusiveScanBuiltin<"unsigned short(unsigned short, int)", ["src0", "src1"]>;
+def __builtin_amdgcn_exclusive_scan_min_i32 : AMDGPUExclusiveScanBuiltin<"int(int, int)", ["src0", "src1"]>;
+def __builtin_amdgcn_exclusive_scan_min_u32 : AMDGPUExclusiveScanBuiltin<"unsigned int(unsigned int, unsigned int)", ["src0", "src1"]>;
+def __builtin_amdgcn_exclusive_scan_max_i16 : AMDGPUExclusiveScanBuiltin<"short(short, int)", ["src0", "src1"]>;
+def __builtin_amdgcn_exclusive_scan_max_u16 : AMDGPUExclusiveScanBuiltin<"unsigned short(unsigned short, int)", ["src0", "src1"]>;
+def __builtin_amdgcn_exclusive_scan_max_i32 : AMDGPUExclusiveScanBuiltin<"int(int, int)", ["src0", "src1"]>;
+def __builtin_amdgcn_exclusive_scan_max_u32 : AMDGPUExclusiveScanBuiltin<"unsigned int(unsigned int, unsigned int)", ["src0", "src1"]>;
+
def __builtin_amdgcn_cvt_scalef32_pk32_fp6_f16 : AMDGPUBuiltin<"_ExtVector<6, unsigned int>(_ExtVector<32, _Float16>, float)", [Const], "f16bf16-to-fp6bf6-cvt-scale-insts">;
def __builtin_amdgcn_cvt_scalef32_pk32_bf6_f16 : AMDGPUBuiltin<"_ExtVector<6, unsigned int>(_ExtVector<32, _Float16>, float)", [Const], "f16bf16-to-fp6bf6-cvt-scale-insts">;
def __builtin_amdgcn_cvt_scalef32_pk32_fp6_bf16 : AMDGPUBuiltin<"_ExtVector<6, unsigned int>(_ExtVector<32, __bf16>, float)", [Const], "f16bf16-to-fp6bf6-cvt-scale-insts">;
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
index bfccd7de6590a..b9b54cce4f28f 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
@@ -653,6 +653,26 @@ nondeterministic value.
}];
}
+def DocExclusiveScan : Documentation {
+ let Category = DocCatWaveDataExchange;
+ let Content = [{
+Performs an exclusive prefix scan of ``src0`` within a subgroup of lanes. Each
+lane receives the running combination of the ``src0`` values of all earlier
+lanes in its subgroup, walking from the lowest to the highest lane index.
+"Exclusive" means a lane's own ``src0`` is not included in its result, so the
+first (lowest) lane of each subgroup receives the operation's identity value.
+
+``src1`` is a 32-bit subgroup mask: bit N is set when lane N belongs to the same
+subgroup as the current lane.
+
+The combining operation is named by the builtin (``sum``, ``xor``, ``or``,
+``and``, ``min``, ``max``).
+
+In wave64 mode the two 32-lane halves of the wave are scanned independently. The
+builtins are convergent: all lanes of a subgroup must execute them together.
+}];
+}
+
def DocWMMA_scale16_GFX1250 : Documentation {
let Category = DocCatWMMA;
let Content = [{
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl
index d8fb243865f76..8bc07bf034128 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl
@@ -3,6 +3,7 @@
// RUN: %clang_cc1 -triple amdgpu12.00-unknown-unknown -verify -emit-llvm -o - %s
typedef unsigned int uint;
+typedef unsigned short ushort;
#pragma OPENCL EXTENSION cl_khr_fp64:enable
@@ -32,6 +33,7 @@ void builtin_test_unsupported(double a_double, float a_float,
v2i a_v2i, v4i a_v4i, v16i a_v16i, v32i a_v32i,
v2f a_v2f, v4f a_v4f, v16f a_v16f, v32f a_v32f,
v4h a_v4h, v8h a_v8h,
+ short a_short, ushort a_ushort,
uint a, uint b) {
@@ -94,4 +96,18 @@ void builtin_test_unsupported(double a_double, float a_float,
a_v16f = __builtin_amdgcn_smfmac_f32_32x32x32_bf8_fp8(a_v2i, a_v4i, a_v16f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_32x32x32_bf8_fp8' needs target feature fp8-insts}}
a_v16f = __builtin_amdgcn_smfmac_f32_32x32x32_fp8_bf8(a_v2i, a_v4i, a_v16f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_32x32x32_fp8_bf8' needs target feature fp8-insts}}
a_v16f = __builtin_amdgcn_smfmac_f32_32x32x32_fp8_fp8(a_v2i, a_v4i, a_v16f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_32x32x32_fp8_fp8' needs target feature fp8-insts}}
+
+ a_int = __builtin_amdgcn_exclusive_scan_sum_i32(a_int, a_int, true); // expected-error {{'__builtin_amdgcn_exclusive_scan_sum_i32' needs target feature exclusive-scan-insts}}
+ a = __builtin_amdgcn_exclusive_scan_sum_u32(a, a, true); // expected-error {{'__builtin_amdgcn_exclusive_scan_sum_u32' needs target feature exclusive-scan-insts}}
+ a_int = __builtin_amdgcn_exclusive_scan_xor_b32(a_int, a_int); // expected-error {{'__builtin_amdgcn_exclusive_scan_xor_b32' needs target feature exclusive-scan-insts}}
+ a_int = __builtin_amdgcn_exclusive_scan_or_b32(a_int, a_int); // expected-error {{'__builtin_amdgcn_exclusive_scan_or_b32' needs target feature exclusive-scan-insts}}
+ a_int = __builtin_amdgcn_exclusive_scan_and_b32(a_int, a_int); // expected-error {{'__builtin_amdgcn_exclusive_scan_and_b32' needs target feature exclusive-scan-insts}}
+ a_short = __builtin_amdgcn_exclusive_scan_min_i16(a_short, a_int); // expected-error {{'__builtin_amdgcn_exclusive_scan_min_i16' needs target feature exclusive-scan-insts}}
+ a_ushort = __builtin_amdgcn_exclusive_scan_min_u16(a_ushort, a_int); // expected-error {{'__builtin_amdgcn_exclusive_scan_min_u16' needs target feature exclusive-scan-insts}}
+ a_int = __builtin_amdgcn_exclusive_scan_min_i32(a_int, a_int); // expected-error {{'__builtin_amdgcn_exclusive_scan_min_i32' needs target feature exclusive-scan-insts}}
+ a = __builtin_amdgcn_exclusive_scan_min_u32(a, a); // expected-error {{'__builtin_amdgcn_exclusive_scan_min_u32' needs target feature exclusive-scan-insts}}
+ a_short = __builtin_amdgcn_exclusive_scan_max_i16(a_short, a_int); // expected-error {{'__builtin_amdgcn_exclusive_scan_max_i16' needs target feature exclusive-scan-insts}}
+ a_ushort = __builtin_amdgcn_exclusive_scan_max_u16(a_ushort, a_int); // expected-error {{'__builtin_amdgcn_exclusive_scan_max_u16' needs target feature exclusive-scan-insts}}
+ a_int = __builtin_amdgcn_exclusive_scan_max_i32(a_int, a_int); // expected-error {{'__builtin_amdgcn_exclusive_scan_max_i32' needs target feature exclusive-scan-insts}}
+ a = __builtin_amdgcn_exclusive_scan_max_u32(a, a); // expected-error {{'__builtin_amdgcn_exclusive_scan_max_u32' needs target feature exclusive-scan-insts}}
}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
index 00e65563a1868..21306cb3a083a 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
@@ -8,6 +8,7 @@ typedef unsigned int __attribute__((ext_vector_type(6))) uint6;
typedef float __attribute__((ext_vector_type(32))) float32;
typedef __bf16 __attribute__((ext_vector_type(2))) bfloat2;
typedef unsigned int uint;
+typedef unsigned short ushort;
// CHECK-LABEL: @test_cvt_scalef32_pk32_bf6_f32(
// CHECK-NEXT: entry:
@@ -211,3 +212,262 @@ void test_s_prefetch_data(global float *gp, unsigned int len)
__builtin_amdgcn_s_prefetch_data(gp, len);
}
+
+// CHECK-LABEL: @test_exclusive_scan_sum_i32(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
+// CHECK-NEXT: [[SRC0_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[SRC1_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[SRC0:%.*]], ptr addrspace(5) [[SRC0_ADDR]], align 4
+// CHECK-NEXT: store i32 [[SRC1:%.*]], ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[SRC0_ADDR]], align 4
+// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.exclusive.scan.sum.i32(i32 [[TMP0]], i32 [[TMP1]], i1 true)
+// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[TMP3]], align 4
+// CHECK-NEXT: [[TMP4:%.*]] = load i32, ptr addrspace(5) [[SRC0_ADDR]], align 4
+// CHECK-NEXT: [[TMP5:%.*]] = load i32, ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.exclusive.scan.sum.i32(i32 [[TMP4]], i32 [[TMP5]], i1 false)
+// CHECK-NEXT: [[TMP7:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP6]], ptr addrspace(1) [[TMP7]], align 4
+// CHECK-NEXT: ret void
+//
+void test_exclusive_scan_sum_i32(global int* out, int src0, uint src1) {
+ *out = __builtin_amdgcn_exclusive_scan_sum_i32(src0, src1, true);
+ *out = __builtin_amdgcn_exclusive_scan_sum_i32(src0, src1, false);
+}
+
+// CHECK-LABEL: @test_exclusive_scan_sum_u32(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
+// CHECK-NEXT: [[SRC0_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[SRC1_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[SRC0:%.*]], ptr addrspace(5) [[SRC0_ADDR]], align 4
+// CHECK-NEXT: store i32 [[SRC1:%.*]], ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[SRC0_ADDR]], align 4
+// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.exclusive.scan.sum.u32(i32 [[TMP0]], i32 [[TMP1]], i1 true)
+// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[TMP3]], align 4
+// CHECK-NEXT: [[TMP4:%.*]] = load i32, ptr addrspace(5) [[SRC0_ADDR]], align 4
+// CHECK-NEXT: [[TMP5:%.*]] = load i32, ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.exclusive.scan.sum.u32(i32 [[TMP4]], i32 [[TMP5]], i1 false)
+// CHECK-NEXT: [[TMP7:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP6]], ptr addrspace(1) [[TMP7]], align 4
+// CHECK-NEXT: ret void
+//
+void test_exclusive_scan_sum_u32(global uint* out, uint src0, uint src1) {
+ *out = __builtin_amdgcn_exclusive_scan_sum_u32(src0, src1, true);
+ *out = __builtin_amdgcn_exclusive_scan_sum_u32(src0, src1, false);
+}
+
+// CHECK-LABEL: @test_exclusive_scan_xor_b32(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
+// CHECK-NEXT: [[SRC0_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[SRC1_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[SRC0:%.*]], ptr addrspace(5) [[SRC0_ADDR]], align 4
+// CHECK-NEXT: store i32 [[SRC1:%.*]], ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[SRC0_ADDR]], align 4
+// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.exclusive.scan.xor.b32(i32 [[TMP0]], i32 [[TMP1]])
+// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[TMP3]], align 4
+// CHECK-NEXT: ret void
+//
+void test_exclusive_scan_xor_b32(global uint* out, uint src0, uint src1) {
+ *out = __builtin_amdgcn_exclusive_scan_xor_b32(src0, src1);
+}
+
+// CHECK-LABEL: @test_exclusive_scan_or_b32(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
+// CHECK-NEXT: [[SRC0_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[SRC1_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[SRC0:%.*]], ptr addrspace(5) [[SRC0_ADDR]], align 4
+// CHECK-NEXT: store i32 [[SRC1:%.*]], ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[SRC0_ADDR]], align 4
+// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.exclusive.scan.or.b32(i32 [[TMP0]], i32 [[TMP1]])
+// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[TMP3]], align 4
+// CHECK-NEXT: ret void
+//
+void test_exclusive_scan_or_b32(global uint* out, uint src0, uint src1) {
+ *out = __builtin_amdgcn_exclusive_scan_or_b32(src0, src1);
+}
+
+// CHECK-LABEL: @test_exclusive_scan_and_b32(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
+// CHECK-NEXT: [[SRC0_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[SRC1_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[SRC0:%.*]], ptr addrspace(5) [[SRC0_ADDR]], align 4
+// CHECK-NEXT: store i32 [[SRC1:%.*]], ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[SRC0_ADDR]], align 4
+// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.exclusive.scan.and.b32(i32 [[TMP0]], i32 [[TMP1]])
+// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[TMP3]], align 4
+// CHECK-NEXT: ret void
+//
+void test_exclusive_scan_and_b32(global uint* out, uint src0, uint src1) {
+ *out = __builtin_amdgcn_exclusive_scan_and_b32(src0, src1);
+}
+
+// CHECK-LABEL: @test_exclusive_scan_min_i16(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
+// CHECK-NEXT: [[SRC0_ADDR:%.*]] = alloca i16, align 2, addrspace(5)
+// CHECK-NEXT: [[SRC1_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i16 [[SRC0:%.*]], ptr addrspace(5) [[SRC0_ADDR]], align 2
+// CHECK-NEXT: store i32 [[SRC1:%.*]], ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load i16, ptr addrspace(5) [[SRC0_ADDR]], align 2
+// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = call i16 @llvm.amdgcn.exclusive.scan.min.i16(i16 [[TMP0]], i32 [[TMP1]])
+// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i16 [[TMP2]], ptr addrspace(1) [[TMP3]], align 2
+// CHECK-NEXT: ret void
+//
+void test_exclusive_scan_min_i16(global short* out, short src0, uint src1) {
+ *out = __builtin_amdgcn_exclusive_scan_min_i16(src0, src1);
+}
+
+// CHECK-LABEL: @test_exclusive_scan_min_u16(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
+// CHECK-NEXT: [[SRC0_ADDR:%.*]] = alloca i16, align 2, addrspace(5)
+// CHECK-NEXT: [[SRC1_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i16 [[SRC0:%.*]], ptr addrspace(5) [[SRC0_ADDR]], align 2
+// CHECK-NEXT: store i32 [[SRC1:%.*]], ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load i16, ptr addrspace(5) [[SRC0_ADDR]], align 2
+// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = call i16 @llvm.amdgcn.exclusive.scan.min.u16(i16 [[TMP0]], i32 [[TMP1]])
+// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i16 [[TMP2]], ptr addrspace(1) [[TMP3]], align 2
+// CHECK-NEXT: ret void
+//
+void test_exclusive_scan_min_u16(global ushort* out, ushort src0, uint src1) {
+ *out = __builtin_amdgcn_exclusive_scan_min_u16(src0, src1);
+}
+
+// CHECK-LABEL: @test_exclusive_scan_min_i32(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
+// CHECK-NEXT: [[SRC0_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[SRC1_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[SRC0:%.*]], ptr addrspace(5) [[SRC0_ADDR]], align 4
+// CHECK-NEXT: store i32 [[SRC1:%.*]], ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[SRC0_ADDR]], align 4
+// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.exclusive.scan.min.i32(i32 [[TMP0]], i32 [[TMP1]])
+// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[TMP3]], align 4
+// CHECK-NEXT: ret void
+//
+void test_exclusive_scan_min_i32(global int* out, int src0, uint src1) {
+ *out = __builtin_amdgcn_exclusive_scan_min_i32(src0, src1);
+}
+
+// CHECK-LABEL: @test_exclusive_scan_min_u32(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
+// CHECK-NEXT: [[SRC0_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[SRC1_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[SRC0:%.*]], ptr addrspace(5) [[SRC0_ADDR]], align 4
+// CHECK-NEXT: store i32 [[SRC1:%.*]], ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[SRC0_ADDR]], align 4
+// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.exclusive.scan.min.u32(i32 [[TMP0]], i32 [[TMP1]])
+// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[TMP3]], align 4
+// CHECK-NEXT: ret void
+//
+void test_exclusive_scan_min_u32(global uint* out, uint src0, uint src1) ...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/226078
More information about the llvm-branch-commits
mailing list