1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals 2 // RUN: %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu \ 3 // RUN: --offload-new-driver -emit-llvm -o - -x cuda %s | FileCheck \ 4 // RUN: --check-prefix=HOST %s 5 6 #include "Inputs/cuda.h" 7 8 //. 9 // HOST: @x = internal global i32 undef, align 4 10 // HOST: @.omp_offloading.entry_name = internal unnamed_addr constant [8 x i8] c"_Z3foov\00" 11 // HOST: @.omp_offloading.entry._Z3foov = weak constant %struct.__tgt_offload_entry { ptr @_Z18__device_stub__foov, ptr @.omp_offloading.entry_name, i64 0, i32 0, i32 0 }, section "cuda_offloading_entries", align 1 12 // HOST: @.omp_offloading.entry_name.1 = internal unnamed_addr constant [8 x i8] c"_Z3barv\00" 13 // HOST: @.omp_offloading.entry._Z3barv = weak constant %struct.__tgt_offload_entry { ptr @_Z18__device_stub__barv, ptr @.omp_offloading.entry_name.1, i64 0, i32 0, i32 0 }, section "cuda_offloading_entries", align 1 14 // HOST: @.omp_offloading.entry_name.2 = internal unnamed_addr constant [2 x i8] c"x\00" 15 // HOST: @.omp_offloading.entry.x = weak constant %struct.__tgt_offload_entry { ptr @x, ptr @.omp_offloading.entry_name.2, i64 4, i32 0, i32 0 }, section "cuda_offloading_entries", align 1 16 //. 17 // HOST-LABEL: @_Z18__device_stub__foov( 18 // HOST-NEXT: entry: 19 // HOST-NEXT: [[TMP0:%.*]] = call i32 @cudaLaunch(ptr @_Z18__device_stub__foov) 20 // HOST-NEXT: br label [[SETUP_END:%.*]] 21 // HOST: setup.end: 22 // HOST-NEXT: ret void 23 // 24 __global__ void foo() {} 25 // HOST-LABEL: @_Z18__device_stub__barv( 26 // HOST-NEXT: entry: 27 // HOST-NEXT: [[TMP0:%.*]] = call i32 @cudaLaunch(ptr @_Z18__device_stub__barv) 28 // HOST-NEXT: br label [[SETUP_END:%.*]] 29 // HOST: setup.end: 30 // HOST-NEXT: ret void 31 // 32 __global__ void bar() {} 33 __device__ int x = 1; 34