[OpenCL] Blocks cannot capture/reference another block

[Anastasia]

Summary

Adding the last restriction from s6.12.5 OpenCL C v2.0.

"A Block cannot reference or capture another Block variable declared in the outer scope".

[Nicola]

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.

[yaxunl]

In 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.

[Anastasia]

In D29764#673679, @yaxunl wrote:

In 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 D29764#673679, @yaxunl wrote:

In 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.

test/SemaOpenCL/invalid-block.cl

71

v0 is not used!

[yaxunl]

Please fix the test. Otherwise LGTM. Thanks!

[Anastasia]

Committed in r295307!

include/clang/Basic/DiagnosticSemaKinds.td

 
 // OpenCL v2.0 s6.12.5 Blocks restrictions
 def err_opencl_block_storage_type : Error<
   "the __block storage type is not permitted">;
 def err_opencl_invalid_block_declaration : Error<
   "invalid block variable declaration - must be %select{const qualified|initialized}0">;
 def err_opencl_extern_block_declaration : Error<
   "invalid block variable declaration - using 'extern' storage class is disallowed">;
+def err_opencl_block_ref_block : Error<
+  "cannot refer to a block inside block">;
 
 // OpenCL v2.0 s6.13.9 - Address space qualifier functions. 
 def err_opencl_builtin_to_addr_arg_num : Error<
   "invalid number of arguments to function: %0">;
 def err_opencl_builtin_to_addr_invalid_arg : Error<
   "invalid argument %0 to function: %1, expecting a generic pointer argument">;
 
 // OpenCL v2.0 s6.13.17 Enqueue kernel restrictions.

lib/Sema/SemaExpr.cpp

if (HasBlocksAttr && (IsLambda || isa<CapturedRegionScopeInfo>(CSI))) {
     if (Diagnose) {
       S.Diag(Loc, diag::err_capture_block_variable)
         << Var->getDeclName() << !IsLambda;
       S.Diag(Var->getLocation(), diag::note_previous_decl)
         << Var->getDeclName();
     }
     return false;
   }
+  // OpenCL v2.0 s6.12.5: Blocks cannot reference/capture other blocks
+  if (S.getLangOpts().OpenCL && IsBlock &&
+      Var->getType()->isBlockPointerType()) {
+    if (Diagnose)
+      S.Diag(Loc, diag::err_opencl_block_ref_block);
+    return false;
+  }
 
   return true;
 }
 
 // Returns true if the capture by block was successful.
 static bool captureInBlock(BlockScopeInfo *BSI, VarDecl *Var, 
                                  SourceLocation Loc, 
                                  const bool BuildAndDiagnose, 

test/SemaOpenCL/invalid-block.cl

 void f6(bl2_t *bl_ptr) { // expected-error{{pointer to type '__generic bl2_t' (aka 'int (__generic ^const __generic)(int)') is invalid in OpenCL}}
   bl2_t bl = ^(int i) {
     return 1;
   };
   bl2_t *p; // expected-error {{pointer to type '__generic bl2_t' (aka 'int (__generic ^const __generic)(int)') is invalid in OpenCL}}
   *bl;      // expected-error {{invalid argument type 'bl2_t' (aka 'int (__generic ^const)(int)') to unary expression}}
   &bl;      // expected-error {{invalid argument type 'bl2_t' (aka 'int (__generic ^const)(int)') to unary expression}}
 }
+// A block can't reference another block
+kernel void foobars() {
+  int v0;
+  bl2_t bl1 = ^(int i) {
+    return 1;
+  };
+  void (^bl2)(void) = ^{
+    int i = bl1(1); // expected-error {{cannot refer to a block inside block}}
+  };
+  void (^bl3)(void) = ^{
+  };
+  void (^bl4)(void) = ^{
+    bl3(); // expected-error {{cannot refer to a block inside block}}
+  };
+  return;
+}