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

Alexey Bader via cfe-commits cfe-commits at lists.llvm.org
Tue May 24 07:40:45 PDT 2016


bader added inline comments.

================
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 =
----------------
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);
  }


================
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*
----------------
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?


http://reviews.llvm.org/D20249





More information about the cfe-commits mailing list