NULL not allowed as pointer argument for zeKernelSetArgumentValue?

[fwinter]

When I try setting a pointer argument to zero using zeKernelSetArgumentValue the call fails with error code 0x78000004.

One could argue why would anyone set a pointer argument to a value that's definitely illegal to use? Well, consider a simplified kernel that appears in a stencil kernel in a multi MPI rank job:

__kernel void multi_node_stencil(__global int* v, __global float* out , __global float* a , __global float* b)
{
  int i = get_global_id(0);

  if (v[i] < 0)
    out[i] = b[ -v[i] - 1 ];
  else
    out[i] = a[ v[i] ];
}

Based on the sign of the index a the source value comes either from the local grid (a) or an MPI receive buffer (b). In circumstances where there's no off-node communication (a runtime choice) this kernel must still work. In this case I'd pass NULL as value for the last parameter, and I know the kernel executes still correctly since it would never read from b. However, Level Zero API seem to not like this and throws the above error.

This is how I use it to set the last parameter:

void* ptr = nullptr;
zeKernelSetArgumentValue( function , 3 , sizeof(ptr) , &ptr );

The API call zeKernelSetArgumentValue should accept a pointer to NULL, right?

[bmyates]

Hi @fwinter,

Your use case looks valid to me. A pointer to null should be okay.

The return code you are getting is ZE_RESULT_ERROR_INVALID_ARGUMENT. This is being returned from the runtime driver. I assume you are using the compute-runtime GPU driver. Can you move this issue to that github? https://github.com/intel/compute-runtime

[jandres742]

@bmyates should we close this since it was moved to driver?

[bmyates]

Yep, closing