| 1 | // Test for offload registration code for two targets |
|---|
| 2 | // RUN: %clang_cc1 -verify -fopenmp -x c -triple x86_64-unknown-linux-gnu -fopenmp-targets=x86_64-pc-linux-gnu,powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s |
|---|
| 3 | // expected-no-diagnostics |
|---|
| 4 | |
|---|
| 5 | void foo() { |
|---|
| 6 | #pragma omp target |
|---|
| 7 | {} |
|---|
| 8 | } |
|---|
| 9 | |
|---|
| 10 | // CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 } |
|---|
| 11 | // CHECK-DAG: [[DEVTY:%.+]] = type { i8*, i8*, [[ENTTY]]*, [[ENTTY]]* } |
|---|
| 12 | // CHECK-DAG: [[DSCTY:%.+]] = type { i32, [[DEVTY]]*, [[ENTTY]]*, [[ENTTY]]* } |
|---|
| 13 | |
|---|
| 14 | // Comdat key for the offload registration code. Should have sorted offload |
|---|
| 15 | // target triples encoded into the name. |
|---|
| 16 | // CHECK-DAG: $[[REGFN:\.omp_offloading\..+\.powerpc64le-ibm-linux-gnu\.x86_64-pc-linux-gnu+]] = comdat any |
|---|
| 17 | |
|---|
| 18 | // Check if offloading descriptor is created. |
|---|
| 19 | // CHECK: [[ENTBEGIN:@.+]] = external constant [[ENTTY]] |
|---|
| 20 | // CHECK: [[ENTEND:@.+]] = external constant [[ENTTY]] |
|---|
| 21 | // CHECK: [[DEV1BEGIN:@.+]] = extern_weak constant i8 |
|---|
| 22 | // CHECK: [[DEV1END:@.+]] = extern_weak constant i8 |
|---|
| 23 | // CHECK: [[DEV2BEGIN:@.+]] = extern_weak constant i8 |
|---|
| 24 | // CHECK: [[DEV2END:@.+]] = extern_weak constant i8 |
|---|
| 25 | // CHECK: [[IMAGES:@.+]] = internal unnamed_addr constant [2 x [[DEVTY]]] [{{.+}} { i8* [[DEV1BEGIN]], i8* [[DEV1END]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }, {{.+}} { i8* [[DEV2BEGIN]], i8* [[DEV2END]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }], comdat($[[REGFN]]) |
|---|
| 26 | // CHECK: [[DESC:@.+]] = internal constant [[DSCTY]] { i32 2, [[DEVTY]]* getelementptr inbounds ([2 x [[DEVTY]]], [2 x [[DEVTY]]]* [[IMAGES]], i32 0, i32 0), [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }, comdat($[[REGFN]]) |
|---|
| 27 | |
|---|
| 28 | // Check target registration is registered as a Ctor. |
|---|
| 29 | // CHECK: appending global [1 x { i32, void ()*, i8* }] [{ i32, void ()*, i8* } { i32 0, void ()* @[[REGFN]], i8* bitcast (void ()* @[[REGFN]] to i8*) }] |
|---|
| 30 | |
|---|
| 31 | // Check presence of foo() and the outlined target region |
|---|
| 32 | // CHECK: define void [[FOO:@.+]]() |
|---|
| 33 | // CHECK: define internal void [[OUTLINEDTARGET:@.+]]() |
|---|
| 34 | |
|---|
| 35 | // Check registration and unregistration code. |
|---|
| 36 | |
|---|
| 37 | // CHECK: define internal void @[[UNREGFN:.+]](i8*) |
|---|
| 38 | // CHECK-SAME: comdat($[[REGFN]]) { |
|---|
| 39 | // CHECK: call i32 @__tgt_unregister_lib([[DSCTY]]* [[DESC]]) |
|---|
| 40 | // CHECK: ret void |
|---|
| 41 | // CHECK: declare i32 @__tgt_unregister_lib([[DSCTY]]*) |
|---|
| 42 | |
|---|
| 43 | // CHECK: define linkonce hidden void @[[REGFN]]() |
|---|
| 44 | // CHECK-SAME: comdat { |
|---|
| 45 | // CHECK: call i32 @__tgt_register_lib([[DSCTY]]* [[DESC]]) |
|---|
| 46 | // CHECK: call i32 @__cxa_atexit(void (i8*)* @[[UNREGFN]], i8* bitcast ([[DSCTY]]* [[DESC]] to i8*), |
|---|
| 47 | // CHECK: ret void |
|---|
| 48 | // CHECK: declare i32 @__tgt_register_lib([[DSCTY]]*) |
|---|
| 49 | |
|---|
| 50 | |
|---|