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 <[email protected]> Date: Wed, 4 Feb 2026 01:12:39 +0000 Subject: [PATCH 1/4] 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 <[email protected]> Date: Wed, 4 Feb 2026 01:49:02 +0000 Subject: [PATCH 2/4] 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(); >From 66863f57ee4772cc4dac9b0d5247908d21182ce2 Mon Sep 17 00:00:00 2001 From: Alex Voicu <[email protected]> Date: Wed, 4 Feb 2026 01:49:15 +0000 Subject: [PATCH 3/4] Fix formatting. --- clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index e47eaef3651c6..97af9a91eca05 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -370,8 +370,8 @@ static Intrinsic::ID getIntrinsicIDforWaveReduction(unsigned BuiltinID, bool IsFP, bool IsSigned) { switch (BuiltinID) { case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_add: - return IsFP ? - Intrinsic::amdgcn_wave_reduce_fadd : Intrinsic::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_max: @@ -385,8 +385,8 @@ static Intrinsic::ID getIntrinsicIDforWaveReduction(unsigned BuiltinID, case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_or: return Intrinsic::amdgcn_wave_reduce_or; case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_sub: - return IsFP ? - Intrinsic::amdgcn_wave_reduce_fsub : Intrinsic::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: >From e96f87a4b7c598f1013c8775c20e05a96d9e061f Mon Sep 17 00:00:00 2001 From: Alex Voicu <[email protected]> Date: Wed, 4 Feb 2026 01:49:55 +0000 Subject: [PATCH 4/4] Fix other noise. --- clang/tools/c-index-test/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/tools/c-index-test/CMakeLists.txt b/clang/tools/c-index-test/CMakeLists.txt index 81bccfc3c29f0..41e80e66ffa7a 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() _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
