32 lines · plain
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 22// REQUIRES: amdgpu-registered-target3 4// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx803 -emit-llvm -o - %s | FileCheck %s5// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -emit-llvm -o - %s | FileCheck %s6// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx90a -emit-llvm -o - %s | FileCheck %s7// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx90c -emit-llvm -o - %s | FileCheck %s8// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx942 -emit-llvm -o - %s | FileCheck %s9// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -emit-llvm -o - %s | FileCheck %s10// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1030 -emit-llvm -o - %s | FileCheck %s11// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -emit-llvm -o - %s | FileCheck %s12 13typedef unsigned int uint;14 15// CHECK-LABEL: define dso_local amdgpu_kernel void @test_builtins_amdgcn_gws_insts16// CHECK-SAME: (i32 noundef [[A:%.*]], i32 noundef [[B:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !kernel_arg_addr_space !8 !kernel_arg_access_qual !9 !kernel_arg_type !10 !kernel_arg_base_type !10 !kernel_arg_type_qual !11 {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 void24//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