mirror of
https://github.com/llvm/llvm-project.git
synced 2025-04-28 03:46:06 +00:00
b8cbc5c02c
The KernelEnvironment is for compile time information about a kernel. It allows the compiler to feed information to the runtime. The KernelLaunchEnvironment is for dynamic information *per* kernel launch. It allows the rutime to feed information to the kernel that is not shared with other invocations of the kernel. The first use case is to replace the globals that synchronize teams reductions with per-launch versions. This allows concurrent teams reductions. More uses cases will follow, e.g., per launch memory pools. Fixes: https://github.com/llvm/llvm-project/issues/70249
140 lines
8.3 KiB
C++
140 lines
8.3 KiB
C++
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _
|
|
// Test device global memory data sharing codegen.
|
|
///==========================================================================///
|
|
|
|
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
|
|
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -disable-llvm-optzns | FileCheck %s --check-prefix=CHECK
|
|
|
|
// expected-no-diagnostics
|
|
|
|
#ifndef HEADER
|
|
#define HEADER
|
|
|
|
void test_ds(){
|
|
#pragma omp target
|
|
{
|
|
int a = 10;
|
|
#pragma omp parallel
|
|
{
|
|
a = 1000;
|
|
}
|
|
int b = 100;
|
|
int c = 1000;
|
|
#pragma omp parallel private(c)
|
|
{
|
|
int *c1 = &c;
|
|
b = a + 10000;
|
|
}
|
|
}
|
|
}
|
|
|
|
#endif
|
|
|
|
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z7test_dsv_l14
|
|
// CHECK-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0:[0-9]+]] {
|
|
// CHECK-NEXT: entry:
|
|
// CHECK-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8
|
|
// CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 8
|
|
// CHECK-NEXT: [[C:%.*]] = alloca i32, align 4
|
|
// CHECK-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [2 x ptr], align 8
|
|
// CHECK-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8
|
|
// CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z7test_dsv_l14_kernel_environment, ptr [[DYN_PTR]])
|
|
// CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
|
|
// CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
|
|
// CHECK: user_code.entry:
|
|
// CHECK-NEXT: [[A:%.*]] = call align 8 ptr @__kmpc_alloc_shared(i64 4)
|
|
// CHECK-NEXT: [[B:%.*]] = call align 8 ptr @__kmpc_alloc_shared(i64 4)
|
|
// CHECK-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])
|
|
// CHECK-NEXT: store i32 10, ptr [[A]], align 4
|
|
// CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 0
|
|
// CHECK-NEXT: store ptr [[A]], ptr [[TMP2]], align 8
|
|
// CHECK-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z7test_dsv_l14_omp_outlined, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z7test_dsv_l14_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS]], i64 1)
|
|
// CHECK-NEXT: store i32 100, ptr [[B]], align 4
|
|
// CHECK-NEXT: store i32 1000, ptr [[C]], align 4
|
|
// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS1]], i64 0, i64 0
|
|
// CHECK-NEXT: store ptr [[B]], ptr [[TMP3]], align 8
|
|
// CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS1]], i64 0, i64 1
|
|
// CHECK-NEXT: store ptr [[A]], ptr [[TMP4]], align 8
|
|
// CHECK-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z7test_dsv_l14_omp_outlined1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z7test_dsv_l14_omp_outlined1_wrapper, ptr [[CAPTURED_VARS_ADDRS1]], i64 2)
|
|
// CHECK-NEXT: call void @__kmpc_free_shared(ptr [[B]], i64 4)
|
|
// CHECK-NEXT: call void @__kmpc_free_shared(ptr [[A]], i64 4)
|
|
// CHECK-NEXT: call void @__kmpc_target_deinit()
|
|
// CHECK-NEXT: ret void
|
|
// CHECK: worker.exit:
|
|
// CHECK-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z7test_dsv_l14_omp_outlined
|
|
// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2:[0-9]+]] {
|
|
// CHECK-NEXT: entry:
|
|
// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
|
|
// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
|
|
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
|
|
// CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
|
|
// CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
|
|
// CHECK-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
|
|
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
|
|
// CHECK-NEXT: store i32 1000, ptr [[TMP0]], align 4
|
|
// CHECK-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z7test_dsv_l14_omp_outlined_wrapper
|
|
// CHECK-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR3:[0-9]+]] {
|
|
// CHECK-NEXT: entry:
|
|
// CHECK-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
|
|
// CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
|
|
// CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
|
|
// CHECK-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 8
|
|
// CHECK-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2
|
|
// CHECK-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4
|
|
// CHECK-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4
|
|
// CHECK-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])
|
|
// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GLOBAL_ARGS]], align 8
|
|
// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 0
|
|
// CHECK-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP3]], align 8
|
|
// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z7test_dsv_l14_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]], ptr [[TMP4]]) #[[ATTR4:[0-9]+]]
|
|
// CHECK-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z7test_dsv_l14_omp_outlined1
|
|
// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[B:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR2]] {
|
|
// CHECK-NEXT: entry:
|
|
// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
|
|
// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
|
|
// CHECK-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 8
|
|
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
|
|
// CHECK-NEXT: [[C:%.*]] = alloca i32, align 4
|
|
// CHECK-NEXT: [[C1:%.*]] = alloca ptr, align 8
|
|
// CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
|
|
// CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
|
|
// CHECK-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 8
|
|
// CHECK-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
|
|
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8
|
|
// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A_ADDR]], align 8
|
|
// CHECK-NEXT: store ptr [[C]], ptr [[C1]], align 8
|
|
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4
|
|
// CHECK-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP2]], 10000
|
|
// CHECK-NEXT: store i32 [[ADD]], ptr [[TMP0]], align 4
|
|
// CHECK-NEXT: ret void
|
|
//
|
|
//
|
|
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z7test_dsv_l14_omp_outlined1_wrapper
|
|
// CHECK-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR3]] {
|
|
// CHECK-NEXT: entry:
|
|
// CHECK-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
|
|
// CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
|
|
// CHECK-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
|
|
// CHECK-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 8
|
|
// CHECK-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2
|
|
// CHECK-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4
|
|
// CHECK-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4
|
|
// CHECK-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])
|
|
// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GLOBAL_ARGS]], align 8
|
|
// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 0
|
|
// CHECK-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP3]], align 8
|
|
// CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 1
|
|
// CHECK-NEXT: [[TMP6:%.*]] = load ptr, ptr [[TMP5]], align 8
|
|
// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z7test_dsv_l14_omp_outlined1(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]], ptr [[TMP4]], ptr [[TMP6]]) #[[ATTR4]]
|
|
// CHECK-NEXT: ret void
|
|
//
|