[clang] ce880ae - [AMDGPU] Implement builtin for raw-buffer F64 atomic add (#219258)
via cfe-commits
cfe-commits at lists.llvm.org
Thu Aug 27 12:57:36 PDT 2026
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,
More information about the cfe-commits
mailing list