Device-side-enqueue throws exception during runtime.

[harald-lang]

Hi,
I have trouble running kernels which are using the OpenCL 2.0 feature enqueue_kernel. The offline compilation runs without errors, but dispatching raises a runtime exception:
kfd kfd: Invalid PPR device 0:1.0 pasid 1 address 0x0 flags 0x124

Sourcecode:

__kernel void storeGlobalId(__global size_t* output, const size_t n) {
   const size_t gid=get_global_id(0);
   if (gid<n) output[gid]=gid;
}

__kernel void enqueueStoreGlobalId(__global size_t* output, const size_t n) {
   const size_t gid=get_global_id(0);
   if (gid == 0) {
      queue_t queue = get_default_queue();
      ndrange_t nd = ndrange_1D(n,128);

      void (^storeGlobalId_block)(void) = ^ {
         storeGlobalId(output,n);
      };

      enqueue_kernel(queue, CLK_ENQUEUE_FLAGS_NO_WAIT, nd, storeGlobalId_block);
   }
}

System setup:

• Kaveri APU (no dGPU)
• Kernel 4.0.0-100002-generic #201511031149 SMP
• kfd-v1.6.1 (7fb04c4 from git repo HSA-Drivers-Linux-AMD)
• HSA-Runtime 1.0.3 (fa0ef7e from git repo HSA-Runtime-AMD)
• CL offline compiler (CLOC) v0.9.8

[bensander]

Device-enqueue requires OpenCL runtime which is not available with the CLOC flow. Recall you are compiling CL code to HSAIL, and then runnign that on HSA runtime. Practically this means that most OpenCL language features are available (ie local memory, atomics, instructions) but features that require an OpenCL runtime are not (two examples are device enqueue and printf).

[harald-lang]

Hi,
HSAIL provides instructions to work with user-mode-queues, such as addqueuewriteindex, casqueuewriteindex etc. Is there a reason for CLOC not to use these instructions? Semantics should be the same for HSAIL and OCL (at least for CLK_ENQUEUE_FLAGS_NO_WAIT).

On the other hand, CLOC does not produce any warning message. Instead, it produces a lot of HSAIL (which I haven't reverse-engineered yet).

[gregrodgers]

Look at the cloc readme to see how to use the snack.sh command to generate the hsa API code to launch a kernel compiled by cloc. As Ben said earlier, we don't have the host side opencl API for HSA or ROCM machines.