[clang] [AMDGPU] Use generic builtins for `wave_reduce` ops (PR #179589)
Alex Voicu via cfe-commits
cfe-commits at lists.llvm.org
Tue Feb 3 17:49:10 PST 2026
https://github.com/AlexVlx updated https://github.com/llvm/llvm-project/pull/179589
>From 24d5fa6eba0855a3044125003960b0240135a0b2 Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu at amd.com>
Date: Wed, 4 Feb 2026 01:12:39 +0000
Subject: [PATCH 1/2] Use generic builtins for wave_reduce ops.
---
clang/include/clang/Basic/BuiltinsAMDGPU.td | 34 +---
clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 102 ++++------
clang/lib/Sema/SemaAMDGPU.cpp | 36 ++++
clang/lib/Sema/SemaChecking.cpp | 5 +-
clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 180 +++++++++---------
.../wave-reduce-builtins-validate-amdgpu.cl | 142 +++++++-------
clang/tools/c-index-test/CMakeLists.txt | 2 +-
clang/tools/libclang/CIndexDiagnostic.cpp | 40 ++--
8 files changed, 268 insertions(+), 273 deletions(-)
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index 1950757097fc6..41d4013a4c4a7 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -386,32 +386,14 @@ def __builtin_amdgcn_set_fpenv : AMDGPUBuiltin<"void(uint64_t)">;
//===----------------------------------------------------------------------===//
-def __builtin_amdgcn_wave_reduce_add_u32 : AMDGPUBuiltin<"uint32_t(uint32_t, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_sub_u32 : AMDGPUBuiltin<"uint32_t(uint32_t, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_min_i32 : AMDGPUBuiltin<"int32_t(int32_t, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_min_u32 : AMDGPUBuiltin<"uint32_t(uint32_t, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_max_i32 : AMDGPUBuiltin<"int32_t(int32_t, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_max_u32 : AMDGPUBuiltin<"uint32_t(uint32_t, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_and_b32 : AMDGPUBuiltin<"int32_t(int32_t, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_or_b32 : AMDGPUBuiltin<"int32_t(int32_t, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_xor_b32 : AMDGPUBuiltin<"int32_t(int32_t, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_add_u64 : AMDGPUBuiltin<"uint64_t(uint64_t, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_sub_u64 : AMDGPUBuiltin<"uint64_t(uint64_t, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_min_i64 : AMDGPUBuiltin<"int64_t(int64_t, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_min_u64 : AMDGPUBuiltin<"uint64_t(uint64_t, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_max_i64 : AMDGPUBuiltin<"int64_t(int64_t, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_max_u64 : AMDGPUBuiltin<"uint64_t(uint64_t, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_and_b64 : AMDGPUBuiltin<"int64_t(int64_t, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_or_b64 : AMDGPUBuiltin<"int64_t(int64_t, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_xor_b64 : AMDGPUBuiltin<"int64_t(int64_t, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_fadd_f32 : AMDGPUBuiltin<"float(float, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_fsub_f32 : AMDGPUBuiltin<"float(float, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_fmin_f32 : AMDGPUBuiltin<"float(float, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_fmax_f32 : AMDGPUBuiltin<"float(float, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_fadd_f64 : AMDGPUBuiltin<"double(double, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_fsub_f64 : AMDGPUBuiltin<"double(double, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_fmin_f64 : AMDGPUBuiltin<"double(double, _Constant int32_t)", [Const]>;
-def __builtin_amdgcn_wave_reduce_fmax_f64 : AMDGPUBuiltin<"double(double, _Constant int32_t)", [Const]>;
+// These are overloaded builtins modelled after the atomic ones
+def __builtin_amdgcn_wave_reduce_add : AMDGPUBuiltin<"void(...)", [Const, CustomTypeChecking]>;
+def __builtin_amdgcn_wave_reduce_sub : AMDGPUBuiltin<"void(...)", [Const, CustomTypeChecking]>;
+def __builtin_amdgcn_wave_reduce_min : AMDGPUBuiltin<"void(...)", [Const, CustomTypeChecking]>;
+def __builtin_amdgcn_wave_reduce_max : AMDGPUBuiltin<"void(...)", [Const, CustomTypeChecking]>;
+def __builtin_amdgcn_wave_reduce_and : AMDGPUBuiltin<"void(...)", [Const, CustomTypeChecking]>;
+def __builtin_amdgcn_wave_reduce_or : AMDGPUBuiltin<"void(...)", [Const, CustomTypeChecking]>;
+def __builtin_amdgcn_wave_reduce_xor : AMDGPUBuiltin<"void(...)", [Const, CustomTypeChecking]>;
//===----------------------------------------------------------------------===//
// R600-NI only builtins.
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index a096ed27a788e..e47eaef3651c6 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -366,49 +366,31 @@ void CodeGenFunction::AddAMDGPUFenceAddressSpaceMMRA(llvm::Instruction *Inst,
Inst->setMetadata(LLVMContext::MD_mmra, MMRAMetadata::getMD(Ctx, MMRAs));
}
-static Intrinsic::ID getIntrinsicIDforWaveReduction(unsigned BuiltinID) {
+static Intrinsic::ID getIntrinsicIDforWaveReduction(unsigned BuiltinID,
+ bool IsFP, bool IsSigned) {
switch (BuiltinID) {
- default:
- llvm_unreachable("Unknown BuiltinID for wave reduction");
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u32:
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u64:
- return Intrinsic::amdgcn_wave_reduce_add;
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_fadd_f32:
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_fadd_f64:
- return Intrinsic::amdgcn_wave_reduce_fadd;
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u32:
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u64:
- return Intrinsic::amdgcn_wave_reduce_sub;
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_fsub_f32:
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_fsub_f64:
- return Intrinsic::amdgcn_wave_reduce_fsub;
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i32:
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i64:
- return Intrinsic::amdgcn_wave_reduce_min;
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_fmin_f32:
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_fmin_f64:
- return Intrinsic::amdgcn_wave_reduce_fmin;
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u32:
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u64:
- return Intrinsic::amdgcn_wave_reduce_umin;
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i32:
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i64:
- return Intrinsic::amdgcn_wave_reduce_max;
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_fmax_f32:
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_fmax_f64:
- return Intrinsic::amdgcn_wave_reduce_fmax;
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u32:
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u64:
- return Intrinsic::amdgcn_wave_reduce_umax;
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b32:
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b64:
+ case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_add:
+ return IsFP ?
+ Intrinsic::amdgcn_wave_reduce_fadd : Intrinsic::amdgcn_wave_reduce_add;
+ case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_and:
return Intrinsic::amdgcn_wave_reduce_and;
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b32:
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b64:
+ case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_max:
+ return IsFP ? Intrinsic::amdgcn_wave_reduce_fmax
+ : (IsSigned ? Intrinsic::amdgcn_wave_reduce_max
+ : Intrinsic::amdgcn_wave_reduce_umax);
+ case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_min:
+ return IsFP ? Intrinsic::amdgcn_wave_reduce_fmin
+ : (IsSigned ? Intrinsic::amdgcn_wave_reduce_min
+ : Intrinsic::amdgcn_wave_reduce_umin);
+ case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_or:
return Intrinsic::amdgcn_wave_reduce_or;
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b32:
- case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b64:
+ case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_sub:
+ return IsFP ?
+ Intrinsic::amdgcn_wave_reduce_fsub : Intrinsic::amdgcn_wave_reduce_sub;
+ case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_xor:
return Intrinsic::amdgcn_wave_reduce_xor;
+ default:
+ llvm_unreachable("Unknown BuiltinID for wave reduction");
}
}
@@ -417,37 +399,17 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
llvm::AtomicOrdering AO = llvm::AtomicOrdering::SequentiallyConsistent;
llvm::SyncScope::ID SSID;
switch (BuiltinID) {
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u32:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_fadd_f32:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_fadd_f64:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u32:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_fsub_f32:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_fsub_f64:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i32:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u32:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_fmin_f32:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_fmin_f64:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i32:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u32:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_fmax_f32:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_fmax_f64:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b32:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b32:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b32:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u64:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u64:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i64:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u64:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i64:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u64:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b64:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b64:
- case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b64: {
- Intrinsic::ID IID = getIntrinsicIDforWaveReduction(BuiltinID);
- llvm::Value *Value = EmitScalarExpr(E->getArg(0));
- llvm::Value *Strategy = EmitScalarExpr(E->getArg(1));
- llvm::Function *F = CGM.getIntrinsic(IID, {Value->getType()});
- return Builder.CreateCall(F, {Value, Strategy});
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_add:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_and:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_max:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_min:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_or:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor: {
+ bool IsFP = E->getArg(0)->getType()->isRealFloatingType();
+ bool IsSigned = E->getArg(0)->getType()->isSignedIntegerType();
+ return emitBuiltinWithOneOverloadedType<2>(
+ *this, E, getIntrinsicIDforWaveReduction(BuiltinID, IsFP, IsSigned));
}
case AMDGPU::BI__builtin_amdgcn_div_scale:
case AMDGPU::BI__builtin_amdgcn_div_scalef: {
diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp
index 4261e1849133f..80949d33b794a 100644
--- a/clang/lib/Sema/SemaAMDGPU.cpp
+++ b/clang/lib/Sema/SemaAMDGPU.cpp
@@ -296,6 +296,42 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
}
return false;
}
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_add:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_and:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_max:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_min:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_or:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub:
+ case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor: {
+ Expr *Val = TheCall->getArg(0);
+ if (Val->isTypeDependent())
+ return true;
+
+ QualType ValTy = Val->getType();
+
+ if (!ValTy->isIntegerType() && !ValTy->isFloatingType())
+ return Diag(Val->getExprLoc(), diag::err_builtin_invalid_arg_type)
+ << Val << /*scalar=*/1 << /*'int'=*/1 << /*floating point=*/1
+ << ValTy;
+
+ Expr *Strategy = TheCall->getArg(1);
+ if (Strategy->isTypeDependent() || Strategy->isValueDependent())
+ return true;
+
+ llvm::APSInt SVal;
+ if (!SemaRef.VerifyIntegerConstantExpression(Strategy, &SVal).isUsable())
+ return true;
+
+ if (SVal.getZExtValue() > 2)
+ return Diag(Strategy->getExprLoc(), diag::err_argument_invalid_range)
+ << SVal.getZExtValue() << 0 << 2;
+
+ // Resolve the overload here, now that we know that the invocation is
+ // correct: the intrinsic returns the type of the value argument.
+ TheCall->setType(ValTy);
+
+ return false;
+ }
default:
return false;
}
diff --git a/clang/lib/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp
index 3f19b6ad16c13..9364e63e37cce 100644
--- a/clang/lib/Sema/SemaChecking.cpp
+++ b/clang/lib/Sema/SemaChecking.cpp
@@ -2126,9 +2126,10 @@ bool Sema::CheckTSBuiltinFunctionCall(const TargetInfo &TI, unsigned BuiltinID,
case llvm::Triple::spirv:
case llvm::Triple::spirv32:
case llvm::Triple::spirv64:
- if (TI.getTriple().getOS() != llvm::Triple::OSType::AMDHSA)
+ if (TI.getTriple().getVendor() != llvm::Triple::VendorType::AMD)
return SPIRV().CheckSPIRVBuiltinFunctionCall(TI, BuiltinID, TheCall);
- return false;
+ else
+ return AMDGPU().CheckAMDGCNBuiltinFunctionCall(BuiltinID, TheCall);
case llvm::Triple::systemz:
return SystemZ().CheckSystemZBuiltinFunctionCall(BuiltinID, TheCall);
case llvm::Triple::x86:
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index 376105cb6594c..bba487cff2544 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -395,546 +395,546 @@ void test_s_sendmsghalt_var(int in)
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.add.i32(
void test_wave_reduce_add_u32_default(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_add_u32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_add(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_add_u64_default
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.add.i64(
void test_wave_reduce_add_u64_default(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_add_u64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_add(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_fadd_f32_default
// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fadd.f32(
void test_wave_reduce_fadd_f32_default(global float* out, float in)
{
- *out = __builtin_amdgcn_wave_reduce_fadd_f32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_add(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_fadd_f64_default
// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.wave.reduce.fadd.f64(
void test_wave_reduce_fadd_f64_default(global double* out, double in)
{
- *out = __builtin_amdgcn_wave_reduce_fadd_f64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_add(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_add_u32_iterative
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.add.i32(
void test_wave_reduce_add_u32_iterative(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_add_u32(in, 1);
+ *out = __builtin_amdgcn_wave_reduce_add(in, 1);
}
// CHECK-LABEL: @test_wave_reduce_add_u64_iterative
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.add.i64(
void test_wave_reduce_add_u64_iterative(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_add_u64(in, 1);
+ *out = __builtin_amdgcn_wave_reduce_add(in, 1);
}
// CHECK-LABEL: @test_wave_reduce_fadd_f32_iterative
// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fadd.f32(
void test_wave_reduce_fadd_f32_iterative(global float* out, float in)
{
- *out = __builtin_amdgcn_wave_reduce_fadd_f32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_add(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_fadd_f64_iterative
// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.wave.reduce.fadd.f64(
void test_wave_reduce_fadd_f64_iterative(global double* out, double in)
{
- *out = __builtin_amdgcn_wave_reduce_fadd_f64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_add(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_add_u32_dpp
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.add.i32(
void test_wave_reduce_add_u32_dpp(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_add_u32(in, 2);
+ *out = __builtin_amdgcn_wave_reduce_add(in, 2);
}
// CHECK-LABEL: @test_wave_reduce_add_u64_dpp
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.add.i64(
void test_wave_reduce_add_u64_dpp(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_add_u64(in, 2);
+ *out = __builtin_amdgcn_wave_reduce_add(in, 2);
}
// CHECK-LABEL: @test_wave_reduce_fadd_f32_dpp
// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fadd.f32(
void test_wave_reduce_fadd_f32_dpp(global float* out, float in)
{
- *out = __builtin_amdgcn_wave_reduce_fadd_f32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_add(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_fadd_f64_dpp
// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.wave.reduce.fadd.f64(
void test_wave_reduce_fadd_f64_dpp(global double* out, double in)
{
- *out = __builtin_amdgcn_wave_reduce_fadd_f64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_add(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_sub_u32_default
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.sub.i32(
void test_wave_reduce_sub_u32_default(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_sub_u32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_sub(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_sub_u64_default
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.sub.i64(
void test_wave_reduce_sub_u64_default(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_sub_u64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_sub(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_fsub_f32_default
// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fsub.f32(
void test_wave_reduce_fsub_f32_default(global float* out, float in)
{
- *out = __builtin_amdgcn_wave_reduce_fsub_f32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_sub(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_fsub_f64_default
// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.wave.reduce.fsub.f64(
void test_wave_reduce_fsub_f64_default(global double* out, double in)
{
- *out = __builtin_amdgcn_wave_reduce_fsub_f64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_sub(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_sub_u32_iterative
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.sub.i32(
void test_wave_reduce_sub_u32_iterative(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_sub_u32(in, 1);
+ *out = __builtin_amdgcn_wave_reduce_sub(in, 1);
}
// CHECK-LABEL: @test_wave_reduce_sub_u64_iterative
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.sub.i64(
void test_wave_reduce_sub_u64_iterative(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_sub_u64(in, 1);
+ *out = __builtin_amdgcn_wave_reduce_sub(in, 1);
}
// CHECK-LABEL: @test_wave_reduce_fsub_f32_iterative
// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fsub.f32(
void test_wave_reduce_fsub_f32_iterative(global float* out, float in)
{
- *out = __builtin_amdgcn_wave_reduce_fsub_f32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_sub(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_fsub_f64_iterative
// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.wave.reduce.fsub.f64(
void test_wave_reduce_fsub_f64_iterative(global double* out, double in)
{
- *out = __builtin_amdgcn_wave_reduce_fsub_f64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_sub(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_sub_u32_dpp
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.sub.i32(
void test_wave_reduce_sub_u32_dpp(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_sub_u32(in, 2);
+ *out = __builtin_amdgcn_wave_reduce_sub(in, 2);
}
// CHECK-LABEL: @test_wave_reduce_sub_u64_dpp
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.sub.i64(
void test_wave_reduce_sub_u64_dpp(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_sub_u64(in, 2);
+ *out = __builtin_amdgcn_wave_reduce_sub(in, 2);
}
// CHECK-LABEL: @test_wave_reduce_fsub_f32_dpp
// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fsub.f32(
void test_wave_reduce_fsub_f32_dpp(global float* out, float in)
{
- *out = __builtin_amdgcn_wave_reduce_fsub_f32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_sub(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_fsub_f64_dpp
// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.wave.reduce.fsub.f64(
void test_wave_reduce_fsub_f64_dpp(global double* out, double in)
{
- *out = __builtin_amdgcn_wave_reduce_fsub_f64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_sub(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_and_b32_default
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.and.i32(
void test_wave_reduce_and_b32_default(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_and_b32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_and(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_and_b64_default
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.and.i64(
void test_wave_reduce_and_b64_default(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_and_b64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_and(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_and_b32_iterative
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.and.i32(
void test_wave_reduce_and_b32_iterative(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_and_b32(in, 1);
+ *out = __builtin_amdgcn_wave_reduce_and(in, 1);
}
// CHECK-LABEL: @test_wave_reduce_and_b64_iterative
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.and.i64(
void test_wave_reduce_and_b64_iterative(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_and_b64(in, 1);
+ *out = __builtin_amdgcn_wave_reduce_and(in, 1);
}
// CHECK-LABEL: @test_wave_reduce_and_b32_dpp
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.and.i32(
void test_wave_reduce_and_b32_dpp(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_and_b32(in, 2);
+ *out = __builtin_amdgcn_wave_reduce_and(in, 2);
}
// CHECK-LABEL: @test_wave_reduce_and_b64_dpp
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.and.i64(
void test_wave_reduce_and_b64_dpp(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_and_b64(in, 2);
+ *out = __builtin_amdgcn_wave_reduce_and(in, 2);
}
// CHECK-LABEL: @test_wave_reduce_or_b32_default
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.or.i32(
void test_wave_reduce_or_b32_default(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_or_b32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_or(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_or_b64_default
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.or.i64(
void test_wave_reduce_or_b64_default(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_or_b64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_or(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_or_b32_iterative
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.or.i32(
void test_wave_reduce_or_b32_iterative(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_or_b32(in, 1);
+ *out = __builtin_amdgcn_wave_reduce_or(in, 1);
}
// CHECK-LABEL: @test_wave_reduce_or_b64_iterative
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.or.i64(
void test_wave_reduce_or_b64_iterative(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_or_b64(in, 1);
+ *out = __builtin_amdgcn_wave_reduce_or(in, 1);
}
// CHECK-LABEL: @test_wave_reduce_or_b32_dpp
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.or.i32(
void test_wave_reduce_or_b32_dpp(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_or_b32(in, 2);
+ *out = __builtin_amdgcn_wave_reduce_or(in, 2);
}
// CHECK-LABEL: @test_wave_reduce_or_b64_dpp
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.or.i64(
void test_wave_reduce_or_b64_dpp(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_or_b64(in, 2);
+ *out = __builtin_amdgcn_wave_reduce_or(in, 2);
}
// CHECK-LABEL: @test_wave_reduce_xor_b32_default
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.xor.i32(
void test_wave_reduce_xor_b32_default(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_xor_b32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_xor(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_xor_b64_default
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.xor.i64(
void test_wave_reduce_xor_b64_default(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_xor_b64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_xor(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_xor_b32_iterative
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.xor.i32(
void test_wave_reduce_xor_b32_iterative(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_xor_b32(in, 1);
+ *out = __builtin_amdgcn_wave_reduce_xor(in, 1);
}
// CHECK-LABEL: @test_wave_reduce_xor_b64_iterative
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.xor.i64(
void test_wave_reduce_xor_b64_iterative(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_xor_b64(in, 1);
+ *out = __builtin_amdgcn_wave_reduce_xor(in, 1);
}
// CHECK-LABEL: @test_wave_reduce_xor_b32_dpp
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.xor.i32(
void test_wave_reduce_xor_b32_dpp(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_xor_b32(in, 2);
+ *out = __builtin_amdgcn_wave_reduce_xor(in, 2);
}
// CHECK-LABEL: @test_wave_reduce_xor_b64_dpp
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.xor.i64(
void test_wave_reduce_xor_b64_dpp(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_xor_b64(in, 2);
+ *out = __builtin_amdgcn_wave_reduce_xor(in, 2);
}
// CHECK-LABEL: @test_wave_reduce_min_i32_default
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.min.i32(
void test_wave_reduce_min_i32_default(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_min_i32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_min(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_min_i64_default
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.min.i64(
void test_wave_reduce_min_i64_default(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_min_i64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_min(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_fmin_f32_default
// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fmin.f32(
void test_wave_reduce_fmin_f32_default(global float* out, float in)
{
- *out = __builtin_amdgcn_wave_reduce_fmin_f32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_min(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_fmin_f64_default
// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.wave.reduce.fmin.f64(
void test_wave_reduce_fmin_f64_default(global double* out, double in)
{
- *out = __builtin_amdgcn_wave_reduce_fmin_f64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_min(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_min_i32_iterative
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.min.i32(
void test_wave_reduce_min_i32_iterative(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_min_i32(in, 1);
+ *out = __builtin_amdgcn_wave_reduce_min(in, 1);
}
// CHECK-LABEL: @test_wave_reduce_min_i64_iterative
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.min.i64(
void test_wave_reduce_min_i64_iterative(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_min_i64(in, 1);
+ *out = __builtin_amdgcn_wave_reduce_min(in, 1);
}
// CHECK-LABEL: @test_wave_reduce_fmin_f32_iterative
// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fmin.f32(
void test_wave_reduce_fmin_f32_iterative(global float* out, float in)
{
- *out = __builtin_amdgcn_wave_reduce_fmin_f32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_min(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_fmin_f64_iterative
// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.wave.reduce.fmin.f64(
void test_wave_reduce_fmin_f64_iterative(global double* out, double in)
{
- *out = __builtin_amdgcn_wave_reduce_fmin_f64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_min(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_min_i32_dpp
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.min.i32(
void test_wave_reduce_min_i32_dpp(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_min_i32(in, 2);
+ *out = __builtin_amdgcn_wave_reduce_min(in, 2);
}
// CHECK-LABEL: @test_wave_reduce_min_i64_dpp
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.min.i64(
void test_wave_reduce_min_i64_dpp(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_min_i64(in, 2);
+ *out = __builtin_amdgcn_wave_reduce_min(in, 2);
}
// CHECK-LABEL: @test_wave_reduce_fmin_f32_dpp
// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fmin.f32(
void test_wave_reduce_fmin_f32_dpp(global float* out, float in)
{
- *out = __builtin_amdgcn_wave_reduce_fmin_f32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_min(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_fmin_f64_dpp
// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.wave.reduce.fmin.f64(
void test_wave_reduce_fmin_f64_dpp(global double* out, double in)
{
- *out = __builtin_amdgcn_wave_reduce_fmin_f64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_min(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_min_u32_default
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.umin.i32(
-void test_wave_reduce_min_u32_default(global int* out, int in)
+void test_wave_reduce_min_u32_default(global int* out, unsigned int in)
{
- *out = __builtin_amdgcn_wave_reduce_min_u32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_min(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_min_u64_default
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.umin.i64(
-void test_wave_reduce_min_u64_default(global int* out, long in)
+void test_wave_reduce_min_u64_default(global int* out, unsigned long in)
{
- *out = __builtin_amdgcn_wave_reduce_min_u64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_min(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_min_u32_iterative
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.umin.i32(
-void test_wave_reduce_min_u32_iterative(global int* out, int in)
+void test_wave_reduce_min_u32_iterative(global int* out, unsigned int in)
{
- *out = __builtin_amdgcn_wave_reduce_min_u32(in, 1);
+ *out = __builtin_amdgcn_wave_reduce_min(in, 1);
}
// CHECK-LABEL: @test_wave_reduce_min_u64_iterative
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.umin.i64(
-void test_wave_reduce_min_u64_iterative(global int* out, long in)
+void test_wave_reduce_min_u64_iterative(global int* out, unsigned long in)
{
- *out = __builtin_amdgcn_wave_reduce_min_u64(in, 1);
+ *out = __builtin_amdgcn_wave_reduce_min(in, 1);
}
// CHECK-LABEL: @test_wave_reduce_min_u32_dpp
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.umin.i32(
-void test_wave_reduce_min_u32_dpp(global int* out, int in)
+void test_wave_reduce_min_u32_dpp(global int* out, unsigned int in)
{
- *out = __builtin_amdgcn_wave_reduce_min_u32(in, 2);
+ *out = __builtin_amdgcn_wave_reduce_min(in, 2);
}
// CHECK-LABEL: @test_wave_reduce_min_u64_dpp
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.umin.i64(
-void test_wave_reduce_min_u64_dpp(global int* out, long in)
+void test_wave_reduce_min_u64_dpp(global int* out, unsigned long in)
{
- *out = __builtin_amdgcn_wave_reduce_min_u64(in, 2);
+ *out = __builtin_amdgcn_wave_reduce_min(in, 2);
}
// CHECK-LABEL: @test_wave_reduce_max_i32_default
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.max.i32(
void test_wave_reduce_max_i32_default(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_max_i32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_max(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_max_i64_default
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.max.i64(
void test_wave_reduce_max_i64_default(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_max_i64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_max(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_fmax_f32_default
// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fmax.f32(
void test_wave_reduce_fmax_f32_default(global float* out, float in)
{
- *out = __builtin_amdgcn_wave_reduce_fmax_f32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_max(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_fmax_f64_default
// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.wave.reduce.fmax.f64(
void test_wave_reduce_fmax_f64_default(global double* out, double in)
{
- *out = __builtin_amdgcn_wave_reduce_fmax_f64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_max(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_max_i32_iterative
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.max.i32(
void test_wave_reduce_max_i32_iterative(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_max_i32(in, 1);
+ *out = __builtin_amdgcn_wave_reduce_max(in, 1);
}
// CHECK-LABEL: @test_wave_reduce_max_i64_iterative
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.max.i64(
void test_wave_reduce_max_i64_iterative(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_max_i64(in, 1);
+ *out = __builtin_amdgcn_wave_reduce_max(in, 1);
}
// CHECK-LABEL: @test_wave_reduce_fmax_f32_iterative
// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fmax.f32(
void test_wave_reduce_fmax_f32_iterative(global float* out, float in)
{
- *out = __builtin_amdgcn_wave_reduce_fmax_f32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_max(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_fmax_f64_iterative
// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.wave.reduce.fmax.f64(
void test_wave_reduce_fmax_f64_iterative(global double* out, double in)
{
- *out = __builtin_amdgcn_wave_reduce_fmax_f64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_max(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_max_i32_dpp
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.max.i32(
void test_wave_reduce_max_i32_dpp(global int* out, int in)
{
- *out = __builtin_amdgcn_wave_reduce_max_i32(in, 2);
+ *out = __builtin_amdgcn_wave_reduce_max(in, 2);
}
// CHECK-LABEL: @test_wave_reduce_max_i64_dpp
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.max.i64(
void test_wave_reduce_max_i64_dpp(global int* out, long in)
{
- *out = __builtin_amdgcn_wave_reduce_max_i64(in, 2);
+ *out = __builtin_amdgcn_wave_reduce_max(in, 2);
}
// CHECK-LABEL: @test_wave_reduce_fmax_f32_dpp
// CHECK: {{.*}}call{{.*}} float @llvm.amdgcn.wave.reduce.fmax.f32(
void test_wave_reduce_fmax_f32_dpp(global float* out, float in)
{
- *out = __builtin_amdgcn_wave_reduce_fmax_f32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_max(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_fmax_f64_dpp
// CHECK: {{.*}}call{{.*}} double @llvm.amdgcn.wave.reduce.fmax.f64(
void test_wave_reduce_fmax_f64_dpp(global double* out, double in)
{
- *out = __builtin_amdgcn_wave_reduce_fmax_f64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_max(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_max_u32_default
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.umax.i32(
-void test_wave_reduce_max_u32_default(global int* out, int in)
+void test_wave_reduce_max_u32_default(global int* out, unsigned int in)
{
- *out = __builtin_amdgcn_wave_reduce_max_u32(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_max(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_max_u64_default
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.umax.i64(
-void test_wave_reduce_max_u64_default(global int* out, long in)
+void test_wave_reduce_max_u64_default(global int* out, unsigned long in)
{
- *out = __builtin_amdgcn_wave_reduce_max_u64(in, 0);
+ *out = __builtin_amdgcn_wave_reduce_max(in, 0);
}
// CHECK-LABEL: @test_wave_reduce_max_u32_iterative
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.umax.i32(
-void test_wave_reduce_max_u32_iterative(global int* out, int in)
+void test_wave_reduce_max_u32_iterative(global int* out, unsigned int in)
{
- *out = __builtin_amdgcn_wave_reduce_max_u32(in, 1);
+ *out = __builtin_amdgcn_wave_reduce_max(in, 1);
}
// CHECK-LABEL: @test_wave_reduce_max_u64_iterative
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.umax.i64(
-void test_wave_reduce_max_u64_iterative(global int* out, long in)
+void test_wave_reduce_max_u64_iterative(global int* out, unsigned long in)
{
- *out = __builtin_amdgcn_wave_reduce_max_u64(in, 1);
+ *out = __builtin_amdgcn_wave_reduce_max(in, 1);
}
// CHECK-LABEL: @test_wave_reduce_max_u32_dpp
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.reduce.umax.i32(
-void test_wave_reduce_max_u32_dpp(global int* out, int in)
+void test_wave_reduce_max_u32_dpp(global int* out, unsigned int in)
{
- *out = __builtin_amdgcn_wave_reduce_max_u32(in, 2);
+ *out = __builtin_amdgcn_wave_reduce_max(in, 2);
}
// CHECK-LABEL: @test_wave_reduce_max_u64_dpp
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.wave.reduce.umax.i64(
-void test_wave_reduce_max_u64_dpp(global int* out, long in)
+void test_wave_reduce_max_u64_dpp(global int* out, unsigned long in)
{
- *out = __builtin_amdgcn_wave_reduce_max_u64(in, 2);
+ *out = __builtin_amdgcn_wave_reduce_max(in, 2);
}
// CHECK-LABEL: @test_s_barrier
diff --git a/clang/test/Sema/wave-reduce-builtins-validate-amdgpu.cl b/clang/test/Sema/wave-reduce-builtins-validate-amdgpu.cl
index 0f1565f1272c1..8a5eb53872ea8 100644
--- a/clang/test/Sema/wave-reduce-builtins-validate-amdgpu.cl
+++ b/clang/test/Sema/wave-reduce-builtins-validate-amdgpu.cl
@@ -4,83 +4,97 @@
// Test that the second argument (strategy) must be a constant integer
void test_wave_reduce_u32(unsigned int val, int strategy) {
- (void)__builtin_amdgcn_wave_reduce_add_u32(val, 0);
- (void)__builtin_amdgcn_wave_reduce_sub_u32(val, 1);
- (void)__builtin_amdgcn_wave_reduce_min_u32(val, 0);
- (void)__builtin_amdgcn_wave_reduce_max_u32(val, 0);
- (void)__builtin_amdgcn_wave_reduce_and_b32(val, 0);
- (void)__builtin_amdgcn_wave_reduce_or_b32(val, 0);
- (void)__builtin_amdgcn_wave_reduce_xor_b32(val, 0);
-
- (void)__builtin_amdgcn_wave_reduce_add_u32(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_add_u32' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_sub_u32(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_sub_u32' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_min_u32(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_min_u32' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_max_u32(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_max_u32' must be a constant integer}}
+ (void)__builtin_amdgcn_wave_reduce_add(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_sub(val, 1);
+ (void)__builtin_amdgcn_wave_reduce_min(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_max(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_and(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_or(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_xor(val, 0);
+
+ (void)__builtin_amdgcn_wave_reduce_add(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_sub(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_min(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_max(val, strategy); // expected-error {{expression is not an integer constant expression}}
}
void test_wave_reduce_i32(int val, int strategy) {
- (void)__builtin_amdgcn_wave_reduce_min_i32(val, 0);
- (void)__builtin_amdgcn_wave_reduce_max_i32(val, 0);
- (void)__builtin_amdgcn_wave_reduce_and_b32(val, 0);
- (void)__builtin_amdgcn_wave_reduce_or_b32(val, 0);
- (void)__builtin_amdgcn_wave_reduce_xor_b32(val, 0);
-
- (void)__builtin_amdgcn_wave_reduce_min_i32(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_min_i32' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_max_i32(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_max_i32' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_and_b32(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_and_b32' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_or_b32(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_or_b32' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_xor_b32(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_xor_b32' must be a constant integer}}
+ (void)__builtin_amdgcn_wave_reduce_min(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_max(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_and(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_or(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_xor(val, 0);
+
+ (void)__builtin_amdgcn_wave_reduce_min(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_max(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_and(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_or(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_xor(val, strategy); // expected-error {{expression is not an integer constant expression}}
}
void test_wave_reduce_u64(unsigned long val, int strategy) {
- (void)__builtin_amdgcn_wave_reduce_add_u64(val, 0);
- (void)__builtin_amdgcn_wave_reduce_sub_u64(val, 1);
- (void)__builtin_amdgcn_wave_reduce_min_u64(val, 0);
- (void)__builtin_amdgcn_wave_reduce_max_u64(val, 0);
- (void)__builtin_amdgcn_wave_reduce_and_b64(val, 0);
- (void)__builtin_amdgcn_wave_reduce_or_b64(val, 0);
- (void)__builtin_amdgcn_wave_reduce_xor_b64(val, 0);
-
- (void)__builtin_amdgcn_wave_reduce_add_u64(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_add_u64' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_sub_u64(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_sub_u64' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_min_u64(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_min_u64' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_max_u64(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_max_u64' must be a constant integer}}
+ (void)__builtin_amdgcn_wave_reduce_add(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_sub(val, 1);
+ (void)__builtin_amdgcn_wave_reduce_min(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_max(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_and(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_or(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_xor(val, 0);
+
+ (void)__builtin_amdgcn_wave_reduce_add(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_sub(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_min(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_max(val, strategy); // expected-error {{expression is not an integer constant expression}}
}
void test_wave_reduce_i64(long val, int strategy) {
- (void)__builtin_amdgcn_wave_reduce_min_i64(val, 0);
- (void)__builtin_amdgcn_wave_reduce_max_i64(val, 0);
- (void)__builtin_amdgcn_wave_reduce_and_b64(val, 0);
- (void)__builtin_amdgcn_wave_reduce_or_b64(val, 0);
- (void)__builtin_amdgcn_wave_reduce_xor_b64(val, 0);
-
- (void)__builtin_amdgcn_wave_reduce_min_i64(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_min_i64' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_max_i64(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_max_i64' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_and_b64(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_and_b64' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_or_b64(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_or_b64' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_xor_b64(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_xor_b64' must be a constant integer}}
+ (void)__builtin_amdgcn_wave_reduce_min(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_max(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_and(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_or(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_xor(val, 0);
+
+ (void)__builtin_amdgcn_wave_reduce_min(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_max(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_and(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_or(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_xor(val, strategy); // expected-error {{expression is not an integer constant expression}}
}
void test_wave_reduce_f32(float val, int strategy) {
- (void)__builtin_amdgcn_wave_reduce_fadd_f32(val, 0);
- (void)__builtin_amdgcn_wave_reduce_fsub_f32(val, 1);
- (void)__builtin_amdgcn_wave_reduce_fmin_f32(val, 0);
- (void)__builtin_amdgcn_wave_reduce_fmax_f32(val, 0);
-
- (void)__builtin_amdgcn_wave_reduce_fadd_f32(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_fadd_f32' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_fsub_f32(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_fsub_f32' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_fmin_f32(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_fmin_f32' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_fmax_f32(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_fmax_f32' must be a constant integer}}
+ (void)__builtin_amdgcn_wave_reduce_add(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_sub(val, 1);
+ (void)__builtin_amdgcn_wave_reduce_min(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_max(val, 0);
+
+ (void)__builtin_amdgcn_wave_reduce_add(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_sub(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_min(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_max(val, strategy); // expected-error {{expression is not an integer constant expression}}
}
void test_wave_reduce_f64(double val, int strategy) {
- (void)__builtin_amdgcn_wave_reduce_fadd_f64(val, 0);
- (void)__builtin_amdgcn_wave_reduce_fsub_f64(val, 1);
- (void)__builtin_amdgcn_wave_reduce_fmin_f64(val, 0);
- (void)__builtin_amdgcn_wave_reduce_fmax_f64(val, 0);
-
- (void)__builtin_amdgcn_wave_reduce_fadd_f64(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_fadd_f64' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_fsub_f64(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_fsub_f64' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_fmin_f64(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_fmin_f64' must be a constant integer}}
- (void)__builtin_amdgcn_wave_reduce_fmax_f64(val, strategy); // expected-error {{argument to '__builtin_amdgcn_wave_reduce_fmax_f64' must be a constant integer}}
+ (void)__builtin_amdgcn_wave_reduce_add(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_sub(val, 1);
+ (void)__builtin_amdgcn_wave_reduce_min(val, 0);
+ (void)__builtin_amdgcn_wave_reduce_max(val, 0);
+
+ (void)__builtin_amdgcn_wave_reduce_add(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_sub(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_min(val, strategy); // expected-error {{expression is not an integer constant expression}}
+ (void)__builtin_amdgcn_wave_reduce_max(val, strategy); // expected-error {{expression is not an integer constant expression}}
+}
+
+struct S { double x; long long y; };
+
+void test_wave_reduce_struct(struct S val, int strategy) {
+ (void)__builtin_amdgcn_wave_reduce_add(val, 0); // expected-error {{'val' argument must be a scalar integer or floating-point type (was '__private struct S')}}
+ (void)__builtin_amdgcn_wave_reduce_sub(val, 1); // expected-error {{'val' argument must be a scalar integer or floating-point type (was '__private struct S')}}
+ (void)__builtin_amdgcn_wave_reduce_min(val, 0); // expected-error {{'val' argument must be a scalar integer or floating-point type (was '__private struct S')}}
+ (void)__builtin_amdgcn_wave_reduce_max(val, 0); // expected-error {{'val' argument must be a scalar integer or floating-point type (was '__private struct S')}}
+}
+
+void test_wave_reduce_invalid_strategy(int val) {
+ (void)__builtin_amdgcn_wave_reduce_add(val, -1); // expected-error {{argument value 4294967295 is outside the valid range [0, 2]}}
+ (void)__builtin_amdgcn_wave_reduce_sub(val, 3); // expected-error {{argument value 3 is outside the valid range [0, 2]}}
}
diff --git a/clang/tools/c-index-test/CMakeLists.txt b/clang/tools/c-index-test/CMakeLists.txt
index 41e80e66ffa7a..81bccfc3c29f0 100644
--- a/clang/tools/c-index-test/CMakeLists.txt
+++ b/clang/tools/c-index-test/CMakeLists.txt
@@ -24,13 +24,13 @@ if (LLVM_BUILD_STATIC)
else()
target_link_libraries(c-index-test
PRIVATE
- libclang
clangAST
clangBasic
clangDriver
clangFrontend
clangIndex
clangSerialization
+ libclang
)
endif()
diff --git a/clang/tools/libclang/CIndexDiagnostic.cpp b/clang/tools/libclang/CIndexDiagnostic.cpp
index d37597e747a84..aa1a46f470849 100644
--- a/clang/tools/libclang/CIndexDiagnostic.cpp
+++ b/clang/tools/libclang/CIndexDiagnostic.cpp
@@ -77,8 +77,8 @@ class CXDiagnosticCustomNoteImpl : public CXDiagnosticImpl {
*ReplacementRange = clang_getNullRange();
return cxstring::createEmpty();
}
-};
-
+};
+
class CXDiagnosticRenderer : public DiagnosticNoteRenderer {
public:
CXDiagnosticRenderer(const LangOptions &LangOpts, DiagnosticOptions &DiagOpts,
@@ -95,7 +95,7 @@ class CXDiagnosticRenderer : public DiagnosticNoteRenderer {
dyn_cast_if_present<const StoredDiagnostic *>(D);
if (!SD)
return;
-
+
if (Level != DiagnosticsEngine::Note)
CurrentSet = MainSet;
@@ -113,7 +113,7 @@ class CXDiagnosticRenderer : public DiagnosticNoteRenderer {
DiagOrStoredDiag D) override {
if (!D.isNull())
return;
-
+
CXSourceLocation L;
if (Loc.hasManager())
L = translateSourceLocation(Loc.getManager(), LangOpts, Loc);
@@ -143,7 +143,7 @@ class CXDiagnosticRenderer : public DiagnosticNoteRenderer {
CXDiagnosticSetImpl *CurrentSet;
CXDiagnosticSetImpl *MainSet;
-};
+};
}
CXDiagnosticSetImpl *cxdiag::lazyCreateDiags(CXTranslationUnit TU,
@@ -246,7 +246,7 @@ CXString clang_formatDiagnostic(CXDiagnostic Diagnostic, unsigned Options) {
SmallString<256> Str;
llvm::raw_svector_ostream Out(Str);
-
+
if (Options & CXDiagnostic_DisplaySourceLocation) {
// Print source location (file:line), along with optional column
// and source ranges.
@@ -267,7 +267,7 @@ CXString clang_formatDiagnostic(CXDiagnostic Diagnostic, unsigned Options) {
for (unsigned I = 0; I != N; ++I) {
CXFile StartFile, EndFile;
CXSourceRange Range = clang_getDiagnosticRange(Diagnostic, I);
-
+
unsigned StartLine, StartColumn, EndLine, EndColumn;
clang_getSpellingLocation(clang_getRangeStart(Range),
&StartFile, &StartLine, &StartColumn,
@@ -277,7 +277,7 @@ CXString clang_formatDiagnostic(CXDiagnostic Diagnostic, unsigned Options) {
if (StartFile != EndFile || StartFile != File)
continue;
-
+
Out << "{" << StartLine << ":" << StartColumn << "-"
<< EndLine << ":" << EndColumn << "}";
PrintedRange = true;
@@ -285,7 +285,7 @@ CXString clang_formatDiagnostic(CXDiagnostic Diagnostic, unsigned Options) {
if (PrintedRange)
Out << ":";
}
-
+
Out << " ";
}
}
@@ -305,7 +305,7 @@ CXString clang_formatDiagnostic(CXDiagnostic Diagnostic, unsigned Options) {
else
Out << "<no diagnostic text>";
clang_disposeString(Text);
-
+
if (Options & (CXDiagnostic_DisplayOption | CXDiagnostic_DisplayCategoryId |
CXDiagnostic_DisplayCategoryName)) {
bool NeedBracket = true;
@@ -322,8 +322,8 @@ CXString clang_formatDiagnostic(CXDiagnostic Diagnostic, unsigned Options) {
}
clang_disposeString(OptionName);
}
-
- if (Options & (CXDiagnostic_DisplayCategoryId |
+
+ if (Options & (CXDiagnostic_DisplayCategoryId |
CXDiagnostic_DisplayCategoryName)) {
if (unsigned CategoryID = clang_getDiagnosticCategory(Diagnostic)) {
if (Options & CXDiagnostic_DisplayCategoryId) {
@@ -335,7 +335,7 @@ CXString clang_formatDiagnostic(CXDiagnostic Diagnostic, unsigned Options) {
NeedBracket = false;
NeedComma = true;
}
-
+
if (Options & CXDiagnostic_DisplayCategoryName) {
CXString CategoryName = clang_getDiagnosticCategoryText(Diagnostic);
if (NeedBracket)
@@ -354,7 +354,7 @@ CXString clang_formatDiagnostic(CXDiagnostic Diagnostic, unsigned Options) {
if (!NeedBracket)
Out << "]";
}
-
+
return cxstring::createDup(Out.str());
}
@@ -396,18 +396,18 @@ unsigned clang_getDiagnosticCategory(CXDiagnostic Diag) {
return D->getCategory();
return 0;
}
-
+
CXString clang_getDiagnosticCategoryName(unsigned Category) {
// Kept for backward compatibility.
return cxstring::createRef(DiagnosticIDs::getCategoryNameFromID(Category));
}
-
+
CXString clang_getDiagnosticCategoryText(CXDiagnostic Diag) {
if (CXDiagnosticImpl *D = static_cast<CXDiagnosticImpl *>(Diag))
return D->getCategoryText();
return cxstring::createEmpty();
}
-
+
unsigned clang_getDiagnosticNumRanges(CXDiagnostic Diag) {
if (CXDiagnosticImpl *D = static_cast<CXDiagnosticImpl *>(Diag))
return D->getNumRanges();
@@ -415,7 +415,7 @@ unsigned clang_getDiagnosticNumRanges(CXDiagnostic Diag) {
}
CXSourceRange clang_getDiagnosticRange(CXDiagnostic Diag, unsigned Range) {
- CXDiagnosticImpl *D = static_cast<CXDiagnosticImpl *>(Diag);
+ CXDiagnosticImpl *D = static_cast<CXDiagnosticImpl *>(Diag);
if (!D || Range >= D->getNumRanges())
return clang_getNullRange();
return D->getRange(Range);
@@ -444,7 +444,7 @@ void clang_disposeDiagnosticSet(CXDiagnosticSet Diags) {
delete D;
}
}
-
+
CXDiagnostic clang_getDiagnosticInSet(CXDiagnosticSet Diags,
unsigned Index) {
if (CXDiagnosticSetImpl *D = static_cast<CXDiagnosticSetImpl*>(Diags))
@@ -452,7 +452,7 @@ CXDiagnostic clang_getDiagnosticInSet(CXDiagnosticSet Diags,
return D->getDiagnostic(Index);
return nullptr;
}
-
+
CXDiagnosticSet clang_getChildDiagnostics(CXDiagnostic Diag) {
if (CXDiagnosticImpl *D = static_cast<CXDiagnosticImpl *>(Diag)) {
CXDiagnosticSetImpl &ChildDiags = D->getChildDiagnostics();
>From f819fd721a8b926d6a97fc69e9cb921636a1be3c Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu at amd.com>
Date: Wed, 4 Feb 2026 01:49:02 +0000
Subject: [PATCH 2/2] Remove accidental noise.
---
clang/tools/libclang/CIndexDiagnostic.cpp | 40 +++++++++++------------
1 file changed, 20 insertions(+), 20 deletions(-)
diff --git a/clang/tools/libclang/CIndexDiagnostic.cpp b/clang/tools/libclang/CIndexDiagnostic.cpp
index aa1a46f470849..d37597e747a84 100644
--- a/clang/tools/libclang/CIndexDiagnostic.cpp
+++ b/clang/tools/libclang/CIndexDiagnostic.cpp
@@ -77,8 +77,8 @@ class CXDiagnosticCustomNoteImpl : public CXDiagnosticImpl {
*ReplacementRange = clang_getNullRange();
return cxstring::createEmpty();
}
-};
-
+};
+
class CXDiagnosticRenderer : public DiagnosticNoteRenderer {
public:
CXDiagnosticRenderer(const LangOptions &LangOpts, DiagnosticOptions &DiagOpts,
@@ -95,7 +95,7 @@ class CXDiagnosticRenderer : public DiagnosticNoteRenderer {
dyn_cast_if_present<const StoredDiagnostic *>(D);
if (!SD)
return;
-
+
if (Level != DiagnosticsEngine::Note)
CurrentSet = MainSet;
@@ -113,7 +113,7 @@ class CXDiagnosticRenderer : public DiagnosticNoteRenderer {
DiagOrStoredDiag D) override {
if (!D.isNull())
return;
-
+
CXSourceLocation L;
if (Loc.hasManager())
L = translateSourceLocation(Loc.getManager(), LangOpts, Loc);
@@ -143,7 +143,7 @@ class CXDiagnosticRenderer : public DiagnosticNoteRenderer {
CXDiagnosticSetImpl *CurrentSet;
CXDiagnosticSetImpl *MainSet;
-};
+};
}
CXDiagnosticSetImpl *cxdiag::lazyCreateDiags(CXTranslationUnit TU,
@@ -246,7 +246,7 @@ CXString clang_formatDiagnostic(CXDiagnostic Diagnostic, unsigned Options) {
SmallString<256> Str;
llvm::raw_svector_ostream Out(Str);
-
+
if (Options & CXDiagnostic_DisplaySourceLocation) {
// Print source location (file:line), along with optional column
// and source ranges.
@@ -267,7 +267,7 @@ CXString clang_formatDiagnostic(CXDiagnostic Diagnostic, unsigned Options) {
for (unsigned I = 0; I != N; ++I) {
CXFile StartFile, EndFile;
CXSourceRange Range = clang_getDiagnosticRange(Diagnostic, I);
-
+
unsigned StartLine, StartColumn, EndLine, EndColumn;
clang_getSpellingLocation(clang_getRangeStart(Range),
&StartFile, &StartLine, &StartColumn,
@@ -277,7 +277,7 @@ CXString clang_formatDiagnostic(CXDiagnostic Diagnostic, unsigned Options) {
if (StartFile != EndFile || StartFile != File)
continue;
-
+
Out << "{" << StartLine << ":" << StartColumn << "-"
<< EndLine << ":" << EndColumn << "}";
PrintedRange = true;
@@ -285,7 +285,7 @@ CXString clang_formatDiagnostic(CXDiagnostic Diagnostic, unsigned Options) {
if (PrintedRange)
Out << ":";
}
-
+
Out << " ";
}
}
@@ -305,7 +305,7 @@ CXString clang_formatDiagnostic(CXDiagnostic Diagnostic, unsigned Options) {
else
Out << "<no diagnostic text>";
clang_disposeString(Text);
-
+
if (Options & (CXDiagnostic_DisplayOption | CXDiagnostic_DisplayCategoryId |
CXDiagnostic_DisplayCategoryName)) {
bool NeedBracket = true;
@@ -322,8 +322,8 @@ CXString clang_formatDiagnostic(CXDiagnostic Diagnostic, unsigned Options) {
}
clang_disposeString(OptionName);
}
-
- if (Options & (CXDiagnostic_DisplayCategoryId |
+
+ if (Options & (CXDiagnostic_DisplayCategoryId |
CXDiagnostic_DisplayCategoryName)) {
if (unsigned CategoryID = clang_getDiagnosticCategory(Diagnostic)) {
if (Options & CXDiagnostic_DisplayCategoryId) {
@@ -335,7 +335,7 @@ CXString clang_formatDiagnostic(CXDiagnostic Diagnostic, unsigned Options) {
NeedBracket = false;
NeedComma = true;
}
-
+
if (Options & CXDiagnostic_DisplayCategoryName) {
CXString CategoryName = clang_getDiagnosticCategoryText(Diagnostic);
if (NeedBracket)
@@ -354,7 +354,7 @@ CXString clang_formatDiagnostic(CXDiagnostic Diagnostic, unsigned Options) {
if (!NeedBracket)
Out << "]";
}
-
+
return cxstring::createDup(Out.str());
}
@@ -396,18 +396,18 @@ unsigned clang_getDiagnosticCategory(CXDiagnostic Diag) {
return D->getCategory();
return 0;
}
-
+
CXString clang_getDiagnosticCategoryName(unsigned Category) {
// Kept for backward compatibility.
return cxstring::createRef(DiagnosticIDs::getCategoryNameFromID(Category));
}
-
+
CXString clang_getDiagnosticCategoryText(CXDiagnostic Diag) {
if (CXDiagnosticImpl *D = static_cast<CXDiagnosticImpl *>(Diag))
return D->getCategoryText();
return cxstring::createEmpty();
}
-
+
unsigned clang_getDiagnosticNumRanges(CXDiagnostic Diag) {
if (CXDiagnosticImpl *D = static_cast<CXDiagnosticImpl *>(Diag))
return D->getNumRanges();
@@ -415,7 +415,7 @@ unsigned clang_getDiagnosticNumRanges(CXDiagnostic Diag) {
}
CXSourceRange clang_getDiagnosticRange(CXDiagnostic Diag, unsigned Range) {
- CXDiagnosticImpl *D = static_cast<CXDiagnosticImpl *>(Diag);
+ CXDiagnosticImpl *D = static_cast<CXDiagnosticImpl *>(Diag);
if (!D || Range >= D->getNumRanges())
return clang_getNullRange();
return D->getRange(Range);
@@ -444,7 +444,7 @@ void clang_disposeDiagnosticSet(CXDiagnosticSet Diags) {
delete D;
}
}
-
+
CXDiagnostic clang_getDiagnosticInSet(CXDiagnosticSet Diags,
unsigned Index) {
if (CXDiagnosticSetImpl *D = static_cast<CXDiagnosticSetImpl*>(Diags))
@@ -452,7 +452,7 @@ CXDiagnostic clang_getDiagnosticInSet(CXDiagnosticSet Diags,
return D->getDiagnostic(Index);
return nullptr;
}
-
+
CXDiagnosticSet clang_getChildDiagnostics(CXDiagnostic Diag) {
if (CXDiagnosticImpl *D = static_cast<CXDiagnosticImpl *>(Diag)) {
CXDiagnosticSetImpl &ChildDiags = D->getChildDiagnostics();
More information about the cfe-commits
mailing list