[PATCH] D20249: [OpenCL] Hierarchical/dynamic parallelism - enqueue kernel in OpenCL 2.0

[Anastasia Stulova via cfe-commits]

Anastasia marked 2 inline comments as done.

================
Comment at: lib/Sema/SemaChecking.cpp:81
@@ +80,3 @@
+static bool checkBlockArgs(Sema &S, Expr *BlockArg) {
+  const BlockPointerType *BPT = cast<BlockPointerType>(BlockArg->getType());
+  ArrayRef<QualType> Params =
----------------
bader wrote:
> There should be some check before cast.
> Here is the code snippet that will crash the compiler:
> 
>   extern queue_t get_default_queue();
>   extern int get_global_id(int);
>   extern ndrange_t get_ndrange();
>   typedef void (^MyBlock)(local void*, local int*);
> 
>   const MyBlock myBlock = (MyBlock)^(local int *p1, local int *p2) {
>     int id = get_global_id(0);
>     p1[id] += p2[id];
>   };
> 
>   void kernel f2(global int* a, global int* b) {
>     enqueue_kernel(get_default_queue(), 0, get_ndrange(), myBlock, 2U, 1U);
>   }
> 
Good spot! Thanks! I have added this case to the tests!

================
Comment at: test/CodeGenOpenCL/cl20-device-side-enqueue.cl:19
@@ +18,3 @@
+  // CHECK: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange
+  // CHECK: [[BL:%[0-9]+]] = bitcast <{ i8*, i32, i32, i8*, %struct.__block_descriptor*, i32 addrspace(1)*, i32 addrspace(1)*, i32 }>* %block to void ()*
+  // CHECK: [[BL_I8:%[0-9]+]] = bitcast void ()* [[BL]] to i8*
----------------
bader wrote:
> This check is failing on my machine:
> 
>   test\CodeGenOpenCL\cl20-device-side-enqueue.cl:19:12: error: expected string not found in input
>   // CHECK: [[BL:%[0-9]+]] = bitcast <{ i8*, i32, i32, i8*, %struct.__block_descriptor*, i32 addrspace(1)*, i32 addrspace(1)*, i32 }>* %block to void ()*
>   <stdin>:47:60: note: scanning from here
>   %3 = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange, align 4
>   <stdin>:67:2: note: possible intended match here
>   %7 = bitcast <{ i8*, i32, i32, i8*, %struct.__block_descriptor*, i32 addrspace(1)*, i32, i32 addrspace(1)* }>* %block to void ()*
> 
> It looks like the order of captured arguments is different from expected? Is the order specified or test shouldn't check it?
It seems that capture order is enforced by the order in which the parser parses statements, and it isn't really enforced  anywhere by ObjC implementation. Their tests only contain 1 capture or captures of the same type, so this issue couldn't be caught.

I am changing the test now, but I have no idea why the parsing order is different though.

If you tell me your revision, I could try to see if I can reproduce this locally.


http://reviews.llvm.org/D20249