./lib/llvmopencl/Kernel.cc:129: pocl::ParallelRegion* pocl::Kernel::createParallelRegionBefore(llvm::BasicBlock*): Assertion `region_entry_barrier != NULL' failed

[picca]

Hello, I am affected by this bug while preparing the Debian package of silx

https://bugs.debian.org/cgi-bin/bugreport.cgi?bug=1060318

reading a bunch of your issues I found that adding

POCL_WORK_GROUP_METHOD=cbs helps saolve the issue or at leat dos not trigger the core dump.

so my question is what is the right way to solve this issue ?

thanks for considering

[pjaaskel]

Hi. If you can create a reproducer in C/C++ with as small kernel as possible that causes it, it wold be helpful in trying to tackle the underlying problem. Meanwhile you can use the CBS for this test case.

[picca]

Here some information

silx-kit/silx#4073

It seems that this is a llvm issue...

llvm 15 ok
llvm 16 ko

[kif]

I simplified the culprit kernel to a certain extent ... This does not do any sensible thing anymore but crashes the compiler:

__kernel void medfilt2d(__global float *image,  // input image
                        __global float *result, // output array
                        __local  float4 *l_data,// local storage 4x the number of threads
                                 int khs1,      // Kernel half-size along dim1 (nb lines)
                                 int khs2,      // Kernel half-size along dim2 (nb columns)
                                 int height,    // Image size along dim1 (nb lines)
                                 int width)     // Image size along dim2 (nb columns)
{
    int threadid = get_local_id(0);
    int x = get_global_id(1);

    if (x < width)
    {
        union
        {
            float  ary[8];
            float8 vec;
        } output, input;
        input.vec = (float8)(MAXFLOAT, MAXFLOAT, MAXFLOAT, MAXFLOAT, MAXFLOAT, MAXFLOAT, MAXFLOAT, MAXFLOAT);
        int kfs1 = 2 * khs1 + 1; 
        int kfs2 = 2 * khs2 + 1;
        int nbands = (kfs1 + 7) / 8; 
        for (int y=0; y<height; y++)
        {
            //Select only the active threads, some may remain inactive
            int nb_threads =  (nbands * kfs2);
            int band_nr = threadid / kfs2;
            int band_id = threadid % kfs2;
            int pos_x = clamp((int)(x + band_id - khs2), (int) 0, (int) width-1);
            int max_vec = clamp(kfs1 - 8 * band_nr, 0, 8);
            if (y == 0)
            {
                for (int i=0; i<max_vec; i++)
                {
                    if (threadid<nb_threads)
                    {
                        int pos_y = clamp((int)(y + 8 * band_nr + i - khs1), (int) 0, (int) height-1);
                        input.ary[i] = image[pos_x + width * pos_y];
                    }
                }
            }
            else
            {
                //store storage.s0 to some shared memory to retrieve it from another thread.
                l_data[threadid].s0 = input.vec.s0;

                //Offset to the bottom
                input.vec = (float8)(input.vec.s1,
                        input.vec.s2,
                        input.vec.s3,
                        input.vec.s4,
                        input.vec.s5,
                        input.vec.s6,
                        input.vec.s7,
                        MAXFLOAT);

                barrier(CLK_LOCAL_MEM_FENCE);

                int read_from = threadid + kfs2;
                if (read_from < nb_threads)
                    input.vec.s7 = l_data[read_from].s0;
                else if (threadid < nb_threads) //we are on the last band
                {
                    int pos_y = clamp((int)(y + 8 * band_nr + max_vec - 1 - khs1), (int) 0, (int) height-1);
                    input.ary[max_vec - 1] = image[pos_x + width * pos_y];
                }

            }

            barrier(CLK_LOCAL_MEM_FENCE);
        }
    }
}

It is the last barrier which is the cause of the crash.

If one removes the if (y==0) or the else, they do compile.

[pjaaskel]

And this compiles with later LLVMs, but not older than 15?

[kif]

This compiles with LLVM15 and elder (as old as we remember ... the code is from 2017) but it fails with LLVM16.

[anbe42]

This is unrelated to the LLVM version being used.
In Debian bookworm we have PoCL 3.1 built against LLVM 15 not showing the bug.
In Debian sid we have PoCL 5.0 built against LLVM 16 showing the bug.
I just rebuilt PoCL 5.0 locally against LLVM 15 and it still shows the bug.

[anbe42]

I bisected it to commit b165b29 "LLVM 17 support", unfortunately that has
77 files changed, 4410 insertions(+), 2341 deletions(-)
I just rebuilt PoCL 5.0 against LLVM 17 locally and that still shows the bug.

[anbe42]

minimized nonsensical kernel to reproduce the problem (reduced with cvise and some manual postprocessing)

__kernel void a(__local float2 *b) {
  struct {
    int c[1];
    float2 d;
  } e;
  for (int f = 0; f < 2; f++) {
    if (f)
      for (int g; g < b; g++)
        e.c[g] = 0;
    else if (b)
      e.d.s0 = b[0].s0;
    barrier(0);
  }
}

[pjaaskel]

Thanks! Very helpful. Also a great pointer to the cvise tool - looks very useful. @franz this is likely something to do with the moving to the new PM by default? Do we still have the ability to run with the old PM to test?

[franz]

Since SVMOffset pass was added without support for legacy PM, main branch does not support the legacy PM anymore.

PoCL 5.0 should be able to run with old PM, it's enough to change MIN_LLVM_NEW_PASSMANAGER in config.h.in.cmake to some higher number (this will only work up to LLVM 16).

[pjaaskel]

The SVMOffset pass is not used in the client-side kernel compilation pipeline, but only called from command line from the pocld server side, so it shouldn't affect this.

[franz]

@picca @kif - could any of you test the "real" kernel in silx with this branch - it should fix the crash, but i'm interested if it produces correct results.

[kif]

Hi Michal, I tested the default ICD (probaly the nvidia one): ``` (py311) lintaillefer:~/workspace/silx % ./run_tests.py src/silx/image/test/test_medianfilter.py INFO:silx.setup:Install requires: numpy >=1.26.4 Building silx to /users/kieffer/workspace-400/silx/build/lib.linux-x86_64-3.11 INFO:silx.setup:Install requires: numpy >=1.26.4 running build running build_py running build_ext INFO: Disabling color, you really want to install colorlog. INFO:pythran:Disabling color, you really want to install colorlog. Patched sys.path, added: '/users/kieffer/workspace-400/silx/build/lib.linux-x86_64-3.11' ==================================================================================== test session starts ===================================================================================== platform linux -- Python 3.11.0, pytest-7.2.2, pluggy-1.0.0 rootdir: /users/kieffer/workspace-400/silx/build/lib.linux-x86_64-3.11/silx, configfile: ../../../pytest.ini plugins: anyio-3.7.1, xvfb-2.0.0 collected 2 items build/lib.linux-x86_64-3.11/silx/image/test/test_medianfilter.py .. [100%] ===================================================================================== 2 passed in 0.46s ====================================================================================== ``` Then the POCL v5 which is known to crash the process: ``` (py311) lintaillefer:~/workspace/silx % OCL_ICD_VENDORS=/opt/pocl5/etc/OpenCL/vendors/pocl.icd ./run_tests.py src/silx/image/test/test_medianfilter.py INFO:silx.setup:Install requires: numpy >=1.26.4 Building silx to /users/kieffer/workspace-400/silx/build/lib.linux-x86_64-3.11 INFO:silx.setup:Install requires: numpy >=1.26.4 running build running build_py running build_ext INFO: Disabling color, you really want to install colorlog. INFO:pythran:Disabling color, you really want to install colorlog. Patched sys.path, added: '/users/kieffer/workspace-400/silx/build/lib.linux-x86_64-3.11' ==================================================================================== test session starts ===================================================================================== platform linux -- Python 3.11.0, pytest-7.2.2, pluggy-1.0.0 rootdir: /users/kieffer/workspace-400/silx/build/lib.linux-x86_64-3.11/silx, configfile: ../../../pytest.ini plugins: anyio-3.7.1, xvfb-2.0.0 collected 2 items build/lib.linux-x86_64-3.11/silx/image/test/test_medianfilter.py . [crashed here] ``` And finally the patched version from branch fix 1435: ``` (py311) lintaillefer:~/workspace/silx % OCL_ICD_VENDORS=/opt/pocl6/etc/OpenCL/vendors/pocl.icd ./run_tests.py src/silx/image/test/test_medianfilter.py INFO:silx.setup:Install requires: numpy >=1.26.4 Building silx to /users/kieffer/workspace-400/silx/build/lib.linux-x86_64-3.11 INFO:silx.setup:Install requires: numpy >=1.26.4 running build running build_py running build_ext INFO: Disabling color, you really want to install colorlog. INFO:pythran:Disabling color, you really want to install colorlog. Patched sys.path, added: '/users/kieffer/workspace-400/silx/build/lib.linux-x86_64-3.11' ==================================================================================== test session starts ===================================================================================== platform linux -- Python 3.11.0, pytest-7.2.2, pluggy-1.0.0 rootdir: /users/kieffer/workspace-400/silx/build/lib.linux-x86_64-3.11/silx, configfile: ../../../pytest.ini plugins: anyio-3.7.1, xvfb-2.0.0 collected 2 items build/lib.linux-x86_64-3.11/silx/image/test/test_medianfilter.py .. [100%] ===================================================================================== 2 passed in 0.13s ====================================================================================== ``` So apparently the result looks good. Thanks for the patch. It looks like it fixes the issue. Jérôme Kieffer

[franz]

fixed in main branch