1 // REQUIRES: x86-registered-target 2 // REQUIRES: nvptx-registered-target 3 4 // RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm \ 5 // RUN: -disable-llvm-passes -o - %s | FileCheck -check-prefix DEVICE %s 6 7 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \ 8 // RUN: -disable-llvm-passes -o - %s | \ 9 // RUN: FileCheck -check-prefix HOST %s 10 11 #include "Inputs/cuda.h" 12 13 // DEVICE: Function Attrs: 14 // DEVICE-SAME: convergent 15 // DEVICE-NEXT: define void @_Z3foov 16 __device__ void foo() {} 17 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]+]] 27 baz(); 28 // DEVICE: call i32 asm "trap;", "=l"() [[ASM_ATTR:#[0-9]+]] 29 int x; 30 asm ("trap;" : "=l"(x)); 31 // DEVICE: call void asm sideeffect "trap;", ""() [[ASM_ATTR:#[0-9]+]] 32 asm volatile ("trap;"); 33 } 34 35 // DEVICE: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]] 36 // DEVICE: attributes [[BAZ_ATTR]] = { 37 // DEVICE-SAME: convergent 38 // DEVICE-SAME: } 39 // DEVICE: attributes [[CALL_ATTR]] = { convergent } 40 // DEVICE: attributes [[ASM_ATTR]] = { convergent 41 42 // HOST: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]] 43 // HOST: attributes [[BAZ_ATTR]] = { 44 // HOST-NOT: convergent 45 // NOST-SAME: } 46
