clang/test/SemaCUDA/exceptions.cu

   1 // RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -fsyntax-only -verify %s
   2 // RUN: %clang_cc1 -fcxx-exceptions -fsyntax-only -verify %s
   3
   4 #include "Inputs/cuda.h"
   5
   6 void host() {
   7   throw NULL;
   8   try {} catch(void*) {}
   9 }
  10 __device__ void device() {
  11   throw NULL;
  12   // expected-error@-1 {{cannot use 'throw' in __device__ function}}
  13   try {} catch(void*) {}
  14   // expected-error@-1 {{cannot use 'try' in __device__ function}}
  15 }
  16 __global__ void kernel() {
  17   throw NULL;
  18   // expected-error@-1 {{cannot use 'throw' in __global__ function}}
  19   try {} catch(void*) {}
  20   // expected-error@-1 {{cannot use 'try' in __global__ function}}
  21 }
  22
  23 // Check that it's an error to use 'try' and 'throw' from a __host__ __device__
  24 // function if and only if it's codegen'ed for device.
  25
  26 __host__ __device__ void hd1() {
  27   throw NULL;
  28   try {} catch(void*) {}
  29 #ifdef __CUDA_ARCH__
  30   // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}}
  31   // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}}
  32 #endif
  33 }
  34
  35 // No error, never instantiated on device.
  36 inline __host__ __device__ void hd2() {
  37   throw NULL;
  38   try {} catch(void*) {}
  39 }
  40 void call_hd2() { hd2(); }
  41
  42 // Error, instantiated on device.
  43 inline __host__ __device__ void hd3() {
  44   throw NULL;
  45   try {} catch(void*) {}
  46 #ifdef __CUDA_ARCH__
  47   // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}}
  48   // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}}
  49 #endif
  50 }
  51
  52 __device__ void call_hd3() { hd3(); }
  53 #ifdef __CUDA_ARCH__
  54 // expected-note@-2 {{called by 'call_hd3'}}
  55 #endif