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 \
   3 // RUN:     -fcuda-include-gpubinary %t -o - \
   4 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,NORDC,CUDA,CUDANORDC
   5 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
   6 // RUN:     -fcuda-include-gpubinary %t -o -  -DNOGLOBALS \
   7 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s -check-prefixes=NOGLOBALS,CUDANOGLOBALS
   8 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
   9 // RUN:     -fcuda-rdc -fcuda-include-gpubinary %t -o - \
  10 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,RDC,CUDA,CUDARDC
  11 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - \
  12 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=NOGPUBIN
  13
  14 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
  15 // RUN:     -fcuda-include-gpubinary %t -o - -x hip\
  16 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,NORDC,HIP
  17 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
  18 // RUN:     -fcuda-include-gpubinary %t -o -  -DNOGLOBALS -x hip \
  19 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s -check-prefixes=NOGLOBALS,HIPNOGLOBALS
  20 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s \
  21 // RUN:     -fcuda-rdc -fcuda-include-gpubinary %t -o - -x hip \
  22 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,NORDC,HIP
  23 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - -x hip\
  24 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=NOGPUBIN
  25
  26 #include "Inputs/cuda.h"
  27
  28 #ifndef NOGLOBALS
  29 // ALL-DAG: @device_var = internal global i32
  30 __device__ int device_var;
  31
  32 // ALL-DAG: @constant_var = internal global i32
  33 __constant__ int constant_var;
  34
  35 // ALL-DAG: @shared_var = internal global i32
  36 __shared__ int shared_var;
  37
  38 // Make sure host globals don't get internalized...
  39 // ALL-DAG: @host_var = global i32
  40 int host_var;
  41 // ... and that extern vars remain external.
  42 // ALL-DAG: @ext_host_var = external global i32
  43 extern int ext_host_var;
  44
  45 // Shadows for external device-side variables are *definitions* of
  46 // those variables.
  47 // ALL-DAG: @ext_device_var = internal global i32
  48 extern __device__ int ext_device_var;
  49 // ALL-DAG: @ext_device_var = internal global i32
  50 extern __constant__ int ext_constant_var;
  51
  52 void use_pointers() {
  53   int *p;
  54   p = &device_var;
  55   p = &constant_var;
  56   p = &shared_var;
  57   p = &host_var;
  58   p = &ext_device_var;
  59   p = &ext_constant_var;
  60   p = &ext_host_var;
  61 }
  62
  63 // Make sure that all parts of GPU code init/cleanup are there:
  64 // * constant unnamed string with the kernel name
  65 // ALL: private unnamed_addr constant{{.*}}kernelfunc{{.*}}\00"
  66 // * constant unnamed string with GPU binary
  67 // HIP: @[[FATBIN:__hip_fatbin]] = external constant i8, section ".hip_fatbin"
  68 // CUDA: @[[FATBIN:.*]] = private constant{{.*GPU binary would be here.*}}\00",
  69 // CUDANORDC-SAME: section ".nv_fatbin", align 8
  70 // CUDARDC-SAME: section "__nv_relfatbin", align 8
  71 // * constant struct that wraps GPU binary
  72 // ALL: @__[[PREFIX:cuda|hip]]_fatbin_wrapper = internal constant
  73 // ALL-SAME: { i32, i32, i8*, i8* }
  74 // CUDA-SAME: { i32 1180844977, i32 1,
  75 // HIP-SAME: { i32 1212764230, i32 1,
  76 // CUDA-SAME: i8* getelementptr inbounds ({{.*}}@[[FATBIN]], i64 0, i64 0),
  77 // HIP-SAME:  i8* @[[FATBIN]],
  78 // ALL-SAME: i8* null }
  79 // CUDA-SAME: section ".nvFatBinSegment"
  80 // HIP-SAME: section ".hipFatBinSegment"
  81 // * variable to save GPU binary handle after initialization
  82 // CUDANORDC: @__[[PREFIX]]_gpubin_handle = internal global i8** null
  83 // HIP: @__[[PREFIX]]_gpubin_handle = linkonce global i8** null
  84 // * constant unnamed string with NVModuleID
  85 // RDC: [[MODULE_ID_GLOBAL:@.*]] = private constant
  86 // CUDARDC-SAME: c"[[MODULE_ID:.+]]\00", section "__nv_module_id", align 32
  87 // * Make sure our constructor was added to global ctor list.
  88 // ALL: @llvm.global_ctors = appending global {{.*}}@__[[PREFIX]]_module_ctor
  89 // * Alias to global symbol containing the NVModuleID.
  90 // RDC: @__fatbinwrap[[MODULE_ID]] = alias { i32, i32, i8*, i8* }
  91 // RDC-SAME: { i32, i32, i8*, i8* }* @__[[PREFIX]]_fatbin_wrapper
  92
  93 // Test that we build the correct number of calls to cudaSetupArgument followed
  94 // by a call to cudaLaunch.
  95
  96 // ALL: define{{.*}}kernelfunc
  97 // ALL: call{{.*}}[[PREFIX]]SetupArgument
  98 // ALL: call{{.*}}[[PREFIX]]SetupArgument
  99 // ALL: call{{.*}}[[PREFIX]]SetupArgument
 100 // ALL: call{{.*}}[[PREFIX]]Launch
 101 __global__ void kernelfunc(int i, int j, int k) {}
 102
 103 // Test that we've built correct kernel launch sequence.
 104 // ALL: define{{.*}}hostfunc
 105 // ALL: call{{.*}}[[PREFIX]]ConfigureCall
 106 // ALL: call{{.*}}kernelfunc
 107 void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
 108 #endif
 109
 110 // Test that we've built a function to register kernels and global vars.
 111 // ALL: define internal void @__[[PREFIX]]_register_globals
 112 // ALL: call{{.*}}[[PREFIX]]RegisterFunction(i8** %0, {{.*}}kernelfunc
 113 // ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}device_var{{.*}}i32 0, i32 4, i32 0, i32 0
 114 // ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}constant_var{{.*}}i32 0, i32 4, i32 1, i32 0
 115 // ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_device_var{{.*}}i32 1, i32 4, i32 0, i32 0
 116 // ALL-DAG: call{{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_constant_var{{.*}}i32 1, i32 4, i32 1, i32 0
 117 // ALL: ret void
 118
 119 // Test that we've built a constructor.
 120 // ALL: define internal void @__[[PREFIX]]_module_ctor
 121
 122 // In separate mode it calls __[[PREFIX]]RegisterFatBinary(&__[[PREFIX]]_fatbin_wrapper)
 123 // HIP only register fat binary once.
 124 // HIP: load i8**, i8*** @__hip_gpubin_handle
 125 // HIP-NEXT: icmp eq i8** {{.*}}, null
 126 // HIP-NEXT: br i1 {{.*}}, label %if, label %exit
 127 // HIP: if:
 128 // NORDC: call{{.*}}[[PREFIX]]RegisterFatBinary{{.*}}__[[PREFIX]]_fatbin_wrapper
 129 //   .. stores return value in __[[PREFIX]]_gpubin_handle
 130 // NORDC-NEXT: store{{.*}}__[[PREFIX]]_gpubin_handle
 131 //   .. and then calls __[[PREFIX]]_register_globals
 132 // HIP-NEXT: br label %exit
 133 // HIP: exit:
 134 // HIP-NEXT: load i8**, i8*** @__hip_gpubin_handle
 135 // NORDC-NEXT: call void @__[[PREFIX]]_register_globals
 136 // * In separate mode we also register a destructor.
 137 // NORDC-NEXT: call i32 @atexit(void (i8*)* @__[[PREFIX]]_module_dtor)
 138
 139 // With relocatable device code we call __[[PREFIX]]RegisterLinkedBinary%NVModuleID%
 140 // RDC: call{{.*}}__[[PREFIX]]RegisterLinkedBinary[[MODULE_ID]](
 141 // RDC-SAME: __[[PREFIX]]_register_globals, {{.*}}__[[PREFIX]]_fatbin_wrapper
 142 // RDC-SAME: [[MODULE_ID_GLOBAL]]
 143
 144 // Test that we've created destructor.
 145 // NORDC: define internal void @__[[PREFIX]]_module_dtor
 146 // NORDC: load{{.*}}__[[PREFIX]]_gpubin_handle
 147 // CUDANORDC-NEXT: call void @__[[PREFIX]]UnregisterFatBinary
 148 // HIP-NEXT: icmp ne i8** {{.*}}, null
 149 // HIP-NEXT: br i1 {{.*}}, label %if, label %exit
 150 // HIP: if:
 151 // HIP-NEXT: call void @__[[PREFIX]]UnregisterFatBinary
 152 // HIP-NEXT: store i8** null, i8*** @__hip_gpubin_handle
 153 // HIP-NEXT: br label %exit
 154 // HIP: exit:
 155
 156 // There should be no __[[PREFIX]]_register_globals if we have no
 157 // device-side globals, but we still need to register GPU binary.
 158 // Skip GPU binary string first.
 159 // CUDANOGLOBALS: @{{.*}} = private constant{{.*}}
 160 // HIPNOGLOBALS: @{{.*}} = external constant{{.*}}
 161 // NOGLOBALS-NOT: define internal void @__{{.*}}_register_globals
 162 // NOGLOBALS: define internal void @__[[PREFIX:cuda|hip]]_module_ctor
 163 // NOGLOBALS: call{{.*}}[[PREFIX]]RegisterFatBinary{{.*}}__[[PREFIX]]_fatbin_wrapper
 164 // NOGLOBALS-NOT: call void @__[[PREFIX]]_register_globals
 165 // NOGLOBALS: define internal void @__[[PREFIX]]_module_dtor
 166 // NOGLOBALS: call void @__[[PREFIX]]UnregisterFatBinary
 167
 168 // There should be no constructors/destructors if we have no GPU binary.
 169 // NOGPUBIN-NOT: define internal void @__[[PREFIX]]_register_globals
 170 // NOGPUBIN-NOT: define internal void @__[[PREFIX]]_module_ctor
 171 // NOGPUBIN-NOT: define internal void @__[[PREFIX]]_module_dtor