Hi, I had a brief look at this issue. Continuing the discussion from the GitHub issue:
I don't think the call to enable CL_KERNEL_EXEC_INFO_INDIRECT_SHARED_ACCESS_INTEL is required. The test app isn't making any SHARED allocations, so this call will have no affect. Enabling CL_KERNEL_EXEC_INFO_INDIRECT_DEVICE_ACCESS_INTEL does have an affect, but because the set of USM pointers is explicitly provided using CL_KERNEL_EXEC_INFO_USM_PTRS_INTEL, enabling indirect access globally for all DEVICE allocations is redundant.
Note: enabling CL_KERNEL_EXEC_INFO_INDIRECT_DEVICE_ACCESS_INTEL does "fix" this issue, but I think this is overkill and shouldn't be required. Said another way, if the global indirect access flag is enabled then there is no need to pass any set of USM pointers explicitly.
I don't see anything wrong with the call to provide the set of USM pointers with CL_KERNEL_EXEC_INFO_USM_PTRS_INTEL, and it works with several other USM implementations I've tried.
Perhaps a good next step would be for somebody on the CPU OpenCL runtime team to report why the call to clEnqueueNDRangeKernel is returning CL_INVALID_OPERATION when the set of USM pointers is provided explicitly by CL_KERNEL_EXEC_INFO_USM_PTRS_INTEL and indirect access is NOT enabled globally.
Thanks!
