xref: /llvm-project/clang/test/CodeGenOpenCL/builtins-amdgcn-gws-insts.cl (revision c5de4dd1eab00df76c1a68c5f397304ceacb71f2) 
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 2
2// REQUIRES: amdgpu-registered-target
3
4// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx803 -emit-llvm -o - %s | FileCheck %s
5// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -emit-llvm -o - %s | FileCheck %s
6// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx90a -emit-llvm -o - %s | FileCheck %s
7// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx90c -emit-llvm -o - %s | FileCheck %s
8// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx940 -emit-llvm -o - %s | FileCheck %s
9// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -emit-llvm -o - %s | FileCheck %s
10// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1030 -emit-llvm -o - %s | FileCheck %s
11// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -emit-llvm -o - %s | FileCheck %s
12
13typedef unsigned int uint;
14
15// CHECK-LABEL: define dso_local amdgpu_kernel void @test_builtins_amdgcn_gws_insts
16// CHECK-SAME: (i32 noundef [[A:%.*]], i32 noundef [[B:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 {
17// CHECK-NEXT:  entry:
18// CHECK-NEXT:    tail call void @llvm.amdgcn.ds.gws.init(i32 [[A]], i32 [[B]])
19// CHECK-NEXT:    tail call void @llvm.amdgcn.ds.gws.barrier(i32 [[A]], i32 [[B]])
20// CHECK-NEXT:    tail call void @llvm.amdgcn.ds.gws.sema.v(i32 [[A]])
21// CHECK-NEXT:    tail call void @llvm.amdgcn.ds.gws.sema.br(i32 [[A]], i32 [[B]])
22// CHECK-NEXT:    tail call void @llvm.amdgcn.ds.gws.sema.p(i32 [[A]])
23// CHECK-NEXT:    ret void
24//
25kernel void test_builtins_amdgcn_gws_insts(uint a, uint b) {
26  __builtin_amdgcn_ds_gws_init(a, b);
27  __builtin_amdgcn_ds_gws_barrier(a, b);
28  __builtin_amdgcn_ds_gws_sema_v(a);
29  __builtin_amdgcn_ds_gws_sema_br(a, b);
30  __builtin_amdgcn_ds_gws_sema_p(a);
31}
32