Bug 109617 - [oland, clover, llvm5] While-If Problem with Booleans
Summary: [oland, clover, llvm5] While-If Problem with Booleans
Status: NEW
Alias: None
Product: Mesa
Classification: Unclassified
Component: Other (show other bugs)
Version: unspecified
Hardware: x86-64 (AMD64) Linux (All)
: medium normal
Assignee: mesa-dev
QA Contact: mesa-dev
URL:
Whiteboard:
Keywords:
Depends on:
Blocks: 99553
  Show dependency treegraph
 
Reported: 2019-02-12 22:44 UTC by Natalia
Modified: 2019-02-27 06:45 UTC (History)
0 users

See Also:
i915 platform:
i915 features:


Attachments

Note You need to log in before you can comment on or make changes to this bug.
Description Natalia 2019-02-12 22:44:04 UTC
OpenCL kernel containing While loop and If statement does not work correctly with boolean variable, while everething is good with integer variable and with For loop.

The kernel is the following: https://textuploader.com/15szh.

The results are the following: https://textuploader.com/15szu.

This is tested at openSUSE 15.0 notebook with AMD Radeon HD 8750M GPU.
Comment 1 Jan Vesely 2019-02-12 22:49:00 UTC
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
Comment 2 Jan Vesely 2019-02-12 22:49:37 UTC
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
Comment 3 Jan Vesely 2019-02-12 23:05:08 UTC
This looks like clang/llvm kernel miscompile.
Can you reproduce using more recent llvm version (ideally git)?
Comment 4 Natalia 2019-02-15 08:29:06 UTC
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?
Comment 5 Aaron Watry 2019-02-27 01:36:06 UTC
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)
Comment 6 Natalia 2019-02-27 06:45:31 UTC
Unfortunately, I will not have access to this machine at the nearest time. When I have it, I think, I will check this.


Use of freedesktop.org services, including Bugzilla, is subject to our Code of Conduct. How we collect and use information is described in our Privacy Policy.