blob: 1e0e81079fe5018d3133a36f715623c9f872cf54 [file] [log] [blame] [edit]
// RUN: %clang_cc1 -std=c++17 -triple nvptx64-nvidia-cuda -fsyntax-only \
// RUN: -fcuda-is-device -verify %s
// RUN: %clang_cc1 -std=c++17 -triple nvptx64-nvidia-cuda -fsyntax-only \
// RUN: -verify %s
#include "Inputs/cuda.h"
template <typename T>
struct S {
__host__ __device__ S(T);
};
// A host+device deduction guide is allowed and participates in CTAD, but its
// explicit target attributes are deprecated and will be rejected in a future
// Clang version.
template <typename T>
__host__ __device__ S(T) -> S<T>; // expected-warning {{use of CUDA/HIP target attributes on deduction guides is deprecated; they will be rejected in a future version of Clang}}
__host__ __device__ void use_hd_guide() {
S s(42); // uses the explicit __host__ __device__ deduction guide above
}
// CUDA/HIP target attributes on deduction guides are rejected when they make
// the guide host-only, device-only, or a kernel.
template <typename U>
__host__ S(U) -> S<U>; // expected-error {{in CUDA/HIP, deduction guides may only be annotated with '__host__ __device__'; '__host__'-only, '__device__'-only, or '__global__' deduction guides are not allowed}}
template <typename V>
__device__ S(V) -> S<V>; // expected-error {{in CUDA/HIP, deduction guides may only be annotated with '__host__ __device__'; '__host__'-only, '__device__'-only, or '__global__' deduction guides are not allowed}}
template <typename W>
__global__ S(W) -> S<W>; // expected-error {{in CUDA/HIP, deduction guides may only be annotated with '__host__ __device__'; '__host__'-only, '__device__'-only, or '__global__' deduction guides are not allowed}}