Scan kernel warning

[yves-surrel]

Simple scan kernel from the doc:

import os
import numpy as np
import pyopencl as cl

os.environ['PYOPENCL_CTX'] = ':0'
os.environ['PYOPENCL_COMPILER_OUTPUT'] = '1'

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

from pyopencl.scan import GenericScanKernel
knl = GenericScanKernel(
        context, np.int32,
        arguments="__global int *ary",
        input_expr="ary[i]",
        scan_expr="a+b", neutral="0",
        output_statement="ary[i+1] = item;")

a = cl.array.arange(queue, 10000, dtype=np.int32)
knl(a, queue=queue)

This works OK. If I switch to one of my my GPUs (MacBookPro) by os.environ['PYOPENCL_CTX'] = ':1' or os.environ['PYOPENCL_CTX'] = ':2', I get a host of warnings:

<program source>:1105:41: warning: comparison of integers of different signs: 'unsigned long' and 'const psc_index_type' (aka 'const int')
                    if (psc_K*psc_LID_0 < psc_offset_end)
                        ~~~~~~~~~~~~~~~ ^ ~~~~~~~~~~~~~~
...
<program source>:1395:39: warning: comparison of integers of different signs: 'unsigned long' and 'int'
                    if (0 + psc_LID_0 <
                        ~~~~~~~~~~~~~ ^
...
etc.

[inducer]

Thanks for reporting this. Unfortunately, I don't seem to have access to a CL implementation that reports these, so I wouldn't be able to effectively fix these. I use pocl most days, and I went as far as their options processing to see if this could be fixed:

https://github.com/pocl/pocl/blob/72b4fe46b14286364ff7a212242c1fd1a69c5663/lib/CL/pocl_build.c#L209-L210=

I also tried passing -Werror, all to no avail. If you've got a recipe for reproducing these on a commonly available ICD, I'd be all ears.

[yves-surrel]

If you can pick up the source lines where these warnings come from (not easy as they seem to be mako outputs), I could try to make some (int) casts...

[inducer]

Maybe try these for a start.

[yves-surrel]

I do not understand: I changed the _cluda.py file (finally putting nonsense text so it should raise an error), removed the maximum __pycache__ directories in pyopencl, set `os.environ['PYOPENCL_NO_CACHE'] = '1' to ensure compiling from scratch, but the kernel goes on compiling ok, with always the same warnings...

[inducer]

Sorry about that exercise in frustration! There's a separate cache in the scan code generation that does not respond to PYOPENCL_NO_CACHE. #573 fixes that.

[yves-surrel]

Glad to confirm that

#define LID_0 (int) get_local_id(0)
#define LID_1 (int) get_local_id(1)
#define LID_2 (int) get_local_id(2)
               
#define GID_0 (int) get_group_id(0)
#define GID_1 (int) get_group_id(1)
#define GID_2 (int) get_group_id(2)

#define LDIM_0 (int) get_local_size(0)
#define LDIM_1 (int) get_local_size(1)
#define LDIM_2 (int) get_local_size(2)

#define GDIM_0 (int) get_num_groups(0)
#define GDIM_1 (int) get_num_groups(1)
#define GDIM_2 (int) get_num_groups(2)

in _cluda.py removes the warnings for all platforms.

[inducer]

Thanks! #574