clang/test/CodeGenCUDA/device-stub.cu

   1 // RUN: echo -n "GPU binary would be here." > %t
   2 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
   3 // RUN:     -target-sdk-version=8.0 -fcuda-include-gpubinary %t -o - \
   4 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s \
   5 // RUN:       --check-prefixes=ALL,LNX,NORDC,CUDA,CUDANORDC,CUDA-OLD
   6 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
   7 // RUN:     -target-sdk-version=8.0  -fcuda-include-gpubinary %t \
   8 // RUN:     -o - -DNOGLOBALS \
   9 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s \
  10 // RUN:     -check-prefixes=NOGLOBALS,CUDANOGLOBALS
  11 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
  12 // RUN:     -target-sdk-version=8.0 -fgpu-rdc -fcuda-include-gpubinary %t \
  13 // RUN:     -o - \
  14 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s \
  15 // RUN:       --check-prefixes=ALL,LNX,RDC,CUDA,CUDARDC,CUDA-OLD
  16 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
  17 // RUN:     -target-sdk-version=8.0 -o - \
  18 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=NOGPUBIN
  19
  20 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s       \
  21 // RUN:     -target-sdk-version=9.2 -fcuda-include-gpubinary %t -o - \
  22 // RUN:   | FileCheck %s -allow-deprecated-dag-overlap \
  23 // RUN:       --check-prefixes=ALL,LNX,NORDC,CUDA,CUDANORDC,CUDA-NEW
  24 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
  25 // RUN:     -target-sdk-version=9.2 -fcuda-include-gpubinary %t -o -  -DNOGLOBALS \
  26 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s \
  27 // RUN:       --check-prefixes=NOGLOBALS,CUDANOGLOBALS
  28 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
  29 // RUN:     -target-sdk-version=9.2 -fgpu-rdc -fcuda-include-gpubinary %t -o - \
  30 // RUN:   | FileCheck %s -allow-deprecated-dag-overlap \
  31 // RUN:       --check-prefixes=ALL,LNX,RDC,CUDA,CUDARDC,CUDA-NEW
  32 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s -std=c++17 \
  33 // RUN:     -target-sdk-version=9.2 -fcuda-include-gpubinary %t -o - \
  34 // RUN:   | FileCheck %s -allow-deprecated-dag-overlap \
  35 // RUN:       --check-prefixes=ALL,LNX,NORDC,CUDA,CUDANORDC,CUDA-NEW,LNX_17,NORDC17
  36 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s -std=c++17 \
  37 // RUN:     -target-sdk-version=9.2 -fgpu-rdc -fcuda-include-gpubinary %t -o - \
  38 // RUN:   | FileCheck %s -allow-deprecated-dag-overlap \
  39 // RUN:       --check-prefixes=ALL,LNX,RDC,CUDA,CUDARDC,CUDA-NEW,LNX_17,RDC17
  40 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
  41 // RUN:     -target-sdk-version=9.2 -o - \
  42 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s -check-prefix=NOGPUBIN
  43
  44 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
  45 // RUN:     -fcuda-include-gpubinary %t -o - -x hip\
  46 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,LNX,NORDC,HIP,HIPEF
  47 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
  48 // RUN:     -fcuda-include-gpubinary %t -o -  -DNOGLOBALS -x hip \
  49 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s -check-prefixes=NOGLOBALS,HIPNOGLOBALS
  50 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s \
  51 // RUN:     -fgpu-rdc -fcuda-include-gpubinary %t -o - -x hip \
  52 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,LNX,RDC,HIP,HIPEF
  53 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-linux-gnu -emit-llvm %s -o - -x hip\
  54 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s -check-prefixes=ALL,LNX,NORDC,HIP,HIPNEF
  55
  56 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-pc-windows-msvc -aux-triple amdgcn -emit-llvm %s \
  57 // RUN:     -fcuda-include-gpubinary %t -o - -x hip\
  58 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,WIN
  59
  60 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-pc-windows-msvc -aux-triple amdgcn -emit-llvm %s \
  61 // RUN:     -o - -x hip\
  62 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,WIN,HIP,HIPNEF
  63
  64 #include "Inputs/cuda.h"
  65
  66 // HIPNEF: $__hip_gpubin_handle = comdat any
  67
  68 #ifndef NOGLOBALS
  69 // NORDC-DAG: @device_var = internal global i32
  70 // RDC-DAG: @device_var = global i32
  71 // WIN-DAG: @"?device_var@@3HA" = internal global i32
  72 __device__ int device_var;
  73
  74 // NORDC-DAG: @constant_var = internal global i32
  75 // RDC-DAG: @constant_var = global i32
  76 // WIN-DAG: @"?constant_var@@3HA" = internal global i32
  77 __constant__ int constant_var;
  78
  79 // NORDC-DAG: @shared_var = internal global i32
  80 // RDC-DAG: @shared_var = global i32
  81 // WIN-DAG: @"?shared_var@@3HA" = internal global i32
  82 __shared__ int shared_var;
  83
  84 // Make sure host globals don't get internalized...
  85 // LNX-DAG: @host_var ={{.*}} global i32
  86 // WIN-DAG: @"?host_var@@3HA" = dso_local global i32
  87 int host_var;
  88 // ... and that extern vars remain external.
  89 // LNX-DAG: @ext_host_var = external global i32
  90 // WIN-DAG: @"?ext_host_var@@3HA" = external dso_local global i32
  91 extern int ext_host_var;
  92
  93 // external device-side variables -> extern references to their shadows.
  94 // LNX-DAG: @ext_device_var = external global i32
  95 // WIN-DAG: @"?ext_device_var@@3HA" = external dso_local global i32
  96 extern __device__ int ext_device_var;
  97 // LNX-DAG: @ext_device_var = external global i32
  98 // WIN-DAG: @"?ext_constant_var@@3HA" = external dso_local global i32
  99 extern __constant__ int ext_constant_var;
 100
 101 // external device-side variables with definitions should generate
 102 // definitions for the shadows.
 103 // NORDC-DAG: @ext_device_var_def = internal global i32 undef,
 104 // RDC-DAG: @ext_device_var_def = global i32 undef,
 105 // WIN-DAG: @"?ext_device_var_def@@3HA" = internal global i32 undef
 106 extern __device__ int ext_device_var_def;
 107 __device__ int ext_device_var_def = 1;
 108 // NORDC-DAG: @ext_device_var_def = internal global i32 undef,
 109 // RDC-DAG: @ext_device_var_def = global i32 undef,
 110 // WIN-DAG: @"?ext_constant_var_def@@3HA" = internal global i32 undef
 111 __constant__ int ext_constant_var_def = 2;
 112
 113 #if __cplusplus > 201402L
 114 // NORDC17: @inline_var = internal global i32 undef, comdat, align 4{{$}}
 115 // RDC17: @inline_var = linkonce_odr global i32 undef, comdat, align 4{{$}}
 116 // NORDC17-NOT: @inline_var2 =
 117 // RDC17-NOT: @inline_var2 =
 118 // NORDC17: @_ZN1C17member_inline_varE = internal constant i32 undef, comdat, align 4{{$}}
 119 // RDC17: @_ZN1C17member_inline_varE = linkonce_odr constant i32 undef, comdat, align 4{{$}}
 120 // Check inline variable ODR-used by host is emitted on host and registered.
 121 __device__ inline int inline_var = 3;
 122 // Check inline variable not ODR-used by host is not emitted on host or registered.
 123 __device__ inline int inline_var2 = 5;
 124 struct C {
 125   __device__ static constexpr int member_inline_var = 4;
 126 };
 127 #endif
 128
 129 void use_pointers() {
 130   const int *p;
 131   p = &device_var;
 132   p = &constant_var;
 133   p = &shared_var;
 134   p = &host_var;
 135   p = &ext_device_var;
 136   p = &ext_constant_var;
 137   p = &ext_host_var;
 138 #if __cplusplus > 201402L
 139   p = &inline_var;
 140   decltype(inline_var2) tmp;
 141   p = &C::member_inline_var;
 142 #endif
 143 }
 144
 145 __device__ void device_use() {
 146 #if __cplusplus > 201402L
 147   const int *p = &inline_var2;
 148 #endif
 149 }
 150
 151 // Make sure that all parts of GPU code init/cleanup are there:
 152 // * constant unnamed string with the device-side kernel name to be passed to
 153 //   __hipRegisterFunction/__cudaRegisterFunction.
 154 // ALL: @0 = private unnamed_addr constant [18 x i8] c"_Z10kernelfunciii\00"
 155 // * constant unnamed string with the device-side kernel name to be passed to
 156 //   __hipRegisterVar/__cudaRegisterVar.
 157 // ALL: @1 = private unnamed_addr constant [11 x i8] c"device_var\00"
 158 // ALL: @2 = private unnamed_addr constant [13 x i8] c"constant_var\00"
 159 // ALL: @3 = private unnamed_addr constant [19 x i8] c"ext_device_var_def\00"
 160 // ALL: @4 = private unnamed_addr constant [21 x i8] c"ext_constant_var_def\00"
 161 // * constant unnamed string with GPU binary
 162 // CUDA: @[[FATBIN:.*]] = private constant{{.*}} c"GPU binary would be here.",
 163 // HIPEF: @[[FATBIN:.*]] = private constant{{.*}} c"GPU binary would be here.",{{.*}}align 4096
 164 // HIPNEF: @[[FATBIN:__hip_fatbin]] = external constant i8, section ".hip_fatbin"
 165 // CUDANORDC-SAME: section ".nv_fatbin", align 8
 166 // CUDARDC-SAME: section "__nv_relfatbin", align 8
 167 // * constant struct that wraps GPU binary
 168 // ALL: @__[[PREFIX:cuda|hip]]_fatbin_wrapper = internal constant
 169 // LNX-SAME: { i32, i32, i8*, i8* }
 170 // CUDA-SAME: { i32 1180844977, i32 1,
 171 // HIP-SAME: { i32 1212764230, i32 1,
 172 // CUDA-SAME: i8* getelementptr inbounds ({{.*}}@[[FATBIN]], i64 0, i64 0),
 173 // HIPEF-SAME: i8* getelementptr inbounds ({{.*}}@[[FATBIN]], i64 0, i64 0),
 174 // HIPNEF-SAME:  i8* @[[FATBIN]],
 175 // LNX-SAME: i8* null }
 176 // CUDA-SAME: section ".nvFatBinSegment"
 177 // HIP-SAME: section ".hipFatBinSegment"
 178 // * variable to save GPU binary handle after initialization
 179 // CUDANORDC: @__[[PREFIX]]_gpubin_handle = internal global i8** null
 180 // HIPNEF: @__[[PREFIX]]_gpubin_handle = linkonce hidden global i8** null
 181 // * constant unnamed string with NVModuleID
 182 // CUDARDC: [[MODULE_ID_GLOBAL:@.*]] = private constant
 183 // CUDARDC-SAME: c"[[MODULE_ID:.+]]\00", section "__nv_module_id", align 32
 184 // * Make sure our constructor was added to global ctor list.
 185 // LNX: @llvm.global_ctors = appending global {{.*}}@__[[PREFIX]]_module_ctor
 186 // * Alias to global symbol containing the NVModuleID.
 187 // CUDARDC: @__fatbinwrap[[MODULE_ID]] ={{.*}} alias { i32, i32, i8*, i8* }
 188 // CUDARDC-SAME: { i32, i32, i8*, i8* }* @__[[PREFIX]]_fatbin_wrapper
 189
 190 // Test that we build the correct number of calls to cudaSetupArgument followed
 191 // by a call to cudaLaunch.
 192
 193 // LNX: define{{.*}}kernelfunc
 194
 195 // New launch sequence stores arguments into local buffer and passes array of
 196 // pointers to them directly to cudaLaunchKernel
 197 // CUDA-NEW: alloca
 198 // CUDA-NEW: store
 199 // CUDA-NEW: store
 200 // CUDA-NEW: store
 201 // CUDA-NEW: call{{.*}}__cudaPopCallConfiguration
 202 // CUDA-NEW: call{{.*}}cudaLaunchKernel
 203
 204 // Legacy style launch sequence sets up arguments by passing them to
 205 // [cuda|hip]SetupArgument.
 206 // CUDA-OLD: call{{.*}}[[PREFIX]]SetupArgument
 207 // CUDA-OLD: call{{.*}}[[PREFIX]]SetupArgument
 208 // CUDA-OLD: call{{.*}}[[PREFIX]]SetupArgument
 209 // CUDA-OLD: call{{.*}}[[PREFIX]]Launch
 210
 211 // HIP: call{{.*}}[[PREFIX]]SetupArgument
 212 // HIP: call{{.*}}[[PREFIX]]SetupArgument
 213 // HIP: call{{.*}}[[PREFIX]]SetupArgument
 214 // HIP: call{{.*}}[[PREFIX]]Launch
 215 __global__ void kernelfunc(int i, int j, int k) {}
 216
 217 // Test that we've built correct kernel launch sequence.
 218 // LNX: define{{.*}}hostfunc
 219 // CUDA-OLD: call{{.*}}[[PREFIX]]ConfigureCall
 220 // CUDA-NEW: call{{.*}}__cudaPushCallConfiguration
 221 // HIP: call{{.*}}[[PREFIX]]ConfigureCall
 222 // LNX: call{{.*}}kernelfunc
 223 void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
 224 #endif
 225
 226 // Test that we've built a function to register kernels and global vars.
 227 // ALL: define internal void @__[[PREFIX]]_register_globals
 228 // ALL: call{{.*}}[[PREFIX]]RegisterFunction(i8** %0, {{.*}}kernelfunc{{[^,]*}}, {{[^@]*}}@0
 229 // ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}device_var{{[^,]*}}, {{[^@]*}}@1, {{.*}}i32 0, {{i32|i64}} 4, i32 0, i32 0
 230 // ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}constant_var{{[^,]*}}, {{[^@]*}}@2, {{.*}}i32 0, {{i32|i64}} 4, i32 1, i32 0
 231 // ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_device_var_def{{[^,]*}}, {{[^@]*}}@3, {{.*}}i32 0, {{i32|i64}} 4, i32 0, i32 0
 232 // ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(i8** %0, {{.*}}ext_constant_var_def{{[^,]*}}, {{[^@]*}}@4, {{.*}}i32 0, {{i32|i64}} 4, i32 1, i32 0
 233 // LNX_17-DAG: [[PREFIX]]RegisterVar(i8** %0, {{.*}}inline_var
 234 // LNX_17-NOT: [[PREFIX]]RegisterVar(i8** %0, {{.*}}inline_var2
 235 // ALL: ret void
 236
 237 // Test that we've built a constructor.
 238 // LNX: define internal void @__[[PREFIX]]_module_ctor
 239
 240 // In separate mode it calls __[[PREFIX]]RegisterFatBinary(&__[[PREFIX]]_fatbin_wrapper)
 241 // HIP only register fat binary once.
 242 // HIP: load i8**, i8*** @__hip_gpubin_handle
 243 // HIP-NEXT: icmp eq i8** {{.*}}, null
 244 // HIP-NEXT: br i1 {{.*}}, label %if, label %exit
 245 // HIP: if:
 246 // CUDANORDC: call{{.*}}[[PREFIX]]RegisterFatBinary{{.*}}__[[PREFIX]]_fatbin_wrapper
 247 //   .. stores return value in __[[PREFIX]]_gpubin_handle
 248 // CUDANORDC-NEXT: store{{.*}}__[[PREFIX]]_gpubin_handle
 249 //   .. and then calls __[[PREFIX]]_register_globals
 250 // HIP: call{{.*}}[[PREFIX]]RegisterFatBinary{{.*}}__[[PREFIX]]_fatbin_wrapper
 251 //   .. stores return value in __[[PREFIX]]_gpubin_handle
 252 // HIP-NEXT: store{{.*}}__[[PREFIX]]_gpubin_handle
 253 //   .. and then calls __[[PREFIX]]_register_globals
 254 // HIP-NEXT: br label %exit
 255 // HIP: exit:
 256 // HIP-NEXT: load i8**, i8*** @__hip_gpubin_handle
 257 // CUDANORDC-NEXT: call void @__[[PREFIX]]_register_globals
 258 // HIP-NEXT: call void @__[[PREFIX]]_register_globals
 259 // * In separate mode we also register a destructor.
 260 // CUDANORDC-NEXT: call i32 @atexit(void ()* @__[[PREFIX]]_module_dtor)
 261 // HIP-NEXT: call i32 @atexit(void ()* @__[[PREFIX]]_module_dtor)
 262
 263 // With relocatable device code we call __[[PREFIX]]RegisterLinkedBinary%NVModuleID%
 264 // CUDARDC: call{{.*}}__[[PREFIX]]RegisterLinkedBinary[[MODULE_ID]](
 265 // CUDARDC-SAME: __[[PREFIX]]_register_globals, {{.*}}__[[PREFIX]]_fatbin_wrapper
 266 // CUDARDC-SAME: [[MODULE_ID_GLOBAL]]
 267
 268 // Test that we've created destructor.
 269 // CUDANORDC: define internal void @__[[PREFIX]]_module_dtor
 270 // HIP: define internal void @__[[PREFIX]]_module_dtor
 271 // CUDANORDC: load{{.*}}__[[PREFIX]]_gpubin_handle
 272 // HIP: load{{.*}}__[[PREFIX]]_gpubin_handle
 273 // CUDANORDC-NEXT: call void @__[[PREFIX]]UnregisterFatBinary
 274 // HIP-NEXT: icmp ne i8** {{.*}}, null
 275 // HIP-NEXT: br i1 {{.*}}, label %if, label %exit
 276 // HIP: if:
 277 // HIP-NEXT: call void @__[[PREFIX]]UnregisterFatBinary
 278 // HIP-NEXT: store i8** null, i8*** @__hip_gpubin_handle
 279 // HIP-NEXT: br label %exit
 280 // HIP: exit:
 281
 282 // There should be no __[[PREFIX]]_register_globals if we have no
 283 // device-side globals, but we still need to register GPU binary.
 284 // Skip GPU binary string first.
 285 // CUDANOGLOBALS-NOT: @{{.*}} = private constant{{.*}}
 286 // HIPNOGLOBALS-NOT: @{{.*}} = internal constant{{.*}}
 287 // NOGLOBALS-NOT: define internal void @__{{.*}}_register_globals
 288 // NOGLOBALS-NOT: define internal void @__{{cuda|hip}}_module_ctor
 289 // NOGLOBALS-NOT: call{{.*}}{{cuda|hip}}RegisterFatBinary{{.*}}__{{cuda|hip}}_fatbin_wrapper
 290 // NOGLOBALS-NOT: call void @__{{cuda|hip}}_register_globals
 291 // NOGLOBALS-NOT: define internal void @__{{cuda|hip}}_module_dtor
 292 // NOGLOBALS-NOT: call void @__{{cuda|hip}}UnregisterFatBinary
 293
 294 // There should be no constructors/destructors if we have no GPU binary.
 295 // NOGPUBIN-NOT: define internal void @__{{cuda|hip}}_register_globals
 296 // NOGPUBIN-NOT: define internal void @__{{cuda|hip}}_module_ctor
 297 // NOGPUBIN-NOT: define internal void @__{{cuda|hip}}_module_dtor