Clarification needed in regards to allocation sizes and constant * arguments in clSetKernelArg and clSetKernelArgSVMPointer

[karolherbst]

At the moment clSetKernelArg states:

If the argument is declared with the constant qualifier, the size in bytes of the memory object cannot exceed CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE

but no such statement is made for clSetKernelArgSVMPointer.

But as far as I can see clSetKernelArg doesn't declare an error if the application would do it anyway. I think we should define what expectations to have here. Should clSetKernelArg return an error if the allocation is too big? Shouldn't be a problem for applications, because they can always create SubBuffers.

What should we do about clSetKernelArgSVMPointer? Doing sub allocation isn't doable with the current API and not sure if unified SVM is going to change that. One could use CL_MEM_USE_HOST_PTR with clCreateBuffer, but it's also not forbidden to use SVM allocations for constant * arguments.

Should application expect that only up to CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE is addressable here? Or should the call fail?

@bashbaug for awareness.

[bashbaug]

Thinking out loud:

• If clSetKernelArg is called to set an argument that is a pointer to the constant address space, and the argument value points to an OpenCL buffer that exceeds CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, this could be a CL_INVALID_MEM_OBJECT error.

  • Other error code possibilities are: CL_INVALID_ARG_VALUE (but this isn't very descriptive) or CL_MAX_SIZE_RESTRICTION_EXCEEDED (but this wasn't added until OpenCL 2.2).
  • At least one implementation I've tested defers the error behavior to clEnqueueNDRangeKernel, which isn't as user-friendly, but matches behavior in other cases that are device-specific.

• clSetKernelArgSVMPointer is trickier. Maybe we could say that it is an error (exact error value TBD) if the size of the SVM allocation is greater than CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE? This is a bit weird, but I'm not sure how to do better.

  • I'm not sure how this would work for system SVM, since the size of a system SVM allocation is unknown.

See also:

• Add missing error case for clSetKernelArgDevicePointerEXT #1492 (comment)
• Add new error code when clSetKernelArgSVMPointer is called for an unsuitable argument #1495

Is this something we can fix with unified SVM?

[karolherbst]

I think making use of CL_MAX_SIZE_RESTRICTION_EXCEEDED does make sense as it's similar to the current error condition.

I think for SVM it's probably fine to leave it undefined behavior, mentioning that some runtimes might misbehave if the access exceeds CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE.

Alternatively, we could also tell runtimes to make it work regardless, but that would imply that in the worst case implementations have to fall back to global addressing (if they have special paths for constant memory) if they can't prove it for a given pointer. I think this was partly the reason for the (no deprecated) nosvm attribute?

I do wonder what's the opinion across vendors here and how many even care these days.