This is an archive of the discontinued LLVM Phabricator instance.

Differential D90248

[libomptarget][nvptx] Undef, weak shared variables
ClosedPublic

Authored by JonChesterfield on Oct 27 2020, 9:50 AM.

Details

Reviewers
jdoerfert
ABataev
grokos
tianshilei1992
ye-luo
Commits
rG5d02ca49a294: [libomptarget][nvptx] Undef, weak shared variables
Summary

[libomptarget][nvptx] Undef, weak shared variables

Shared variables on nvptx, and LDS on amdgcn, are uninitialized at
the start of kernel execution. Therefore create the variables with
undef instead of zeros, motivated in part by the amdgcn back end
rejecting LDS+initializer.

Common is zero initialized, which seems incompatible with shared. Thus
change them to weak, following the direction of
https://reviews.llvm.org/rG7b3eabdcd215

Diff Detail

Unit TestsFailed

TimeTest
400 mslinux > HWAddressSanitizer-x86_64.TestCases::sizes.cpp

Event Timeline

JonChesterfield created this revision.Oct 27 2020, 9:50 AM
Herald added a project: Restricted Project. · View Herald TranscriptOct 27 2020, 9:50 AM
Herald added a subscriber: cfe-commits. · View Herald Transcript
JonChesterfield requested review of this revision.Oct 27 2020, 9:50 AM
JonChesterfield mentioned this in D89994: [libomptarget][nvptx] Undef, internal shared variables.Oct 27 2020, 9:51 AM
Harbormaster completed remote builds in B76598: Diff 301041.Oct 27 2020, 11:33 AM
jdoerfert accepted this revision.Oct 28 2020, 7:04 AM

LGTM, not totally convinced about weak linkage but ok

This revision is now accepted and ready to land.Oct 28 2020, 7:04 AM
This revision was landed with ongoing or failed builds.Oct 28 2020, 7:25 AM
Closed by commit rG5d02ca49a294: [libomptarget][nvptx] Undef, weak shared variables (authored by JonChesterfield). · Explain Why
This revision was automatically updated to reflect the committed changes.
JonChesterfield added a commit: rG5d02ca49a294: [libomptarget][nvptx] Undef, weak shared variables.
saiislam mentioned this in D86097: [OpenMP][AMDGCN] Generate global variables and attributes for AMDGCN.Nov 23 2020, 9:54 AM

Revision Contents

Diff 301041

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

