Skip to content

Domain tree inferences gets confused when sub-domains are not adjacently listed

@inducer, is this intended behavior? If it is, how do I add a sub-domain to a kernel and tell it under which tree it should be nested?

>>> import loopy as lp
>>> knl = lp.make_kernel(
... ["{[i]: 0 <= i <= 10}",
...  "[i] -> {[a]: 0 <= a <= i}",
...  "{[b]: 0 <= b <= 10}",
...  "[i] -> {[c]: 0 <= c <= i}"],
... "<>tmp = 1 {inames=i:c}",
... "...")
>>> print(knl)
---------------------------------------------------------------------------
KERNEL: loopy_kernel
---------------------------------------------------------------------------
ARGUMENTS:
---------------------------------------------------------------------------
DOMAINS:
{ [i] : 0 <= i <= 10 }
  [i] -> { [a] : 0 <= a <= i }
{ [b] : 0 <= b <= 10 }
[i] -> { [c] : 0 <= c <= i }
---------------------------------------------------------------------------
INAME IMPLEMENTATION TAGS:
a: None
b: None
c: None
i: None
---------------------------------------------------------------------------
TEMPORARIES:
tmp: type: <auto>, shape: () scope:auto
---------------------------------------------------------------------------
INSTRUCTIONS:
 [c,i]                                tmp <- 1   # insn
---------------------------------------------------------------------------
>>> import pyopencl as cl
>>> c = cl._csc()
No protocol specified
Choose platform:
[0] <pyopencl.Platform 'Intel(R) OpenCL' at 0x1df1b20>
[1] <pyopencl.Platform 'Portable Computing Language' at 0x7f36ab1e4400>
[2] <pyopencl.Platform 'AMD Accelerated Parallel Processing' at 0x7f36a6ab6430>
Choice [0]:1
Set the environment variable PYOPENCL_CTX='1' to avoid being asked again.
>>> q = cl.CommandQueue(c)
>>> knl(queue)
Traceback (most recent call last):
  File "<stdin>", line 1, in <module>
NameError: name 'queue' is not defined
>>> lp.generate_code_v2(knl)
/home/matt/src/loopy/loopy/target/pyopencl.py:141: LoopyAdvisory: in kernel loopy_kernel: No device parameter was passed to the PyOpenCLTarget. Perhaps you want to pass a device to benefit from additional checking. (add 'no_device_in_pre_codegen_checks' to silenced_warnings kernel attribute to disable)
  "additional checking.", LoopyAdvisory)
/home/matt/src/loopy/loopy/target/pyopencl.py:431: UserWarning: loopy_kernel: device not supplied to PyOpenCLTarget--workarounds for broken OpenCL implementations (such as those relating to complex numbers) may not be enabled when needed
  .format(knl_name=kernel.name))
-------------------------------------------------------------------------------
CODE:
-------------------------------------------------------------------------------
__kernel void __attribute__ ((reqd_work_group_size(1, 1, 1))) loopy_kernel()
{
  int tmp;

  for (int i = 0; i <= 10; ++i)
    for (int c = 0; c <= i; ++c)
      tmp = 1;
}

-------------------------------------------------------------------------------
Traceback (most recent call last):
  File "<stdin>", line 1, in <module>
  File "/home/matt/src/loopy/loopy/codegen/__init__.py", line 482, in generate_code_v2
    device_code_str)
  File "/home/matt/src/loopy/loopy/check.py", line 742, in check_implemented_domains
    % (insn_id, insn_impl_domain, desired_domain, "\n".join(lines)))
loopy.diagnostic.LoopyError: sanity check failed--implemented and desired domain for instruction 'insn' do not match

implemented: { [i, c] : 0 <= c <= 10 }

desired:{ [i, c] : i <= 10 and 0 <= c <= i }

sample point in implemented but not desired: c=0, i=11
gist of constraints in desired but not implemented: { [i, c] : i <= 10 and c <= i }