clang/test/CodeGenCUDA/function-overload.cu

   1 // REQUIRES: x86-registered-target
   2 // REQUIRES: nvptx-registered-target
   3
   4 // Make sure we handle target overloads correctly.  Most of this is checked in
   5 // sema, but special functions like constructors and destructors are here.
   6 //
   7 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s \
   8 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST %s
   9 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s \
  10 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE %s
  11
  12 #include "Inputs/cuda.h"
  13
  14 // Check constructors/destructors for D/H functions
  15 #ifdef __CUDA_ARCH__
  16 __device__
  17 #endif
  18 int x;
  19 struct s_cd_dh {
  20   __host__ s_cd_dh() { x = 11; }
  21   __device__ s_cd_dh() { x = 12; }
  22 };
  23
  24 struct s_cd_hd {
  25   __host__ __device__ s_cd_hd() { x = 31; }
  26   __host__ __device__ ~s_cd_hd() { x = 32; }
  27 };
  28
  29 // CHECK-BOTH: define{{.*}} void @_Z7wrapperv
  30 #if defined(__CUDA_ARCH__)
  31 __device__
  32 #else
  33 __host__
  34 #endif
  35 void wrapper() {
  36   s_cd_dh scddh;
  37   // CHECK-BOTH: call void @_ZN7s_cd_dhC1Ev(
  38   s_cd_hd scdhd;
  39   // CHECK-BOTH: call void @_ZN7s_cd_hdC1Ev
  40
  41   // CHECK-BOTH: call void @_ZN7s_cd_hdD1Ev(
  42 }
  43 // CHECK-BOTH: ret void
  44
  45 // Now it's time to check what's been generated for the methods we used.
  46
  47 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhC2Ev(
  48 // CHECK-HOST:   store i32 11,
  49 // CHECK-DEVICE: store i32 12,
  50 // CHECK-BOTH: ret void
  51
  52 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdC2Ev(
  53 // CHECK-BOTH:   store i32 31,
  54 // CHECK-BOTH: ret void
  55
  56 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2Ev(
  57 // CHECK-BOTH: store i32 32,
  58 // CHECK-BOTH: ret void