Show First 20 Lines • Show All 1,096 Lines • ▼ Show 20 Linesvoid CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D,
CodeGen.setAction(Action); CodeGen.setAction(Action);
IsInTTDRegion = true; IsInTTDRegion = true;
// Reserve place for the globalized memory. // Reserve place for the globalized memory.
GlobalizedRecords.emplace_back(); GlobalizedRecords.emplace_back();
if (!KernelStaticGlobalized) { if (!KernelStaticGlobalized) {
KernelStaticGlobalized = new llvm::GlobalVariable( KernelStaticGlobalized = new llvm::GlobalVariable(
CGM.getModule(), CGM.VoidPtrTy, /*isConstant=*/false, CGM.getModule(), CGM.VoidPtrTy, /*isConstant=*/false,
llvm::GlobalValue::InternalLinkage, llvm::GlobalValue::InternalLinkage,
llvm::ConstantPointerNull::get(CGM.VoidPtrTy), llvm::UndefValue::get(CGM.VoidPtrTy),
"_openmp_kernel_static_glob_rd$ptr", /*InsertBefore=*/nullptr, "_openmp_kernel_static_glob_rd$ptr", /*InsertBefore=*/nullptr,
llvm::GlobalValue::NotThreadLocal, llvm::GlobalValue::NotThreadLocal,
CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared)); CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared));
} }
emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID, emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
IsOffloadEntry, CodeGen); IsOffloadEntry, CodeGen);
IsInTTDRegion = false; IsInTTDRegion = false;
▲ Show 20 Lines • Show All 115 Lines • ▼ Show 20 Linesvoid CGOpenMPRuntimeGPU::emitSPMDKernel(const OMPExecutableDirective &D,
CodeGen.setAction(Action); CodeGen.setAction(Action);
IsInTTDRegion = true; IsInTTDRegion = true;
// Reserve place for the globalized memory. // Reserve place for the globalized memory.
GlobalizedRecords.emplace_back(); GlobalizedRecords.emplace_back();
if (!KernelStaticGlobalized) { if (!KernelStaticGlobalized) {
KernelStaticGlobalized = new llvm::GlobalVariable( KernelStaticGlobalized = new llvm::GlobalVariable(
CGM.getModule(), CGM.VoidPtrTy, /*isConstant=*/false, CGM.getModule(), CGM.VoidPtrTy, /*isConstant=*/false,
llvm::GlobalValue::InternalLinkage, llvm::GlobalValue::InternalLinkage,
llvm::ConstantPointerNull::get(CGM.VoidPtrTy), llvm::UndefValue::get(CGM.VoidPtrTy),
"_openmp_kernel_static_glob_rd$ptr", /*InsertBefore=*/nullptr, "_openmp_kernel_static_glob_rd$ptr", /*InsertBefore=*/nullptr,
llvm::GlobalValue::NotThreadLocal, llvm::GlobalValue::NotThreadLocal,
CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared)); CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared));
} }
emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID, emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
IsOffloadEntry, CodeGen); IsOffloadEntry, CodeGen);
IsInTTDRegion = false; IsInTTDRegion = false;
} }
▲ Show 20 Lines • Show All 1,604 Lines • ▼ Show 20 Lines StringRef TransferMediumName =
"__openmp_nvptx_data_transfer_temporary_storage"; "__openmp_nvptx_data_transfer_temporary_storage";
llvm::GlobalVariable *TransferMedium = llvm::GlobalVariable *TransferMedium =
M.getGlobalVariable(TransferMediumName); M.getGlobalVariable(TransferMediumName);
unsigned WarpSize = CGF.getTarget().getGridValue(llvm::omp::GV_Warp_Size); unsigned WarpSize = CGF.getTarget().getGridValue(llvm::omp::GV_Warp_Size);
if (!TransferMedium) { if (!TransferMedium) {
auto *Ty = llvm::ArrayType::get(CGM.Int32Ty, WarpSize); auto *Ty = llvm::ArrayType::get(CGM.Int32Ty, WarpSize);
unsigned SharedAddressSpace = C.getTargetAddressSpace(LangAS::cuda_shared); unsigned SharedAddressSpace = C.getTargetAddressSpace(LangAS::cuda_shared);
TransferMedium = new llvm::GlobalVariable( TransferMedium = new llvm::GlobalVariable(
M, Ty, /*isConstant=*/false, llvm::GlobalVariable::CommonLinkage, M, Ty, /*isConstant=*/false, llvm::GlobalVariable::WeakAnyLinkage,
llvm::Constant::getNullValue(Ty), TransferMediumName, llvm::UndefValue::get(Ty), TransferMediumName,
/*InsertBefore=*/nullptr, llvm::GlobalVariable::NotThreadLocal, /*InsertBefore=*/nullptr, llvm::GlobalVariable::NotThreadLocal,
SharedAddressSpace); SharedAddressSpace);
CGM.addCompilerUsedGlobal(TransferMedium); CGM.addCompilerUsedGlobal(TransferMedium);
} }
auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime()); auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
// Get the CUDA thread id of the current OpenMP thread on the GPU. // Get the CUDA thread id of the current OpenMP thread on the GPU.
llvm::Value *ThreadID = RT.getGPUThreadID(CGF); llvm::Value *ThreadID = RT.getGPUThreadID(CGF);
▲ Show 20 Lines • Show All 1,918 Lines • ▼ Show 20 Lines if (!SharedStaticRD->field_empty()) {
SharedStaticRD->addDecl(Field); SharedStaticRD->addDecl(Field);
} }
SharedStaticRD->completeDefinition(); SharedStaticRD->completeDefinition();
if (!SharedStaticRD->field_empty()) { if (!SharedStaticRD->field_empty()) {
QualType StaticTy = C.getRecordType(SharedStaticRD); QualType StaticTy = C.getRecordType(SharedStaticRD);
llvm::Type *LLVMStaticTy = CGM.getTypes().ConvertTypeForMem(StaticTy); llvm::Type *LLVMStaticTy = CGM.getTypes().ConvertTypeForMem(StaticTy);
auto *GV = new llvm::GlobalVariable( auto *GV = new llvm::GlobalVariable(
CGM.getModule(), LLVMStaticTy, CGM.getModule(), LLVMStaticTy,
/*isConstant=*/false, llvm::GlobalValue::CommonLinkage, /*isConstant=*/false, llvm::GlobalValue::WeakAnyLinkage,
llvm::Constant::getNullValue(LLVMStaticTy), llvm::UndefValue::get(LLVMStaticTy),
"_openmp_shared_static_glob_rd_$_", /*InsertBefore=*/nullptr, "_openmp_shared_static_glob_rd_$_", /*InsertBefore=*/nullptr,
llvm::GlobalValue::NotThreadLocal, llvm::GlobalValue::NotThreadLocal,
C.getTargetAddressSpace(LangAS::cuda_shared)); C.getTargetAddressSpace(LangAS::cuda_shared));
auto *Replacement = llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast( auto *Replacement = llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(
GV, CGM.VoidPtrTy); GV, CGM.VoidPtrTy);
for (const GlobalPtrSizeRecsTy *Rec : SharedRecs) { for (const GlobalPtrSizeRecsTy *Rec : SharedRecs) {
Rec->Buffer->replaceAllUsesWith(Replacement); Rec->Buffer->replaceAllUsesWith(Replacement);
Rec->Buffer->eraseFromParent(); Rec->Buffer->eraseFromParent();
▲ Show 20 Lines • Show All 64 Lines • Show Last 20 Lines

clang/test/OpenMP/nvptx_data_sharing.cpp

Show All 22 Lines #pragma omp target
#pragma omp parallel private(c) #pragma omp parallel private(c)
{ {
int *c1 = &c; int *c1 = &c;
b = a + 10000; b = a + 10000;
} }
} }
} }
// SEQ: [[MEM_TY:%.+]] = type { [128 x i8] } // SEQ: [[MEM_TY:%.+]] = type { [128 x i8] }
// SEQ-DAG: [[SHARED_GLOBAL_RD:@.+]] = common addrspace(3) global [[MEM_TY]] zeroinitializer // SEQ-DAG: [[SHARED_GLOBAL_RD:@.+]] = weak addrspace(3) global [[MEM_TY]] undef
// SEQ-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* null // SEQ-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* undef
// SEQ-DAG: [[KERNEL_SIZE:@.+]] = internal unnamed_addr constant i64 8 // SEQ-DAG: [[KERNEL_SIZE:@.+]] = internal unnamed_addr constant i64 8
// SEQ-DAG: [[KERNEL_SHARED:@.+]] = internal unnamed_addr constant i16 1 // SEQ-DAG: [[KERNEL_SHARED:@.+]] = internal unnamed_addr constant i16 1
/// ========= In the worker function ========= /// /// ========= In the worker function ========= ///
// CK1: {{.*}}define internal void @__omp_offloading{{.*}}test_ds{{.*}}_worker() // CK1: {{.*}}define internal void @__omp_offloading{{.*}}test_ds{{.*}}_worker()
// CK1: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CK1: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
// CK1-NOT: call void @__kmpc_data_sharing_init_stack // CK1-NOT: call void @__kmpc_data_sharing_init_stack
▲ Show 20 Lines • Show All 79 Lines • Show Last 20 Lines

