https://github.com/shiltian created
https://github.com/llvm/llvm-project/pull/94576
None
>From c874a4d69d472a83dfcd11e8c07518e4d216b725 Mon Sep 17 00:00:00 2001
From: Shilei Tian <i...@tianshilei.me>
Date: Thu, 6 Jun 2024 01:46:17 -0400
Subject: [PATCH] [Clang][AMDGPU] Add builtins for instrinsic
`llvm.amdgcn.raw.buffer.store`
---
clang/include/clang/Basic/BuiltinsAMDGPU.def | 14 +
clang/lib/CodeGen/CGBuiltin.cpp | 23 ++
.../builtins-amdgcn-raw-buffer-store.cl | 264 ++++++++++++++++++
.../builtins-amdgcn-raw-buffer-store-error.cl | 94 +++++++
4 files changed, 395 insertions(+)
create mode 100644 clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store.cl
create mode 100644
clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-store-error.cl
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def
b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 433c7795325f0..d6866304f8b1b 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -148,6 +148,20 @@ BUILTIN(__builtin_amdgcn_qsad_pk_u16_u8, "WUiWUiUiWUi",
"nc")
BUILTIN(__builtin_amdgcn_mqsad_pk_u16_u8, "WUiWUiUiWUi", "nc")
BUILTIN(__builtin_amdgcn_mqsad_u32_u8, "V4UiWUiUiV4Ui", "nc")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_i8, "vcV4iIiIiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_i16, "vsV4iIiIiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_i32, "viV4iIiIiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_f16, "vhV4iIiIiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_f32, "vfV4iIiIiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_v2i16, "vV2sV4iIiIiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_v2i32, "vV2iV4iIiIiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_v2f16, "vV2hV4iIiIiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_v2f32, "vV2fV4iIiIiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_v4i16, "vV4sV4iIiIiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_v4i32, "vV4iV4iIiIiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_v4f16, "vV4hV4iIiIiIi", "n")
+BUILTIN(__builtin_amdgcn_raw_buffer_store_v4f32, "vV4fV4iIiIiIi", "n")
+
//===----------------------------------------------------------------------===//
// Ballot builtins.
//===----------------------------------------------------------------------===//
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 37d0c478e0330..706141cbc85ce 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -19063,6 +19063,29 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned
BuiltinID,
CGM.getIntrinsic(Intrinsic::amdgcn_s_sendmsg_rtn, {ResultType});
return Builder.CreateCall(F, {Arg});
}
+
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_i8:
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_i16:
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_i32:
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_f32:
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_f16:
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_v2i16:
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_v2i32:
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_v2f16:
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_v2f32:
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_v4i16:
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_v4i32:
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_v4f16:
+ case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_v4f32: {
+ llvm::Value *VData = EmitScalarExpr(E->getArg(0));
+ llvm::Value *Rsrc = EmitScalarExpr(E->getArg(1));
+ llvm::Value *Offset = EmitScalarExpr(E->getArg(2));
+ llvm::Value *SOffset = EmitScalarExpr(E->getArg(3));
+ llvm::Value *Aux = EmitScalarExpr(E->getArg(4));
+ Function *F =
+ CGM.getIntrinsic(Intrinsic::amdgcn_raw_buffer_store, VData->getType());
+ return Builder.CreateCall(F, {VData, Rsrc, Offset, SOffset, Aux});
+ }
default:
return nullptr;
}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store.cl
new file mode 100644
index 0000000000000..0553c53ab41ad
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store.cl
@@ -0,0 +1,264 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm
-o - %s | FileCheck %s --check-prefixes=VERDE
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -emit-llvm
-o - %s | FileCheck %s --check-prefixes=GFX8
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100
-emit-llvm -o - %s | FileCheck %s --check-prefixes=GFX11
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+typedef short v2i16 __attribute__((ext_vector_type(2)));
+typedef int v2i32 __attribute__((ext_vector_type(2)));
+typedef half v2f16 __attribute__((ext_vector_type(2)));
+typedef float v2f32 __attribute__((ext_vector_type(2)));
+typedef short v4i16 __attribute__((ext_vector_type(4)));
+typedef int v4i32 __attribute__((ext_vector_type(4)));
+typedef half v4f16 __attribute__((ext_vector_type(4)));
+typedef float v4f32 __attribute__((ext_vector_type(4)));
+
+// VERDE-LABEL: @test_amdgcn_raw_buffer_store_i8(
+// VERDE-NEXT: entry:
+// VERDE-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.i8(i8
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// VERDE-NEXT: ret void
+//
+// GFX8-LABEL: @test_amdgcn_raw_buffer_store_i8(
+// GFX8-NEXT: entry:
+// GFX8-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.i8(i8
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX8-NEXT: ret void
+//
+// GFX11-LABEL: @test_amdgcn_raw_buffer_store_i8(
+// GFX11-NEXT: entry:
+// GFX11-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.i8(i8
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX11-NEXT: ret void
+//
+void test_amdgcn_raw_buffer_store_i8(char vdata, v4i32 rsrc) {
+ __builtin_amdgcn_raw_buffer_store_i8(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, /*aux=*/0);
+}
+
+// VERDE-LABEL: @test_amdgcn_raw_buffer_store_i16(
+// VERDE-NEXT: entry:
+// VERDE-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.i16(i16
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// VERDE-NEXT: ret void
+//
+// GFX8-LABEL: @test_amdgcn_raw_buffer_store_i16(
+// GFX8-NEXT: entry:
+// GFX8-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.i16(i16
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX8-NEXT: ret void
+//
+// GFX11-LABEL: @test_amdgcn_raw_buffer_store_i16(
+// GFX11-NEXT: entry:
+// GFX11-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.i16(i16
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX11-NEXT: ret void
+//
+void test_amdgcn_raw_buffer_store_i16(short vdata, v4i32 rsrc) {
+ __builtin_amdgcn_raw_buffer_store_i16(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, /*aux=*/0);
+}
+
+// VERDE-LABEL: @test_amdgcn_raw_buffer_store_i32(
+// VERDE-NEXT: entry:
+// VERDE-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.i32(i32
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// VERDE-NEXT: ret void
+//
+// GFX8-LABEL: @test_amdgcn_raw_buffer_store_i32(
+// GFX8-NEXT: entry:
+// GFX8-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.i32(i32
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX8-NEXT: ret void
+//
+// GFX11-LABEL: @test_amdgcn_raw_buffer_store_i32(
+// GFX11-NEXT: entry:
+// GFX11-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.i32(i32
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX11-NEXT: ret void
+//
+void test_amdgcn_raw_buffer_store_i32(int vdata, v4i32 rsrc) {
+ __builtin_amdgcn_raw_buffer_store_i32(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, /*aux=*/0);
+}
+
+// VERDE-LABEL: @test_amdgcn_raw_buffer_store_f16(
+// VERDE-NEXT: entry:
+// VERDE-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.f16(half
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// VERDE-NEXT: ret void
+//
+// GFX8-LABEL: @test_amdgcn_raw_buffer_store_f16(
+// GFX8-NEXT: entry:
+// GFX8-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.f16(half
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX8-NEXT: ret void
+//
+// GFX11-LABEL: @test_amdgcn_raw_buffer_store_f16(
+// GFX11-NEXT: entry:
+// GFX11-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.f16(half
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX11-NEXT: ret void
+//
+void test_amdgcn_raw_buffer_store_f16(half vdata, v4i32 rsrc) {
+ __builtin_amdgcn_raw_buffer_store_f16(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, /*aux=*/0);
+}
+
+// VERDE-LABEL: @test_amdgcn_raw_buffer_store_f32(
+// VERDE-NEXT: entry:
+// VERDE-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.f32(float
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// VERDE-NEXT: ret void
+//
+// GFX8-LABEL: @test_amdgcn_raw_buffer_store_f32(
+// GFX8-NEXT: entry:
+// GFX8-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.f32(float
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX8-NEXT: ret void
+//
+// GFX11-LABEL: @test_amdgcn_raw_buffer_store_f32(
+// GFX11-NEXT: entry:
+// GFX11-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.f32(float
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX11-NEXT: ret void
+//
+void test_amdgcn_raw_buffer_store_f32(float vdata, v4i32 rsrc) {
+ __builtin_amdgcn_raw_buffer_store_f32(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, /*aux=*/0);
+}
+
+// VERDE-LABEL: @test_amdgcn_raw_buffer_store_v2i16(
+// VERDE-NEXT: entry:
+// VERDE-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v2i16(<2 x i16>
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// VERDE-NEXT: ret void
+//
+// GFX8-LABEL: @test_amdgcn_raw_buffer_store_v2i16(
+// GFX8-NEXT: entry:
+// GFX8-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v2i16(<2 x i16>
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX8-NEXT: ret void
+//
+// GFX11-LABEL: @test_amdgcn_raw_buffer_store_v2i16(
+// GFX11-NEXT: entry:
+// GFX11-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v2i16(<2 x i16>
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX11-NEXT: ret void
+//
+void test_amdgcn_raw_buffer_store_v2i16(v2i16 vdata, v4i32 rsrc) {
+ __builtin_amdgcn_raw_buffer_store_v2i16(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, /*aux=*/0);
+}
+
+// VERDE-LABEL: @test_amdgcn_raw_buffer_store_v2i32(
+// VERDE-NEXT: entry:
+// VERDE-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v2i32(<2 x i32>
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// VERDE-NEXT: ret void
+//
+// GFX8-LABEL: @test_amdgcn_raw_buffer_store_v2i32(
+// GFX8-NEXT: entry:
+// GFX8-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v2i32(<2 x i32>
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX8-NEXT: ret void
+//
+// GFX11-LABEL: @test_amdgcn_raw_buffer_store_v2i32(
+// GFX11-NEXT: entry:
+// GFX11-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v2i32(<2 x i32>
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX11-NEXT: ret void
+//
+void test_amdgcn_raw_buffer_store_v2i32(v2i32 vdata, v4i32 rsrc) {
+ __builtin_amdgcn_raw_buffer_store_v2i32(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, /*aux=*/0);
+}
+
+// VERDE-LABEL: @test_amdgcn_raw_buffer_store_v2f16(
+// VERDE-NEXT: entry:
+// VERDE-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v2f16(<2 x
half> [[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// VERDE-NEXT: ret void
+//
+// GFX8-LABEL: @test_amdgcn_raw_buffer_store_v2f16(
+// GFX8-NEXT: entry:
+// GFX8-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v2f16(<2 x half>
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX8-NEXT: ret void
+//
+// GFX11-LABEL: @test_amdgcn_raw_buffer_store_v2f16(
+// GFX11-NEXT: entry:
+// GFX11-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v2f16(<2 x
half> [[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX11-NEXT: ret void
+//
+void test_amdgcn_raw_buffer_store_v2f16(v2f16 vdata, v4i32 rsrc) {
+ __builtin_amdgcn_raw_buffer_store_v2f16(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, /*aux=*/0);
+}
+
+// VERDE-LABEL: @test_amdgcn_raw_buffer_store_v2f32(
+// VERDE-NEXT: entry:
+// VERDE-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v2f32(<2 x
float> [[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// VERDE-NEXT: ret void
+//
+// GFX8-LABEL: @test_amdgcn_raw_buffer_store_v2f32(
+// GFX8-NEXT: entry:
+// GFX8-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v2f32(<2 x
float> [[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX8-NEXT: ret void
+//
+// GFX11-LABEL: @test_amdgcn_raw_buffer_store_v2f32(
+// GFX11-NEXT: entry:
+// GFX11-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v2f32(<2 x
float> [[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX11-NEXT: ret void
+//
+void test_amdgcn_raw_buffer_store_v2f32(v2f32 vdata, v4i32 rsrc) {
+ __builtin_amdgcn_raw_buffer_store_v2f32(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, /*aux=*/0);
+}
+
+// VERDE-LABEL: @test_amdgcn_raw_buffer_store_v4i16(
+// VERDE-NEXT: entry:
+// VERDE-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v4i16(<4 x i16>
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// VERDE-NEXT: ret void
+//
+// GFX8-LABEL: @test_amdgcn_raw_buffer_store_v4i16(
+// GFX8-NEXT: entry:
+// GFX8-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v4i16(<4 x i16>
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX8-NEXT: ret void
+//
+// GFX11-LABEL: @test_amdgcn_raw_buffer_store_v4i16(
+// GFX11-NEXT: entry:
+// GFX11-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v4i16(<4 x i16>
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX11-NEXT: ret void
+//
+void test_amdgcn_raw_buffer_store_v4i16(v4i16 vdata, v4i32 rsrc) {
+ __builtin_amdgcn_raw_buffer_store_v4i16(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, /*aux=*/0);
+}
+
+// VERDE-LABEL: @test_amdgcn_raw_buffer_store_v4i32(
+// VERDE-NEXT: entry:
+// VERDE-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v4i32(<4 x i32>
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// VERDE-NEXT: ret void
+//
+// GFX8-LABEL: @test_amdgcn_raw_buffer_store_v4i32(
+// GFX8-NEXT: entry:
+// GFX8-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v4i32(<4 x i32>
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX8-NEXT: ret void
+//
+// GFX11-LABEL: @test_amdgcn_raw_buffer_store_v4i32(
+// GFX11-NEXT: entry:
+// GFX11-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v4i32(<4 x i32>
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX11-NEXT: ret void
+//
+void test_amdgcn_raw_buffer_store_v4i32(v4i32 vdata, v4i32 rsrc) {
+ __builtin_amdgcn_raw_buffer_store_v4i32(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, /*aux=*/0);
+}
+
+// VERDE-LABEL: @test_amdgcn_raw_buffer_store_v4f16(
+// VERDE-NEXT: entry:
+// VERDE-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v4f16(<4 x
half> [[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// VERDE-NEXT: ret void
+//
+// GFX8-LABEL: @test_amdgcn_raw_buffer_store_v4f16(
+// GFX8-NEXT: entry:
+// GFX8-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v4f16(<4 x half>
[[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX8-NEXT: ret void
+//
+// GFX11-LABEL: @test_amdgcn_raw_buffer_store_v4f16(
+// GFX11-NEXT: entry:
+// GFX11-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v4f16(<4 x
half> [[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX11-NEXT: ret void
+//
+void test_amdgcn_raw_buffer_store_v4f16(v4f16 vdata, v4i32 rsrc) {
+ __builtin_amdgcn_raw_buffer_store_v4f16(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, /*aux=*/0);
+}
+
+// VERDE-LABEL: @test_amdgcn_raw_buffer_store_v4f32(
+// VERDE-NEXT: entry:
+// VERDE-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v4f32(<4 x
float> [[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// VERDE-NEXT: ret void
+//
+// GFX8-LABEL: @test_amdgcn_raw_buffer_store_v4f32(
+// GFX8-NEXT: entry:
+// GFX8-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v4f32(<4 x
float> [[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX8-NEXT: ret void
+//
+// GFX11-LABEL: @test_amdgcn_raw_buffer_store_v4f32(
+// GFX11-NEXT: entry:
+// GFX11-NEXT: tail call void @llvm.amdgcn.raw.buffer.store.v4f32(<4 x
float> [[VDATA:%.*]], <4 x i32> [[RSRC:%.*]], i32 0, i32 0, i32 0)
+// GFX11-NEXT: ret void
+//
+void test_amdgcn_raw_buffer_store_v4f32(v4f32 vdata, v4i32 rsrc) {
+ __builtin_amdgcn_raw_buffer_store_v4f32(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, /*aux=*/0);
+}
+
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-store-error.cl
b/clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-store-error.cl
new file mode 100644
index 0000000000000..33e20d6dc771a
--- /dev/null
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-raw-buffer-store-error.cl
@@ -0,0 +1,94 @@
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -S -verify
-o - %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -S -verify
-o - %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -S
-verify -o - %s
+// REQUIRES: amdgpu-registered-target
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+typedef short v2i16 __attribute__((ext_vector_type(2)));
+typedef int v2i32 __attribute__((ext_vector_type(2)));
+typedef half v2f16 __attribute__((ext_vector_type(2)));
+typedef float v2f32 __attribute__((ext_vector_type(2)));
+typedef short v4i16 __attribute__((ext_vector_type(4)));
+typedef int v4i32 __attribute__((ext_vector_type(4)));
+typedef half v4f16 __attribute__((ext_vector_type(4)));
+typedef float v4f32 __attribute__((ext_vector_type(4)));
+
+void test_amdgcn_raw_buffer_store_i8(char vdata, v4i32 rsrc, int offset, int
soffset, int aux) {
+ __builtin_amdgcn_raw_buffer_store_i8(vdata, rsrc, offset, /*soffset=*/0,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_i8' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_i8(vdata, rsrc, /*offset=*/0, soffset,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_i8' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_i8(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, aux); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_i8' must be a constant integer}}
+}
+
+void test_amdgcn_raw_buffer_store_i16(short vdata, v4i32 rsrc, int offset, int
soffset, int aux) {
+ __builtin_amdgcn_raw_buffer_store_i16(vdata, rsrc, offset, /*soffset=*/0,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_i16' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_i16(vdata, rsrc, /*offset=*/0, soffset,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_i16' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_i16(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, aux); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_i16' must be a constant integer}}
+}
+
+void test_amdgcn_raw_buffer_store_i32(int vdata, v4i32 rsrc, int offset, int
soffset, int aux) {
+ __builtin_amdgcn_raw_buffer_store_i32(vdata, rsrc, offset, /*soffset=*/0,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_i32' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_i32(vdata, rsrc, /*offset=*/0, soffset,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_i32' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_i32(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, aux); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_i32' must be a constant integer}}
+}
+
+void test_amdgcn_raw_buffer_store_f16(half vdata, v4i32 rsrc, int offset, int
soffset, int aux) {
+ __builtin_amdgcn_raw_buffer_store_f16(vdata, rsrc, offset, /*soffset=*/0,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_f16' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_f16(vdata, rsrc, /*offset=*/0, soffset,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_f16' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_f16(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, aux); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_f16' must be a constant integer}}
+}
+
+void test_amdgcn_raw_buffer_store_f32(float vdata, v4i32 rsrc, int offset, int
soffset, int aux) {
+ __builtin_amdgcn_raw_buffer_store_f32(vdata, rsrc, offset, /*soffset=*/0,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_f32' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_f32(vdata, rsrc, /*offset=*/0, soffset,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_f32' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_f32(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, aux); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_f32' must be a constant integer}}
+}
+
+void test_amdgcn_raw_buffer_store_v2i16(v2i16 vdata, v4i32 rsrc, int offset,
int soffset, int aux) {
+ __builtin_amdgcn_raw_buffer_store_v2i16(vdata, rsrc, offset, /*soffset=*/0,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v2i16' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_v2i16(vdata, rsrc, /*offset=*/0, soffset,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v2i16' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_v2i16(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, aux); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v2i16' must be a constant integer}}
+}
+
+void test_amdgcn_raw_buffer_store_v2i32(v2i32 vdata, v4i32 rsrc, int offset,
int soffset, int aux) {
+ __builtin_amdgcn_raw_buffer_store_v2i32(vdata, rsrc, offset, /*soffset=*/0,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v2i32' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_v2i32(vdata, rsrc, /*offset=*/0, soffset,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v2i32' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_v2i32(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, aux); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v2i32' must be a constant integer}}
+}
+
+void test_amdgcn_raw_buffer_store_v2f16(v2f16 vdata, v4i32 rsrc, int offset,
int soffset, int aux) {
+ __builtin_amdgcn_raw_buffer_store_v2f16(vdata, rsrc, offset, /*soffset=*/0,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v2f16' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_v2f16(vdata, rsrc, /*offset=*/0, soffset,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v2f16' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_v2f16(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, aux); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v2f16' must be a constant integer}}
+}
+
+void test_amdgcn_raw_buffer_store_v2f32(v2f32 vdata, v4i32 rsrc, int offset,
int soffset, int aux) {
+ __builtin_amdgcn_raw_buffer_store_v2f32(vdata, rsrc, offset, /*soffset=*/0,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v2f32' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_v2f32(vdata, rsrc, /*offset=*/0, soffset,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v2f32' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_v2f32(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, aux); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v2f32' must be a constant integer}}
+}
+
+void test_amdgcn_raw_buffer_store_v4i16(v4i16 vdata, v4i32 rsrc, int offset,
int soffset, int aux) {
+ __builtin_amdgcn_raw_buffer_store_v4i16(vdata, rsrc, offset, /*soffset=*/0,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v4i16' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_v4i16(vdata, rsrc, /*offset=*/0, soffset,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v4i16' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_v4i16(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, aux); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v4i16' must be a constant integer}}
+}
+
+void test_amdgcn_raw_buffer_store_v4i32(v4i32 vdata, v4i32 rsrc, int offset,
int soffset, int aux) {
+ __builtin_amdgcn_raw_buffer_store_v4i32(vdata, rsrc, offset, /*soffset=*/0,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v4i32' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_v4i32(vdata, rsrc, /*offset=*/0, soffset,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v4i32' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_v4i32(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, aux); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v4i32' must be a constant integer}}
+}
+
+void test_amdgcn_raw_buffer_store_v4f16(v4f16 vdata, v4i32 rsrc, int offset,
int soffset, int aux) {
+ __builtin_amdgcn_raw_buffer_store_v4f16(vdata, rsrc, offset, /*soffset=*/0,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v4f16' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_v4f16(vdata, rsrc, /*offset=*/0, soffset,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v4f16' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_v4f16(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, aux); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v4f16' must be a constant integer}}
+}
+
+void test_amdgcn_raw_buffer_store_v4f32(v4f32 vdata, v4i32 rsrc, int offset,
int soffset, int aux) {
+ __builtin_amdgcn_raw_buffer_store_v4f32(vdata, rsrc, offset, /*soffset=*/0,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v4f32' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_v4f32(vdata, rsrc, /*offset=*/0, soffset,
/*aux=*/0); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v4f32' must be a constant integer}}
+ __builtin_amdgcn_raw_buffer_store_v4f32(vdata, rsrc, /*offset=*/0,
/*soffset=*/0, aux); //expected-error{{argument to
'__builtin_amdgcn_raw_buffer_store_v4f32' must be a constant integer}}
+}
+
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
