| // REQUIRES: amdgpu-registered-target | |
| // REQUIRES: nvptx-registered-target | |
| // Implicit-H+D functions reached only via an explicit template instantiation | |
| // get a real body if their device codegen is clean, or a trap body if it | |
| // would otherwise raise a target-mismatch diagnostic. | |
| // RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip \ | |
| // RUN: -emit-llvm -o - %s | FileCheck %s | |
| // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -x hip \ | |
| // RUN: -emit-llvm -o - %s | FileCheck %s | |
| #include "Inputs/cuda.h" | |
| __host__ int host_only() { return 7; } | |
| // Class-template member, safe body. | |
| template <typename T> | |
| struct ETI_Safe { | |
| constexpr T add_one(T x) { return x + (T)1; } | |
| }; | |
| template class ETI_Safe<float>; | |
| // CHECK: define {{.*}} float @_ZN8ETI_SafeIfE7add_oneEf( | |
| // CHECK-NOT: call void @llvm.trap() | |
| // CHECK: ret float | |
| // Class-template member, body calls a __host__-only function. | |
| template <typename T> | |
| struct ETI_BadBody { | |
| constexpr T call_host(T x) { return x + (T)host_only(); } | |
| }; | |
| template class ETI_BadBody<float>; | |
| // CHECK-LABEL: define {{.*}} float @_ZN11ETI_BadBodyIfE9call_hostEf( | |
| // CHECK: call void @llvm.trap() | |
| // CHECK-NEXT: ret float poison | |
| // Function-template form, safe body. | |
| template <typename T> | |
| constexpr T ft_safe(T x) { return x + (T)1; } | |
| template float ft_safe<float>(float); | |
| // CHECK: define {{.*}} float @_Z7ft_safeIfET_S0_( | |
| // CHECK-NOT: call void @llvm.trap() | |
| // CHECK: ret float | |
| // Function-template form, body calls a __host__-only function. | |
| template <typename T> | |
| constexpr T ft_bad(T x) { return x + (T)host_only(); } | |
| template float ft_bad<float>(float); | |
| // CHECK-LABEL: define {{.*}} float @_Z6ft_badIfET_S0_( | |
| // CHECK: call void @llvm.trap() | |
| // CHECK-NEXT: ret float poison | |
| __global__ void kernel(float *p) { *p = 1.0f; } | |
| // CHECK-NOT: __clang_hip_unreachable_dtor | |
| // CHECK-NOT: __clang_cuda_unreachable_dtor |