clEnqueueNDRangeKernel() throws CL_INVALID_KERNEL_ARGS even when all the kernel arguments are set (only occurs on the CPU on macOS)

user_1285261 0
Viewed 44

I've been contributing to an OpenCL program called mfakto that trial factors Mersenne numbers for GIMPS. It uses a modified Sieve of Eratosthenes to create a list of potential factors and then uses modular exponentiation to test the factors. The sieving step can be done on either the GPU or CPU while the modular exponentiation step is done only on the target device.

The program uses different kernels depending on the number's size. I'm able to initialize the 15-bit kernels without issues. However, clEnqueueNDRangeKernel() throws a CL_INVALID_KERNEL_ARGS error for 32-bit kernels even though I've set all the arguments. Here's a sample 32-bit kernel:

__kernel void cl_barrett32_76(__private uint exponent, const int96_t k_base, const __global uint * restrict k_tab, const int shiftcount,
#ifdef WA_FOR_CATALYST11_10_BUG
                           const uint8 b_in,
#else
                           const __private int192_t bb,
#endif
                           __global uint * restrict RES, const int bit_max65
                           MODBASECASE_PAR_DEF         )
{
    ...
}

In normal circumstances, the kernel takes eight arguments. MODBASECASE_PAR_DEF adds a ninth argument that is only used when the application is compiled with certain debug flags. I traced the code and verified that clSetKernelArg() is used to set each argument at least once.

The issue only occurs when the OpenCL code is run on the CPU on macOS. It does not come up when the program is run on an AMD GPU or on any device on Windows.

Apple's OpenCL implementation only supports a kernel work-group size of 128 on the CPU, but I've already added checks to prevent clEnqueueNDRangeKernel() from trying to use more local threads than the kernel allows.

I realize this is a very specific problem in a very complex application, but any advice would be appreciated.

opencl c++
1 Answers

I suspect the issue might be related to the int96_t and int192_t types you are using. Are those typedefs for structs? The OpenCL standard is not clear on the use of non-primitive, non-buffer types as kernel arguments, so you will often find this working for some implementations but failing for others.

I suggest putting the data for these arguments in OpenCL buffers, and declaring them in the kernel function signature as constant int96_t* or global int96_t and similar. You'll obviously need to dereference the pointer in the kernel, either on every use of the value or when copying it into a private variable once.

user_486600
Related