Error dispatching a single thread SPIRV Kernel?

[jjfumero]

Executing a single thread kernel (without specifying the thread-id intrinsic in the kernel) gives me wrong results.

I simplified the kernel as follows:

__kernel void initValues(__global long* input, __global long* output) {
       output[0] = input[0];
} 

After running the SPIRV kernel I get wrong values in the output. However, if I run the following kernel:

__kernel void initValues(__global long* input, __global long* output) {
       uint idx = get_global_id(0);
       output[idx] = input[idx];
} 

I get the correct results when running with a single thread.

This is the kernel dispatch

 uint32_t groupSizeX = 1u;
 uint32_t groupSizeY = 1u;
 uint32_t groupSizeZ = 1u;
 VALIDATECALL(zeKernelSuggestGroupSize(kernel, items, 1U, 1U, &groupSizeX, &groupSizeY, &groupSizeZ));
 VALIDATECALL(zeKernelSetGroupSize(kernel, groupSizeX, groupSizeY, groupSizeZ));

 // Push arguments
 VALIDATECALL(zeKernelSetArgumentValue(kernel, 0, sizeof(inputBuffer), &inputBuffer));
 VALIDATECALL(zeKernelSetArgumentValue(kernel, 1, sizeof(outputBuffer), &outputBuffer));

// Kernel thread-dispatch
ze_group_count_t dispatch;
dispatch.groupCountX = 1u;
dispatch.groupCountY = 1u;
dispatch.groupCountZ = 1u;
VALIDATECALL(zeCommandListAppendLaunchKernel(cmdList, kernel, &dispatch, nullptr, 0, nullptr))

Note, I am working on a tool that generates SPIRV, sometimes from single treads kernels (so we do not expose thread-ids). I wonder if I am missing something. I am not sure if this is an error in Level Zero, but it works in pure OpenCL.

Full program available here: https://gist.github.com/jjfumero/a15f24c567953af57c3a02a963577980

Any ideas/pointers are appreciated.

[jandres742]

@jjfumero Thanks. This seems like an issue more suitable for the driver repo in https://github.com/intel/compute-runtime. Could you move the issue there so appropriate attention is given?

[bmyates]

Closing this issue as I believe it's being handled in compute-runtime issue. Please reopen if needed.