[PATCH] D29764: [OpenCL] Blocks cannot capture/reference another block

[Anastasia Stulova via Phabricator via cfe-commits]

Anastasia added a comment.

In https://reviews.llvm.org/D29764#673679, @yaxunl wrote:

> In https://reviews.llvm.org/D29764#673548, @Nicola wrote:
>
> > Looking at "Example 4" in the standard it looks like this should also be illegal.
> >
> >   int (^block1)(void) = ^int {return 1;};
> >   int foo() { return block1(); }
> >  
> >   __kernel void k(global int *z)
> >   {
> >    int (^block2)(void) = ^int {
> >     return foo(); // expected-error {{cannot refer to a block inside block}}
> >    }; 
> >   }
> >  
> >
> >
> > Unless I missed something it's not erroring in this case.
>
>
> To diagnose this needs to traverse the AST tree. I think it is too much to do it during parsing.
>
> It may be done through static analysis though.




In https://reviews.llvm.org/D29764#673679, @yaxunl wrote:

> In https://reviews.llvm.org/D29764#673548, @Nicola wrote:
>
> > Looking at "Example 4" in the standard it looks like this should also be illegal.
> >
> >   int (^block1)(void) = ^int {return 1;};
> >   int foo() { return block1(); }
> >  
> >   __kernel void k(global int *z)
> >   {
> >    int (^block2)(void) = ^int {
> >     return foo(); // expected-error {{cannot refer to a block inside block}}
> >    }; 
> >   }
> >  
> >
> >
> > Unless I missed something it's not erroring in this case.
>
>
> To diagnose this needs to traverse the AST tree. I think it is too much to do it during parsing.
>
> It may be done through static analysis though.


Yes, it will require quite some expensive checks to implement visiting the AST to trace back the block variables declared in other functions. Additionally this example doesn't seem to cause any issues really. It is mainly in the case with blocks captured from the stack:

  kernel void foo() {
    bl2_t bl1 = ^(int i) {
      return 1;
    };
    void (^bl2)(void) = ^{
      int i = bl1(1); // expected-error {{cannot refer to a block inside block}}
    };
  }

}
for which ObjC implementation creates copy/destroy helpers that use some symbols that are supposed to be defined elsewhere (presumably in the ObjC runtime) which causes issue in OpenCL because we suddenly end up with undefined symbols. I am guessing this is needed in order to promote stack allocated blocks into heap. Although the copy itself doesn't happen in our cases.



================
Comment at: test/SemaOpenCL/invalid-block.cl:71
+kernel void foobars() {
+  int v0;
+  bl2_t bl1 = ^(int i) {
----------------
v0 is not used!


https://reviews.llvm.org/D29764