doe300 / VC4C

Compiler for the VC4CL OpenCL implementation
MIT License
117 stars 37 forks source link

Generating invalid code about vector rotation. #142

Closed long-long-float closed 4 years ago

long-long-float commented 4 years ago

When VC4C compiles the following code, an validation error is occurred.

#define I(x, y) ((y) * width / 4 + (x))
__kernel void stencil(int width, int height, __global uchar *in, __global uchar *out)
{
    for (int x = 0; x < width / 4; x++) {
        for (int y = 1; y < height - 1; y++) {
            uchar4 right = (uchar4)(0);
            uchar16 center = vload16(I(x, y), in);

            uchar16 r = (uchar16)(center.s456789AB, center.sCDEF, right);

            vstore16(r, I(x, y), out);

        }
    }
}
[E] Wed Mar 25 18:04:41 2020: Validation-error 'Should not write to the source r0 of the vector rotation in the pr
evious instruction.' in: v8min r1, r0, r0 << 12 (full)
[E] Wed Mar 25 18:04:41 2020: With reference to instruction: or r0, r0, r0

But, the following code can be compiled without errors.

#define I(x, y) ((y) * width / 4 + (x))
__kernel void stencil(int width, int height, __global uchar *in, __global uchar *out)
{
    for (int x = 0; x < width / 4; x++) {
        for (int y = 1; y < height - 1; y++) {
            uchar4 left  = (uchar4)(0);
            uchar4 right = (uchar4)(0);
            uchar16 up   = vload16(I(x, y - 1), in);
            uchar16 down = vload16(I(x, y + 1), in);
            uchar16 center = vload16(I(x, y), in);

            uchar16 r = (
                up                                               / (uchar16)(5) +
                (uchar16)(left, center.s01234567, center.s89AB)  / (uchar16)(5) +
                center                                           / (uchar16)(5) +
                (uchar16)(center.s456789AB, center.sCDEF, right) / (uchar16)(5) + // HERE is the same code.
                down                                             / (uchar16)(5));

            vstore16(r, I(x, y), out);

        }
    }
}
doe300 commented 4 years ago

Seems like I am unable to reproduce the issue on my host PC...

What version of VC4C and LLVM are you using? Can you please post the output of vc4c --version?

long-long-float commented 4 years ago

Here is output of ./build/src/vc4c --version.

Running VC4C in version: 0.4.9999
Build configuration:
debug mode; multi-threaded optimization; clang 3.9+ OpenCL features; clang in /usr/bin/clang; LLVM library front-end with libLLVM 5; custom VC4CL standard-library path /home/***/vc4c/VC4C-netdist/../VC4CLStdLib/include/; LLVM linker; vc4asm verification
long-long-float commented 4 years ago

The above code is compiled successfully. Thank you!