clang/test/OpenMP/nvptx_distribute_parallel_generic_mode_codegen.cpp

Show All 29 Lines
#pragma omp target teams map(tofrom:a) #pragma omp target teams map(tofrom:a)
#pragma omp distribute parallel for firstprivate(b) lastprivate(c) if(a) #pragma omp distribute parallel for firstprivate(b) lastprivate(c) if(a)
for (int i= 0; i < argc; ++i) for (int i= 0; i < argc; ++i)
a = foo(&i) + foo(&a) + foo(&b[i]) + foo(&c[i]) + foo(&d[i]); a = foo(&i) + foo(&a) + foo(&b[i]) + foo(&c[i]) + foo(&d[i]);
return 0; return 0;
} }
// SEQ: [[MEM_TY:%.+]] = type { [128 x i8] } // SEQ: [[MEM_TY:%.+]] = type { [128 x i8] }
// SEQ-DAG: [[SHARED_GLOBAL_RD:@.+]] = common addrspace(3) global [[MEM_TY]] zeroinitializer // SEQ-DAG: [[SHARED_GLOBAL_RD:@.+]] = weak addrspace(3) global [[MEM_TY]] undef
// SEQ-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* null // SEQ-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* undef
// SEQ-DAG: [[KERNEL_SIZE:@.+]] = internal unnamed_addr constant i{{64|32}} 40 // SEQ-DAG: [[KERNEL_SIZE:@.+]] = internal unnamed_addr constant i{{64|32}} 40
// SEQ-DAG: [[KERNEL_SHARED:@.+]] = internal unnamed_addr constant i16 1 // SEQ-DAG: [[KERNEL_SHARED:@.+]] = internal unnamed_addr constant i16 1
// CHECK-DAG: @__omp_offloading_{{.*}}_main_[[LINE:l.+]]_exec_mode = weak constant i8 0 // CHECK-DAG: @__omp_offloading_{{.*}}_main_[[LINE:l.+]]_exec_mode = weak constant i8 0
// CHECK: define weak void @__omp_offloading_{{.*}}_main_[[LINE]]([10 x i32]* nonnull align 4 dereferenceable(40) %{{.+}}, [10 x i32]* nonnull align 4 dereferenceable(40) %{{.+}}, i32* nonnull align 4 dereferenceable(4) %{{.+}}, i{{64|32}} %{{.+}}, [10 x i32]* nonnull align 4 dereferenceable(40) %{{.+}}) // CHECK: define weak void @__omp_offloading_{{.*}}_main_[[LINE]]([10 x i32]* nonnull align 4 dereferenceable(40) %{{.+}}, [10 x i32]* nonnull align 4 dereferenceable(40) %{{.+}}, i32* nonnull align 4 dereferenceable(4) %{{.+}}, i{{64|32}} %{{.+}}, [10 x i32]* nonnull align 4 dereferenceable(40) %{{.+}})
// SEQ: [[SHARED:%.+]] = load i16, i16* [[KERNEL_SHARED]], // SEQ: [[SHARED:%.+]] = load i16, i16* [[KERNEL_SHARED]],
// SEQ: [[SIZE:%.+]] = load i{{64|32}}, i{{64|32}}* [[KERNEL_SIZE]], // SEQ: [[SIZE:%.+]] = load i{{64|32}}, i{{64|32}}* [[KERNEL_SIZE]],
// SEQ: call void @__kmpc_get_team_static_memory(i16 1, i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([[MEM_TY]], [[MEM_TY]] addrspace(3)* [[SHARED_GLOBAL_RD]], i32 0, i32 0, i32 0) to i8*), i{{64|32}} [[SIZE]], i16 [[SHARED]], i8** addrspacecast (i8* addrspace(3)* [[KERNEL_PTR]] to i8**)) // SEQ: call void @__kmpc_get_team_static_memory(i16 1, i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([[MEM_TY]], [[MEM_TY]] addrspace(3)* [[SHARED_GLOBAL_RD]], i32 0, i32 0, i32 0) to i8*), i{{64|32}} [[SIZE]], i16 [[SHARED]], i8** addrspacecast (i8* addrspace(3)* [[KERNEL_PTR]] to i8**))
Show All 22 Lines

