1 // Test declare target link under unified memory requirement. 2 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK 3 // expected-no-diagnostics 4 5 #ifndef HEADER 6 #define HEADER 7 8 #define N 1000 9 10 double var = 10.0; 11 12 #pragma omp requires unified_shared_memory 13 #pragma omp declare target link(var) 14 15 int bar(int n){ 16 double sum = 0; 17 18 #pragma omp target 19 for(int i = 0; i < n; i++) { 20 sum += var; 21 } 22 23 return sum; 24 } 25 26 // CHECK: [[VAR:@.+]] = global double 1.000000e+01 27 // CHECK: [[VAR_DECL_TGT_LINK_PTR:@.+]] = global double* [[VAR]] 28 29 // CHECK: [[OFFLOAD_SIZES:@.+]] = private unnamed_addr constant [3 x i64] [i64 4, i64 8, i64 8] 30 // CHECK: [[OFFLOAD_MAPTYPES:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 531] 31 32 // CHECK: [[N_CASTED:%.+]] = alloca i64 33 // CHECK: [[SUM_CASTED:%.+]] = alloca i64 34 35 // CHECK: [[OFFLOAD_BASEPTRS:%.+]] = alloca [3 x i8*] 36 // CHECK: [[OFFLOAD_PTRS:%.+]] = alloca [3 x i8*] 37 38 // CHECK: [[LOAD1:%.+]] = load i64, i64* [[N_CASTED]] 39 // CHECK: [[LOAD2:%.+]] = load i64, i64* [[SUM_CASTED]] 40 41 // CHECK: [[BPTR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[OFFLOAD_BASEPTRS]], i32 0, i32 0 42 // CHECK: [[BCAST1:%.+]] = bitcast i8** [[BPTR1]] to i64* 43 // CHECK: store i64 [[LOAD1]], i64* [[BCAST1]] 44 // CHECK: [[BPTR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 0 45 // CHECK: [[BCAST2:%.+]] = bitcast i8** [[BPTR2]] to i64* 46 // CHECK: store i64 [[LOAD1]], i64* [[BCAST2]] 47 48 // CHECK: [[BPTR3:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[OFFLOAD_BASEPTRS]], i32 0, i32 1 49 // CHECK: [[BCAST3:%.+]] = bitcast i8** [[BPTR3]] to i64* 50 // CHECK: store i64 [[LOAD2]], i64* [[BCAST3]] 51 // CHECK: [[BPTR4:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 1 52 // CHECK: [[BCAST4:%.+]] = bitcast i8** [[BPTR4]] to i64* 53 // CHECK: store i64 [[LOAD2]], i64* [[BCAST4]] 54 55 // CHECK: [[BPTR5:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[OFFLOAD_BASEPTRS]], i32 0, i32 2 56 // CHECK: [[BCAST5:%.+]] = bitcast i8** [[BPTR5]] to double*** 57 // CHECK: store double** [[VAR_DECL_TGT_LINK_PTR]], double*** [[BCAST5]] 58 // CHECK: [[BPTR6:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 2 59 // CHECK: [[BCAST6:%.+]] = bitcast i8** [[BPTR6]] to double** 60 // CHECK: store double* [[VAR]], double** [[BCAST6]] 61 62 // CHECK: [[BPTR7:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[OFFLOAD_BASEPTRS]], i32 0, i32 0 63 // CHECK: [[BPTR8:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 0 64 65 // CHECK: call i32 @__tgt_target(i64 -1, i8* @{{.*}}.region_id, i32 3, i8** [[BPTR7]], i8** [[BPTR8]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[OFFLOAD_SIZES]], i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[OFFLOAD_MAPTYPES]], i32 0, i32 0)) 66 67 #endif 68