1 // Test declare target link under unified memory requirement. 2 3 // RUN: %clang_cc1 -no-opaque-pointers -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-HOST 4 5 // RUN: %clang_cc1 -no-opaque-pointers -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc 6 // RUN: %clang_cc1 -no-opaque-pointers -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -target-cpu sm_70 -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CHECK-DEVICE 7 8 // Test declare target link under unified memory requirement. 9 10 // RUN: %clang_cc1 -no-opaque-pointers -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-HOST 11 12 // RUN: %clang_cc1 -no-opaque-pointers -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc 13 // RUN: %clang_cc1 -no-opaque-pointers -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -target-cpu sm_70 -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CHECK-DEVICE 14 15 // expected-no-diagnostics 16 17 #ifndef HEADER 18 #define HEADER 19 20 #define N 1000 21 22 double var = 10.0; 23 double to_var = 20.0; 24 25 #pragma omp requires unified_shared_memory 26 #pragma omp declare target link(var) 27 #pragma omp declare target to(to_var) 28 29 int bar(int n){ 30 double sum = 0; 31 32 #pragma omp target 33 for(int i = 0; i < n; i++) { 34 sum += var + to_var; 35 } 36 37 return sum; 38 } 39 40 // CHECK-HOST: [[VAR:@.+]] ={{.*}} global double 1.000000e+01 41 // CHECK-HOST: [[VAR_DECL_TGT_LINK_PTR:@.+]] = weak{{.*}} global double* [[VAR]] 42 43 // CHECK-HOST: [[TO_VAR:@.+]] ={{.*}} global double 2.000000e+01 44 // CHECK-HOST: [[VAR_DECL_TGT_TO_PTR:@.+]] = weak{{.*}} global double* [[TO_VAR]] 45 46 // CHECK-HOST: [[OFFLOAD_SIZES:@.+]] = private unnamed_addr constant [2 x i64] [i64 4, i64 8] 47 // CHECK-HOST: [[OFFLOAD_MAPTYPES:@.+]] = private unnamed_addr constant [2 x i64] [i64 800, i64 800] 48 49 // CHECK-HOST: [[OMP_OFFLOAD_ENTRY_LINK_VAR_PTR_NAME:@.+]] = internal unnamed_addr constant [21 x i8] 50 // CHECK-HOST: [[OMP_OFFLOAD_ENTRY_LINK_VAR_PTR:@.+]] = weak{{.*}} constant %struct.__tgt_offload_entry { i8* bitcast (double** [[VAR_DECL_TGT_LINK_PTR]] to i8*), i8* getelementptr inbounds ([21 x i8], [21 x i8]* [[OMP_OFFLOAD_ENTRY_LINK_VAR_PTR_NAME]], i32 0, i32 0), i64 8, i32 1, i32 0 }, section "omp_offloading_entries" 51 52 // CHECK-HOST: [[OMP_OFFLOAD_ENTRY_TO_VAR_PTR_NAME:@.+]] = internal unnamed_addr constant [24 x i8] 53 // CHECK-HOST: [[OMP_OFFLOAD_ENTRY_TO_VAR_PTR:@.+]] = weak{{.*}} constant %struct.__tgt_offload_entry { i8* bitcast (double** [[VAR_DECL_TGT_TO_PTR]] to i8*), i8* getelementptr inbounds ([24 x i8], [24 x i8]* [[OMP_OFFLOAD_ENTRY_TO_VAR_PTR_NAME]], i32 0, i32 0), i64 8, i32 0, i32 0 }, section "omp_offloading_entries" 54 55 // CHECK-HOST: [[N_CASTED:%.+]] = alloca i64 56 // CHECK-HOST: [[SUM_CASTED:%.+]] = alloca i64 57 58 // CHECK-HOST: [[OFFLOAD_BASEPTRS:%.+]] = alloca [2 x i8*] 59 // CHECK-HOST: [[OFFLOAD_PTRS:%.+]] = alloca [2 x i8*] 60 61 // CHECK-HOST: [[LOAD1:%.+]] = load i64, i64* [[N_CASTED]] 62 // CHECK-HOST: [[LOAD2:%.+]] = load i64, i64* [[SUM_CASTED]] 63 64 // CHECK-HOST: [[BPTR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[OFFLOAD_BASEPTRS]], i32 0, i32 0 65 // CHECK-HOST: [[BCAST1:%.+]] = bitcast i8** [[BPTR1]] to i64* 66 // CHECK-HOST: store i64 [[LOAD1]], i64* [[BCAST1]] 67 // CHECK-HOST: [[BPTR2:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 0 68 // CHECK-HOST: [[BCAST2:%.+]] = bitcast i8** [[BPTR2]] to i64* 69 // CHECK-HOST: store i64 [[LOAD1]], i64* [[BCAST2]] 70 71 // CHECK-HOST: [[BPTR3:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[OFFLOAD_BASEPTRS]], i32 0, i32 1 72 // CHECK-HOST: [[BCAST3:%.+]] = bitcast i8** [[BPTR3]] to i64* 73 // CHECK-HOST: store i64 [[LOAD2]], i64* [[BCAST3]] 74 // CHECK-HOST: [[BPTR4:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 1 75 // CHECK-HOST: [[BCAST4:%.+]] = bitcast i8** [[BPTR4]] to i64* 76 // CHECK-HOST: store i64 [[LOAD2]], i64* [[BCAST4]] 77 78 // CHECK-HOST: [[BPTR7:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[OFFLOAD_BASEPTRS]], i32 0, i32 0 79 // CHECK-HOST: [[BPTR8:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 0 80 81 // CHECK-HOST: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* %{{.+}}) 82 83 // CHECK-DEVICE: [[VAR_LINK:@.+]] = weak{{.*}} global double* null 84 // CHECK-DEVICE: [[VAR_TO:@.+]] = weak{{.*}} global double* null 85 86 // CHECK-DEVICE: [[VAR_TO_PTR:%.+]] = load double*, double** [[VAR_TO]] 87 // CHECK-DEVICE: load double, double* [[VAR_TO_PTR]] 88 89 #endif 90