clang/test/OpenMP/nvptx_parallel_codegen.cpp

Show First 20 Lines • Show All 70 Lines • ▼ Show 20 Linesint bar(int n){
int a = 0; int a = 0;
a += ftemplate<int>(n); a += ftemplate<int>(n);
return a; return a;
} }
// SEQ: [[MEM_TY:%.+]] = type { [128 x i8] } // SEQ: [[MEM_TY:%.+]] = type { [128 x i8] }
// SEQ-DAG: [[SHARED_GLOBAL_RD:@.+]] = common addrspace(3) global [[MEM_TY]] zeroinitializer // SEQ-DAG: [[SHARED_GLOBAL_RD:@.+]] = weak addrspace(3) global [[MEM_TY]] undef
// SEQ-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* null // SEQ-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* undef
// SEQ-DAG: [[KERNEL_SIZE:@.+]] = internal unnamed_addr constant i{{64|32}} 4 // SEQ-DAG: [[KERNEL_SIZE:@.+]] = internal unnamed_addr constant i{{64|32}} 4
// SEQ-DAG: [[KERNEL_SHARED:@.+]] = internal unnamed_addr constant i16 1 // SEQ-DAG: [[KERNEL_SHARED:@.+]] = internal unnamed_addr constant i16 1
// CHECK-NOT: define {{.*}}void {{@__omp_offloading_.+template.+l20}}_worker() // CHECK-NOT: define {{.*}}void {{@__omp_offloading_.+template.+l20}}_worker()
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l29}}_worker() // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l29}}_worker()
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8, // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*, // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
▲ Show 20 Lines • Show All 294 Lines • Show Last 20 Lines

