| // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --version 5 |
| // REQUIRES: x86-registered-target |
| // REQUIRES: nvptx-registered-target |
| |
| // RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm \ |
| // RUN: -disable-llvm-passes -o - %s | FileCheck -allow-deprecated-dag-overlap -check-prefix DEVICE %s |
| |
| // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \ |
| // RUN: -disable-llvm-passes -o - %s | \ |
| // RUN: FileCheck -allow-deprecated-dag-overlap -check-prefix HOST %s |
| |
| #include "Inputs/cuda.h" |
| |
| // DEVICE-LABEL: define dso_local void @_Z3foov( |
| // DEVICE-SAME: ) #[[ATTR0:[0-9]+]] { |
| // DEVICE-NEXT: [[ENTRY:.*:]] |
| // DEVICE-NEXT: ret void |
| // |
| __device__ void foo() {} |
| // DEVICE-LABEL: define dso_local void @_Z3baxv( |
| // DEVICE-SAME: ) #[[ATTR1:[0-9]+]] { |
| // DEVICE-NEXT: [[ENTRY:.*:]] |
| // DEVICE-NEXT: ret void |
| // |
| [[clang::noconvergent]] __device__ void bax() {} |
| |
| __host__ __device__ void baz(); |
| |
| __host__ __device__ float aliasf0(int) asm("something"); |
| __host__ __device__ [[clang::noconvergent]] float aliasf1(int) asm("somethingelse"); |
| |
| // DEVICE-LABEL: define dso_local void @_Z3barv( |
| // DEVICE-SAME: ) #[[ATTR0]] { |
| // DEVICE-NEXT: [[ENTRY:.*:]] |
| // DEVICE-NEXT: [[X:%.*]] = alloca i32, align 4 |
| // DEVICE-NEXT: call void @_Z3bazv() #[[ATTR4:[0-9]+]] |
| // DEVICE-NEXT: [[TMP0:%.*]] = call i32 asm "trap", "=l"() #[[ATTR5:[0-9]+]], !srcloc [[META3:![0-9]+]] |
| // DEVICE-NEXT: store i32 [[TMP0]], ptr [[X]], align 4 |
| // DEVICE-NEXT: call void asm sideeffect "trap", ""() #[[ATTR4]], !srcloc [[META4:![0-9]+]] |
| // DEVICE-NEXT: call void asm sideeffect "nop", ""() #[[ATTR6:[0-9]+]], !srcloc [[META5:![0-9]+]] |
| // DEVICE-NEXT: [[TMP1:%.*]] = load i32, ptr [[X]], align 4 |
| // DEVICE-NEXT: [[CALL:%.*]] = call contract noundef float @something(i32 noundef [[TMP1]]) #[[ATTR4]] |
| // DEVICE-NEXT: [[TMP2:%.*]] = load i32, ptr [[X]], align 4 |
| // DEVICE-NEXT: [[CALL1:%.*]] = call contract noundef float @somethingelse(i32 noundef [[TMP2]]) #[[ATTR6]] |
| // DEVICE-NEXT: ret void |
| // |
| // HOST-LABEL: define dso_local void @_Z3barv( |
| // HOST-SAME: ) #[[ATTR0:[0-9]+]] { |
| // HOST-NEXT: [[ENTRY:.*:]] |
| // HOST-NEXT: [[X:%.*]] = alloca i32, align 4 |
| // HOST-NEXT: call void @_Z3bazv() |
| // HOST-NEXT: [[TMP0:%.*]] = call i32 asm "trap", "=l,~{dirflag},~{fpsr},~{flags}"() #[[ATTR2:[0-9]+]], !srcloc [[META2:![0-9]+]] |
| // HOST-NEXT: store i32 [[TMP0]], ptr [[X]], align 4 |
| // HOST-NEXT: call void asm sideeffect "trap", "~{dirflag},~{fpsr},~{flags}"() #[[ATTR3:[0-9]+]], !srcloc [[META3:![0-9]+]] |
| // HOST-NEXT: call void asm sideeffect "nop", "~{dirflag},~{fpsr},~{flags}"() #[[ATTR3]], !srcloc [[META4:![0-9]+]] |
| // HOST-NEXT: [[TMP1:%.*]] = load i32, ptr [[X]], align 4 |
| // HOST-NEXT: [[CALL:%.*]] = call contract noundef float @something(i32 noundef [[TMP1]]) |
| // HOST-NEXT: [[TMP2:%.*]] = load i32, ptr [[X]], align 4 |
| // HOST-NEXT: [[CALL1:%.*]] = call contract noundef float @somethingelse(i32 noundef [[TMP2]]) |
| // HOST-NEXT: ret void |
| // |
| __host__ __device__ void bar() { |
| baz(); |
| int x; |
| asm ("trap" : "=l"(x)); |
| asm volatile ("trap"); |
| [[clang::noconvergent]] { asm volatile ("nop"); } |
| aliasf0(x); |
| aliasf1(x); |
| } |
| |
| |
| //. |
| // DEVICE: attributes #[[ATTR0]] = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32" } |
| // DEVICE: attributes #[[ATTR1]] = { mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32" } |
| // DEVICE: attributes #[[ATTR2:[0-9]+]] = { convergent nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32" } |
| // DEVICE: attributes #[[ATTR3:[0-9]+]] = { nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+ptx32" } |
| // DEVICE: attributes #[[ATTR4]] = { convergent nounwind } |
| // DEVICE: attributes #[[ATTR5]] = { convergent nounwind memory(none) } |
| // DEVICE: attributes #[[ATTR6]] = { nounwind } |
| //. |
| // 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" } |
| // HOST: attributes #[[ATTR1:[0-9]+]] = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" } |
| // HOST: attributes #[[ATTR2]] = { nounwind memory(none) } |
| // HOST: attributes #[[ATTR3]] = { nounwind } |
| //. |
| // DEVICE: [[META0:![0-9]+]] = !{i32 1, !"wchar_size", i32 4} |
| // DEVICE: [[META1:![0-9]+]] = !{i32 4, !"nvvm-reflect-ftz", i32 0} |
| // DEVICE: [[META2:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} |
| // DEVICE: [[META3]] = !{i64 3120} |
| // DEVICE: [[META4]] = !{i64 3155} |
| // DEVICE: [[META5]] = !{i64 3206} |
| //. |
| // HOST: [[META0:![0-9]+]] = !{i32 1, !"wchar_size", i32 4} |
| // HOST: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} |
| // HOST: [[META2]] = !{i64 3120} |
| // HOST: [[META3]] = !{i64 3155} |
| // HOST: [[META4]] = !{i64 3206} |
| //. |
