Unable to include a directory located in macOS Documents in program build options

[dhellfeld]

I have a set of header files that define some inline functions that I use in several CL kernels. In the past I have been able to successfully include and retrieve these header files by supplying a compiler flag (namely -I /path/to/headers) to the program build (using the options kwarg). At some point after updating macOS (think this may have started on Catalina, I am now on Big Sur), I am no longer able to include the headers if they are in the macOS "Documents", "Desktop", or "Downloads" folders or in an external drive. If I move the headers somewhere else (say ~/) and update the -I flag, everything works. It's just something about these specific places it seems. I believe Apple has started controlling access to these folders, but what is interesting is that if I compile an example OpenCL program from the command line (clang -framework OpenCL -I /Users/username/Documents/include test.c, where test.c has an #include <test.h> and test.h lives in /Users/username/Documents/include ), it works fine.

MWE to reproduce the behavior:

1. Create and activate fresh anaconda environment (conda create -y -n pyopencl-test python=3.8 && conda activate pyopencl-test)
2. pip install pyopencl
3. Create test.py (see below) and replace "username" in the build options
4. Create test.h (see below) and place in "/Users/username/Documents/include" with appropriate "username"
5. Run python test.py (should fail)
6. Now move test.h to "/Users/username/include" and update the build options in test.py accordingly.
7. Run python test.py (should succeed)

test.py (augmented from documentation example):

#!/usr/bin/env python

import numpy as np
import pyopencl as cl

a_np = np.random.rand(50000).astype(np.float32)
b_np = np.random.rand(50000).astype(np.float32)

ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx)

mf = cl.mem_flags
a_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a_np)
b_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b_np)

prg = cl.Program(ctx, """
#include <test.h>  // header file

__kernel void sum(
    __global const float *a_g, __global const float *b_g, __global float *res_g)
{
  int gid = get_global_id(0);
  res_g[gid] = a_plus_b(a_g[gid], b_g[gid]);   // using function in header file
}
""").build(options="-I /Users/username/Documents/include")  # specify the include directory in the build options

res_g = cl.Buffer(ctx, mf.WRITE_ONLY, a_np.nbytes)
knl = prg.sum  # Use this Kernel object for repeated calls
knl(queue, a_np.shape, None, a_g, b_g, res_g)

res_np = np.empty_like(a_np)
cl.enqueue_copy(queue, res_np, res_g)

# Check on CPU with Numpy:
print(res_np - (a_np + b_np))
print(np.linalg.norm(res_np - (a_np + b_np)))
assert np.allclose(res_np, a_np + b_np)

test.h:

/*Test header.*/

inline float a_plus_b(const float a, const float b){
    float c = a + b;
    return c;
}

Expected behavior
Program should find test.h and use the a_plus_b function in the kernel. After step 5 above the build fails, claiming that it cannot find test.h. Once it is moved out of "Documents" and rerun (step 7) it works.

Output after step 5:

Output after step 7:

Environment

• OS: macOS Big Sur 11.6.1
• CPU device: <pyopencl.Device 'Intel(R) Xeon(R) W-2170B CPU @ 2.50GHz' on 'Apple' at 0xffffffff>
• GPU device: <pyopencl.Device 'AMD Radeon Pro Vega 64X Compute Engine' on 'Apple' at 0x1021d00>
• Python version: 3.8
• PyOpenCL version: 2021.2.10

[inducer]

I suspect this is a behavior of the Apple OpenCL implementation and not something PyOpenCL can control. AFAICT, Apple has some systemwide "compilation server" process for CL, and they may have changed security settings around that server so that it's no longer able to access header code in those locations.

I compile an example OpenCL program from the command line (clang -framework OpenCL -I /Users/username/Documents/include test.c, where test.c has an #include <test.h> and test.h lives in /Users/username/Documents/include ), it works fine.

That's not a good comparison. Instead, use a C program like this to try and compile kernel source code that includes the header. I suspect you'll see the same behavior.

Moving to Troubleshooting discussion.

[dhellfeld]

Ah yes 🤦‍♂️, that was indeed not a fair comparison. If I change my test.c example (or use the example you linked) to compile the kernel source code with the header included I get the same behavior (works fine if the header is not in those specific Apple protected locations, but fails if it is). Easy workaround is to move my project out of these folders, but I wonder if there is some other solution...