1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --check-attributes --check-globals
2 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fno-ident -fcuda-is-device \
3 // RUN: -emit-llvm -o - %s | FileCheck -check-prefix=OPTNONE %s
5 // RUN: %clang_cc1 -O3 -triple amdgcn-amd-amdhsa -x hip -fno-ident -fcuda-is-device \
6 // RUN: -emit-llvm -o - %s | FileCheck -check-prefix=OPT %s
8 #define __device__ __attribute__((device))
9 #define __global__ __attribute__((global))
11 // OPTNONE: Function Attrs: convergent mustprogress noinline nounwind optnone
12 // OPTNONE-LABEL: define {{[^@]+}}@_Z4funcv
13 // OPTNONE-SAME: () #[[ATTR0:[0-9]+]] {
14 // OPTNONE-NEXT: entry:
15 // OPTNONE-NEXT: ret void
17 // OPT: Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none)
18 // OPT-LABEL: define {{[^@]+}}@_Z4funcv
19 // OPT-SAME: () local_unnamed_addr #[[ATTR0:[0-9]+]] {
23 __device__ void func() {
27 // OPTNONE: Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
28 // OPTNONE-LABEL: define {{[^@]+}}@_Z6kernelv
29 // OPTNONE-SAME: () #[[ATTR1:[0-9]+]] {
30 // OPTNONE-NEXT: entry:
31 // OPTNONE-NEXT: ret void
33 // OPT: Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none)
34 // OPT-LABEL: define {{[^@]+}}@_Z6kernelv
35 // OPT-SAME: () local_unnamed_addr #[[ATTR1:[0-9]+]] {
39 __global__ void kernel() {
43 // OPTNONE: attributes #0 = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
44 // OPTNONE: attributes #1 = { convergent mustprogress noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,1024" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }
46 // OPT: attributes #0 = { mustprogress nofree norecurse nosync nounwind willreturn memory(none) "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
47 // OPT: attributes #1 = { mustprogress nofree norecurse nosync nounwind willreturn memory(none) "amdgpu-flat-work-group-size"="1,1024" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }
49 // OPTNONE: !0 = !{i32 1, !"amdgpu_code_object_version", i32 400}
50 // OPTNONE: !1 = !{i32 1, !"wchar_size", i32 4}
52 // OPT: !0 = !{i32 1, !"amdgpu_code_object_version", i32 400}
53 // OPT: !1 = !{i32 1, !"wchar_size", i32 4}