| // RUN: echo "GPU binary would be here" > %t |
| // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -fcuda-include-gpubinary %t -o - | FileCheck %s |
| // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -fcuda-include-gpubinary %t -o - -DNOGLOBALS \ |
| // RUN: | FileCheck %s -check-prefix=NOGLOBALS |
| // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - | FileCheck %s -check-prefix=NOGPUBIN |
| |
| #include "Inputs/cuda.h" |
| |
| #ifndef NOGLOBALS |
| // CHECK-DAG: @device_var = internal global i32 |
| __device__ int device_var; |
| |
| // CHECK-DAG: @constant_var = internal global i32 |
| __constant__ int constant_var; |
| |
| // CHECK-DAG: @shared_var = internal global i32 |
| __shared__ int shared_var; |
| |
| // Make sure host globals don't get internalized... |
| // CHECK-DAG: @host_var = global i32 |
| int host_var; |
| // ... and that extern vars remain external. |
| // CHECK-DAG: @ext_host_var = external global i32 |
| extern int ext_host_var; |
| |
| // Shadows for external device-side variables are *definitions* of |
| // those variables. |
| // CHECK-DAG: @ext_device_var = internal global i32 |
| extern __device__ int ext_device_var; |
| // CHECK-DAG: @ext_device_var = internal global i32 |
| extern __constant__ int ext_constant_var; |
| |
| void use_pointers() { |
| int *p; |
| p = &device_var; |
| p = &constant_var; |
| p = &shared_var; |
| p = &host_var; |
| p = &ext_device_var; |
| p = &ext_constant_var; |
| p = &ext_host_var; |
| } |
| |
| // Make sure that all parts of GPU code init/cleanup are there: |
| // * constant unnamed string with the kernel name |
| // CHECK: private unnamed_addr constant{{.*}}kernelfunc{{.*}}\00" |
| // * constant unnamed string with GPU binary |
| // CHECK: private unnamed_addr constant{{.*GPU binary would be here.*}}\00" |
| // CHECK-SAME: section ".nv_fatbin", align 8 |
| // * constant struct that wraps GPU binary |
| // CHECK: @__cuda_fatbin_wrapper = internal constant { i32, i32, i8*, i8* } |
| // CHECK-SAME: { i32 1180844977, i32 1, {{.*}}, i8* null } |
| // CHECK-SAME: section ".nvFatBinSegment" |
| // * variable to save GPU binary handle after initialization |
| // CHECK: @__cuda_gpubin_handle = internal global i8** null |
| // * Make sure our constructor/destructor was added to global ctor/dtor list. |
| // CHECK: @llvm.global_ctors = appending global {{.*}}@__cuda_module_ctor |
| // CHECK: @llvm.global_dtors = appending global {{.*}}@__cuda_module_dtor |
| |
| // Test that we build the correct number of calls to cudaSetupArgument followed |
| // by a call to cudaLaunch. |
| |
| // CHECK: define{{.*}}kernelfunc |
| // CHECK: call{{.*}}cudaSetupArgument |
| // CHECK: call{{.*}}cudaSetupArgument |
| // CHECK: call{{.*}}cudaSetupArgument |
| // CHECK: call{{.*}}cudaLaunch |
| __global__ void kernelfunc(int i, int j, int k) {} |
| |
| // Test that we've built correct kernel launch sequence. |
| // CHECK: define{{.*}}hostfunc |
| // CHECK: call{{.*}}cudaConfigureCall |
| // CHECK: call{{.*}}kernelfunc |
| void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); } |
| #endif |
| |
| // Test that we've built a function to register kernels and global vars. |
| // CHECK: define internal void @__cuda_register_globals |
| // CHECK: call{{.*}}cudaRegisterFunction(i8** %0, {{.*}}kernelfunc |
| // CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}device_var{{.*}}i32 0, i32 4, i32 0, i32 0 |
| // CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}constant_var{{.*}}i32 0, i32 4, i32 1, i32 0 |
| // CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}ext_device_var{{.*}}i32 1, i32 4, i32 0, i32 0 |
| // CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}ext_constant_var{{.*}}i32 1, i32 4, i32 1, i32 0 |
| // CHECK: ret void |
| |
| // Test that we've built contructor.. |
| // CHECK: define internal void @__cuda_module_ctor |
| // .. that calls __cudaRegisterFatBinary(&__cuda_fatbin_wrapper) |
| // CHECK: call{{.*}}cudaRegisterFatBinary{{.*}}__cuda_fatbin_wrapper |
| // .. stores return value in __cuda_gpubin_handle |
| // CHECK-NEXT: store{{.*}}__cuda_gpubin_handle |
| // .. and then calls __cuda_register_globals |
| // CHECK-NEXT: call void @__cuda_register_globals |
| |
| // Test that we've created destructor. |
| // CHECK: define internal void @__cuda_module_dtor |
| // CHECK: load{{.*}}__cuda_gpubin_handle |
| // CHECK-NEXT: call void @__cudaUnregisterFatBinary |
| |
| // There should be no __cuda_register_globals if we have no |
| // device-side globals, but we still need to register GPU binary. |
| // Skip GPU binary string first. |
| // NOGLOBALS: @0 = private unnamed_addr constant{{.*}} |
| // NOGLOBALS-NOT: define internal void @__cuda_register_globals |
| // NOGLOBALS: define internal void @__cuda_module_ctor |
| // NOGLOBALS: call{{.*}}cudaRegisterFatBinary{{.*}}__cuda_fatbin_wrapper |
| // NOGLOBALS-NOT: call void @__cuda_register_globals |
| // NOGLOBALS: define internal void @__cuda_module_dtor |
| // NOGLOBALS: call void @__cudaUnregisterFatBinary |
| |
| // There should be no constructors/destructors if we have no GPU binary. |
| // NOGPUBIN-NOT: define internal void @__cuda_register_globals |
| // NOGPUBIN-NOT: define internal void @__cuda_module_ctor |
| // NOGPUBIN-NOT: define internal void @__cuda_module_dtor |
