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 15 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 1 24 template<typename T> template_no_overload(T x)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 2 36 // DEV: ret i32 3 37 template<typename T> template_with_overload(T x)38T template_with_overload(T x) { 39 return 2; 40 } 41 42 template<typename T> template_with_overload(T x)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 10 52 template<typename T> template_used_by_host(T x)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 11 62 template<typename T> template_indirectly_used_by_host(T x)63T template_indirectly_used_by_host(T x) { 64 return 11; 65 } 66 67 template<typename T> template_in_middle_by_host(T x)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 21 77 template<typename T> template_indirectly_used_by_device(T x)78T template_indirectly_used_by_device(T x) { 79 return 21; 80 } 81 82 template<typename T> template_in_middle_by_device(T x)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 31 92 template<typename T> template_indirectly_used_by_host_device(T x)93T template_indirectly_used_by_host_device(T x) { 94 return 31; 95 } 96 97 template<typename T> template_in_middle_by_host_device(T x)98T template_in_middle_by_host_device(T x) { 99 template_indirectly_used_by_host_device(x); 100 return 32; 101 } 102 host_fun()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 device_fun()110__device__ void device_fun() { 111 template_no_overload(0); 112 template_with_overload(0); 113 template_in_middle_by_device(0); 114 } 115 host_device_fun()116__host__ __device__ void host_device_fun() { 117 template_in_middle_by_host_device(0); 118 } 119