296 lines · plain
1// RUN: echo -n "GPU binary would be here." > %t2// RUN: %clang_cc1 -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-OLD6// RUN: %clang_cc1 -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,CUDANOGLOBALS11// RUN: %clang_cc1 -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-OLD16// RUN: %clang_cc1 -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=NOGPUBIN19 20// RUN: %clang_cc1 -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-NEW24// RUN: %clang_cc1 -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,CUDANOGLOBALS28// RUN: %clang_cc1 -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-NEW32// RUN: %clang_cc1 -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,NORDC1736// RUN: %clang_cc1 -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,RDC1740// RUN: %clang_cc1 -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=NOGPUBIN43 44// RUN: %clang_cc1 -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,HIPEF47// RUN: %clang_cc1 -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,HIPNOGLOBALS50// RUN: %clang_cc1 -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,HIPEF53// RUN: %clang_cc1 -cuid=123 -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,HIPNEF55 56// RUN: %clang_cc1 -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,WIN59 60// RUN: %clang_cc1 -cuid=123 -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,HIPNEF63 64#include "Inputs/cuda.h"65 66#ifndef NOGLOBALS67// NORDC-DAG: @device_var = internal global i3268// RDC-DAG: @device_var = global i3269// WIN-DAG: @"?device_var@@3HA" = internal global i3270__device__ int device_var;71 72// NORDC-DAG: @constant_var = internal global i3273// RDC-DAG: @constant_var = global i3274// WIN-DAG: @"?constant_var@@3HA" = internal global i3275__constant__ int constant_var;76 77// NORDC-DAG: @shared_var = internal global i3278// RDC-DAG: @shared_var = global i3279// WIN-DAG: @"?shared_var@@3HA" = internal global i3280__shared__ int shared_var;81 82// Make sure host globals don't get internalized...83// LNX-DAG: @host_var ={{.*}} global i3284// WIN-DAG: @"?host_var@@3HA" = dso_local global i3285int host_var;86// ... and that extern vars remain external.87// LNX-DAG: @ext_host_var = external global i3288// WIN-DAG: @"?ext_host_var@@3HA" = external dso_local global i3289extern int ext_host_var;90 91// external device-side variables -> extern references to their shadows.92// LNX-DAG: @ext_device_var = external global i3293// WIN-DAG: @"?ext_device_var@@3HA" = external dso_local global i3294extern __device__ int ext_device_var;95// LNX-DAG: @ext_device_var = external global i3296// WIN-DAG: @"?ext_constant_var@@3HA" = external dso_local global i3297extern __constant__ int ext_constant_var;98 99// external device-side variables with definitions should generate100// definitions for the shadows.101// NORDC-DAG: @ext_device_var_def = internal global i32 undef,102// RDC-DAG: @ext_device_var_def = global i32 undef,103// WIN-DAG: @"?ext_device_var_def@@3HA" = internal global i32 undef104extern __device__ int ext_device_var_def;105__device__ int ext_device_var_def = 1;106// NORDC-DAG: @ext_device_var_def = internal global i32 undef,107// RDC-DAG: @ext_device_var_def = global i32 undef,108// WIN-DAG: @"?ext_constant_var_def@@3HA" = internal global i32 undef109__constant__ int ext_constant_var_def = 2;110 111#if __cplusplus > 201402L112// NORDC17: @inline_var = internal global i32 undef, comdat, align 4{{$}}113// RDC17: @inline_var = linkonce_odr global i32 undef, comdat, align 4{{$}}114// NORDC17-NOT: @inline_var2 =115// RDC17-NOT: @inline_var2 =116// NORDC17: @_ZN1C17member_inline_varE = internal constant i32 undef, comdat, align 4{{$}}117// RDC17: @_ZN1C17member_inline_varE = linkonce_odr constant i32 undef, comdat, align 4{{$}}118// Check inline variable ODR-used by host is emitted on host and registered.119__device__ inline int inline_var = 3;120// Check inline variable not ODR-used by host is not emitted on host or registered.121__device__ inline int inline_var2 = 5;122struct C {123 __device__ static constexpr int member_inline_var = 4;124};125#endif126 127void use_pointers() {128 const int *p;129 p = &device_var;130 p = &constant_var;131 p = &shared_var;132 p = &host_var;133 p = &ext_device_var;134 p = &ext_constant_var;135 p = &ext_host_var;136#if __cplusplus > 201402L137 p = &inline_var;138 decltype(inline_var2) tmp;139 p = &C::member_inline_var;140#endif141}142 143__device__ void device_use() {144#if __cplusplus > 201402L145 const int *p = &inline_var2;146#endif147}148 149// Make sure that all parts of GPU code init/cleanup are there:150// * constant unnamed string with the device-side kernel name to be passed to151// __hipRegisterFunction/__cudaRegisterFunction.152// ALL: @0 = private unnamed_addr constant [18 x i8] c"_Z10kernelfunciii\00"153// * constant unnamed string with the device-side kernel name to be passed to154// __hipRegisterVar/__cudaRegisterVar.155// ALL: @1 = private unnamed_addr constant [11 x i8] c"device_var\00"156// ALL: @2 = private unnamed_addr constant [13 x i8] c"constant_var\00"157// ALL: @3 = private unnamed_addr constant [19 x i8] c"ext_device_var_def\00"158// ALL: @4 = private unnamed_addr constant [21 x i8] c"ext_constant_var_def\00"159// * constant unnamed string with GPU binary160// CUDA: @[[FATBIN:.*]] = private constant{{.*}} c"GPU binary would be here.",161// HIPEF: @[[FATBIN:.*]] = private constant{{.*}} c"GPU binary would be here.",{{.*}}align 4096162// HIPNEF: @[[FATBIN:__hip_fatbin_[0-9a-f]+]] = external constant i8, section ".hip_fatbin"163// CUDANORDC-SAME: section ".nv_fatbin", align 8164// CUDARDC-SAME: section "__nv_relfatbin", align 8165// * constant struct that wraps GPU binary166// ALL: @__[[PREFIX:cuda|hip]]_fatbin_wrapper = internal constant167// LNX-SAME: { i32, i32, ptr, ptr }168// CUDA-SAME: { i32 1180844977, i32 1,169// HIP-SAME: { i32 1212764230, i32 1,170// CUDA-SAME: ptr @[[FATBIN]],171// HIPEF-SAME: ptr @[[FATBIN]],172// HIPNEF-SAME: ptr @[[FATBIN]],173// LNX-SAME: ptr null }174// CUDA-SAME: section ".nvFatBinSegment"175// HIP-SAME: section ".hipFatBinSegment"176// * variable to save GPU binary handle after initialization177// CUDANORDC: @__[[PREFIX]]_gpubin_handle = internal global ptr null178// HIPNEF: @__[[PREFIX]]_gpubin_handle_{{[0-9a-f]+}} = internal global ptr null, align 8179// * constant unnamed string with NVModuleID180// CUDARDC: [[MODULE_ID_GLOBAL:@.*]] = private constant181// CUDARDC-SAME: c"[[MODULE_ID:.+]]\00", section "__nv_module_id", align 32182// * Make sure our constructor was added to global ctor list.183// LNX: @llvm.global_ctors = appending global {{.*}}@__[[PREFIX]]_module_ctor184// * Alias to global symbol containing the NVModuleID.185// CUDARDC: @__fatbinwrap[[MODULE_ID]] ={{.*}} alias { i32, i32, ptr, ptr }186// CUDARDC-SAME: ptr @__[[PREFIX]]_fatbin_wrapper187 188// Test that we build the correct number of calls to cudaSetupArgument followed189// by a call to cudaLaunch.190 191// LNX: define{{.*}}kernelfunc192 193// New launch sequence stores arguments into local buffer and passes array of194// pointers to them directly to cudaLaunchKernel195// CUDA-NEW: alloca196// CUDA-NEW: store197// CUDA-NEW: store198// CUDA-NEW: store199// CUDA-NEW: call{{.*}}__cudaPopCallConfiguration200// CUDA-NEW: call{{.*}}cudaLaunchKernel201 202// Legacy style launch sequence sets up arguments by passing them to203// [cuda|hip]SetupArgument.204// CUDA-OLD: call{{.*}}[[PREFIX]]SetupArgument205// CUDA-OLD: call{{.*}}[[PREFIX]]SetupArgument206// CUDA-OLD: call{{.*}}[[PREFIX]]SetupArgument207// CUDA-OLD: call{{.*}}[[PREFIX]]Launch208 209// HIP: call{{.*}}[[PREFIX]]SetupArgument210// HIP: call{{.*}}[[PREFIX]]SetupArgument211// HIP: call{{.*}}[[PREFIX]]SetupArgument212// HIP: call{{.*}}[[PREFIX]]Launch213__global__ void kernelfunc(int i, int j, int k) {}214 215// Test that we've built correct kernel launch sequence.216// LNX: define{{.*}}hostfunc217// CUDA-OLD: call{{.*}}[[PREFIX]]ConfigureCall218// CUDA-NEW: call{{.*}}__cudaPushCallConfiguration219// HIP: call{{.*}}[[PREFIX]]ConfigureCall220// LNX: call{{.*}}kernelfunc221void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }222#endif223 224// Test that we've built a function to register kernels and global vars.225// ALL: define internal void @__[[PREFIX]]_register_globals226// ALL: call{{.*}}[[PREFIX]]RegisterFunction(ptr %0, {{.*}}kernelfunc{{[^,]*}}, {{[^@]*}}@0227// ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(ptr %0, {{.*}}device_var{{[^,]*}}, {{[^@]*}}@1, {{.*}}i32 0, {{i32|i64}} 4, i32 0, i32 0228// ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(ptr %0, {{.*}}constant_var{{[^,]*}}, {{[^@]*}}@2, {{.*}}i32 0, {{i32|i64}} 4, i32 1, i32 0229// ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(ptr %0, {{.*}}ext_device_var_def{{[^,]*}}, {{[^@]*}}@3, {{.*}}i32 0, {{i32|i64}} 4, i32 0, i32 0230// ALL-DAG: call void {{.*}}[[PREFIX]]RegisterVar(ptr %0, {{.*}}ext_constant_var_def{{[^,]*}}, {{[^@]*}}@4, {{.*}}i32 0, {{i32|i64}} 4, i32 1, i32 0231// LNX_17-DAG: [[PREFIX]]RegisterVar(ptr %0, {{.*}}inline_var232// LNX_17-NOT: [[PREFIX]]RegisterVar(ptr %0, {{.*}}inline_var2233// ALL: ret void234 235// Test that we've built a constructor.236// LNX: define internal void @__[[PREFIX]]_module_ctor237 238// In separate mode it calls __[[PREFIX]]RegisterFatBinary(&__[[PREFIX]]_fatbin_wrapper)239// HIP only register fat binary once.240// HIP: load ptr, ptr @__hip_gpubin_handle241// HIP-NEXT: icmp eq ptr {{.*}}, null242// HIP-NEXT: br i1 {{.*}}, label %if, label %exit243// HIP: if:244// CUDANORDC: call{{.*}}[[PREFIX]]RegisterFatBinary{{.*}}__[[PREFIX]]_fatbin_wrapper245// .. stores return value in __[[PREFIX]]_gpubin_handle246// CUDANORDC-NEXT: store{{.*}}__[[PREFIX]]_gpubin_handle247// .. and then calls __[[PREFIX]]_register_globals248// HIP: call{{.*}}[[PREFIX]]RegisterFatBinary{{.*}}__[[PREFIX]]_fatbin_wrapper249// .. stores return value in __[[PREFIX]]_gpubin_handle250// HIP-NEXT: store{{.*}}__[[PREFIX]]_gpubin_handle251// .. and then calls __[[PREFIX]]_register_globals252// HIP-NEXT: br label %exit253// HIP: exit:254// HIP-NEXT: load ptr, ptr @__hip_gpubin_handle255// CUDANORDC-NEXT: call void @__[[PREFIX]]_register_globals256// HIP-NEXT: call void @__[[PREFIX]]_register_globals257// * In separate mode we also register a destructor.258// CUDANORDC-NEXT: call i32 @atexit(ptr @__[[PREFIX]]_module_dtor)259// HIP-NEXT: call i32 @atexit(ptr @__[[PREFIX]]_module_dtor)260 261// With relocatable device code we call __[[PREFIX]]RegisterLinkedBinary%NVModuleID%262// CUDARDC: call{{.*}}__[[PREFIX]]RegisterLinkedBinary[[MODULE_ID]](263// CUDARDC-SAME: __[[PREFIX]]_register_globals, {{.*}}__[[PREFIX]]_fatbin_wrapper264// CUDARDC-SAME: [[MODULE_ID_GLOBAL]]265 266// Test that we've created destructor.267// CUDANORDC: define internal void @__[[PREFIX]]_module_dtor268// HIP: define internal void @__[[PREFIX]]_module_dtor269// CUDANORDC: load{{.*}}__[[PREFIX]]_gpubin_handle270// HIP: load{{.*}}__[[PREFIX]]_gpubin_handle271// CUDANORDC-NEXT: call void @__[[PREFIX]]UnregisterFatBinary272// HIP-NEXT: icmp ne ptr {{.*}}, null273// HIP-NEXT: br i1 {{.*}}, label %if, label %exit274// HIP: if:275// HIP-NEXT: call void @__[[PREFIX]]UnregisterFatBinary276// HIP-NEXT: store ptr null, ptr @__hip_gpubin_handle277// HIP-NEXT: br label %exit278// HIP: exit:279 280// There should be no __[[PREFIX]]_register_globals if we have no281// device-side globals, but we still need to register GPU binary.282// Skip GPU binary string first.283// CUDANOGLOBALS-NOT: @{{.*}} = private constant{{.*}}284// HIPNOGLOBALS-NOT: @{{.*}} = internal constant{{.*}}285// NOGLOBALS-NOT: define internal void @__{{.*}}_register_globals286// NOGLOBALS-NOT: define internal void @__{{cuda|hip}}_module_ctor287// NOGLOBALS-NOT: call{{.*}}{{cuda|hip}}RegisterFatBinary{{.*}}__{{cuda|hip}}_fatbin_wrapper288// NOGLOBALS-NOT: call void @__{{cuda|hip}}_register_globals289// NOGLOBALS-NOT: define internal void @__{{cuda|hip}}_module_dtor290// NOGLOBALS-NOT: call void @__{{cuda|hip}}UnregisterFatBinary291 292// There should be no constructors/destructors if we have no GPU binary.293// NOGPUBIN-NOT: define internal void @__{{cuda|hip}}_register_globals294// NOGPUBIN-NOT: define internal void @__{{cuda|hip}}_module_ctor295// NOGPUBIN-NOT: define internal void @__{{cuda|hip}}_module_dtor296