49 lines · plain
1// RUN: %clang_cc1 -fsyntax-only -verify %s2 3// Check that it's OK for kernels to call HD functions that call device-only4// functions.5 6#include "Inputs/cuda.h"7 8__device__ void device_fn(int) {}9// expected-note@-1 2 {{declared here}}10 11inline __host__ __device__ int hd1() {12 device_fn(0); // expected-error {{reference to __device__ function}}13 return 0;14}15 16inline __host__ __device__ int hd2() {17 // No error here because hd2 is only referenced from a kernel.18 device_fn(0);19 return 0;20}21 22inline __host__ __device__ void hd3(int) {23 device_fn(0); // expected-error {{reference to __device__ function 'device_fn'}}24}25inline __host__ __device__ void hd3(double) {}26 27inline __host__ __device__ void hd4(int) {}28inline __host__ __device__ void hd4(double) {29 device_fn(0); // No error; this function is never called.30}31 32__global__ void kernel(int) { hd2(); }33 34template <typename T>35void launch_kernel() {36 kernel<<<0, 0>>>(T());37 38 // Notice that these two diagnostics are different: Because the call to hd139 // is not dependent on T, the call to hd1 comes from 'launch_kernel', while40 // the call to hd3, being dependent, comes from 'launch_kernel<int>'.41 hd1(); // expected-note {{called by 'launch_kernel<int>'}}42 hd3(T()); // expected-note {{called by 'launch_kernel<int>'}}43}44 45void host_fn() {46 launch_kernel<int>();47 // expected-note@-1 2 {{called by 'host_fn'}}48}49