[OpenCL, OpenMP] Fix crash when OpenMP used in OpenCL file

[mikerice]

Summary

Compiler crashes when omp simd is used in an OpenCL file:

clang -c -fopenmp omp_simd.cl

kernel void test(global int *data, int size) {

#pragma omp simd
for (int i = 0; i < size; ++i) {
}

}

The problem seems to be the check added to verify block pointers have initializers. An OMPCapturedExprDecl is created to capture ‘size’ but there is no TypeSourceInfo.

The change just uses getType() directly.

[ABataev]

1. Is this allowed to use OpenMP for OpenCL programs at all?
2. If it is allowed, I think it would be better to set the TypeSourceInfo for the artificial variable

[ABataev]

BTW, the test itself does not check anything.

[mikerice]

1. It probably makes sense to allow omp simd with OpenCL. Some people here have been successfully using it anyway.
2. I can give that a try.

The test just tests that the compiler doesn't segfault. What would you prefer instead, something that checks the AST or the codegen?

[Anastasia]

OpenCL C is based on C99, so OpenMP isn't enabled by default. But in your tests you use -fopenmp to activate it.

OpenCL general philosophy is that vectors are written explicitly, but it's not always very easy. In OpenCL C++ we have added an attribute hint for auto vectorization cl::vec_type_hint. It's given to a kernel rather than for a loop though. So I think this simd pragma is useful. My only worry is that other OpenMP features might not work well in OpenCL and we have no way to reject them at the moment.

[Anastasia]

lib/Sema/SemaDecl.cpp

11348

Does simd result in block creation? My original understanding of it was that it will result in vector operations, not extra threads.

[ABataev]

In D46667#1094340, @Anastasia wrote:

OpenCL C is based on C99, so OpenMP isn't enabled by default. But in your tests you use -fopenmp to activate it.

OpenCL general philosophy is that vectors are written explicitly, but it's not always very easy. In OpenCL C++ we have added an attribute hint for auto vectorization cl::vec_type_hint. It's given to a kernel rather than for a loop though. So I think this simd pragma is useful. My only worry is that other OpenMP features might not work well in OpenCL and we have no way to reject them at the moment.

We have special flag '-fopenmp-simd' to support only simd-based constructs. We can allow to use only this subset of OpenMP in OpenCL programs and disable support for '-fopenmp'.

[ABataev]

lib/Sema/SemaDecl.cpp

11348

Currently all of the OpenMP constructs (even simd) create special lambda-like construct CapturedStmt that are kind of a block. It is required for unified processing of the OpenMP constructs.

test/SemaOpenCL/omp_simd.cl

1

It would be good to check that at least -ast-print or -ast-dump works correctly.

[Anastasia]

lib/Sema/SemaDecl.cpp

11348

I see. It does add some extra code for captures and extra BBs but most of it seems to be removed with optimizations.

test/SemaOpenCL/omp_simd.cl

1

Generated IR seems fine. Might be enough to just convert to IR test?

[mikerice]

Sorry for the delay in getting back to this. I've found that we are using many OpenMP directives not just simd. We'd like to continue doing that.

I updated the patch to provide TypeSourceInfo for OMPCapturedExprDecls. I changed the test to check for captures (in the AST dump) for a few OpenMP directives.

[ABataev]

test/SemaOpenCL/with_openmp.cl

1 ↗

(On Diff #147538)

Still the same question, do we really want to support full OpenMP for OpenCL or only simd?

[mikerice]

test/SemaOpenCL/with_openmp.cl

1 ↗

(On Diff #147538)

We think it makes sense to have full support. We are using OpenCL for devices that can run full OpenMP. Is there any known problem with allowing full support?

[Anastasia]

LGTM!

test/SemaOpenCL/with_openmp.cl

1 ↗

(On Diff #147538)

Sounds good! What target do you use though? Would it help to have an IR test as well?

lib/Sema/SemaDecl.cpp

 }
 
 void Sema::CheckCompleteVariableDeclaration(VarDecl *var) {
   if (var->isInvalidDecl()) return;
 
   if (getLangOpts().OpenCL) {
     // OpenCL v2.0 s6.12.5 - Every block variable declaration must have an
     // initialiser
-    if (var->getTypeSourceInfo()->getType()->isBlockPointerType() &&
-        !var->hasInit()) {
+    if (var->getType()->isBlockPointerType() && !var->hasInit()) {
       Diag(var->getLocation(), diag::err_opencl_invalid_block_declaration)
           << 1 /*Init*/;
       var->setInvalidDecl();
       return;
     }
   }
 
   // In Objective-C, don't allow jumps past the implicit initialization of a

test/SemaOpenCL/omp_simd.cl

+// RUN: %clang_cc1 -verify -fopenmp -fsyntax-only -x cl %s
+// expected-no-diagnostics
+
+__kernel void test(__global int *data, int size) {
+  #pragma omp simd
+  for (int i = 0; i < size; ++i) {
+  }
+}