OpenCL says CL_KERNEL_WORK_GROUP_SIZE is 256 - but accepts 1024; why?

On NVIDIA GPUs from recent years, one can run kernels with up to 1024 "threads" per block, or work-items per workgroup in OpenCL parlance - as long as the kernel uses less than 64 registers. Specifically, that holds for a simple kernel such as:

__kernel void vectorAdd(
   __global unsigned char       * __restrict C,
   __global unsigned char const * __restrict A,
   __global unsigned char const * __restrict B,
   unsigned long length)
{
   int i = get_global_id(0);
   if (i < length)
       C[i] = A[i] + B[i] + 2;
}

however, if we build this kernel and run:

built_kernel.getWorkGroupInfo<CL_KERNEL_WORK_GROUP_SIZE>(device);

this yields 256, not 1024 (on a Quadro RTX 6000 with CUDA 11.6.55 for example).

That seems like a bug with the API. Is it? Or - is there some legitimate reason why CL_KERNEL_WORK_GROUP_SIZE should be 256?