From f69e6504d9b12b972343b6f4fe0ec7b1e11711ab Mon Sep 17 00:00:00 2001 From: Alexey Bataev Date: Thu, 3 Jan 2019 16:25:35 +0000 Subject: [PATCH] [OPENMP][NVPTX]Use __kmpc_barrier_simple_spmd(nullptr, 0) instead of nvvm_barrier0. Use runtime functions instead of the direct call to the nvvm intrinsics. It allows to prevent some dangerous LLVM optimizations, that breaks the code for the NVPTX target. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@350328 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp | 41 ++++++++++++----- lib/CodeGen/CGOpenMPRuntimeNVPTX.h | 3 ++ test/OpenMP/nvptx_data_sharing.cpp | 10 ++--- test/OpenMP/nvptx_parallel_codegen.cpp | 24 +++++----- test/OpenMP/nvptx_parallel_for_codegen.cpp | 6 +-- test/OpenMP/nvptx_target_codegen.cpp | 44 +++++++++---------- test/OpenMP/nvptx_target_teams_codegen.cpp | 12 ++--- .../nvptx_target_teams_distribute_codegen.cpp | 6 +-- 8 files changed, 83 insertions(+), 63 deletions(-) diff --git a/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp b/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp index 5e2676bf16..97b8f79a9f 100644 --- a/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp +++ b/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp @@ -96,8 +96,11 @@ enum OpenMPRTLFunctionNVPTX { OMPRTL_NVPTX__kmpc_get_team_static_memory, /// Call to void __kmpc_restore_team_static_memory(int16_t is_shared); OMPRTL_NVPTX__kmpc_restore_team_static_memory, - // Call to void __kmpc_barrier(ident_t *loc, kmp_int32 global_tid); + /// Call to void __kmpc_barrier(ident_t *loc, kmp_int32 global_tid); OMPRTL__kmpc_barrier, + /// Call to void __kmpc_barrier_simple_spmd(ident_t *loc, kmp_int32 + /// global_tid); + OMPRTL__kmpc_barrier_simple_spmd, }; /// Pre(post)-action for different OpenMP constructs specialized for NVPTX. @@ -640,17 +643,6 @@ static llvm::Value *getNVPTXNumThreads(CodeGenFunction &CGF) { "nvptx_num_threads"); } -/// Get barrier to synchronize all threads in a block. -static void getNVPTXCTABarrier(CodeGenFunction &CGF) { - llvm::Function *F = llvm::Intrinsic::getDeclaration( - &CGF.CGM.getModule(), llvm::Intrinsic::nvvm_barrier0); - F->addFnAttr(llvm::Attribute::Convergent); - CGF.EmitRuntimeCall(F); -} - -/// Synchronize all GPU threads in a block. -static void syncCTAThreads(CodeGenFunction &CGF) { getNVPTXCTABarrier(CGF); } - /// Get the value of the thread_limit clause in the teams directive. /// For the 'generic' execution mode, the runtime encodes thread_limit in /// the launch parameters, always starting thread_limit+warpSize threads per @@ -1813,6 +1805,17 @@ CGOpenMPRuntimeNVPTX::createNVPTXRuntimeFunction(unsigned Function) { cast(RTLFn)->addFnAttr(llvm::Attribute::Convergent); break; } + case OMPRTL__kmpc_barrier_simple_spmd: { + // Build void __kmpc_barrier_simple_spmd(ident_t *loc, kmp_int32 + // global_tid); + llvm::Type *TypeParams[] = {getIdentTyPointerTy(), CGM.Int32Ty}; + auto *FnTy = + llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false); + RTLFn = + CGM.CreateRuntimeFunction(FnTy, /*Name*/ "__kmpc_barrier_simple_spmd"); + cast(RTLFn)->addFnAttr(llvm::Attribute::Convergent); + break; + } } return RTLFn; } @@ -2665,6 +2668,20 @@ void CGOpenMPRuntimeNVPTX::emitSPMDParallelCall( } } +void CGOpenMPRuntimeNVPTX::syncCTAThreads(CodeGenFunction &CGF) { + // Always emit simple barriers! + if (!CGF.HaveInsertPoint()) + return; + // Build call __kmpc_barrier_simple_spmd(nullptr, 0); + // This function does not use parameters, so we can emit just default values. + llvm::Value *Args[] = { + llvm::ConstantPointerNull::get( + cast(getIdentTyPointerTy())), + llvm::ConstantInt::get(CGF.Int32Ty, /*V=*/0, /*isSigned=*/true)}; + CGF.EmitRuntimeCall( + createNVPTXRuntimeFunction(OMPRTL__kmpc_barrier_simple_spmd), Args); +} + void CGOpenMPRuntimeNVPTX::emitBarrierCall(CodeGenFunction &CGF, SourceLocation Loc, OpenMPDirectiveKind Kind, bool, diff --git a/lib/CodeGen/CGOpenMPRuntimeNVPTX.h b/lib/CodeGen/CGOpenMPRuntimeNVPTX.h index 8fb3b0a061..6091610c37 100644 --- a/lib/CodeGen/CGOpenMPRuntimeNVPTX.h +++ b/lib/CodeGen/CGOpenMPRuntimeNVPTX.h @@ -58,6 +58,9 @@ private: bool requiresFullRuntime() const { return RequiresFullRuntime; } + /// Get barrier to synchronize all threads in a block. + void syncCTAThreads(CodeGenFunction &CGF); + /// Emit the worker function for the current target region. void emitWorkerFunction(WorkerFunctionState &WST); diff --git a/test/OpenMP/nvptx_data_sharing.cpp b/test/OpenMP/nvptx_data_sharing.cpp index 3baee007fa..df9c3ee83b 100644 --- a/test/OpenMP/nvptx_data_sharing.cpp +++ b/test/OpenMP/nvptx_data_sharing.cpp @@ -34,7 +34,7 @@ void test_ds(){ /// ========= In the worker function ========= /// // CK1: {{.*}}define internal void @__omp_offloading{{.*}}test_ds{{.*}}_worker() -// CK1: call void @llvm.nvvm.barrier0() +// CK1: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CK1-NOT: call void @__kmpc_data_sharing_init_stack /// ========= In the kernel function ========= /// @@ -59,8 +59,8 @@ void test_ds(){ // CK1: [[SHARGSTMP2:%.+]] = getelementptr inbounds i8*, i8** [[SHARGSTMP1]], i64 0 // CK1: [[SHAREDVAR:%.+]] = bitcast i32* [[A]] to i8* // CK1: store i8* [[SHAREDVAR]], i8** [[SHARGSTMP2]] -// CK1: call void @llvm.nvvm.barrier0() -// CK1: call void @llvm.nvvm.barrier0() +// 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: call void @__kmpc_end_sharing_variables() // CK1: store i32 100, i32* [[B]] // CK1: call void @__kmpc_kernel_prepare_parallel({{.*}}, i16 1) @@ -72,8 +72,8 @@ void test_ds(){ // CK1: [[SHARGSTMP12:%.+]] = getelementptr inbounds i8*, i8** [[SHARGSTMP3]], i64 1 // CK1: [[SHAREDVAR2:%.+]] = bitcast i32* [[A]] to i8* // CK1: store i8* [[SHAREDVAR2]], i8** [[SHARGSTMP12]] -// CK1: call void @llvm.nvvm.barrier0() -// CK1: call void @llvm.nvvm.barrier0() +// 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: call void @__kmpc_end_sharing_variables() // CK1: [[SHARED_MEM_FLAG:%.+]] = load i16, i16* [[KERNEL_SHARED]], // CK1: call void @__kmpc_restore_team_static_memory(i16 [[SHARED_MEM_FLAG]]) diff --git a/test/OpenMP/nvptx_parallel_codegen.cpp b/test/OpenMP/nvptx_parallel_codegen.cpp index 821897cc0f..21a616d091 100644 --- a/test/OpenMP/nvptx_parallel_codegen.cpp +++ b/test/OpenMP/nvptx_parallel_codegen.cpp @@ -88,7 +88,7 @@ int bar(int n){ // CHECK: br label {{%?}}[[AWAIT_WORK:.+]] // // CHECK: [[AWAIT_WORK]] -// CHECK: call void @llvm.nvvm.barrier0() +// CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: [[KPR:%.+]] = call i1 @__kmpc_kernel_parallel(i8** [[OMP_WORK_FN]] // CHECK: [[KPRB:%.+]] = zext i1 [[KPR]] to i8 // store i8 [[KPRB]], i8* [[OMP_EXEC_STATUS]], align 1 @@ -127,7 +127,7 @@ int bar(int n){ // CHECK: br label {{%?}}[[BAR_PARALLEL]] // // CHECK: [[BAR_PARALLEL]] -// CHECK: call void @llvm.nvvm.barrier0() +// CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[AWAIT_WORK]] // // CHECK: [[EXIT]] @@ -164,21 +164,21 @@ int bar(int n){ // CHECK: [[MTMP1:%.+]] = sub nuw i32 [[MNTH]], [[MWS]] // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]] // CHECK: call void @__kmpc_kernel_prepare_parallel(i8* bitcast (void (i16, i32)* [[PARALLEL_FN1]]_wrapper to i8*), -// CHECK: call void @llvm.nvvm.barrier0() -// CHECK: call void @llvm.nvvm.barrier0() +// 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 void @__kmpc_serialized_parallel( // CHECK: {{call|invoke}} void [[PARALLEL_FN3:@.+]]( // CHECK: call void @__kmpc_end_serialized_parallel( // CHECK: call void @__kmpc_kernel_prepare_parallel(i8* bitcast (void (i16, i32)* [[PARALLEL_FN2]]_wrapper to i8*), -// CHECK: call void @llvm.nvvm.barrier0() -// CHECK: call void @llvm.nvvm.barrier0() +// 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-64-DAG: load i32, i32* [[REF_A]] // CHECK-32-DAG: load i32, i32* [[LOCAL_A]] // CHECK: br label {{%?}}[[TERMINATE:.+]] // // CHECK: [[TERMINATE]] // CHECK: call void @__kmpc_kernel_deinit( -// CHECK: call void @llvm.nvvm.barrier0() +// CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[EXIT]] // // CHECK: [[EXIT]] @@ -207,7 +207,7 @@ int bar(int n){ // CHECK: br label {{%?}}[[AWAIT_WORK:.+]] // // CHECK: [[AWAIT_WORK]] -// CHECK: call void @llvm.nvvm.barrier0() +// CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: [[KPR:%.+]] = call i1 @__kmpc_kernel_parallel(i8** [[OMP_WORK_FN]], // CHECK: [[KPRB:%.+]] = zext i1 [[KPR]] to i8 // store i8 [[KPRB]], i8* [[OMP_EXEC_STATUS]], align 1 @@ -237,7 +237,7 @@ int bar(int n){ // CHECK: br label {{%?}}[[BAR_PARALLEL]] // // CHECK: [[BAR_PARALLEL]] -// CHECK: call void @llvm.nvvm.barrier0() +// CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[AWAIT_WORK]] // // CHECK: [[EXIT]] @@ -289,8 +289,8 @@ int bar(int n){ // // CHECK: [[IF_THEN]] // CHECK: call void @__kmpc_kernel_prepare_parallel(i8* bitcast (void (i16, i32)* [[PARALLEL_FN4]]_wrapper to i8*), -// CHECK: call void @llvm.nvvm.barrier0() -// CHECK: call void @llvm.nvvm.barrier0() +// 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: br label {{%?}}[[IF_END:.+]] // // CHECK: [[IF_ELSE]] @@ -309,7 +309,7 @@ int bar(int n){ // // CHECK: [[TERMINATE]] // CHECK: call void @__kmpc_kernel_deinit( -// CHECK: call void @llvm.nvvm.barrier0() +// CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[EXIT]] // // CHECK: [[EXIT]] diff --git a/test/OpenMP/nvptx_parallel_for_codegen.cpp b/test/OpenMP/nvptx_parallel_for_codegen.cpp index 6a8601f3ae..92783d6085 100644 --- a/test/OpenMP/nvptx_parallel_for_codegen.cpp +++ b/test/OpenMP/nvptx_parallel_for_codegen.cpp @@ -37,7 +37,7 @@ int bar(int n){ // CHECK-DAG: [[KERNEL_SHARED:@.+]] = internal unnamed_addr constant i16 1 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l12}}_worker() -// CHECK: call void @llvm.nvvm.barrier0() +// CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: call i1 @__kmpc_kernel_parallel( // CHECK: call void @__omp_outlined___wrapper( @@ -52,8 +52,8 @@ int bar(int n){ // CHECK: [[STACK:%.+]] = getelementptr inbounds i8, i8* [[KERNEL_RD]], i{{64|32}} 0 // CHECK: call void @__kmpc_kernel_prepare_parallel( // CHECK: call void @__kmpc_begin_sharing_variables({{.*}}, i64 2) -// CHECK: call void @llvm.nvvm.barrier0() -// CHECK: call void @llvm.nvvm.barrier0() +// 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 void @__kmpc_end_sharing_variables() // CHECK: [[IS_SHARED:%.+]] = load i16, i16* [[KERNEL_SHARED]], // CHECK: call void @__kmpc_restore_team_static_memory(i16 [[IS_SHARED]]) diff --git a/test/OpenMP/nvptx_target_codegen.cpp b/test/OpenMP/nvptx_target_codegen.cpp index 033ed1e684..ff44c0e8fb 100644 --- a/test/OpenMP/nvptx_target_codegen.cpp +++ b/test/OpenMP/nvptx_target_codegen.cpp @@ -52,8 +52,8 @@ struct TT{ // CHECK: [[ARG_PTR2:%.+]] = getelementptr inbounds i8*, i8** [[ARG_PTRS]], i{{[0-9]+}} 1 // CHECK: [[BC:%.+]] = bitcast i32** [[PTR2_REF]] to i8* // CHECK: store i8* [[BC]], i8** [[ARG_PTR2]], -// CHECK: call void @llvm.nvvm.barrier0() -// CHECK: call void @llvm.nvvm.barrier0() +// 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 void @__kmpc_end_sharing_variables() void targetBar(int *Ptr1, int *Ptr2) { #pragma omp target map(Ptr1[:0], Ptr2) @@ -78,7 +78,7 @@ int foo(int n) { // CHECK: br label {{%?}}[[AWAIT_WORK:.+]] // // CHECK: [[AWAIT_WORK]] - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]], // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]] @@ -95,7 +95,7 @@ int foo(int n) { // CHECK: br label {{%?}}[[BAR_PARALLEL]] // // CHECK: [[BAR_PARALLEL]] - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[AWAIT_WORK]] // // CHECK: [[EXIT]] @@ -129,7 +129,7 @@ int foo(int n) { // // CHECK: [[TERMINATE]] // CHECK: call void @__kmpc_kernel_deinit( - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[EXIT]] // // CHECK: [[EXIT]] @@ -151,7 +151,7 @@ int foo(int n) { // CHECK: br label {{%?}}[[AWAIT_WORK:.+]] // // CHECK: [[AWAIT_WORK]] - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]], // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]] @@ -168,7 +168,7 @@ int foo(int n) { // CHECK: br label {{%?}}[[BAR_PARALLEL]] // // CHECK: [[BAR_PARALLEL]] - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[AWAIT_WORK]] // // CHECK: [[EXIT]] @@ -206,7 +206,7 @@ int foo(int n) { // // CHECK: [[TERMINATE]] // CHECK: call void @__kmpc_kernel_deinit( - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[EXIT]] // // CHECK: [[EXIT]] @@ -225,7 +225,7 @@ int foo(int n) { // CHECK: br label {{%?}}[[AWAIT_WORK:.+]] // // CHECK: [[AWAIT_WORK]] - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]], // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]] @@ -242,7 +242,7 @@ int foo(int n) { // CHECK: br label {{%?}}[[BAR_PARALLEL]] // // CHECK: [[BAR_PARALLEL]] - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[AWAIT_WORK]] // // CHECK: [[EXIT]] @@ -316,7 +316,7 @@ int foo(int n) { // // CHECK: [[TERMINATE]] // CHECK: call void @__kmpc_kernel_deinit( - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[EXIT]] // // CHECK: [[EXIT]] @@ -417,7 +417,7 @@ int baz(int f, double &a) { // CHECK: br label {{%?}}[[AWAIT_WORK:.+]] // // CHECK: [[AWAIT_WORK]] - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]], // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]] @@ -434,7 +434,7 @@ int baz(int f, double &a) { // CHECK: br label {{%?}}[[BAR_PARALLEL]] // // CHECK: [[BAR_PARALLEL]] - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[AWAIT_WORK]] // // CHECK: [[EXIT]] @@ -487,7 +487,7 @@ int baz(int f, double &a) { // // CHECK: [[TERMINATE]] // CHECK: call void @__kmpc_kernel_deinit( - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[EXIT]] // // CHECK: [[EXIT]] @@ -503,7 +503,7 @@ int baz(int f, double &a) { // CHECK: br label {{%?}}[[AWAIT_WORK:.+]] // // CHECK: [[AWAIT_WORK]] - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]], // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]] @@ -523,7 +523,7 @@ int baz(int f, double &a) { // CHECK: br label {{%?}}[[BAR_PARALLEL]] // // CHECK: [[BAR_PARALLEL]] - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[AWAIT_WORK]] // // CHECK: [[EXIT]] @@ -581,7 +581,7 @@ int baz(int f, double &a) { // // CHECK: [[TERMINATE]] // CHECK: call void @__kmpc_kernel_deinit( - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[EXIT]] // // CHECK: [[EXIT]] @@ -633,8 +633,8 @@ int baz(int f, double &a) { // CHECK: [[REF:%.+]] = getelementptr inbounds i8*, i8** [[SHARED]], i{{64|32}} 0 // CHECK: [[F_REF:%.+]] = bitcast i32* [[F_PTR]] to i8* // CHECK: store i8* [[F_REF]], i8** [[REF]], - // CHECK: call void @llvm.nvvm.barrier0() - // CHECK: call void @llvm.nvvm.barrier0() + // 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 void @__kmpc_end_sharing_variables() // CHECK: br label @@ -656,7 +656,7 @@ int baz(int f, double &a) { // CHECK: br label {{%?}}[[AWAIT_WORK:.+]] // // CHECK: [[AWAIT_WORK]] - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]], // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]] @@ -673,7 +673,7 @@ int baz(int f, double &a) { // CHECK: br label {{%?}}[[BAR_PARALLEL]] // // CHECK: [[BAR_PARALLEL]] - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[AWAIT_WORK]] // // CHECK: [[EXIT]] @@ -725,7 +725,7 @@ int baz(int f, double &a) { // // CHECK: [[TERMINATE]] // CHECK: call void @__kmpc_kernel_deinit( - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[EXIT]] // // CHECK: [[EXIT]] diff --git a/test/OpenMP/nvptx_target_teams_codegen.cpp b/test/OpenMP/nvptx_target_teams_codegen.cpp index 33c8b065d3..069eecb607 100644 --- a/test/OpenMP/nvptx_target_teams_codegen.cpp +++ b/test/OpenMP/nvptx_target_teams_codegen.cpp @@ -67,7 +67,7 @@ int bar(int n){ // CHECK: br label {{%?}}[[AWAIT_WORK:.+]] // // CHECK: [[AWAIT_WORK]] - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: [[KPR:%.+]] = call i1 @__kmpc_kernel_parallel(i8** [[OMP_WORK_FN]], i16 1) // CHECK: [[KPRB:%.+]] = zext i1 [[KPR]] to i8 // store i8 [[KPRB]], i8* [[OMP_EXEC_STATUS]], align 1 @@ -88,7 +88,7 @@ int bar(int n){ // CHECK: br label {{%?}}[[BAR_PARALLEL]] // // CHECK: [[BAR_PARALLEL]] - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[AWAIT_WORK]] // // CHECK: [[EXIT]] @@ -134,7 +134,7 @@ int bar(int n){ // // CHECK: [[TERMINATE]] // CHECK: call void @__kmpc_kernel_deinit( - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[EXIT]] // // CHECK: [[EXIT]] @@ -153,7 +153,7 @@ int bar(int n){ // CHECK: br label {{%?}}[[AWAIT_WORK:.+]] // // CHECK: [[AWAIT_WORK]] - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: [[KPR:%.+]] = call i1 @__kmpc_kernel_parallel(i8** [[OMP_WORK_FN]], i16 1) // CHECK: [[KPRB:%.+]] = zext i1 [[KPR]] to i8 // store i8 [[KPRB]], i8* [[OMP_EXEC_STATUS]], align 1 @@ -174,7 +174,7 @@ int bar(int n){ // CHECK: br label {{%?}}[[BAR_PARALLEL]] // // CHECK: [[BAR_PARALLEL]] - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[AWAIT_WORK]] // // CHECK: [[EXIT]] @@ -220,7 +220,7 @@ int bar(int n){ // // CHECK: [[TERMINATE]] // CHECK: call void @__kmpc_kernel_deinit( - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[EXIT]] // // CHECK: [[EXIT]] diff --git a/test/OpenMP/nvptx_target_teams_distribute_codegen.cpp b/test/OpenMP/nvptx_target_teams_distribute_codegen.cpp index a13762549c..cf007427ec 100644 --- a/test/OpenMP/nvptx_target_teams_distribute_codegen.cpp +++ b/test/OpenMP/nvptx_target_teams_distribute_codegen.cpp @@ -73,15 +73,15 @@ int bar(int n){ // CHECK: [[SHARED_VARS_BUF:%.+]] = load i8**, i8*** [[SHARED_VARS_PTR]], // CHECK: [[I_ADDR_BC:%.+]] = bitcast i32* [[I_ADDR]] to i8* // CHECK: store i8* [[I_ADDR_BC]], i8** [[SHARED_VARS_BUF]], - // CHECK: call void @llvm.nvvm.barrier0() - // CHECK: call void @llvm.nvvm.barrier0() + // 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 void @__kmpc_end_sharing_variables() // CHECK: call void @__kmpc_for_static_fini( // CHECK: br label {{%?}}[[TERMINATE:.+]] // // CHECK: [[TERMINATE]] // CHECK: call void @__kmpc_kernel_deinit( - // CHECK: call void @llvm.nvvm.barrier0() + // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0) // CHECK: br label {{%?}}[[EXIT]] // // CHECK: [[EXIT]] -- 2.40.0