Author: Alexey Bataev Date: 2019-12-11T13:18:39-05:00 New Revision: d8c31d41989b0748e2e5b8d7fa9cf7e7023bcbff
URL: https://github.com/llvm/llvm-project/commit/d8c31d41989b0748e2e5b8d7fa9cf7e7023bcbff DIFF: https://github.com/llvm/llvm-project/commit/d8c31d41989b0748e2e5b8d7fa9cf7e7023bcbff.diff LOG: [OPENMP50]Fix capturing of if condition in target parallel for simd directive. Fixed capturing of the if condition if no modifer was specified in this condition. Previously could capture it only in outer region and it could lead to a compiler crash. Added: Modified: clang/lib/Sema/SemaOpenMP.cpp clang/test/OpenMP/target_parallel_for_simd_codegen.cpp Removed: ################################################################################ diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index 4f4c4d2e887d..91d76e648613 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -10705,8 +10705,10 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause( switch (DKind) { case OMPD_target_parallel_for_simd: if (OpenMPVersion >= 50 && - (NameModifier == OMPD_unknown || NameModifier == OMPD_simd)) + (NameModifier == OMPD_unknown || NameModifier == OMPD_simd)) { CaptureRegion = OMPD_parallel; + break; + } LLVM_FALLTHROUGH; case OMPD_target_parallel: case OMPD_target_parallel_for: diff --git a/clang/test/OpenMP/target_parallel_for_simd_codegen.cpp b/clang/test/OpenMP/target_parallel_for_simd_codegen.cpp index be62c6f7bbf0..5e1ed2d6275b 100644 --- a/clang/test/OpenMP/target_parallel_for_simd_codegen.cpp +++ b/clang/test/OpenMP/target_parallel_for_simd_codegen.cpp @@ -539,7 +539,7 @@ struct S1 { short int c[2][n]; #ifdef OMP5 - #pragma omp target parallel for simd if(target: n>60) if(simd:n) + #pragma omp target parallel for simd if(n>60) #else #pragma omp target parallel for simd if(target: n>60) #endif // OMP5 @@ -576,20 +576,22 @@ int bar(int n){ // CHECK: define {{.*}}[[FS1]] // // CHECK: i8* @llvm.stacksave() +// CHECK-32: store i32 %{{.+}}, i32* %__vla_expr +// OMP50: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 60 // CHECK-64: [[B_ADDR:%.+]] = bitcast i[[SZ]]* [[B_CADDR:%.+]] to i32* // CHECK-64: store i32 %{{.+}}, i32* [[B_ADDR]], // CHECK-64: [[B_CVAL:%.+]] = load i[[SZ]], i[[SZ]]* [[B_CADDR]], -// CHECK-32: store i32 %{{.+}}, i32* %__vla_expr // CHECK-32: store i32 %{{.+}}, i32* [[B_ADDR:%.+]], // CHECK-32: [[B_CVAL:%.+]] = load i[[SZ]], i[[SZ]]* [[B_ADDR]], +// OMP45: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 60 // OMP50: [[TOBOOL:%.+]] = trunc i8 %{{.+}} to i1 // OMP50: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAP:%.+]] to i8* // OMP50: [[FROMBOOL:%.+]] = zext i1 [[TOBOOL]] to i8 // OMP50: store i8 [[FROMBOOL]], i8* [[CONV]], // OMP50: [[SIMD_COND:%.+]] = load i[[SZ]], i[[SZ]]* [[CAP]], -// CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 60 +// OMP50: [[IF:%.+]] = trunc i8 %{{.+}} to i1 // CHECK: br i1 [[IF]], label %[[TRY:[^,]+]], label %[[FAIL:[^,]+]] // CHECK: [[TRY]] // We capture 2 VLA sizes in this target region @@ -599,7 +601,7 @@ int bar(int n){ // CHECK-32: [[CSIZE:%.+]] = sext i32 [[CSZSIZE]] to i64 // OMP45-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 6, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MAPT7]], i32 0, i32 0), i32 1, i32 0) -// OMP50-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 7, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([7 x i64], [7 x i64]* [[MAPT7]], i32 0, i32 0), i32 1, i32 0) +// OMP50-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 7, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([7 x i64], [7 x i64]* [[MAPT7]], i32 0, i32 0), i32 1, i32 %{{.+}}) // OMP45-DAG: [[BPR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP:%.+]], i32 0, i32 0 // OMP45-DAG: [[PR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P:%.+]], i32 0, i32 0 // OMP45-DAG: [[SR]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S:%.+]], i32 0, i32 0 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
