1 // REQUIRES: x86-registered-target
2 // REQUIRES: nvptx-registered-target
4 // RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm \
5 // RUN: -disable-llvm-passes -o - %s | FileCheck -allow-deprecated-dag-overlap -check-prefix DEVICE %s
7 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \
8 // RUN: -disable-llvm-passes -o - %s | \
9 // RUN: FileCheck -allow-deprecated-dag-overlap -check-prefix HOST %s
11 #include "Inputs/cuda.h"
13 // DEVICE: Function Attrs:
14 // DEVICE-SAME: convergent
15 // DEVICE-NEXT: define{{.*}} void @_Z3foov
16 __device__ void foo() {}
18 // HOST: Function Attrs:
19 // HOST-NOT: convergent
20 // HOST-NEXT: define{{.*}} void @_Z3barv
21 // DEVICE: Function Attrs:
22 // DEVICE-SAME: convergent
23 // DEVICE-NEXT: define{{.*}} void @_Z3barv
24 __host__ __device__ void baz();
25 __host__ __device__ void bar() {
26 // DEVICE: call void @_Z3bazv() [[CALL_ATTR:#[0-9]+]]
28 // DEVICE: call i32 asm "trap;", "=l"() [[ASM_ATTR:#[0-9]+]]
30 asm ("trap;" : "=l"(x));
31 // DEVICE: call void asm sideeffect "trap;", ""() [[ASM_ATTR:#[0-9]+]]
32 asm volatile ("trap;");
35 // DEVICE: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]]
36 // DEVICE: attributes [[BAZ_ATTR]] = {
37 // DEVICE-SAME: convergent
39 // DEVICE-DAG: attributes [[CALL_ATTR]] = { convergent
40 // DEVICE-DAG: attributes [[ASM_ATTR]] = { convergent
42 // HOST: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]]
43 // HOST: attributes [[BAZ_ATTR]] = {
44 // HOST-NOT: convergent