mirror of
https://github.com/llvm/llvm-project.git
synced 2025-04-28 00:46:06 +00:00
9774d0ce5f
Added option -foffload-implicit-host-device-templates which is off by default. When the option is on, template functions and specializations without host/device attributes have implicit host device attributes. They can be overridden by device template functions with the same signagure. They are emitted on device side only if they are used on device side. This feature is added as an extension. `__has_extension(cuda_implicit_host_device_templates)` can be used to check whether it is enabled. This is to facilitate using standard C++ headers for device. Fixes: https://github.com/llvm/llvm-project/issues/69956 Fixes: SWDEV-428314
119 lines
3.3 KiB
Plaintext
119 lines
3.3 KiB
Plaintext
// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu \
|
|
// RUN: -foffload-implicit-host-device-templates \
|
|
// RUN: -emit-llvm -o - -x hip %s 2>&1 | \
|
|
// RUN: FileCheck -check-prefixes=COMM,HOST %s
|
|
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
|
|
// RUN: -target-cpu gfx1100 \
|
|
// RUN: -foffload-implicit-host-device-templates \
|
|
// RUN: -emit-llvm -o - -x hip %s 2>&1 | \
|
|
// RUN: FileCheck -check-prefixes=COMM,DEV %s
|
|
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
|
|
// RUN: -target-cpu gfx1100 \
|
|
// RUN: -foffload-implicit-host-device-templates \
|
|
// RUN: -emit-llvm -o - -x hip %s 2>&1 | \
|
|
// RUN: FileCheck -check-prefixes=DEV-NEG %s
|
|
|
|
#include "Inputs/cuda.h"
|
|
|
|
// Implicit host device template not overloaded by device template.
|
|
// Used by both device and host function.
|
|
// Emitted on both host and device.
|
|
|
|
// COMM-LABEL: define {{.*}}@_Z20template_no_overloadIiET_S0_(
|
|
// COMM: ret i32 1
|
|
template<typename T>
|
|
T template_no_overload(T x) {
|
|
return 1;
|
|
}
|
|
|
|
// Implicit host device template overloaded by device template.
|
|
// Used by both device and host function.
|
|
// Implicit host device template emitted on host.
|
|
// Device template emitted on device.
|
|
|
|
// COMM-LABEL: define {{.*}}@_Z22template_with_overloadIiET_S0_(
|
|
// HOST: ret i32 2
|
|
// DEV: ret i32 3
|
|
template<typename T>
|
|
T template_with_overload(T x) {
|
|
return 2;
|
|
}
|
|
|
|
template<typename T>
|
|
__device__ T template_with_overload(T x) {
|
|
return 3;
|
|
}
|
|
|
|
// Implicit host device template used by host function only.
|
|
// Emitted on host only.
|
|
// HOST-LABEL: define {{.*}}@_Z21template_used_by_hostIiET_S0_(
|
|
// DEV-NEG-NOT: define {{.*}}@_Z21template_used_by_hostIiET_S0_(
|
|
// HOST: ret i32 10
|
|
template<typename T>
|
|
T template_used_by_host(T x) {
|
|
return 10;
|
|
}
|
|
|
|
// Implicit host device template indirectly used by host function only.
|
|
// Emitted on host only.
|
|
// HOST-LABEL: define {{.*}}@_Z32template_indirectly_used_by_hostIiET_S0_(
|
|
// DEV-NEG-NOT: define {{.*}}@_Z32template_indirectly_used_by_hostIiET_S0_(
|
|
// HOST: ret i32 11
|
|
template<typename T>
|
|
T template_indirectly_used_by_host(T x) {
|
|
return 11;
|
|
}
|
|
|
|
template<typename T>
|
|
T template_in_middle_by_host(T x) {
|
|
template_indirectly_used_by_host(x);
|
|
return 12;
|
|
}
|
|
|
|
// Implicit host device template indirectly used by device function only.
|
|
// Emitted on device.
|
|
// DEVICE-LABEL: define {{.*}}@_Z34template_indirectly_used_by_deviceIiET_S0_(
|
|
// DEVICE: ret i32 21
|
|
template<typename T>
|
|
T template_indirectly_used_by_device(T x) {
|
|
return 21;
|
|
}
|
|
|
|
template<typename T>
|
|
T template_in_middle_by_device(T x) {
|
|
template_indirectly_used_by_device(x);
|
|
return 22;
|
|
}
|
|
|
|
// Implicit host device template indirectly used by host device function only.
|
|
// Emitted on host and device.
|
|
// COMMON-LABEL: define {{.*}}@_Z39template_indirectly_used_by_host_deviceIiET_S0_(
|
|
// COMMON: ret i32 31
|
|
template<typename T>
|
|
T template_indirectly_used_by_host_device(T x) {
|
|
return 31;
|
|
}
|
|
|
|
template<typename T>
|
|
T template_in_middle_by_host_device(T x) {
|
|
template_indirectly_used_by_host_device(x);
|
|
return 32;
|
|
}
|
|
|
|
void host_fun() {
|
|
template_no_overload(0);
|
|
template_with_overload(0);
|
|
template_used_by_host(0);
|
|
template_in_middle_by_host(0);
|
|
}
|
|
|
|
__device__ void device_fun() {
|
|
template_no_overload(0);
|
|
template_with_overload(0);
|
|
template_in_middle_by_device(0);
|
|
}
|
|
|
|
__host__ __device__ void host_device_fun() {
|
|
template_in_middle_by_host_device(0);
|
|
}
|