linking printf

`clLinkProgram` has issues with `printf`.

Consider this pyopencl example:
```python
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: https://github.com/KhronosGroup/OpenCL-CLHPP/issues/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.

Start with the minimal Python reproduction in the issue and the two related C++ examples in OpenCL-CLHPP issue 316. Add coverage for one and two kernels containing printf, including an empty second source, checking compilation and linking and, where execution is available, the output. Done means the cases are represented in the conformance tests and their expected behavior is explicit.

Written by the indexing model from the issue text.