BUILTIN(__builtin_amdgcn_buffer_wbinvl1, "v", "n")
BUILTIN(__builtin_amdgcn_ds_gws_init, "vUiUi", "n")
BUILTIN(__builtin_amdgcn_ds_gws_barrier, "vUiUi", "n")
+BUILTIN(__builtin_amdgcn_ds_gws_sema_v, "vUi", "n")
+BUILTIN(__builtin_amdgcn_ds_gws_sema_br, "vUiUi", "n")
+BUILTIN(__builtin_amdgcn_ds_gws_sema_p, "vUi", "n")
// FIXME: Need to disallow constant address space.
BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n")
//===----------------------------------------------------------------------===//
TARGET_BUILTIN(__builtin_amdgcn_s_dcache_inv_vol, "v", "n", "ci-insts")
TARGET_BUILTIN(__builtin_amdgcn_buffer_wbinvl1_vol, "v", "n", "ci-insts")
+TARGET_BUILTIN(__builtin_amdgcn_ds_gws_sema_release_all, "vUi", "n", "ci-insts")
//===----------------------------------------------------------------------===//
// Interpolation builtins.
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu fiji -S -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck %s
+typedef unsigned int uint;
+
// CHECK-LABEL: @test_s_dcache_inv_vol
// CHECK: call void @llvm.amdgcn.s.dcache.inv.vol(
void test_s_dcache_inv_vol()
__builtin_amdgcn_buffer_wbinvl1_vol();
}
+// CHECK-LABEL: @test_gws_sema_release_all(
+// CHECK: call void @llvm.amdgcn.ds.gws.sema.release.all(i32 %id)
+void test_gws_sema_release_all(uint id)
+{
+ __builtin_amdgcn_ds_gws_sema_release_all(id);
+}
__builtin_amdgcn_ds_gws_barrier(value, id);
}
+// CHECK-LABEL: @test_gws_sema_v(
+// CHECK: call void @llvm.amdgcn.ds.gws.sema.v(i32 %id)
+kernel void test_gws_sema_v(uint id) {
+ __builtin_amdgcn_ds_gws_sema_v(id);
+}
+
+// CHECK-LABEL: @test_gws_sema_br(
+// CHECK: call void @llvm.amdgcn.ds.gws.sema.br(i32 %value, i32 %id)
+kernel void test_gws_sema_br(uint value, uint id) {
+ __builtin_amdgcn_ds_gws_sema_br(value, id);
+}
+
+// CHECK-LABEL: @test_gws_sema_p(
+// CHECK: call void @llvm.amdgcn.ds.gws.sema.p(i32 %id)
+kernel void test_gws_sema_p(uint id) {
+ __builtin_amdgcn_ds_gws_sema_p(id);
+}
+
// CHECK-DAG: [[$WI_RANGE]] = !{i32 0, i32 1024}
// CHECK-DAG: attributes #[[$NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly }
// CHECK-DAG: attributes #[[$READ_EXEC_ATTRS]] = { convergent }
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
-void test_ci_biltins()
+void test_ci_builtins()
{
__builtin_amdgcn_s_dcache_inv_vol(); // expected-error {{'__builtin_amdgcn_s_dcache_inv_vol' needs target feature ci-insts}}
__builtin_amdgcn_buffer_wbinvl1_vol(); // expected-error {{'__builtin_amdgcn_buffer_wbinvl1_vol' needs target feature ci-insts}}
+ __builtin_amdgcn_ds_gws_sema_release_all(0); // expected-error {{'__builtin_amdgcn_ds_gws_sema_release_all' needs target feature ci-insts}}
}