Bug 109617

please don't use text upload sites. either post in a comment or add as an attachment.

kernel code:

#ifndef TEST_KERNEL_CL
#define TEST_KERNEL_CL

__kernel void
test_kernel_1(__global float4 *result)
{
    bool flag = false;
    //uint flag = false;
    uint gid = get_global_id(0);
    float rnd;
    uint i = 0;
    
#define VARIANT 1
#if (VARIANT == 1)
    while ((i < 2) && (flag == false)){
        rnd = sin((float)(gid + i)) + 0.1;
    
        if ((rnd * rnd)<=0.5){flag=true;}
        i++;
    }
#elif (VARIANT == 2)
    for (i = 0; i < 2; i++){
        if (flag == false){
        rnd = sin((float)(gid + i)) + 0.1;
    
        if ((rnd * rnd)<=0.5){flag=true; i++; break;}
        }
    }
#endif
    
    if(flag){    
        result[gid].x = rnd;
    }
            
    result[gid].y = (float)i;
}

#endif

results:

xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx
Correct execution    xxxxxxxxxxxxxxxxxxxxxx
xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx
-------------------------------------------
bool flag = false;
#define VARIANT 2
---
:::::::::::::::::::::::::::::::::::::::::::::::::::::
Platform: Clover
Device: AMD OLAND (DRM 2.50.0 / 4.12.14-lp150.12.45-default, LLVM 5.0.1) (GPU)
:::::::::::::::::::::::::::::::::::::::::::::::::::::
Number of test_kernels: 1
----------
[ 0.100000,  1.000000,  0.000000,  0.000000]
[ 0.000000,  2.000000,  0.000000,  0.000000]
[ 0.241120,  2.000000,  0.000000,  0.000000]

-------------------------------------------
uint flag = false;
#define VARIANT 1
---
:::::::::::::::::::::::::::::::::::::::::::::::::::::
Platform: Clover
Device: AMD OLAND (DRM 2.50.0 / 4.12.14-lp150.12.45-default, LLVM 5.0.1) (GPU)
:::::::::::::::::::::::::::::::::::::::::::::::::::::
Number of test_kernels: 1
----------
[ 0.100000,  1.000000,  0.000000,  0.000000]
[ 0.000000,  2.000000,  0.000000,  0.000000]
[ 0.241120,  2.000000,  0.000000,  0.000000]

xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx
Incorrect execution    xxxxxxxxxxxxxxxxxxxx
xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx
-------------------------------------------
bool flag = false;
#define VARIANT 1
---
:::::::::::::::::::::::::::::::::::::::::::::::::::::
Platform: Clover
Device: AMD OLAND (DRM 2.50.0 / 4.12.14-lp150.12.45-default, LLVM 5.0.1) (GPU)
:::::::::::::::::::::::::::::::::::::::::::::::::::::
Number of test_kernels: 1
----------
[ 0.000000,  1.000000,  0.000000,  0.000000]
[ 0.000000,  2.000000,  0.000000,  0.000000]
[ 0.241120,  2.000000,  0.000000,  0.000000]

xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx

This looks like clang/llvm kernel miscompile.
Can you reproduce using more recent llvm version (ideally git)?

I have downloaded LLVM from https://github.com/llvm-mirror/llvm, produced Makefile with CMake, did 
# make
# make install
# reboot

so that

