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 115