1 // RUN: %clang_cc1 -triple x86_64-pc-linux-gnu \
2 // RUN: -foffload-implicit-host-device-templates \
3 // RUN: -emit-llvm -o - -x hip %s 2>&1 | \
4 // RUN: FileCheck -check-prefixes=COMM,HOST %s
5 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
6 // RUN: -target-cpu gfx1100 \
7 // RUN: -foffload-implicit-host-device-templates \
8 // RUN: -emit-llvm -o - -x hip %s 2>&1 | \
9 // RUN: FileCheck -check-prefixes=COMM,DEV %s
10 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
11 // RUN: -target-cpu gfx1100 \
12 // RUN: -foffload-implicit-host-device-templates \
13 // RUN: -emit-llvm -o - -x hip %s 2>&1 | \
14 // RUN: FileCheck -check-prefixes=DEV-NEG %s
16 #include "Inputs/cuda.h"
18 // Implicit host device template not overloaded by device template.
19 // Used by both device and host function.
20 // Emitted on both host and device.
22 // COMM-LABEL: define {{.*}}@_Z20template_no_overloadIiET_S0_(
25 T template_no_overload(T x) {
29 // Implicit host device template overloaded by device template.
30 // Used by both device and host function.
31 // Implicit host device template emitted on host.
32 // Device template emitted on device.
34 // COMM-LABEL: define {{.*}}@_Z22template_with_overloadIiET_S0_(
38 T template_with_overload(T x) {
43 __device__ T template_with_overload(T x) {
47 // Implicit host device template used by host function only.
48 // Emitted on host only.
49 // HOST-LABEL: define {{.*}}@_Z21template_used_by_hostIiET_S0_(
50 // DEV-NEG-NOT: define {{.*}}@_Z21template_used_by_hostIiET_S0_(
53 T template_used_by_host(T x) {
57 // Implicit host device template indirectly used by host function only.
58 // Emitted on host only.
59 // HOST-LABEL: define {{.*}}@_Z32template_indirectly_used_by_hostIiET_S0_(
60 // DEV-NEG-NOT: define {{.*}}@_Z32template_indirectly_used_by_hostIiET_S0_(
63 T template_indirectly_used_by_host(T x) {
68 T template_in_middle_by_host(T x) {
69 template_indirectly_used_by_host(x);
73 // Implicit host device template indirectly used by device function only.
75 // DEVICE-LABEL: define {{.*}}@_Z34template_indirectly_used_by_deviceIiET_S0_(
78 T template_indirectly_used_by_device(T x) {
83 T template_in_middle_by_device(T x) {
84 template_indirectly_used_by_device(x);
88 // Implicit host device template indirectly used by host device function only.
89 // Emitted on host and device.
90 // COMMON-LABEL: define {{.*}}@_Z39template_indirectly_used_by_host_deviceIiET_S0_(
93 T template_indirectly_used_by_host_device(T x) {
98 T template_in_middle_by_host_device(T x) {
99 template_indirectly_used_by_host_device(x);
104 template_no_overload(0);
105 template_with_overload(0);
106 template_used_by_host(0);
107 template_in_middle_by_host(0);
110 __device__ void device_fun() {
111 template_no_overload(0);
112 template_with_overload(0);
113 template_in_middle_by_device(0);
116 __host__ __device__ void host_device_fun() {
117 template_in_middle_by_host_device(0);