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