test/CodeGenCUDA/device-stub.cu

   1 // RUN: echo "GPU binary would be here" > %t
   2 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -fcuda-include-gpubinary %t -o - | FileCheck %s
   3 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -fcuda-include-gpubinary %t -o -  -DNOGLOBALS \
   4 // RUN:   | FileCheck %s -check-prefix=NOGLOBALS
   5 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - | FileCheck %s -check-prefix=NOGPUBIN
   6
   7 #include "Inputs/cuda.h"
   8
   9 #ifndef NOGLOBALS
  10 // CHECK-DAG: @device_var = internal global i32
  11 __device__ int device_var;
  12
  13 // CHECK-DAG: @constant_var = internal global i32
  14 __constant__ int constant_var;
  15
  16 // CHECK-DAG: @shared_var = internal global i32
  17 __shared__ int shared_var;
  18
  19 // Make sure host globals don't get internalized...
  20 // CHECK-DAG: @host_var = global i32
  21 int host_var;
  22 // ... and that extern vars remain external.
  23 // CHECK-DAG: @ext_host_var = external global i32
  24 extern int ext_host_var;
  25
  26 // Shadows for external device-side variables are *definitions* of
  27 // those variables.
  28 // CHECK-DAG: @ext_device_var = internal global i32
  29 extern __device__ int ext_device_var;
  30 // CHECK-DAG: @ext_device_var = internal global i32
  31 extern __constant__ int ext_constant_var;
  32
  33 void use_pointers() {
  34   int *p;
  35   p = &device_var;
  36   p = &constant_var;
  37   p = &shared_var;
  38   p = &host_var;
  39   p = &ext_device_var;
  40   p = &ext_constant_var;
  41   p = &ext_host_var;
  42 }
  43
  44 // Make sure that all parts of GPU code init/cleanup are there:
  45 // * constant unnamed string with the kernel name
  46 // CHECK: private unnamed_addr constant{{.*}}kernelfunc{{.*}}\00"
  47 // * constant unnamed string with GPU binary
  48 // CHECK: private unnamed_addr constant{{.*GPU binary would be here.*}}\00"
  49 // CHECK-SAME: section ".nv_fatbin", align 8
  50 // * constant struct that wraps GPU binary
  51 // CHECK: @__cuda_fatbin_wrapper = internal constant { i32, i32, i8*, i8* }
  52 // CHECK-SAME: { i32 1180844977, i32 1, {{.*}}, i8* null }
  53 // CHECK-SAME: section ".nvFatBinSegment"
  54 // * variable to save GPU binary handle after initialization
  55 // CHECK: @__cuda_gpubin_handle = internal global i8** null
  56 // * Make sure our constructor/destructor was added to global ctor/dtor list.
  57 // CHECK: @llvm.global_ctors = appending global {{.*}}@__cuda_module_ctor
  58 // CHECK: @llvm.global_dtors = appending global {{.*}}@__cuda_module_dtor
  59
  60 // Test that we build the correct number of calls to cudaSetupArgument followed
  61 // by a call to cudaLaunch.
  62
  63 // CHECK: define{{.*}}kernelfunc
  64 // CHECK: call{{.*}}cudaSetupArgument
  65 // CHECK: call{{.*}}cudaSetupArgument
  66 // CHECK: call{{.*}}cudaSetupArgument
  67 // CHECK: call{{.*}}cudaLaunch
  68 __global__ void kernelfunc(int i, int j, int k) {}
  69
  70 // Test that we've built correct kernel launch sequence.
  71 // CHECK: define{{.*}}hostfunc
  72 // CHECK: call{{.*}}cudaConfigureCall
  73 // CHECK: call{{.*}}kernelfunc
  74 void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
  75 #endif
  76
  77 // Test that we've built a function to register kernels and global vars.
  78 // CHECK: define internal void @__cuda_register_globals
  79 // CHECK: call{{.*}}cudaRegisterFunction(i8** %0, {{.*}}kernelfunc
  80 // CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}device_var{{.*}}i32 0, i32 4, i32 0, i32 0
  81 // CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}constant_var{{.*}}i32 0, i32 4, i32 1, i32 0
  82 // CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}ext_device_var{{.*}}i32 1, i32 4, i32 0, i32 0
  83 // CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}ext_constant_var{{.*}}i32 1, i32 4, i32 1, i32 0
  84 // CHECK: ret void
  85
  86 // Test that we've built contructor..
  87 // CHECK: define internal void @__cuda_module_ctor
  88 //   .. that calls __cudaRegisterFatBinary(&__cuda_fatbin_wrapper)
  89 // CHECK: call{{.*}}cudaRegisterFatBinary{{.*}}__cuda_fatbin_wrapper
  90 //   .. stores return value in __cuda_gpubin_handle
  91 // CHECK-NEXT: store{{.*}}__cuda_gpubin_handle
  92 //   .. and then calls __cuda_register_globals
  93 // CHECK-NEXT: call void @__cuda_register_globals
  94
  95 // Test that we've created destructor.
  96 // CHECK: define internal void @__cuda_module_dtor
  97 // CHECK: load{{.*}}__cuda_gpubin_handle
  98 // CHECK-NEXT: call void @__cudaUnregisterFatBinary
  99
 100 // There should be no __cuda_register_globals if we have no
 101 // device-side globals, but we still need to register GPU binary.
 102 // Skip GPU binary string first.
 103 // NOGLOBALS: @0 = private unnamed_addr constant{{.*}}
 104 // NOGLOBALS-NOT: define internal void @__cuda_register_globals
 105 // NOGLOBALS: define internal void @__cuda_module_ctor
 106 // NOGLOBALS: call{{.*}}cudaRegisterFatBinary{{.*}}__cuda_fatbin_wrapper
 107 // NOGLOBALS-NOT: call void @__cuda_register_globals
 108 // NOGLOBALS: define internal void @__cuda_module_dtor
 109 // NOGLOBALS: call void @__cudaUnregisterFatBinary
 110
 111 // There should be no constructors/destructors if we have no GPU binary.
 112 // NOGPUBIN-NOT: define internal void @__cuda_register_globals
 113 // NOGPUBIN-NOT: define internal void @__cuda_module_ctor
 114 // NOGPUBIN-NOT: define internal void @__cuda_module_dtor