119 lines · plain
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 %s10// 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 %s15 16#include "Inputs/cuda.h"17 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.21 22// COMM-LABEL: define {{.*}}@_Z20template_no_overloadIiET_S0_(23// COMM: ret i32 124template<typename T>25T template_no_overload(T x) {26 return 1;27}28 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.33 34// COMM-LABEL: define {{.*}}@_Z22template_with_overloadIiET_S0_(35// HOST: ret i32 236// DEV: ret i32 337template<typename T>38T template_with_overload(T x) {39 return 2;40}41 42template<typename T>43__device__ T template_with_overload(T x) {44 return 3;45}46 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_(51// HOST: ret i32 1052template<typename T>53T template_used_by_host(T x) {54 return 10;55}56 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_(61// HOST: ret i32 1162template<typename T>63T template_indirectly_used_by_host(T x) {64 return 11;65}66 67template<typename T>68T template_in_middle_by_host(T x) {69 template_indirectly_used_by_host(x);70 return 12;71}72 73// Implicit host device template indirectly used by device function only.74// Emitted on device.75// DEVICE-LABEL: define {{.*}}@_Z34template_indirectly_used_by_deviceIiET_S0_(76// DEVICE: ret i32 2177template<typename T>78T template_indirectly_used_by_device(T x) {79 return 21;80}81 82template<typename T>83T template_in_middle_by_device(T x) {84 template_indirectly_used_by_device(x);85 return 22;86}87 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_(91// COMMON: ret i32 3192template<typename T>93T template_indirectly_used_by_host_device(T x) {94 return 31;95}96 97template<typename T>98T template_in_middle_by_host_device(T x) {99 template_indirectly_used_by_host_device(x);100 return 32;101}102 103void host_fun() {104 template_no_overload(0);105 template_with_overload(0);106 template_used_by_host(0);107 template_in_middle_by_host(0);108}109 110__device__ void device_fun() {111 template_no_overload(0);112 template_with_overload(0);113 template_in_middle_by_device(0);114}115 116__host__ __device__ void host_device_fun() {117 template_in_middle_by_host_device(0);118}119