| 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 -allow-deprecated-dag-overlap -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 -allow-deprecated-dag-overlap -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-DAG: attributes [[CALL_ATTR]] = { convergent |
|---|
| 40 | // DEVICE-DAG: 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 | |
|---|