clang/test/OpenMP/nvptx_parallel_for_codegen.cpp

Show All 26 Linesint bar(int n){
int a = 0; int a = 0;
a += ftemplate<int>(n); a += ftemplate<int>(n);
return a; return a;
} }
// SEQ: [[MEM_TY:%.+]] = type { [128 x i8] } // SEQ: [[MEM_TY:%.+]] = type { [128 x i8] }
// SEQ-DAG: [[SHARED_GLOBAL_RD:@.+]] = common addrspace(3) global [[MEM_TY]] zeroinitializer // SEQ-DAG: [[SHARED_GLOBAL_RD:@.+]] = weak addrspace(3) global [[MEM_TY]] undef
// SEQ-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* null // SEQ-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* undef
// SEQ-DAG: [[KERNEL_SIZE:@.+]] = internal unnamed_addr constant i{{64|32}} 4 // SEQ-DAG: [[KERNEL_SIZE:@.+]] = internal unnamed_addr constant i{{64|32}} 4
// SEQ-DAG: [[KERNEL_SHARED:@.+]] = internal unnamed_addr constant i16 1 // SEQ-DAG: [[KERNEL_SHARED:@.+]] = internal unnamed_addr constant i16 1
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l13}}_worker() // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l13}}_worker()
// CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
// CHECK: call i1 @__kmpc_kernel_parallel( // CHECK: call i1 @__kmpc_kernel_parallel(
// CHECK: call void @__omp_outlined___wrapper( // CHECK: call void @__omp_outlined___wrapper(
▲ Show 20 Lines • Show All 93 Lines • Show Last 20 Lines

clang/test/OpenMP/nvptx_target_parallel_reduction_codegen.cpp