# l /usr/local/bin/ | grep "llvm"
-rwxr-xr-x  1 root root  174338016 фев 15 02:21 llvm-ar*
-rwxr-xr-x  1 root root  141089976 фев 15 03:32 llvm-as*
-rwxr-xr-x  1 root root    7825712 фев 15 03:32 llvm-bcanalyzer*
-rwxr-xr-x  1 root root  141473560 фев 15 03:38 llvm-cat*
-rwxr-xr-x  1 root root  213087112 фев 15 03:39 llvm-cfi-verify*
-rwxr-xr-x  1 root root    6324176 фев 15 02:21 llvm-config*
-rwxr-xr-x  1 root root  106796744 фев 15 03:39 llvm-cov*
-rwxr-xr-x  1 root root 1138615128 фев 15 03:38 llvm-c-test*
-rwxr-xr-x  1 root root   93546704 фев 15 03:40 llvm-cvtres*
-rwxr-xr-x  1 root root   96751584 фев 15 03:40 llvm-cxxdump*
-rwxr-xr-x  1 root root    7013808 фев 15 03:40 llvm-cxxfilt*
-rwxr-xr-x  1 root root    8395432 фев 15 03:40 llvm-cxxmap*
-rwxr-xr-x  1 root root   69219032 фев 15 03:40 llvm-diff*
-rwxr-xr-x  1 root root   62994736 фев 15 03:40 llvm-dis*
lrwxrwxrwx  1 root root          7 фев 15 08:20 llvm-dlltool -> llvm-ar*
-rwxr-xr-x  1 root root  167804064 фев 15 03:40 llvm-dwarfdump*
-rwxr-xr-x  1 root root 1208619744 фев 15 03:57 llvm-dwp*
-rwxr-xr-x  1 root root   93318584 фев 15 03:57 llvm-elfabi*
-rwxr-xr-x  1 root root  538498472 фев 15 03:59 llvm-exegesis*
-rwxr-xr-x  1 root root  171478416 фев 15 04:00 llvm-extract*
lrwxrwxrwx  1 root root          7 фев 15 08:20 llvm-lib -> llvm-ar*
-rwxr-xr-x  1 root root  171099032 фев 15 04:15 llvm-link*
-rwxr-xr-x  1 root root 1275040600 фев 15 02:37 llvm-lto*
-rwxr-xr-x  1 root root 1268033504 фев 15 04:32 llvm-lto2*
-rwxr-xr-x  1 root root  153453544 фев 15 04:33 llvm-mc*
-rwxr-xr-x  1 root root  147927736 фев 15 04:33 llvm-mca*
-rwxr-xr-x  1 root root  136001376 фев 15 04:34 llvm-modextract*
-rwxr-xr-x  1 root root    7430880 фев 15 04:34 llvm-mt*
-rwxr-xr-x  1 root root  173515424 фев 15 04:34 llvm-nm*
-rwxr-xr-x  1 root root  100393360 фев 15 04:34 llvm-objcopy*
-rwxr-xr-x  1 root root  198631080 фев 15 04:35 llvm-objdump*
-rwxr-xr-x  1 root root    8992552 фев 15 04:52 llvm-opt-report*
-rwxr-xr-x  1 root root  141847808 фев 15 04:53 llvm-pdbutil*
-rwxr-xr-x  1 root root   53838256 фев 15 02:37 llvm-profdata*
lrwxrwxrwx  1 root root          7 фев 15 08:20 llvm-ranlib -> llvm-ar*
-rwxr-xr-x  1 root root   10510000 фев 15 04:53 llvm-rc*
lrwxrwxrwx  1 root root         12 фев 15 08:26 llvm-readelf -> llvm-readobj*
-rwxr-xr-x  1 root root  109591424 фев 15 04:54 llvm-readobj*
-rwxr-xr-x  1 root root  182818080 фев 15 04:54 llvm-rtdyld*
-rwxr-xr-x  1 root root   92759976 фев 15 04:54 llvm-size*
-rwxr-xr-x  1 root root  166803512 фев 15 04:55 llvm-split*
-rwxr-xr-x  1 root root   56764944 фев 15 04:55 llvm-stress*
-rwxr-xr-x  1 root root    6193784 фев 15 04:55 llvm-strings*
lrwxrwxrwx  1 root root         12 фев 15 08:26 llvm-strip -> llvm-objcopy*
-rwxr-xr-x  1 root root  120943912 фев 15 04:55 llvm-symbolizer*
-rwxr-xr-x  1 root root   57481560 фев 14 23:40 llvm-tblgen*
-rwxr-xr-x  1 root root    6472472 фев 15 04:55 llvm-undname*
-rwxr-xr-x  1 root root  131050384 фев 15 04:56 llvm-xray*

# l /usr/local/include/ | grep "llvm"
drwxr-xr-x 35 root root 4096 фев 15 08:19 llvm/
drwxr-xr-x  3 root root 4096 фев 15 08:19 llvm-c/

"фев" means "Feb".

However, the result is the same:
> ./program 
PlatformID: 0
DeviceID: 0
:::::::::::::::::::::::::::::::::::::::::::::::::::::
Platform: Clover
Device: AMD OLAND (DRM 2.50.0 / 4.12.14-lp150.12.45-default, LLVM 5.0.1) (GPU)
:::::::::::::::::::::::::::::::::::::::::::::::::::::
Number of test_kernels: 1
----------
[ 0.000000,  1.000000,  0.000000,  0.000000]
[ 0.000000,  2.000000,  0.000000,  0.000000]
[ 0.241120,  2.000000,  0.000000,  0.000000]

I am confused that there still stays "LLVM 5.0.1". Is it correct?

For where ever you have clover installed, can you check that it's linked against the right LLVM? If you've just installed an updated LLVM, it's likely that clover is still running against the old version.

E.g.:
#Determine name of mesa opencl driver library
$ cat /etc/OpenCL/vendors/mesa.icd
 libMesaOpenCL.so.1

#If you know where your clover install is, great, otherwise:
$ find / -type f -name "libMesaOpenCL.so.1" -exec ldd {} \;
$ ldd ${PATH_TO_MESA_OPENCL_SO}

If it says it's linked against LLVM 5, then you'll probably have to rebuild mesa against the newer LLVM.

My output (trimmed a bit):
find /usr/local/lib -type f -name "libMesaOpenCL.so.1*" -exec ldd {} \;
 libdl.so.2 => /lib/x86_64-linux-gnu/libdl.so.2 (0x00007f8eee84f000)
 libelf.so.1 => /usr/lib/x86_64-linux-gnu/libelf.so.1 (0x00007f8eee834000)
 libdrm.so.2 => /usr/local/lib/libdrm.so.2 (0x00007f8eee7e3000)
 libLLVM-9svn.so => /usr/local/lib/libLLVM-9svn.so (0x00007f8eeafad000)

Unfortunately, I will not have access to this machine at the nearest time. When I have it, I think, I will check this.

-- GitLab Migration Automatic Message --

This bug has been migrated to freedesktop.org's GitLab instance and has been closed from further activity.

You can subscribe and participate further through the new bug through this link to our GitLab instance: https://gitlab.freedesktop.org/mesa/mesa/issues/933.