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.
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.
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.