Index: clang/test/PCH/pragma-cuda-force-host-device.cu =================================================================== --- clang/test/PCH/pragma-cuda-force-host-device.cu +++ clang/test/PCH/pragma-cuda-force-host-device.cu @@ -1,5 +1,5 @@ // RUN: %clang_cc1 -emit-pch %s -o %t -// RUN: %clang_cc1 -verify -verify-ignore-unexpected=note -include-pch %t -S -o /dev/null %s +// RUN: %clang_cc1 -verify -verify-ignore-unexpected=note -include-pch %t -fsyntax-only %s #ifndef HEADER #define HEADER Index: clang/test/Parser/cuda-force-host-device-templates.cu =================================================================== --- clang/test/Parser/cuda-force-host-device-templates.cu +++ clang/test/Parser/cuda-force-host-device-templates.cu @@ -1,8 +1,7 @@ -// RUN: %clang_cc1 -std=c++14 -S -verify -fcuda-is-device %s -o /dev/null +// RUN: %clang_cc1 -std=c++14 -fsyntax-only -verify -fcuda-is-device %s // Check how the force_cuda_host_device pragma interacts with template -// instantiations. The errors here are emitted at codegen, so we can't do -// -fsyntax-only. +// instantiations. template auto foo() { // expected-note {{declared here}} Index: clang/test/SemaCUDA/device-var-init.cu =================================================================== --- clang/test/SemaCUDA/device-var-init.cu +++ clang/test/SemaCUDA/device-var-init.cu @@ -213,3 +213,15 @@ static int v; // expected-error@-1 {{within a __device__ function, only __shared__ variables may be marked 'static'}} } + +__host__ __device__ void hd_sema() { + static int x = 42; +#ifdef __CUDA_ARCH__ + // expected-error@-2 {{within a __host__ __device__ function, only __shared__ variables may be marked 'static'}} +#endif +} + +inline __host__ __device__ void hd_emitted_host_only() { + static int x = 42; // no error on device because this is never codegen'ed there. +} +void call_hd_emitted_host_only() { hd_emitted_host_only(); } Index: clang/test/SemaCUDA/exceptions-host-device.cu =================================================================== --- clang/test/SemaCUDA/exceptions-host-device.cu +++ /dev/null @@ -1,38 +0,0 @@ -// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -verify %s -S -o /dev/null -// RUN: %clang_cc1 -fcxx-exceptions -verify -DHOST %s -S -o /dev/null - -#include "Inputs/cuda.h" - -// Check that it's an error to use 'try' and 'throw' from a __host__ __device__ -// function if and only if it's codegen'ed for device. - -#ifdef HOST -// expected-no-diagnostics -#endif - -__host__ __device__ void hd1() { - throw NULL; - try {} catch(void*) {} -#ifndef HOST - // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}} - // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}} -#endif -} - -// No error, never instantiated on device. -inline __host__ __device__ void hd2() { - throw NULL; - try {} catch(void*) {} -} -void call_hd2() { hd2(); } - -// Error, instantiated on device. -inline __host__ __device__ void hd3() { - throw NULL; - try {} catch(void*) {} -#ifndef HOST - // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}} - // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}} -#endif -} -__device__ void call_hd3() { hd3(); } Index: clang/test/SemaCUDA/exceptions.cu =================================================================== --- clang/test/SemaCUDA/exceptions.cu +++ clang/test/SemaCUDA/exceptions.cu @@ -19,3 +19,37 @@ try {} catch(void*) {} // expected-error@-1 {{cannot use 'try' in __global__ function}} } + +// Check that it's an error to use 'try' and 'throw' from a __host__ __device__ +// function if and only if it's codegen'ed for device. + +__host__ __device__ void hd1() { + throw NULL; + try {} catch(void*) {} +#ifdef __CUDA_ARCH__ + // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}} + // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}} +#endif +} + +// No error, never instantiated on device. +inline __host__ __device__ void hd2() { + throw NULL; + try {} catch(void*) {} +} +void call_hd2() { hd2(); } + +// Error, instantiated on device. +inline __host__ __device__ void hd3() { + throw NULL; + try {} catch(void*) {} +#ifdef __CUDA_ARCH__ + // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}} + // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}} +#endif +} + +__device__ void call_hd3() { hd3(); } +#ifdef __CUDA_ARCH__ +// expected-note@-2 {{called by 'call_hd3'}} +#endif Index: clang/test/SemaCUDA/function-overload-hd.cu =================================================================== --- clang/test/SemaCUDA/function-overload-hd.cu +++ /dev/null @@ -1,31 +0,0 @@ -// REQUIRES: x86-registered-target -// REQUIRES: nvptx-registered-target - -// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -S -o /dev/null -verify \ -// RUN: -verify-ignore-unexpected=note %s -// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -S -o /dev/null -fcuda-is-device \ -// RUN: -verify -verify-ignore-unexpected=note %s - -#include "Inputs/cuda.h" - -// FIXME: Merge into function-overload.cu once deferred errors can be emitted -// when non-deferred errors are present. - -#if !defined(__CUDA_ARCH__) -//expected-no-diagnostics -#endif - -typedef void (*GlobalFnPtr)(); // __global__ functions must return void. - -__global__ void g() {} - -__host__ __device__ void hd() { - GlobalFnPtr fp_g = g; -#if defined(__CUDA_ARCH__) - // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}} -#endif - g<<<0,0>>>(); -#if defined(__CUDA_ARCH__) - // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}} -#endif // __CUDA_ARCH__ -} Index: clang/test/SemaCUDA/function-overload.cu =================================================================== --- clang/test/SemaCUDA/function-overload.cu +++ clang/test/SemaCUDA/function-overload.cu @@ -193,12 +193,22 @@ CurrentFnPtr fp_cdh = cdh; CurrentReturnTy ret_cdh = cdh(); + GlobalFnPtr fp_g = g; +#if defined(__CUDA_ARCH__) + // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}} +#endif + g(); #if defined (__CUDA_ARCH__) // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}} #else // expected-error@-4 {{call to global function g not configured}} #endif + + g<<<0,0>>>(); +#if defined(__CUDA_ARCH__) + // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}} +#endif } // Test for address of overloaded function resolution in the global context. Index: clang/test/SemaCUDA/implicit-device-lambda-hd.cu =================================================================== --- clang/test/SemaCUDA/implicit-device-lambda-hd.cu +++ /dev/null @@ -1,27 +0,0 @@ -// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -verify -verify-ignore-unexpected=note \ -// RUN: -S -o /dev/null %s -// RUN: %clang_cc1 -std=c++11 -verify -fsyntax-only -verify-ignore-unexpected=note \ -// RUN: -DHOST -S -o /dev/null %s -#include "Inputs/cuda.h" - -__host__ __device__ void hd_fn() { - auto f1 = [&] {}; - f1(); // implicitly __host__ __device__ - - auto f2 = [&] __device__ {}; - f2(); -#ifdef HOST - // expected-error@-2 {{reference to __device__ function}} -#endif - - auto f3 = [&] __host__ {}; - f3(); -#ifndef HOST - // expected-error@-2 {{reference to __host__ function}} -#endif - - auto f4 = [&] __host__ __device__ {}; - f4(); -} - - Index: clang/test/SemaCUDA/implicit-device-lambda.cu =================================================================== --- clang/test/SemaCUDA/implicit-device-lambda.cu +++ clang/test/SemaCUDA/implicit-device-lambda.cu @@ -76,6 +76,26 @@ f4(); } +__host__ __device__ void hd_fn() { + auto f1 = [&] {}; + f1(); // implicitly __host__ __device__ + + auto f2 = [&] __device__ {}; + f2(); +#ifndef __CUDA_ARCH__ + // expected-error@-2 {{reference to __device__ function}} +#endif + + auto f3 = [&] __host__ {}; + f3(); +#ifdef __CUDA_ARCH__ + // expected-error@-2 {{reference to __host__ function}} +#endif + + auto f4 = [&] __host__ __device__ {}; + f4(); +} + // The special treatment above only applies to lambdas. __device__ void foo() { struct X { Index: clang/test/SemaCUDA/static-vars-hd.cu =================================================================== --- clang/test/SemaCUDA/static-vars-hd.cu +++ /dev/null @@ -1,20 +0,0 @@ -// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -S -o /dev/null -verify %s -// RUN: %clang_cc1 -fcxx-exceptions -S -o /dev/null -D HOST -verify %s - -#include "Inputs/cuda.h" - -#ifdef HOST -// expected-no-diagnostics -#endif - -__host__ __device__ void f() { - static int x = 42; -#ifndef HOST - // expected-error@-2 {{within a __host__ __device__ function, only __shared__ variables may be marked 'static'}} -#endif -} - -inline __host__ __device__ void g() { - static int x = 42; // no error on device because this is never codegen'ed there. -} -void call_g() { g(); } Index: clang/test/SemaCUDA/vla-host-device.cu =================================================================== --- clang/test/SemaCUDA/vla-host-device.cu +++ /dev/null @@ -1,21 +0,0 @@ -// RUN: %clang_cc1 -fcuda-is-device -verify -S %s -o /dev/null -// RUN: %clang_cc1 -verify -DHOST %s -S -o /dev/null - -#include "Inputs/cuda.h" - -#ifdef HOST -// expected-no-diagnostics -#endif - -__host__ __device__ void hd(int n) { - int x[n]; -#ifndef HOST - // expected-error@-2 {{cannot use variable-length arrays in __host__ __device__ functions}} -#endif -} - -// No error because never codegen'ed for device. -__host__ __device__ inline void hd_inline(int n) { - int x[n]; -} -void call_hd_inline() { hd_inline(42); } Index: clang/test/SemaCUDA/vla.cu =================================================================== --- clang/test/SemaCUDA/vla.cu +++ clang/test/SemaCUDA/vla.cu @@ -10,3 +10,16 @@ __device__ void device(int n) { int x[n]; // expected-error {{cannot use variable-length arrays in __device__ functions}} } + +__host__ __device__ void hd(int n) { + int x[n]; +#ifdef __CUDA_ARCH__ + // expected-error@-2 {{cannot use variable-length arrays in __host__ __device__ functions}} +#endif +} + +// No error because never codegen'ed for device. +__host__ __device__ inline void hd_inline(int n) { + int x[n]; +} +void call_hd_inline() { hd_inline(42); }