This is an archive of the discontinued LLVM Phabricator instance.

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

Authored by mikerice on May 9 2018, 5:28 PM.

Details

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.

Diff Detail

Repository
rC Clang

Event Timeline

mikerice created this revision.May 9 2018, 5:28 PM
  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

BTW, the test itself does not check anything.

  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?

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 added inline comments.May 10 2018, 5:13 AM
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.

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 added inline comments.May 10 2018, 5:35 AM
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 added inline comments.May 10 2018, 7:57 AM
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 updated this revision to Diff 147538.May 18 2018, 10:04 AM

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 added inline comments.May 18 2018, 10:23 AM
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 added inline comments.May 18 2018, 11:24 AM
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 accepted this revision.May 23 2018, 4:03 AM

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?

This revision is now accepted and ready to land.May 23 2018, 4:03 AM
This revision was automatically updated to reflect the committed changes.