linking printf

[Richardk2n]

clLinkProgram has issues with printf.

Consider this pyopencl example:

import pyopencl as cl

plattform = cl.get_platforms()[0]
device = plattform.get_devices()[0]

ctx = cl.Context([device])
queue = cl.CommandQueue(ctx)

src = r"""
kernel void tet() {
    const long m = get_global_id(0);
	printf("ID %ld\n", m);
}
"""

src2 = r"""
kernel void other() {
    const long n = get_global_id(0);
    printf("Do not print %ld\n", n);
}
"""

p = cl.Program(ctx, src).compile()
p2 = cl.Program(ctx, src2).compile()

program = cl.link_program(ctx, [p, p2])

program.tet(queue, (2,), None)

This should return

ID 0
ID 1

And using pocl with my CPU it does.
However, using Nvidia the linking segfaults.
Using AMD, I get the following:

ID 0
Do not print 4294967297
(null)

These problems are related to the presence of a printf in both kernels. If I get rid of one of them the issues disappear as well.
I have two related C++ examples here: KhronosGroup/OpenCL-CLHPP#316
For some conditions linking “just” fails under Nvidia instead of a segfault. Nvidia can sometimes segfault with a printf in a single kernel as well though (see example in linked issue).

A test covering this (One kernel with a printf, two kernels with a printf) would be appreciated.
While part of the issue can be captured by testing if the kernels compile and can be linked, to capture what AMD is doing wrong, execution is required. And I do not know how to test that.

By the way, while testing this I removed the contents of src2. This is how I discovered, that passing an empty string (which should be completely valid, and can happen with generated code) also creates issues.