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 _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