// Test target codegen - host bc file has to be created first. // Test target codegen - host bc file has to be created first.
// 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.bc // 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.bc
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc // RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
// expected-no-diagnostics // expected-no-diagnostics
#ifndef HEADER #ifndef HEADER
#define HEADER #define HEADER
// Check for the data transfer medium in shared memory to transfer the reduction list to the first warp. // Check for the data transfer medium in shared memory to transfer the reduction list to the first warp.
// CHECK-DAG: [[TRANSFER_STORAGE:@.+]] = common addrspace([[SHARED_ADDRSPACE:[0-9]+]]) global [32 x i32] // CHECK-DAG: [[TRANSFER_STORAGE:@.+]] = weak addrspace([[SHARED_ADDRSPACE:[0-9]+]]) global [32 x i32]
// Check that the execution mode of all 3 target regions is set to Spmd Mode. // Check that the execution mode of all 3 target regions is set to Spmd Mode.
// CHECK-DAG: {{@__omp_offloading_.+l27}}_exec_mode = weak constant i8 0 // CHECK-DAG: {{@__omp_offloading_.+l27}}_exec_mode = weak constant i8 0
// CHECK-DAG: {{@__omp_offloading_.+l32}}_exec_mode = weak constant i8 0 // CHECK-DAG: {{@__omp_offloading_.+l32}}_exec_mode = weak constant i8 0
// CHECK-DAG: {{@__omp_offloading_.+l38}}_exec_mode = weak constant i8 0 // CHECK-DAG: {{@__omp_offloading_.+l38}}_exec_mode = weak constant i8 0
template<typename tx> template<typename tx>
tx ftemplate(int n) { tx ftemplate(int n) {
▲ Show 20 Lines • Show All 787 Lines • Show Last 20 Lines

clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_codegen.cpp

Show First 20 Lines • Show All 87 Lines • ▼ Show 20 Linesint bar(int n){
int a = 0; int a = 0;
a += ftemplate<int>(n); a += ftemplate<int>(n);
return a; return a;
} }
// SEQ-DAG: [[MEM_TY:%.+]] = type { [128 x i8] } // SEQ-DAG: [[MEM_TY:%.+]] = type { [128 x i8] }
// SEQ-DAG: [[SHARED_GLOBAL_RD:@.+]] = common addrspace(3) global [[MEM_TY]] zeroinitializer // SEQ-DAG: [[SHARED_GLOBAL_RD:@.+]] = weak addrspace(3) global [[MEM_TY]] undef
// SEQ-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* null // SEQ-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* undef
// SEQ-DAG: [[KERNEL_SIZE:@.+]] = internal unnamed_addr constant i{{64|32}} 4 // SEQ-DAG: [[KERNEL_SIZE:@.+]] = internal unnamed_addr constant i{{64|32}} 4
// SEQ-DAG: [[KERNEL_SHARED:@.+]] = internal unnamed_addr constant i16 1 // SEQ-DAG: [[KERNEL_SHARED:@.+]] = internal unnamed_addr constant i16 1
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+}}_l50( // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+}}_l50(
// CHECK-DAG: [[THREAD_LIMIT:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() // CHECK-DAG: [[THREAD_LIMIT:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
// CHECK: call void @__kmpc_spmd_kernel_init(i32 [[THREAD_LIMIT]], i16 0) // CHECK: call void @__kmpc_spmd_kernel_init(i32 [[THREAD_LIMIT]], i16 0)
// CHECK: call void [[PARALLEL:@.+]]( // CHECK: call void [[PARALLEL:@.+]](
// CHECK: call void @__kmpc_spmd_kernel_deinit_v2(i16 0) // CHECK: call void @__kmpc_spmd_kernel_deinit_v2(i16 0)
▲ Show 20 Lines • Show All 159 Lines • Show Last 20 Lines

clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_simd_codegen.cpp

Show First 20 Lines • Show All 70 Lines • ▼ Show 20 Linesint bar(int n){
int a = 0; int a = 0;
a += ftemplate<int>(n); a += ftemplate<int>(n);
return a; return a;
} }
// SEQ-DAG: [[MEM_TY:%.+]] = type { [128 x i8] } // SEQ-DAG: [[MEM_TY:%.+]] = type { [128 x i8] }
// SEQ-DAG: [[SHARED_GLOBAL_RD:@.+]] = common addrspace(3) global [[MEM_TY]] zeroinitializer // SEQ-DAG: [[SHARED_GLOBAL_RD:@.+]] = weak addrspace(3) global [[MEM_TY]] undef
// SEQ-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* null // SEQ-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* undef
// SEQ-DAG: [[KERNEL_SIZE:@.+]] = internal unnamed_addr constant i{{64|32}} 4 // SEQ-DAG: [[KERNEL_SIZE:@.+]] = internal unnamed_addr constant i{{64|32}} 4
// SEQ-DAG: [[KERNEL_SHARED:@.+]] = internal unnamed_addr constant i16 1 // SEQ-DAG: [[KERNEL_SHARED:@.+]] = internal unnamed_addr constant i16 1
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+}}_l43( // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+}}_l43(
// CHECK-DAG: [[THREAD_LIMIT:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() // CHECK-DAG: [[THREAD_LIMIT:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
// CHECK: call void @__kmpc_spmd_kernel_init(i32 [[THREAD_LIMIT]], i16 0) // CHECK: call void @__kmpc_spmd_kernel_init(i32 [[THREAD_LIMIT]], i16 0)
// CHECK: call void @__kmpc_spmd_kernel_deinit_v2(i16 0) // CHECK: call void @__kmpc_spmd_kernel_deinit_v2(i16 0)
▲ Show 20 Lines • Show All 69 Lines • Show Last 20 Lines

clang/test/OpenMP/nvptx_teams_codegen.cpp

Show All 24 Lines
#pragma omp teams #pragma omp teams
{ {
argc = 0; argc = 0;
} }
return tmain(argv); return tmain(argv);
} }
// SEQ: [[MEM_TY:%.+]] = type { [128 x i8] } // SEQ: [[MEM_TY:%.+]] = type { [128 x i8] }
// SEQ-DAG: [[SHARED_GLOBAL_RD:@.+]] = common addrspace(3) global [[MEM_TY]] zeroinitializer // SEQ-DAG: [[SHARED_GLOBAL_RD:@.+]] = weak addrspace(3) global [[MEM_TY]] undef
// SEQ-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* null // SEQ-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* undef
// SEQ-DAG: [[KERNEL_SIZE1:@.+]] = internal unnamed_addr constant i{{64|32}} 4 // SEQ-DAG: [[KERNEL_SIZE1:@.+]] = internal unnamed_addr constant i{{64|32}} 4
// SEQ-DAG: [[KERNEL_SIZE2:@.+]] = internal unnamed_addr constant i{{64|32}} {{8|4}} // SEQ-DAG: [[KERNEL_SIZE2:@.+]] = internal unnamed_addr constant i{{64|32}} {{8|4}}
// SEQ-DAG: [[KERNEL_SHARED1:@.+]] = internal unnamed_addr constant i16 1 // SEQ-DAG: [[KERNEL_SHARED1:@.+]] = internal unnamed_addr constant i16 1
// SEQ-DAG: [[KERNEL_SHARED2:@.+]] = internal unnamed_addr constant i16 1 // SEQ-DAG: [[KERNEL_SHARED2:@.+]] = internal unnamed_addr constant i16 1
// only nvptx side: do not outline teams region and do not call fork_teams // only nvptx side: do not outline teams region and do not call fork_teams
// CK1: define {{.*}}void @{{[^,]+}}(i{{[0-9]+}} [[ARGC:%.+]]) // CK1: define {{.*}}void @{{[^,]+}}(i{{[0-9]+}} [[ARGC:%.+]])
// CK1: [[ARGCADDR:%.+]] = alloca i{{[0-9]+}}, // CK1: [[ARGCADDR:%.+]] = alloca i{{[0-9]+}},
▲ Show 20 Lines • Show All 71 Lines • ▼ Show 20 Lines
#pragma omp teams num_teams(a) thread_limit(b) #pragma omp teams num_teams(a) thread_limit(b)
{ {
argc = 0; argc = 0;
} }
return tmain(argv); return tmain(argv);
} }
// SEQ2: [[MEM_TY:%.+]] = type { [128 x i8] } // SEQ2: [[MEM_TY:%.+]] = type { [128 x i8] }
// SEQ2-DAG: [[SHARED_GLOBAL_RD:@.+]] = common addrspace(3) global [[MEM_TY]] zeroinitializer // SEQ2-DAG: [[SHARED_GLOBAL_RD:@.+]] = weak addrspace(3) global [[MEM_TY]] undef
// SEQ2-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* null // SEQ2-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* undef
// SEQ2-DAG: [[KERNEL_SIZE1:@.+]] = internal unnamed_addr constant i{{64|32}} 4 // SEQ2-DAG: [[KERNEL_SIZE1:@.+]] = internal unnamed_addr constant i{{64|32}} 4
// SEQ2-DAG: [[KERNEL_SIZE2:@.+]] = internal unnamed_addr constant i{{64|32}} {{8|4}} // SEQ2-DAG: [[KERNEL_SIZE2:@.+]] = internal unnamed_addr constant i{{64|32}} {{8|4}}
// SEQ2-DAG: [[KERNEL_SHARED1:@.+]] = internal unnamed_addr constant i16 1 // SEQ2-DAG: [[KERNEL_SHARED1:@.+]] = internal unnamed_addr constant i16 1
// SEQ2-DAG: [[KERNEL_SHARED2:@.+]] = internal unnamed_addr constant i16 1 // SEQ2-DAG: [[KERNEL_SHARED2:@.+]] = internal unnamed_addr constant i16 1
// CK2: define {{.*}}void @{{[^,]+}}(i{{[0-9]+}} [[A_IN:%.+]], i{{[0-9]+}} [[B_IN:%.+]], i{{[0-9]+}} [[ARGC_IN:.+]]) // CK2: define {{.*}}void @{{[^,]+}}(i{{[0-9]+}} [[A_IN:%.+]], i{{[0-9]+}} [[B_IN:%.+]], i{{[0-9]+}} [[ARGC_IN:.+]])
// CK2: [[AADDR:%.+]] = alloca i{{[0-9]+}}, // CK2: [[AADDR:%.+]] = alloca i{{[0-9]+}},
// CK2: [[BADDR:%.+]] = alloca i{{[0-9]+}}, // CK2: [[BADDR:%.+]] = alloca i{{[0-9]+}},
▲ Show 20 Lines • Show All 53 Lines • Show Last 20 Lines

