1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --version 5
2 // REQUIRES: x86-registered-target
3 // REQUIRES: nvptx-registered-target
5 // RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm \
6 // RUN: -disable-llvm-passes -o - %s | FileCheck -allow-deprecated-dag-overlap -check-prefix DEVICE %s
8 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \
9 // RUN: -disable-llvm-passes -o - %s | \
10 // RUN: FileCheck -allow-deprecated-dag-overlap -check-prefix HOST %s
12 #include "Inputs/cuda.h"
14 // DEVICE-LABEL: define dso_local void @_Z3foov(
15 // DEVICE-SAME: ) #[[ATTR0:[0-9]+]] {
16 // DEVICE-NEXT: [[ENTRY:.*:]]
17 // DEVICE-NEXT: ret void
19 __device__ void foo() {}
20 // DEVICE-LABEL: define dso_local void @_Z3baxv(
21 // DEVICE-SAME: ) #[[ATTR1:[0-9]+]] {
22 // DEVICE-NEXT: [[ENTRY:.*:]]
23 // DEVICE-NEXT: ret void
25 [[clang::noconvergent]] __device__ void bax() {}
27 __host__ __device__ void baz();
29 __host__ __device__ float aliasf0(int) asm("something");
30 __host__ __device__ [[clang::noconvergent]] float aliasf1(int) asm("somethingelse");
32 // DEVICE-LABEL: define dso_local void @_Z3barv(
33 // DEVICE-SAME: ) #[[ATTR0]] {
34 // DEVICE-NEXT: [[ENTRY:.*:]]
35 // DEVICE-NEXT: [[X:%.*]] = alloca i32, align 4
36 // DEVICE-NEXT: call void @_Z3bazv() #[[ATTR4:[0-9]+]]
37 // DEVICE-NEXT: [[TMP0:%.*]] = call i32 asm "trap", "=l"() #[[ATTR5:[0-9]+]], !srcloc [[META3:![0-9]+]]
38 // DEVICE-NEXT: store i32 [[TMP0]], ptr [[X]], align 4
39 // DEVICE-NEXT: call void asm sideeffect "trap", ""() #[[ATTR4]], !srcloc [[META4:![0-9]+]]
40 // DEVICE-NEXT: call void asm sideeffect "nop", ""() #[[ATTR6:[0-9]+]], !srcloc [[META5:![0-9]+]]
41 // DEVICE-NEXT: [[TMP1:%.*]] = load i32, ptr [[X]], align 4
42 // DEVICE-NEXT: [[CALL:%.*]] = call contract noundef float @something(i32 noundef [[TMP1]]) #[[ATTR4]]
43 // DEVICE-NEXT: [[TMP2:%.*]] = load i32, ptr [[X]], align 4
44 // DEVICE-NEXT: [[CALL1:%.*]] = call contract noundef float @somethingelse(i32 noundef [[TMP2]]) #[[ATTR6]]
45 // DEVICE-NEXT: ret void
47 // HOST-LABEL: define dso_local void @_Z3barv(
48 // HOST-SAME: ) #[[ATTR0:[0-9]+]] {
49 // HOST-NEXT: [[ENTRY:.*:]]
50 // HOST-NEXT: [[X:%.*]] = alloca i32, align 4
51 // HOST-NEXT: call void @_Z3bazv()
52 // HOST-NEXT: [[TMP0:%.*]] = call i32 asm "trap", "=l,~{dirflag},~{fpsr},~{flags}"() #[[ATTR2:[0-9]+]], !srcloc [[META2:![0-9]+]]
53 // HOST-NEXT: store i32 [[TMP0]], ptr [[X]], align 4
54 // HOST-NEXT: call void asm sideeffect "trap", "~{dirflag},~{fpsr},~{flags}"() #[[ATTR3:[0-9]+]], !srcloc [[META3:![0-9]+]]
55 // HOST-NEXT: call void asm sideeffect "nop", "~{dirflag},~{fpsr},~{flags}"() #[[ATTR3]], !srcloc [[META4:![0-9]+]]
56 // HOST-NEXT: [[TMP1:%.*]] = load i32, ptr [[X]], align 4
57 // HOST-NEXT: [[CALL:%.*]] = call contract noundef float @something(i32 noundef [[TMP1]])
58 // HOST-NEXT: [[TMP2:%.*]] = load i32, ptr [[X]], align 4
59 // HOST-NEXT: [[CALL1:%.*]] = call contract noundef float @somethingelse(i32 noundef [[TMP2]])
60 // HOST-NEXT: ret void
62 __host__ __device__ void bar() {
65 asm ("trap" : "=l"(x));
66 asm volatile ("trap");
67 [[clang::noconvergent]] { asm volatile ("nop"); }
74 // DEVICE: attributes #[[ATTR0]] = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32" }
75 // DEVICE: attributes #[[ATTR1]] = { mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32" }
76 // DEVICE: attributes #[[ATTR2:[0-9]+]] = { convergent nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32" }
77 // DEVICE: attributes #[[ATTR3:[0-9]+]] = { nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32" }
78 // DEVICE: attributes #[[ATTR4]] = { convergent nounwind }
79 // DEVICE: attributes #[[ATTR5]] = { convergent nounwind memory(none) }
80 // DEVICE: attributes #[[ATTR6]] = { nounwind }
82 // HOST: attributes #[[ATTR0]] = { mustprogress noinline nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" }
83 // HOST: attributes #[[ATTR1:[0-9]+]] = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" }
84 // HOST: attributes #[[ATTR2]] = { nounwind memory(none) }
85 // HOST: attributes #[[ATTR3]] = { nounwind }
87 // DEVICE: [[META0:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
88 // DEVICE: [[META1:![0-9]+]] = !{i32 4, !"nvvm-reflect-ftz", i32 0}
89 // DEVICE: [[META2:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
90 // DEVICE: [[META3]] = !{i64 3120}
91 // DEVICE: [[META4]] = !{i64 3155}
92 // DEVICE: [[META5]] = !{i64 3206}
94 // HOST: [[META0:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
95 // HOST: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
96 // HOST: [[META2]] = !{i64 3120}
97 // HOST: [[META3]] = !{i64 3155}
98 // HOST: [[META4]] = !{i64 3206}