https://github.com/zGoldthorpe updated
https://github.com/llvm/llvm-project/pull/149216
>From c82bff38f45c65d8fbac77bd4ef3369a16b768d9 Mon Sep 17 00:00:00 2001
From: Zach Goldthorpe <zach.goldtho...@amd.com>
Date: Wed, 16 Jul 2025 17:45:04 -0500
Subject: [PATCH 1/2] Exposed some AMDGCN raw buffer atomic intrinsics to
clang.
---
clang/include/clang/Basic/BuiltinsAMDGPU.def | 9 +++
clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 46 +++++++++++
...cn-raw-buffer-atomics-gfx908-target-err.cl | 16 ++++
...ns-amdgcn-raw-buffer-atomics-gfx90a-err.cl | 16 ++++
...iltins-amdgcn-raw-buffer-atomics-gfx90a.cl | 76 +++++++++++++++++++
5 files changed, 163 insertions(+)
create mode 100644
clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx908-target-err.cl
create mode 100644
clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx90a-err.cl
create mode 100644
clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx90a.cl
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def
b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 313c0e640d240..58e26f0a2c458 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -252,6 +252,15 @@ TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fmax_f64,
"dd*0d", "t", "gfx90a-inst
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f64, "dd*3d", "t",
"gfx90a-insts")
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f32, "ff*3f", "t", "gfx8-insts")
+TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32, "iiQbiiIi",
"t", "gfx90a-insts")
+TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32, "ffQbiiIi",
"t", "gfx90a-insts")
+TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f64, "ddQbiiIi",
"t", "gfx90a-insts")
+TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16,
"V2hV2hQbiiIi", "t", "gfx90a-insts")
+
+TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32, "ffQbiiIi",
"t", "gfx90a-insts")
+TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64, "ddQbiiIi",
"t", "gfx90a-insts")
+TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_v2f16,
"V2hV2hQbiiIi", "t", "gfx90a-insts")
+
TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_f32, "ff*0f", "t",
"gfx940-insts")
TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2f16, "V2hV2h*0V2h", "t",
"atomic-flat-pk-add-16-insts")
TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2bf16, "V2sV2s*0V2s", "t",
"atomic-flat-pk-add-16-insts")
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index dcfdea648e93c..4ed6439c0d2d4 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -1439,6 +1439,52 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned
BuiltinID,
F, {EmitScalarExpr(E->getArg(0)), EmitScalarExpr(E->getArg(1)),
EmitScalarExpr(E->getArg(2)), EmitScalarExpr(E->getArg(3))});
}
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32:
+ 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:
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32:
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64:
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_v2f16: {
+ llvm::Type *RetTy;
+ switch (BuiltinID) {
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32:
+ RetTy = Int32Ty;
+ break;
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32:
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32:
+ RetTy = FloatTy;
+ break;
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f64:
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64:
+ RetTy = DoubleTy;
+ break;
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16:
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_v2f16:
+ RetTy = llvm::FixedVectorType::get(HalfTy, 2);
+ }
+ unsigned IID;
+ switch (BuiltinID) {
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32:
+ IID = Intrinsic::amdgcn_raw_ptr_buffer_atomic_add;
+ break;
+ 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:
+ IID = Intrinsic::amdgcn_raw_ptr_buffer_atomic_fadd;
+ break;
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32:
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64:
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_v2f16:
+ IID = Intrinsic::amdgcn_raw_ptr_buffer_atomic_fmax;
+ break;
+ }
+ llvm::Function *F = CGM.getIntrinsic(IID, RetTy);
+ return Builder.CreateCall(
+ F, {EmitScalarExpr(E->getArg(0)), EmitScalarExpr(E->getArg(1)),
+ EmitScalarExpr(E->getArg(2)), EmitScalarExpr(E->getArg(3)),
+ EmitScalarExpr(E->getArg(4))});
+ }
case AMDGPU::BI__builtin_amdgcn_s_prefetch_data:
return emitBuiltinWithOneOverloadedType<2>(
*this, E, Intrinsic::amdgcn_s_prefetch_data);
diff --git
a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx908-target-err.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx908-target-err.cl
new file mode 100644
index 0000000000000..3e72ef98de2c1
--- /dev/null
+++
b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx908-target-err.cl
@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx908 -S
-verify=gfx908,expected -o - %s
+
+// REQUIRES: amdgpu-registered-target
+
+typedef half __attribute__((ext_vector_type(2))) float16x2_t;
+
+void test_raw_ptr_atomics(__amdgpu_buffer_rsrc_t rsrc, int i32, float f32,
double f64, float16x2_t v2f16, int offset, int soffset) {
+ i32 = __builtin_amdgcn_raw_ptr_buffer_atomic_add_i32(i32, rsrc, offset,
soffset, 0); //
expected-error{{'__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32' needs target
feature gfx90a-insts}}
+ 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 gfx90a-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 gfx90a-insts}}
+ 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 gfx90a-insts}}
+
+ f32 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32(f32, rsrc, offset,
soffset, 0); //
expected-error{{'__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32' needs target
feature gfx90a-insts}}
+ f64 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64(f64, rsrc, offset,
soffset, 0); //
expected-error{{'__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64' needs target
feature gfx90a-insts}}
+ v2f16 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_v2f16(v2f16, rsrc,
offset, soffset, 0); //
expected-error{{'__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_v2f16' needs
target feature gfx90a-insts}}
+}
diff --git
a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx90a-err.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx90a-err.cl
new file mode 100644
index 0000000000000..b32c9cc620896
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx90a-err.cl
@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx90a -S
-verify=gfx90a,expected -o - %s
+
+// REQUIRES: amdgpu-registered-target
+
+typedef half __attribute__((ext_vector_type(2))) float16x2_t;
+
+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}}
+
+ f32 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32(f32, rsrc, offset,
soffset, x); // expected-error{{argument to
'__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32' must be a constant integer}}
+ f64 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64(f64, rsrc, offset,
soffset, x); // expected-error{{argument to
'__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64' must be a constant integer}}
+ v2f16 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_v2f16(v2f16, rsrc,
offset, soffset, x); // expected-error{{argument to
'__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_v2f16' must be a constant integer}}
+}
diff --git
a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx90a.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx90a.cl
new file mode 100644
index 0000000000000..f13d4f3459d58
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx90a.cl
@@ -0,0 +1,76 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
UTC_ARGS: --version 5
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx90a
-emit-llvm -o - %s | FileCheck %s
+
+// REQUIRES: amdgpu-registered-target
+
+typedef half __attribute__((ext_vector_type(2))) float16x2_t;
+
+// CHECK-LABEL: define dso_local i32 @test_atomic_add_i32(
+// CHECK-SAME: ptr addrspace(8) captures(none) [[RSRC:%.*]], i32 noundef
[[X:%.*]], i32 noundef [[OFFSET:%.*]], i32 noundef [[SOFFSET:%.*]])
local_unnamed_addr #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[TMP0:%.*]] = tail call i32
@llvm.amdgcn.raw.ptr.buffer.atomic.add.i32(i32 [[X]], ptr addrspace(8)
[[RSRC]], i32 [[OFFSET]], i32 [[SOFFSET]], i32 0)
+// CHECK-NEXT: ret i32 [[TMP0]]
+//
+int test_atomic_add_i32(__amdgpu_buffer_rsrc_t rsrc, int x, int offset, int
soffset) {
+ return __builtin_amdgcn_raw_ptr_buffer_atomic_add_i32(x, rsrc, offset,
soffset, 0);
+}
+
+// CHECK-LABEL: define dso_local float @test_atomic_fadd_f32(
+// CHECK-SAME: ptr addrspace(8) captures(none) [[RSRC:%.*]], float noundef
[[X:%.*]], i32 noundef [[OFFSET:%.*]], i32 noundef [[SOFFSET:%.*]])
local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[TMP0:%.*]] = tail call float
@llvm.amdgcn.raw.ptr.buffer.atomic.fadd.f32(float [[X]], ptr addrspace(8)
[[RSRC]], i32 [[OFFSET]], i32 [[SOFFSET]], i32 0)
+// CHECK-NEXT: ret float [[TMP0]]
+//
+float test_atomic_fadd_f32(__amdgpu_buffer_rsrc_t rsrc, float x, int offset,
int soffset) {
+ 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) 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) captures(none) [[RSRC:%.*]], <2 x half>
noundef [[X:%.*]], i32 noundef [[OFFSET:%.*]], i32 noundef [[SOFFSET:%.*]])
local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x half>
@llvm.amdgcn.raw.ptr.buffer.atomic.fadd.v2f16(<2 x half> [[X]], ptr
addrspace(8) [[RSRC]], i32 [[OFFSET]], i32 [[SOFFSET]], i32 0)
+// CHECK-NEXT: ret <2 x half> [[TMP0]]
+//
+float16x2_t test_atomic_fadd_v2f16(__amdgpu_buffer_rsrc_t rsrc, float16x2_t x,
int offset, int soffset) {
+ return __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16(x, rsrc, offset,
soffset, 0);
+}
+
+// CHECK-LABEL: define dso_local float @test_atomic_fmax_f32(
+// CHECK-SAME: ptr addrspace(8) captures(none) [[RSRC:%.*]], float noundef
[[X:%.*]], i32 noundef [[OFFSET:%.*]], i32 noundef [[SOFFSET:%.*]])
local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[TMP0:%.*]] = tail call float
@llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f32(float [[X]], ptr addrspace(8)
[[RSRC]], i32 [[OFFSET]], i32 [[SOFFSET]], i32 0)
+// CHECK-NEXT: ret float [[TMP0]]
+//
+float test_atomic_fmax_f32(__amdgpu_buffer_rsrc_t rsrc, float x, int offset,
int soffset) {
+ return __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32(x, rsrc, offset,
soffset, 0);
+}
+
+// CHECK-LABEL: define dso_local double @test_atomic_fmax_f64(
+// CHECK-SAME: ptr addrspace(8) 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.fmax.f64(double [[X]], ptr addrspace(8)
[[RSRC]], i32 [[OFFSET]], i32 [[SOFFSET]], i32 0)
+// CHECK-NEXT: ret double [[TMP0]]
+//
+double test_atomic_fmax_f64(__amdgpu_buffer_rsrc_t rsrc, double x, int offset,
int soffset) {
+ return __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64(x, rsrc, offset,
soffset, 0);
+}
+
+// CHECK-LABEL: define dso_local <2 x half> @test_atomic_fmax_v2f16(
+// CHECK-SAME: ptr addrspace(8) captures(none) [[RSRC:%.*]], <2 x half>
noundef [[X:%.*]], i32 noundef [[OFFSET:%.*]], i32 noundef [[SOFFSET:%.*]])
local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x half>
@llvm.amdgcn.raw.ptr.buffer.atomic.fmax.v2f16(<2 x half> [[X]], ptr
addrspace(8) [[RSRC]], i32 [[OFFSET]], i32 [[SOFFSET]], i32 0)
+// CHECK-NEXT: ret <2 x half> [[TMP0]]
+//
+float16x2_t test_atomic_fmax_v2f16(__amdgpu_buffer_rsrc_t rsrc, float16x2_t x,
int offset, int soffset) {
+ return __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_v2f16(x, rsrc, offset,
soffset, 0);
+}
>From 59e1b303486a6c051deeeec6a91773bdfb82aa3e Mon Sep 17 00:00:00 2001
From: Zach Goldthorpe <zach.goldtho...@amd.com>
Date: Thu, 17 Jul 2025 12:29:01 -0500
Subject: [PATCH 2/2] Refined choice of target features.
---
clang/include/clang/Basic/BuiltinsAMDGPU.def | 16 ++--
clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 9 +--
...iltins-amdgcn-raw-buffer-atomic-add-err.cl | 10 +++
...amdgcn-raw-buffer-atomic-add-target-err.cl | 9 +++
.../builtins-amdgcn-raw-buffer-atomic-add.cl | 36 +++++++++
...iltins-amdgcn-raw-buffer-atomic-max-err.cl | 7 ++
...amdgcn-raw-buffer-atomic-max-target-err.cl | 7 ++
.../builtins-amdgcn-raw-buffer-atomic-max.cl | 24 ++++++
...cn-raw-buffer-atomics-gfx908-target-err.cl | 16 ----
...ns-amdgcn-raw-buffer-atomics-gfx90a-err.cl | 16 ----
...iltins-amdgcn-raw-buffer-atomics-gfx90a.cl | 76 -------------------
11 files changed, 102 insertions(+), 124 deletions(-)
create mode 100644
clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-add-err.cl
create mode 100644
clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-add-target-err.cl
create mode 100644
clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-add.cl
create mode 100644
clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max-err.cl
create mode 100644
clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max-target-err.cl
create mode 100644
clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max.cl
delete mode 100644
clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx908-target-err.cl
delete mode 100644
clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx90a-err.cl
delete mode 100644
clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx90a.cl
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def
b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 58e26f0a2c458..9b1e618a490a5 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -163,6 +163,13 @@ BUILTIN(__builtin_amdgcn_raw_buffer_load_b64,
"V2UiQbiiIi", "n")
BUILTIN(__builtin_amdgcn_raw_buffer_load_b96, "V3UiQbiiIi", "n")
BUILTIN(__builtin_amdgcn_raw_buffer_load_b128, "V4UiQbiiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32, "iiQbiiIi", "t")
+TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32, "ffQbiiIi",
"t", "atomic-fadd-rtn-insts")
+TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16,
"V2hV2hQbiiIi", "t", "atomic-buffer-global-pk-add-f16-insts")
+
+TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32, "ffQbiiIi",
"t", "atomic-fmin-fmax-global-f32")
+TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64, "ddQbiiIi",
"t", "atomic-fmin-fmax-global-f64")
+
TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_load_lds, "vQbv*3IUiiiIiIi",
"t", "vmem-to-lds-load-insts")
TARGET_BUILTIN(__builtin_amdgcn_struct_ptr_buffer_load_lds,
"vQbv*3IUiiiiIiIi", "t", "vmem-to-lds-load-insts")
@@ -252,15 +259,6 @@ TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fmax_f64,
"dd*0d", "t", "gfx90a-inst
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f64, "dd*3d", "t",
"gfx90a-insts")
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f32, "ff*3f", "t", "gfx8-insts")
-TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32, "iiQbiiIi",
"t", "gfx90a-insts")
-TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32, "ffQbiiIi",
"t", "gfx90a-insts")
-TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f64, "ddQbiiIi",
"t", "gfx90a-insts")
-TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16,
"V2hV2hQbiiIi", "t", "gfx90a-insts")
-
-TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32, "ffQbiiIi",
"t", "gfx90a-insts")
-TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64, "ddQbiiIi",
"t", "gfx90a-insts")
-TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_v2f16,
"V2hV2hQbiiIi", "t", "gfx90a-insts")
-
TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_f32, "ff*0f", "t",
"gfx940-insts")
TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2f16, "V2hV2h*0V2h", "t",
"atomic-flat-pk-add-16-insts")
TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2bf16, "V2sV2s*0V2s", "t",
"atomic-flat-pk-add-16-insts")
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index 4ed6439c0d2d4..0c9b8fafc8cd2 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -1441,11 +1441,9 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned
BuiltinID,
}
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32:
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:
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32:
- case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64:
- case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_v2f16: {
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64: {
llvm::Type *RetTy;
switch (BuiltinID) {
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32:
@@ -1455,13 +1453,12 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned
BuiltinID,
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32:
RetTy = FloatTy;
break;
- case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f64:
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64:
RetTy = DoubleTy;
break;
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16:
- case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_v2f16:
RetTy = llvm::FixedVectorType::get(HalfTy, 2);
+ break;
}
unsigned IID;
switch (BuiltinID) {
@@ -1469,13 +1466,11 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned
BuiltinID,
IID = Intrinsic::amdgcn_raw_ptr_buffer_atomic_add;
break;
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:
IID = Intrinsic::amdgcn_raw_ptr_buffer_atomic_fadd;
break;
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32:
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64:
- case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_v2f16:
IID = Intrinsic::amdgcn_raw_ptr_buffer_atomic_fmax;
break;
}
diff --git
a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-add-err.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-add-err.cl
new file mode 100644
index 0000000000000..5c5ae937797e8
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-add-err.cl
@@ -0,0 +1,10 @@
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx90a -S
-verify=gfx90a,expected -o - %s
+// REQUIRES: amdgpu-registered-target
+
+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) {
+ 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}}
+ 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/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-add-target-err.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-add-target-err.cl
new file mode 100644
index 0000000000000..509a498af5efa
--- /dev/null
+++
b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-add-target-err.cl
@@ -0,0 +1,9 @@
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx908 -S
-verify=gfx908,expected -o - %s
+// REQUIRES: amdgpu-registered-target
+
+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) {
+ 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}}
+ 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/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-add.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-add.cl
new file mode 100644
index 0000000000000..7641e10c89c04
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-add.cl
@@ -0,0 +1,36 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
UTC_ARGS: --version 5
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx90a
-emit-llvm -o - %s | FileCheck %s
+
+// REQUIRES: amdgpu-registered-target
+
+typedef half __attribute__((ext_vector_type(2))) float16x2_t;
+
+// CHECK-LABEL: define dso_local i32 @test_atomic_add_i32(
+// CHECK-SAME: ptr addrspace(8) captures(none) [[RSRC:%.*]], i32 noundef
[[X:%.*]], i32 noundef [[OFFSET:%.*]], i32 noundef [[SOFFSET:%.*]])
local_unnamed_addr #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[TMP0:%.*]] = tail call i32
@llvm.amdgcn.raw.ptr.buffer.atomic.add.i32(i32 [[X]], ptr addrspace(8)
[[RSRC]], i32 [[OFFSET]], i32 [[SOFFSET]], i32 0)
+// CHECK-NEXT: ret i32 [[TMP0]]
+//
+int test_atomic_add_i32(__amdgpu_buffer_rsrc_t rsrc, int x, int offset, int
soffset) {
+ return __builtin_amdgcn_raw_ptr_buffer_atomic_add_i32(x, rsrc, offset,
soffset, 0);
+}
+
+// CHECK-LABEL: define dso_local float @test_atomic_fadd_f32(
+// CHECK-SAME: ptr addrspace(8) captures(none) [[RSRC:%.*]], float noundef
[[X:%.*]], i32 noundef [[OFFSET:%.*]], i32 noundef [[SOFFSET:%.*]])
local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[TMP0:%.*]] = tail call float
@llvm.amdgcn.raw.ptr.buffer.atomic.fadd.f32(float [[X]], ptr addrspace(8)
[[RSRC]], i32 [[OFFSET]], i32 [[SOFFSET]], i32 0)
+// CHECK-NEXT: ret float [[TMP0]]
+//
+float test_atomic_fadd_f32(__amdgpu_buffer_rsrc_t rsrc, float x, int offset,
int soffset) {
+ return __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32(x, rsrc, offset,
soffset, 0);
+}
+
+// CHECK-LABEL: define dso_local <2 x half> @test_atomic_fadd_v2f16(
+// CHECK-SAME: ptr addrspace(8) captures(none) [[RSRC:%.*]], <2 x half>
noundef [[X:%.*]], i32 noundef [[OFFSET:%.*]], i32 noundef [[SOFFSET:%.*]])
local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x half>
@llvm.amdgcn.raw.ptr.buffer.atomic.fadd.v2f16(<2 x half> [[X]], ptr
addrspace(8) [[RSRC]], i32 [[OFFSET]], i32 [[SOFFSET]], i32 0)
+// CHECK-NEXT: ret <2 x half> [[TMP0]]
+//
+float16x2_t test_atomic_fadd_v2f16(__amdgpu_buffer_rsrc_t rsrc, float16x2_t x,
int offset, int soffset) {
+ return __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16(x, rsrc, offset,
soffset, 0);
+}
diff --git
a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max-err.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max-err.cl
new file mode 100644
index 0000000000000..aeb7426177d6d
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max-err.cl
@@ -0,0 +1,7 @@
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-feature
+atomic-fmin-fmax-global-f32 -target-feature +atomic-fmin-fmax-global-f64 -S
-verify=expected -o - %s
+// REQUIRES: amdgpu-registered-target
+
+void test_raw_ptr_atomics(__amdgpu_buffer_rsrc_t rsrc, float f32, double f64,
int offset, int soffset, int x) {
+ f32 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32(f32, rsrc, offset,
soffset, x); // expected-error{{argument to
'__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32' must be a constant integer}}
+ f64 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64(f64, rsrc, offset,
soffset, x); // expected-error{{argument to
'__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64' must be a constant integer}}
+}
diff --git
a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max-target-err.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max-target-err.cl
new file mode 100644
index 0000000000000..3f1ad416d132f
--- /dev/null
+++
b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max-target-err.cl
@@ -0,0 +1,7 @@
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -verify=expected -o - %s
+// REQUIRES: amdgpu-registered-target
+
+void test_raw_ptr_atomics(__amdgpu_buffer_rsrc_t rsrc, float f32, double f64,
int offset, int soffset) {
+ f32 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32(f32, rsrc, offset,
soffset, 0); //
expected-error{{'__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32' needs target
feature atomic-fmin-fmax-global-f32}}
+ f64 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64(f64, rsrc, offset,
soffset, 0); //
expected-error{{'__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64' needs target
feature atomic-fmin-fmax-global-f64}}
+}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max.cl
new file mode 100644
index 0000000000000..d2a9314906a65
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max.cl
@@ -0,0 +1,24 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
UTC_ARGS: --version 5
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-feature
+atomic-fmin-fmax-global-f32 -target-feature +atomic-fmin-fmax-global-f64
-emit-llvm -o - %s | FileCheck %s
+
+// REQUIRES: amdgpu-registered-target
+
+// CHECK-LABEL: define dso_local float @test_atomic_fmax_f32(
+// CHECK-SAME: ptr addrspace(8) captures(none) [[RSRC:%.*]], float noundef
[[X:%.*]], i32 noundef [[OFFSET:%.*]], i32 noundef [[SOFFSET:%.*]])
local_unnamed_addr #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[TMP0:%.*]] = tail call float
@llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f32(float [[X]], ptr addrspace(8)
[[RSRC]], i32 [[OFFSET]], i32 [[SOFFSET]], i32 0)
+// CHECK-NEXT: ret float [[TMP0]]
+//
+float test_atomic_fmax_f32(__amdgpu_buffer_rsrc_t rsrc, float x, int offset,
int soffset) {
+ return __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32(x, rsrc, offset,
soffset, 0);
+}
+
+// CHECK-LABEL: define dso_local double @test_atomic_fmax_f64(
+// CHECK-SAME: ptr addrspace(8) 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.fmax.f64(double [[X]], ptr addrspace(8)
[[RSRC]], i32 [[OFFSET]], i32 [[SOFFSET]], i32 0)
+// CHECK-NEXT: ret double [[TMP0]]
+//
+double test_atomic_fmax_f64(__amdgpu_buffer_rsrc_t rsrc, double x, int offset,
int soffset) {
+ return __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64(x, rsrc, offset,
soffset, 0);
+}
diff --git
a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx908-target-err.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx908-target-err.cl
deleted file mode 100644
index 3e72ef98de2c1..0000000000000
---
a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx908-target-err.cl
+++ /dev/null
@@ -1,16 +0,0 @@
-// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx908 -S
-verify=gfx908,expected -o - %s
-
-// REQUIRES: amdgpu-registered-target
-
-typedef half __attribute__((ext_vector_type(2))) float16x2_t;
-
-void test_raw_ptr_atomics(__amdgpu_buffer_rsrc_t rsrc, int i32, float f32,
double f64, float16x2_t v2f16, int offset, int soffset) {
- i32 = __builtin_amdgcn_raw_ptr_buffer_atomic_add_i32(i32, rsrc, offset,
soffset, 0); //
expected-error{{'__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32' needs target
feature gfx90a-insts}}
- 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 gfx90a-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 gfx90a-insts}}
- 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 gfx90a-insts}}
-
- f32 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32(f32, rsrc, offset,
soffset, 0); //
expected-error{{'__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32' needs target
feature gfx90a-insts}}
- f64 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64(f64, rsrc, offset,
soffset, 0); //
expected-error{{'__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64' needs target
feature gfx90a-insts}}
- v2f16 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_v2f16(v2f16, rsrc,
offset, soffset, 0); //
expected-error{{'__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_v2f16' needs
target feature gfx90a-insts}}
-}
diff --git
a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx90a-err.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx90a-err.cl
deleted file mode 100644
index b32c9cc620896..0000000000000
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx90a-err.cl
+++ /dev/null
@@ -1,16 +0,0 @@
-// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx90a -S
-verify=gfx90a,expected -o - %s
-
-// REQUIRES: amdgpu-registered-target
-
-typedef half __attribute__((ext_vector_type(2))) float16x2_t;
-
-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}}
-
- f32 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32(f32, rsrc, offset,
soffset, x); // expected-error{{argument to
'__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32' must be a constant integer}}
- f64 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64(f64, rsrc, offset,
soffset, x); // expected-error{{argument to
'__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64' must be a constant integer}}
- v2f16 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_v2f16(v2f16, rsrc,
offset, soffset, x); // expected-error{{argument to
'__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_v2f16' must be a constant integer}}
-}
diff --git
a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx90a.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx90a.cl
deleted file mode 100644
index f13d4f3459d58..0000000000000
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomics-gfx90a.cl
+++ /dev/null
@@ -1,76 +0,0 @@
-// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
UTC_ARGS: --version 5
-// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx90a
-emit-llvm -o - %s | FileCheck %s
-
-// REQUIRES: amdgpu-registered-target
-
-typedef half __attribute__((ext_vector_type(2))) float16x2_t;
-
-// CHECK-LABEL: define dso_local i32 @test_atomic_add_i32(
-// CHECK-SAME: ptr addrspace(8) captures(none) [[RSRC:%.*]], i32 noundef
[[X:%.*]], i32 noundef [[OFFSET:%.*]], i32 noundef [[SOFFSET:%.*]])
local_unnamed_addr #[[ATTR0:[0-9]+]] {
-// CHECK-NEXT: [[ENTRY:.*:]]
-// CHECK-NEXT: [[TMP0:%.*]] = tail call i32
@llvm.amdgcn.raw.ptr.buffer.atomic.add.i32(i32 [[X]], ptr addrspace(8)
[[RSRC]], i32 [[OFFSET]], i32 [[SOFFSET]], i32 0)
-// CHECK-NEXT: ret i32 [[TMP0]]
-//
-int test_atomic_add_i32(__amdgpu_buffer_rsrc_t rsrc, int x, int offset, int
soffset) {
- return __builtin_amdgcn_raw_ptr_buffer_atomic_add_i32(x, rsrc, offset,
soffset, 0);
-}
-
-// CHECK-LABEL: define dso_local float @test_atomic_fadd_f32(
-// CHECK-SAME: ptr addrspace(8) captures(none) [[RSRC:%.*]], float noundef
[[X:%.*]], i32 noundef [[OFFSET:%.*]], i32 noundef [[SOFFSET:%.*]])
local_unnamed_addr #[[ATTR0]] {
-// CHECK-NEXT: [[ENTRY:.*:]]
-// CHECK-NEXT: [[TMP0:%.*]] = tail call float
@llvm.amdgcn.raw.ptr.buffer.atomic.fadd.f32(float [[X]], ptr addrspace(8)
[[RSRC]], i32 [[OFFSET]], i32 [[SOFFSET]], i32 0)
-// CHECK-NEXT: ret float [[TMP0]]
-//
-float test_atomic_fadd_f32(__amdgpu_buffer_rsrc_t rsrc, float x, int offset,
int soffset) {
- 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) 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) captures(none) [[RSRC:%.*]], <2 x half>
noundef [[X:%.*]], i32 noundef [[OFFSET:%.*]], i32 noundef [[SOFFSET:%.*]])
local_unnamed_addr #[[ATTR0]] {
-// CHECK-NEXT: [[ENTRY:.*:]]
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x half>
@llvm.amdgcn.raw.ptr.buffer.atomic.fadd.v2f16(<2 x half> [[X]], ptr
addrspace(8) [[RSRC]], i32 [[OFFSET]], i32 [[SOFFSET]], i32 0)
-// CHECK-NEXT: ret <2 x half> [[TMP0]]
-//
-float16x2_t test_atomic_fadd_v2f16(__amdgpu_buffer_rsrc_t rsrc, float16x2_t x,
int offset, int soffset) {
- return __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16(x, rsrc, offset,
soffset, 0);
-}
-
-// CHECK-LABEL: define dso_local float @test_atomic_fmax_f32(
-// CHECK-SAME: ptr addrspace(8) captures(none) [[RSRC:%.*]], float noundef
[[X:%.*]], i32 noundef [[OFFSET:%.*]], i32 noundef [[SOFFSET:%.*]])
local_unnamed_addr #[[ATTR0]] {
-// CHECK-NEXT: [[ENTRY:.*:]]
-// CHECK-NEXT: [[TMP0:%.*]] = tail call float
@llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f32(float [[X]], ptr addrspace(8)
[[RSRC]], i32 [[OFFSET]], i32 [[SOFFSET]], i32 0)
-// CHECK-NEXT: ret float [[TMP0]]
-//
-float test_atomic_fmax_f32(__amdgpu_buffer_rsrc_t rsrc, float x, int offset,
int soffset) {
- return __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32(x, rsrc, offset,
soffset, 0);
-}
-
-// CHECK-LABEL: define dso_local double @test_atomic_fmax_f64(
-// CHECK-SAME: ptr addrspace(8) 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.fmax.f64(double [[X]], ptr addrspace(8)
[[RSRC]], i32 [[OFFSET]], i32 [[SOFFSET]], i32 0)
-// CHECK-NEXT: ret double [[TMP0]]
-//
-double test_atomic_fmax_f64(__amdgpu_buffer_rsrc_t rsrc, double x, int offset,
int soffset) {
- return __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64(x, rsrc, offset,
soffset, 0);
-}
-
-// CHECK-LABEL: define dso_local <2 x half> @test_atomic_fmax_v2f16(
-// CHECK-SAME: ptr addrspace(8) captures(none) [[RSRC:%.*]], <2 x half>
noundef [[X:%.*]], i32 noundef [[OFFSET:%.*]], i32 noundef [[SOFFSET:%.*]])
local_unnamed_addr #[[ATTR0]] {
-// CHECK-NEXT: [[ENTRY:.*:]]
-// CHECK-NEXT: [[TMP0:%.*]] = tail call <2 x half>
@llvm.amdgcn.raw.ptr.buffer.atomic.fmax.v2f16(<2 x half> [[X]], ptr
addrspace(8) [[RSRC]], i32 [[OFFSET]], i32 [[SOFFSET]], i32 0)
-// CHECK-NEXT: ret <2 x half> [[TMP0]]
-//
-float16x2_t test_atomic_fmax_v2f16(__amdgpu_buffer_rsrc_t rsrc, float16x2_t x,
int offset, int soffset) {
- return __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_v2f16(x, rsrc, offset,
soffset, 0);
-}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