clang/test/OpenMP/nvptx_teams_reduction_codegen.cpp

Show All 11 Lines
#define HEADER #define HEADER
// CHECK-DAG: [[TEAM1_REDUCE_TY:%.+]] = type { [{{1024|2048}} x double] } // CHECK-DAG: [[TEAM1_REDUCE_TY:%.+]] = type { [{{1024|2048}} x double] }
// CHECK-DAG: [[TEAM2_REDUCE_TY:%.+]] = type { [{{1024|2048}} x i8], [{{1024|2048}} x float] } // CHECK-DAG: [[TEAM2_REDUCE_TY:%.+]] = type { [{{1024|2048}} x i8], [{{1024|2048}} x float] }
// CHECK-DAG: [[TEAM3_REDUCE_TY:%.+]] = type { [{{1024|2048}} x i32], [{{1024|2048}} x i16] } // CHECK-DAG: [[TEAM3_REDUCE_TY:%.+]] = type { [{{1024|2048}} x i32], [{{1024|2048}} x i16] }
// CHECK-DAG: [[TEAMS_REDUCE_UNION_TY:%.+]] = type { [[TEAM1_REDUCE_TY]] } // CHECK-DAG: [[TEAMS_REDUCE_UNION_TY:%.+]] = type { [[TEAM1_REDUCE_TY]] }
// SEQ-DAG: [[MAP_TY:%.+]] = type { [128 x i8] } // SEQ-DAG: [[MAP_TY:%.+]] = type { [128 x i8] }
// SEQ-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* null // SEQ-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* undef
// SEQ-DAG: [[KERNEL_SHARED1:@.+]] = internal unnamed_addr constant i16 1 // SEQ-DAG: [[KERNEL_SHARED1:@.+]] = internal unnamed_addr constant i16 1
// SEQ-DAG: [[KERNEL_SHARED2:@.+]] = internal unnamed_addr constant i16 1 // SEQ-DAG: [[KERNEL_SHARED2:@.+]] = internal unnamed_addr constant i16 1
// SEQ-DAG: [[KERNEL_SIZE1:@.+]] = internal unnamed_addr constant i{{64|32}} {{16|8}} // SEQ-DAG: [[KERNEL_SIZE1:@.+]] = internal unnamed_addr constant i{{64|32}} {{16|8}}
// SEQ-DAG: [[KERNEL_SIZE2:@.+]] = internal unnamed_addr constant i{{64|32}} 16 // SEQ-DAG: [[KERNEL_SIZE2:@.+]] = internal unnamed_addr constant i{{64|32}} 16
// Check for the data transfer medium in shared memory to transfer the reduction list to the first warp. // Check for the data transfer medium in shared memory to transfer the reduction list to the first warp.
// CHECK-DAG: [[TRANSFER_STORAGE:@.+]] = common addrspace([[SHARED_ADDRSPACE:[0-9]+]]) global [32 x i32] // CHECK-DAG: [[TRANSFER_STORAGE:@.+]] = weak addrspace([[SHARED_ADDRSPACE:[0-9]+]]) global [32 x i32]
// Check that the execution mode of 2 target regions is set to Non-SPMD and the 3rd is in SPMD. // Check that the execution mode of 2 target regions is set to Non-SPMD and the 3rd is in SPMD.
// CHECK-DAG: {{@__omp_offloading_.+l44}}_exec_mode = weak constant i8 1 // CHECK-DAG: {{@__omp_offloading_.+l44}}_exec_mode = weak constant i8 1
// CHECK-DAG: {{@__omp_offloading_.+l50}}_exec_mode = weak constant i8 1 // CHECK-DAG: {{@__omp_offloading_.+l50}}_exec_mode = weak constant i8 1
// CHECK-DAG: {{@__omp_offloading_.+l57}}_exec_mode = weak constant i8 0 // CHECK-DAG: {{@__omp_offloading_.+l57}}_exec_mode = weak constant i8 0
// CHECK-DAG: [[TEAMS_RED_BUFFER:@.+]] = internal global [[TEAMS_REDUCE_UNION_TY]] zeroinitializer // CHECK-DAG: [[TEAMS_RED_BUFFER:@.+]] = internal global [[TEAMS_REDUCE_UNION_TY]] zeroinitializer
▲ Show 20 Lines • Show All 1,365 Lines • Show Last 20 Lines