86 lines · cpp
1// Test declare target link under unified memory requirement.2 3// 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-HOST4 5// RUN: %clang_cc1 -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.bc6// RUN: %clang_cc1 -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-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CHECK-DEVICE7 8// Test declare target link under unified memory requirement.9 10// 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-HOST11 12// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc13// RUN: %clang_cc1 -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-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CHECK-DEVICE14 15// expected-no-diagnostics16 17#ifndef HEADER18#define HEADER19 20#define N 100021 22double var = 10.0;23double to_var = 20.0;24 25#pragma omp requires unified_shared_memory26#pragma omp declare target link(var)27#pragma omp declare target to(to_var)28 29int bar(int n){30 double sum = 0;31 32#pragma omp target33 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+0141// CHECK-HOST: [[VAR_DECL_TGT_LINK_PTR:@.+]] = weak{{.*}} global ptr [[VAR]]42 43// CHECK-HOST: [[TO_VAR:@.+]] ={{.*}} global double 2.000000e+0144// CHECK-HOST: [[VAR_DECL_TGT_TO_PTR:@.+]] = weak{{.*}} global ptr [[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 { i64 0, i16 1, i16 1, i32 1, ptr [[VAR_DECL_TGT_LINK_PTR]], ptr [[OMP_OFFLOAD_ENTRY_LINK_VAR_PTR_NAME]], i64 8, i64 0, ptr null }, section "llvm_offload_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 { i64 0, i16 1, i16 1, i32 0, ptr [[VAR_DECL_TGT_TO_PTR]], ptr @.offloading.entry_name.1, i64 8, i64 0, ptr null }, section "llvm_offload_entries"54 55// CHECK-HOST: [[N_CASTED:%.+]] = alloca i6456// CHECK-HOST: [[SUM_CASTED:%.+]] = alloca i6457 58// CHECK-HOST: [[OFFLOAD_BASEPTRS:%.+]] = alloca [2 x ptr]59// CHECK-HOST: [[OFFLOAD_PTRS:%.+]] = alloca [2 x ptr]60 61// CHECK-HOST: [[LOAD1:%.+]] = load i64, ptr [[N_CASTED]]62// CHECK-HOST: [[LOAD2:%.+]] = load i64, ptr [[SUM_CASTED]]63 64// CHECK-HOST: [[BPTR1:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOAD_BASEPTRS]], i32 0, i32 065// CHECK-HOST: store i64 [[LOAD1]], ptr [[BPTR1]]66// CHECK-HOST: [[BPTR2:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOAD_PTRS]], i32 0, i32 067// CHECK-HOST: store i64 [[LOAD1]], ptr [[BPTR2]]68 69// CHECK-HOST: [[BPTR3:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOAD_BASEPTRS]], i32 0, i32 170// CHECK-HOST: store i64 [[LOAD2]], ptr [[BPTR3]]71// CHECK-HOST: [[BPTR4:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOAD_PTRS]], i32 0, i32 172// CHECK-HOST: store i64 [[LOAD2]], ptr [[BPTR4]]73 74// CHECK-HOST: [[BPTR7:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOAD_BASEPTRS]], i32 0, i32 075// CHECK-HOST: [[BPTR8:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOAD_PTRS]], i32 0, i32 076 77// CHECK-HOST: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr %{{.+}})78 79// CHECK-DEVICE: [[VAR_LINK:@.+]] = weak{{.*}} global ptr null80// CHECK-DEVICE: [[VAR_TO:@.+]] = weak{{.*}} global ptr null81 82// CHECK-DEVICE: [[VAR_TO_PTR:%.+]] = load ptr, ptr [[VAR_TO]]83// CHECK-DEVICE: load double, ptr [[VAR_TO_PTR]]84 85#endif86