Author: Changpeng Fang Date: 2026-08-27T12:57:29-07:00 New Revision: ce880aebe23a1462e2290db13c468ebad1af16d4
URL: https://github.com/llvm/llvm-project/commit/ce880aebe23a1462e2290db13c468ebad1af16d4 DIFF: https://github.com/llvm/llvm-project/commit/ce880aebe23a1462e2290db13c468ebad1af16d4.diff LOG: [AMDGPU] Implement builtin for raw-buffer F64 atomic add (#219258) Fixes: LCOMPILER-2659 Added: Modified: clang/include/clang/Basic/BuiltinsAMDGPU.td clang/include/clang/Basic/BuiltinsAMDGPUDocs.td clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp clang/test/CodeGen/link-builtin-bitcode.c clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-add.cl clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-atomic-add-err.cl clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-atomic-add-target-err.cl llvm/lib/Target/AMDGPU/AMDGPU.td Removed: ################################################################################ diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td index 794fa517016a6..4e4e0ee275a16 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td @@ -272,6 +272,10 @@ def __builtin_amdgcn_struct_buffer_store_format_v4f16 : AMDGPUBuiltin<"void(_Ext def __builtin_amdgcn_raw_ptr_buffer_atomic_add_i32 : AMDGPUBuiltin<"int(int, __amdgpu_buffer_rsrc_t, int, int, _Constant int)">; def __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32 : AMDGPUBuiltin<"float(float, __amdgpu_buffer_rsrc_t, int, int, _Constant int)", [], "atomic-fadd-rtn-insts">; +def __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f64 : AMDGPUBuiltin<"double(double, __amdgpu_buffer_rsrc_t, int, int, _Constant int)", [], "flat-buffer-global-fadd-f64-inst"> { + let Documentation = [DocRawPtrBufferAtomicFAddF64]; + let ArgNames = ["vdata", "rsrc", "offset", "soffset", "aux"]; +} def __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16 : AMDGPUBuiltin<"_ExtVector<2, _Float16>(_ExtVector<2, _Float16>, __amdgpu_buffer_rsrc_t, int, int, _Constant int)", [], "atomic-buffer-global-pk-add-f16-insts">; def __builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f32 : AMDGPUBuiltin<"float(float, __amdgpu_buffer_rsrc_t, int, int, _Constant int)", [], "atomic-fmin-fmax-global-f32">; diff --git a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td index a835f0e010fdc..ad390e8d024c1 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td @@ -746,6 +746,45 @@ value gives an undefined result. }]; } +//===----------------------------------------------------------------------===// +// Raw Buffer Atomic Builtins +//===----------------------------------------------------------------------===// + +def DocCatRawBufferAtomic : DocumentationCategory<"Raw Buffer Atomic Builtins"> { + let Content = [{ +These builtins perform an atomic read-modify-write on a location addressed +through a buffer resource, and return the value that was in memory before the +operation. + +The buffer resource is passed as ``__amdgpu_buffer_rsrc_t``, which can be +created with ``__builtin_amdgcn_make_buffer_rsrc``. The address is formed from +the resource base plus ``offset`` and ``soffset``; the ``raw`` addressing form +does not take a record index. + +Each builtin requires the target feature listed in its description, and is only +available when that feature is enabled for the target. +}]; +} + +def DocRawPtrBufferAtomicFAddF64 : Documentation { + let Category = DocCatRawBufferAtomic; + let Content = [{ +Atomically adds a double-precision value to a location in a buffer and returns +the value previously stored there. + +- ``vdata``: value added to the value in memory. +- ``rsrc``: buffer resource descriptor. +- ``offset``: byte offset from the base of the buffer, held in a VGPR. +- ``soffset``: additional byte offset, held in an SGPR. +- ``aux``: cache-policy and control flags. Must be a compile-time constant. + The encoding is target-dependent; see the corresponding field of + ``__builtin_amdgcn_s_buffer_load_*`` for the common fields. + +Requires target feature ``flat-buffer-global-fadd-f64-inst``. Maps to the +``buffer_atomic_add_f64`` instruction. +}]; +} + //===----------------------------------------------------------------------===// // S-Buffer Load Builtins //===----------------------------------------------------------------------===// diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index ef40039eb6800..e3a12d825434a 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -990,6 +990,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, return mlir::Value{}; } case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32: + case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f64: case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16: { cgm.errorNYI(expr->getSourceRange(), std::string("unimplemented AMDGPU builtin call: ") + diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index 667c6508040ae..53e9315c85539 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -2185,6 +2185,7 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, return emitBuiltinWithOneOverloadedType<5>( *this, E, Intrinsic::amdgcn_raw_ptr_buffer_atomic_add); case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32: + case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f64: case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16: return emitBuiltinWithOneOverloadedType<5>( *this, E, Intrinsic::amdgcn_raw_ptr_buffer_atomic_fadd); diff --git a/clang/test/CodeGen/link-builtin-bitcode.c b/clang/test/CodeGen/link-builtin-bitcode.c index 8a6aebdbd66dd..2cae7c027a4f3 100644 --- a/clang/test/CodeGen/link-builtin-bitcode.c +++ b/clang/test/CodeGen/link-builtin-bitcode.c @@ -44,6 +44,6 @@ int bar() { return no_attr() + attr_in_target() + attr_not_in_target() + attr_in // CHECK-SAME: () #[[ATTR_INCOMPATIBLE:[0-9]+]] { // CHECK: attributes #[[ATTR_BAR]] = { {{.*}} "no-trapping-math"="true" {{.*}} } -// CHECK: attributes #[[ATTR_COMPATIBLE]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-global-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" } -// CHECK: attributes #[[ATTR_EXTEND]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dpp,+flat-global-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" } -// CHECK: attributes #[[ATTR_INCOMPATIBLE]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-global-insts,+gfx8-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64,-gfx9-insts" } +// CHECK: attributes #[[ATTR_COMPATIBLE]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-buffer-global-fadd-f64-inst,+flat-global-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" } +// CHECK: attributes #[[ATTR_EXTEND]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dpp,+flat-buffer-global-fadd-f64-inst,+flat-global-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" } +// CHECK: attributes #[[ATTR_INCOMPATIBLE]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-buffer-global-fadd-f64-inst,+flat-global-insts,+gfx8-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64,-gfx9-insts" } diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-add.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-add.cl index 95abeb92e59e7..2d038366f8c29 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-add.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-add.cl @@ -2,6 +2,7 @@ // RUN: %clang_cc1 -triple amdgpu9.0a-unknown-unknown -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -triple amdgpu9.42-unknown-unknown -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -triple amdgpu9.50-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu12.50-unknown-unknown -emit-llvm -o - %s | FileCheck %s // REQUIRES: amdgpu-registered-target @@ -27,6 +28,16 @@ float test_atomic_fadd_f32(__amdgpu_buffer_rsrc_t rsrc, float x, int offset, int return __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32(x, rsrc, offset, soffset, 0); } +// CHECK-LABEL: define dso_local double @test_atomic_fadd_f64( +// CHECK-SAME: ptr addrspace(8) nofree captures(none) [[RSRC:%.*]], double noundef [[X:%.*]], i32 noundef [[OFFSET:%.*]], i32 noundef [[SOFFSET:%.*]]) local_unnamed_addr #[[ATTR0]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[TMP0:%.*]] = tail call double @llvm.amdgcn.raw.ptr.buffer.atomic.fadd.f64(double [[X]], ptr addrspace(8) [[RSRC]], i32 [[OFFSET]], i32 [[SOFFSET]], i32 0) +// CHECK-NEXT: ret double [[TMP0]] +// +double test_atomic_fadd_f64(__amdgpu_buffer_rsrc_t rsrc, double x, int offset, int soffset) { + return __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f64(x, rsrc, offset, soffset, 0); +} + // CHECK-LABEL: define dso_local <2 x half> @test_atomic_fadd_v2f16( // CHECK-SAME: ptr addrspace(8) nofree captures(none) [[RSRC:%.*]], <2 x half> noundef [[X:%.*]], i32 noundef [[OFFSET:%.*]], i32 noundef [[SOFFSET:%.*]]) local_unnamed_addr #[[ATTR0]] { // CHECK-NEXT: [[ENTRY:.*:]] diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-atomic-add-err.cl b/clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-atomic-add-err.cl index 1fc23e75ceafb..76b9f7dfa8d2f 100644 --- a/clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-atomic-add-err.cl +++ b/clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-atomic-add-err.cl @@ -3,8 +3,9 @@ typedef half __attribute__((ext_vector_type(2))) float16x2_t; -void test_raw_ptr_atomics(__amdgpu_buffer_rsrc_t rsrc, int i32, float f32, float16x2_t v2f16, int offset, int soffset, int x) { +void test_raw_ptr_atomics(__amdgpu_buffer_rsrc_t rsrc, int i32, float f32, double f64, float16x2_t v2f16, int offset, int soffset, int x) { i32 = __builtin_amdgcn_raw_ptr_buffer_atomic_add_i32(i32, rsrc, offset, soffset, x); // expected-error{{argument to '__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32' must be a constant integer}} f32 = __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32(f32, rsrc, offset, soffset, x); // expected-error{{argument to '__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32' must be a constant integer}} + f64 = __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f64(f64, rsrc, offset, soffset, x); // expected-error{{argument to '__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f64' must be a constant integer}} v2f16 = __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16(v2f16, rsrc, offset, soffset, x); // expected-error{{argument to '__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16' must be a constant integer}} } diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-atomic-add-target-err.cl b/clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-atomic-add-target-err.cl index 6e5b9754fccbe..4b393aa88f5ab 100644 --- a/clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-atomic-add-target-err.cl +++ b/clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-atomic-add-target-err.cl @@ -3,7 +3,8 @@ typedef half __attribute__((ext_vector_type(2))) float16x2_t; -void test_raw_ptr_atomics(__amdgpu_buffer_rsrc_t rsrc, float f32, float16x2_t v2f16, int offset, int soffset) { +void test_raw_ptr_atomics(__amdgpu_buffer_rsrc_t rsrc, float f32, double f64, float16x2_t v2f16, int offset, int soffset) { f32 = __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32(f32, rsrc, offset, soffset, 0); // expected-error{{'__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32' needs target feature atomic-fadd-rtn-insts}} + f64 = __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f64(f64, rsrc, offset, soffset, 0); // expected-error{{'__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f64' needs target feature flat-buffer-global-fadd-f64-inst}} v2f16 = __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16(v2f16, rsrc, offset, soffset, 0); // expected-error{{'__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16' needs target feature atomic-buffer-global-pk-add-f16-insts}} } diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td index ed2ee1ff70f4f..e0473880aa7da 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPU.td +++ b/llvm/lib/Target/AMDGPU/AMDGPU.td @@ -3200,7 +3200,8 @@ def AMDGPUFrontendVisibleFeatures { FeatureDot9Insts, FeatureExtendedImageInsts, FeatureF16BF16ToFP6BF6ConversionScaleInsts, FeatureF32ToF16BF16ConversionSRInsts, FeatureF32ToFP6BF6ConversionScaleInsts, FeatureFP4ConversionScaleInsts, FeatureFP6BF6ConversionScaleInsts, FeatureFP8ConversionInsts, FeatureFP8ConversionScaleInsts, - FeatureFP8E5M3Insts, FeatureFP8Insts, FeatureFlatGlobalInsts, + FeatureFP8E5M3Insts, FeatureFP8Insts, + FeatureFlatBufferGlobalAtomicFaddF64Inst, FeatureFlatGlobalInsts, FeatureGFX10Insts, FeatureGFX10_3Insts, FeatureGFX11Insts, FeatureGFX1250Insts, FeatureGFX1251GEMMInsts, FeatureGFX12Insts, FeatureGFX13Insts, FeatureGFX8Insts, FeatureGFX90AInsts, _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
