{"action":"created","issue":{"url":"https://api.github.com/repos/magnumripper/JohnTheRipper/issues/1571","labels_url":"https://api.github.com/repos/magnumripper/JohnTheRipper/issues/1571/labels{/name}","comments_url":"https://api.github.com/repos/magnumripper/JohnTheRipper/issues/1571/comments","events_url":"https://api.github.com/repos/magnumripper/JohnTheRipper/issues/1571/events","html_url":"https://github.com/magnumripper/JohnTheRipper/issues/1571","id":96686904,"number":1571,"title":"Some failing opencl format tests with beignet on haswell GPU","user":{"login":"frank-dittrich","id":2026819,"avatar_url":"https://avatars.githubusercontent.com/u/2026819?v=3","gravatar_id":"","url":"https://api.github.com/users/frank-dittrich","html_url":"https://github.com/frank-dittrich","followers_url":"https://api.github.com/users/frank-dittrich/followers","following_url":"https://api.github.com/users/frank-dittrich/following{/other_user}","gists_url":"https://api.github.com/users/frank-dittrich/gists{/gist_id}","starred_url":"https://api.github.com/users/frank-dittrich/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/frank-dittrich/subscriptions","organizations_url":"https://api.github.com/users/frank-dittrich/orgs","repos_url":"https://api.github.com/users/frank-dittrich/repos","events_url":"https://api.github.com/users/frank-dittrich/events{/privacy}","received_events_url":"https://api.github.com/users/frank-dittrich/received_events","type":"User","site_admin":false},"labels":[{"url":"https://api.github.com/repos/magnumripper/JohnTheRipper/labels/portability","name":"portability","color":"fbca04"}],"state":"open","locked":false,"assignee":null,"milestone":null,"comments":15,"created_at":"2015-07-22T23:05:44Z","updated_at":"2015-09-22T18:11:03Z","closed_at":null,"body":"During my experiments to get the Haswell GPU working, I probably found some beignet / libdrm / i915 driver issues and may be some bleeding-jumbo issues.\r\n\r\nOn my 64bit Fedora 22 with Haswell CPU Intel(R) Core(TM) i5-4570 CPU @ 3.20GHz, after installing beignet (I tried both the fedora package and the latest http://anongit.freedesktop.org/git/beignet.git commit):\r\n\r\nHere, I am not sure whether the `multiple unsequenced modifications to 'p'` warning indicates a real problem:\r\n```\r\n$ OCL_IGNORE_SELF_TEST=1 ./john --test=0 --format=oldoffice-opencl\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\ndrm_intel_gem_bo_context_exec() failed: Invalid argument\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nBeignet: self-test failed: (3, 7, 5) + (5, 7, 3) returned (3, 7, 5)\r\nSee README.md or http://www.freedesktop.org/wiki/Software/Beignet/\r\nBeignet: warning - disable atomic in L3 feature.\r\nBeignet: Warning - overriding self-test failure\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nBeignet: Warning - overriding self-test failure\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nTesting: oldoffice-opencl, MS Office <= 2003 [MD5/SHA1 RC4 OpenCL]... Build log: /tmp/fxgfQc.cl:224:24: warning: multiple unsequenced modifications to 'p'\r\n/tmp/fxgfQc.cl:441:15: warning: multiple unsequenced modifications to 'p'\r\n\r\nFAILED (cmp_all(1))\r\n\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\n```\r\n\r\nThe only statements that may trigger that `multiple unsequenced modifications to 'p'` warning IMHO are\r\n```\r\noldoffice_kernel.cl:194:                        W[i >> 1] = (uint)*p++ | (*p++ << 16U);\r\noldoffice_kernel.cl:197:                        W[i >> 1] = (uint)*p++ | (*p++ << 16U);\r\noldoffice_kernel.cl:210:                        W[i >> 1] = (uint)*p++ | (*p++ << 16U);\r\noldoffice_kernel.cl:224:                        W[i >> 1] = (uint)*p++ | (*p++ << 16U);\r\noldoffice_kernel.cl:408:                        uint u = *p++ | (*p++ << 16U);\r\noldoffice_kernel.cl:413:                        uint u = *p++ | (*p++ << 16U);\r\noldoffice_kernel.cl:426:                        uint u = *p++ | (*p++ << 16U);\r\noldoffice_kernel.cl:441:                        uint u = *p++ | (*p++ << 16U);\r\n```\r\nCan any of the above be a problem?\r\n\r\n\r\n\r\nI hope that these warnings\r\n```\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\n```\r\nwill disappear after a future libdrm or linux kernel upgrade and that the\r\n```\r\ndrm_intel_gem_bo_context_exec() failed: Invalid argument\r\n```\r\nproblem will be fixed in future beignet versions.\r\n\r\nMay be when I get rid of the `disable atomic in L3 feature` due to the failed self test and when I disable the debug CFLAGS from my beignet build, the performance will improve...\r\n\r\nFor now, I just use OCL_IGNORE_SELF_TEST=1, otherwise the Haswell GPU will be skipped:\r\n```\r\n$ ./john --list=opencl-devices \r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\ndrm_intel_gem_bo_context_exec() failed: Invalid argument\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nBeignet: self-test failed: (3, 7, 5) + (5, 7, 3) returned (3, 7, 5)\r\nSee README.md or http://www.freedesktop.org/wiki/Software/Beignet/\r\nBeignet: warning - disable atomic in L3 feature.\r\nBeignet: disabling non-working device\r\nOpenCL error (CL_DEVICE_NOT_FOUND) in file (common-opencl.c) at line (385) - (No OpenCL device of that type exist)\r\n```\r\n\r\n```\r\n$ OCL_IGNORE_SELF_TEST=1 ./john --list=opencl-devices             \r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\ndrm_intel_gem_bo_context_exec() failed: Invalid argument\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nBeignet: self-test failed: (3, 7, 5) + (5, 7, 3) returned (3, 7, 5)\r\nSee README.md or http://www.freedesktop.org/wiki/Software/Beignet/\r\nBeignet: warning - disable atomic in L3 feature.\r\nBeignet: Warning - overriding self-test failure\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nBeignet: Warning - overriding self-test failure\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nPlatform #0 name: Intel Gen OCL Driver\r\nPlatform version: OpenCL 1.2 beignet 1.1 (git-834d0ae)\r\n\tDevice #0 (0) name:\tIntel(R) HD Graphics Haswell GT2 Desktop\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\n\tDevice vendor:\t\tIntel\r\n\tDevice type:\t\tGPU (LE)\r\n\tDevice version:\t\tOpenCL 1.2 beignet 1.1 (git-834d0ae)\r\n\tDriver version:\t\t1.1\r\n\tNative vector widths:\tchar 8, short 8, int 4, long 2\r\n\tPreferred vector width:\tchar 16, short 8, int 4, long 2\r\n\tGlobal Memory:\t\t2.0 GB\r\n\tGlobal Memory Cache:\t8.0 KB\r\n\tLocal Memory:\t\t64.0 KB (Global)\r\n\tMax memory alloc. size:\t1024.0 MB\r\n\tMax clock (MHz):\t1000\r\n\tProfiling timer res.:\t80 ns\r\n\tMax Work Group Size:\t512\r\n\tParallel compute cores:\t20\r\n\r\n```\r\n\r\n```\r\n$ OCL_IGNORE_SELF_TEST=1 ./john --list=opencl-devices --verbosity=5\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\ndrm_intel_gem_bo_context_exec() failed: Invalid argument\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nBeignet: self-test failed: (3, 7, 5) + (5, 7, 3) returned (3, 7, 5)\r\nSee README.md or http://www.freedesktop.org/wiki/Software/Beignet/\r\nBeignet: warning - disable atomic in L3 feature.\r\nBeignet: Warning - overriding self-test failure\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nBeignet: Warning - overriding self-test failure\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nPlatform #0 name: Intel Gen OCL Driver\r\nPlatform version: OpenCL 1.2 beignet 1.1 (git-834d0ae)\r\n\tPlatform extensions:\tcl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_byte_addressable_store cl_khr_spir cl_khr_icd\r\n\tDevice #0 (0) name:\tIntel(R) HD Graphics Haswell GT2 Desktop\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\n\tDevice vendor:\t\tIntel\r\n\tDevice type:\t\tGPU (LE)\r\n\tDevice version:\t\tOpenCL 1.2 beignet 1.1 (git-834d0ae)\r\n\tDriver version:\t\t1.1\r\n\tNative vector widths:\tchar 8, short 8, int 4, long 2\r\n\tPreferred vector width:\tchar 16, short 8, int 4, long 2\r\n\tGlobal Memory:\t\t2.0 GB\r\n\tDevice extensions:\tcl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_byte_addressable_store cl_khr_spir cl_khr_icd\r\n\tGlobal Memory Cache:\t8.0 KB\r\n\tLocal Memory:\t\t64.0 KB (Global)\r\n\tMax memory alloc. size:\t1024.0 MB\r\n\tMax clock (MHz):\t1000\r\n\tProfiling timer res.:\t80 ns\r\n\tMax Work Group Size:\t512\r\n\tParallel compute cores:\t20\r\n```\r\n\r\nMost opencl formats pass self test, with c/s rate between 10% and 90% of the CPU (4*OMP).\r\nFor rar, opencl gets closest:\r\n```\r\n$ ./john --test --format=rar\r\nWill run 4 OpenMP threads\r\nBenchmarking: rar, RAR3 (4 characters) [SHA1 AES 32/64]... (4xOMP) DONE\r\nRaw:\t403 c/s real, 100 c/s virtual\r\n\r\n$ OCL_IGNORE_SELF_TEST=1 ./john --test --format=rar-opencl --verbosity=5\r\ninitUnicode(UNICODE, ASCII/ASCII)\r\nASCII -> ASCII -> ASCII\r\nWill run 4 OpenMP threads\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\ndrm_intel_gem_bo_context_exec() failed: Invalid argument\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nBeignet: self-test failed: (3, 7, 5) + (5, 7, 3) returned (3, 7, 5)\r\nSee README.md or http://www.freedesktop.org/wiki/Software/Beignet/\r\nBeignet: warning - disable atomic in L3 feature.\r\nBeignet: Warning - overriding self-test failure\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nBeignet: Warning - overriding self-test failure\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nBenchmarking: rar-opencl, RAR3 (length 5) [SHA1 OpenCL AES]... (4xOMP) Options used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=1 -D_OPENCL_COMPILER -DPLAINTEXT_LENGTH=22\r\nCalculating best global worksize (GWS); max. 10s total for crypt_all()\r\nRaw speed figures including buffer transfers:\r\nkey xfer: 115.517ms*, len xfer: 115.533ms*, init: 2.800us, loop: 16x52.082ms, final: 19.920us, key xfer: 101.292ms*, iv xfer: 107.760us*\r\ngws:       256\t       219 c/s    57409536 rounds/s    1.165s per crypt_all()!\r\nkey xfer: 199.327ms*, len xfer: 199.329ms*, init: 1.680us, loop: 16x99.704ms, final: 19.520us, key xfer: 199.272ms*, iv xfer: 88.480us*\r\ngws:       512\t       233 c/s    61079552 rounds/s    2.193s per crypt_all()+\r\nkey xfer: 364.582ms*, len xfer: 364.586ms*, init: 2.320us, loop: 16x182.919ms, final: 35.040us, key xfer: 364.498ms*, iv xfer: 73.920us*\r\ngws:      1024\t       254 c/s    66584576 rounds/s    4.020s per crypt_all()+\r\nkey xfer: 1.440us, len xfer: 1.280us, init: 2.800us, loop: 16x349.420ms, final: 36us, key xfer: 1.280us, iv xfer: 1.200us\r\ngws:      2048\t       366 c/s    95944704 rounds/s    5.590s per crypt_all()+\r\nkey xfer: 1.440us, len xfer: 1.120us, init: 5.760us, loop: 16x687.015ms, final: 65.600us, key xfer: 1.120us, iv xfer: 10.575ms*\r\ngws:      4096\t       372 c/s    97517568 rounds/s   11.002s per crypt_all() - too slow\r\nLocal worksize (LWS) 64, global worksize (GWS) 2048\r\nDONE\r\nRaw:\t364 c/s real, 102400 c/s virtual\r\n\r\n```\r\n\r\nBTW: why is the CPU format using length 4, but the GPU length 5?\r\n\r\nAnd these are all the failing self tests:\r\n```\r\ninitUnicode(UNICODE, ASCII/ASCII)\r\nASCII -> ASCII -> ASCII\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nBenchmarking: oldoffice-opencl, MS Office <= 2003 [MD5/SHA1 RC4 OpenCL]... Options used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=1 -D_OPENCL_COMPILER -DASCII -DPLAINTEXT_LENGTH=19\r\nBuild log: /tmp/DAFnPU.cl:224:24: warning: multiple unsequenced modifications to 'p'\r\n/tmp/DAFnPU.cl:441:15: warning: multiple unsequenced modifications to 'p'\r\n\r\nCalculating best global worksize (GWS); max. 1s total for crypt_all()\r\nRaw speed figures including buffer transfers:\r\nxP: 1.746ms*, xI: 1.747ms*, enc: 53.760us, md5+rc4: 1.466ms, xR: 1.543ms*\r\ngws:      1024\t    156160 c/s      156160 rounds/s   6.557ms per crypt_all()!\r\nxP: 3.068ms*, xI: 3.072ms*, enc: 39.360us, md5+rc4: 2.868ms, xR: 3.004ms*\r\ngws:      2048\t    169902 c/s      169902 rounds/s  12.054ms per crypt_all()+\r\nxP: 5.818ms*, xI: 5.805ms*, enc: 75.280us, md5+rc4: 5.614ms, xR: 5.729ms*\r\ngws:      4096\t    177750 c/s      177750 rounds/s  23.043ms per crypt_all()+\r\nxP: 10.606ms*, xI: 10.581ms*, enc: 146.080us, md5+rc4: 10.284ms, xR: 10.468ms*\r\ngws:      8192\t    194643 c/s      194643 rounds/s  42.087ms per crypt_all()+\r\nxP: 20.173ms*, xI: 20.117ms*, enc: 285.600us, md5+rc4: 19.524ms, xR: 19.923ms*\r\ngws:     16384\t    204735 c/s      204735 rounds/s  80.025ms per crypt_all()+\r\nxP: 38.848ms*, xI: 38.835ms*, enc: 595.600us, md5+rc4: 37.806ms, xR: 38.401ms*\r\ngws:     32768\t    212107 c/s      212107 rounds/s 154.487ms per crypt_all()+\r\nxP: 75.759ms*, xI: 75.554ms*, enc: 1.200ms, md5+rc4: 73.755ms, xR: 75.026ms*\r\ngws:     65536\t    217513 c/s      217513 rounds/s 301.295ms per crypt_all()+\r\nxP: 149.298ms*, xI: 148.891ms*, enc: 2.294ms, md5+rc4: 145.494ms, xR: 147.853ms*\r\ngws:    131072\t    220722 c/s      220722 rounds/s 593.831ms per crypt_all()+\r\nxP: 1.440us, xI: 298.903ms*, enc: 4.523ms, md5+rc4: 291.022ms, xR: 295.578ms*\r\ngws:    262144\t    294534 c/s      294534 rounds/s 890.028ms per crypt_all()+\r\nxP: 1.920us, xI: 1.520us, enc: 9.227ms, md5+rc4: 576.541ms, xR: 1.360us\r\ngws:    524288\t    895035 c/s      895035 rounds/s 585.773ms per crypt_all()!\r\nxP: 2.240us, xI: 1.280us, enc: 18.194ms, md5+rc4: 1.155s, xR: 1.440us\r\ngws:   1048576\t    893037 c/s      893037 rounds/s    1.174s per crypt_all() - too slow\r\nLocal worksize (LWS) 64, global worksize (GWS) 524288\r\nFAILED (cmp_all(1))\r\n\r\ninitUnicode(UNICODE, ASCII/ASCII)\r\nASCII -> ASCII -> ASCII\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nBenchmarking: lotus5-opencl, Lotus Notes/Domino 5 [OpenCL]... Options used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=1 -D_OPENCL_COMPILER \r\nCalculating best global worksize (GWS); max. 1s single kernel invocation.\r\nRaw speed figures including buffer transfers:\r\nxfer: 833.920us*, crypt: 487.520us, xfer: 550.880us*\r\ngws:       256\t    136728 c/s      136728 rounds/s   1.872ms per crypt_all()!\r\nxfer: 886.800us*, crypt: 752.320us, xfer: 861.200us*\r\ngws:       512\t    204773 c/s      204773 rounds/s   2.500ms per crypt_all()+\r\nxfer: 1.773ms*, crypt: 1.641ms, xfer: 1.738ms*\r\ngws:      1024\t    198683 c/s      198683 rounds/s   5.153ms per crypt_all()\r\nxfer: 3.115ms*, crypt: 2.980ms, xfer: 3.087ms*\r\ngws:      2048\t    223010 c/s      223010 rounds/s   9.183ms per crypt_all()+\r\nxfer: 6.011ms*, crypt: 5.858ms, xfer: 5.957ms*\r\ngws:      4096\t    229759 c/s      229759 rounds/s  17.827ms per crypt_all()+\r\nxfer: 11.924ms*, crypt: 11.736ms, xfer: 11.839ms*\r\ngws:      8192\t    230756 c/s      230756 rounds/s  35.500ms per crypt_all()\r\nxfer: 23.686ms*, crypt: 23.426ms, xfer: 23.538ms*\r\ngws:     16384\t    231898 c/s      231898 rounds/s  70.651ms per crypt_all()\r\nxfer: 46.946ms*, crypt: 46.512ms, xfer: 46.656ms*\r\ngws:     32768\t    233863 c/s      233863 rounds/s 140.115ms per crypt_all()+\r\nxfer: 1.040us, crypt: 92.609ms, xfer: 92.877ms*\r\ngws:     65536\t    353316 c/s      353316 rounds/s 185.487ms per crypt_all()+\r\nxfer: 1.360us, crypt: 185.871ms, xfer: 186.244ms*\r\ngws:    131072\t    352232 c/s      352232 rounds/s 372.117ms per crypt_all()\r\nxfer: 1.360us, crypt: 370.840ms, xfer: 371.800ms*\r\ngws:    262144\t    352988 c/s      352988 rounds/s 742.642ms per crypt_all()\r\nxfer: 1.040us, crypt: 743.818ms, xfer: 745.505ms*\r\ngws:    524288\t    352030 c/s      352030 rounds/s    1.489s per crypt_all()\r\nxfer: 2.400us, crypt: 1.484s (exceeds 1s)\r\nLocal worksize (LWS) 64, global worksize (GWS) 65536\r\nFAILED (cmp_all(1))\r\n\r\ninitUnicode(UNICODE, ASCII/ASCII)\r\nASCII -> ASCII -> ASCII\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nOptions used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=1 -D_OPENCL_COMPILER -DWORK_GROUP_SIZE=8\r\nLocal worksize (LWS) 8, Global worksize (GWS) 1024\r\nBenchmarking: bcrypt-opencl (\"$2a$05\", 32 iterations) [Blowfish OpenCL]... FAILED (cmp_all(1))\r\n\r\ninitUnicode(UNICODE, ASCII/ASCII)\r\nASCII -> ASCII -> ASCII\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nOptions used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=1 -D_OPENCL_COMPILER \r\nASSERTION FAILED: 0\r\n  at file /home/fd/git/beignet/backend/src/backend/gen_context.cpp, function virtual void gbe::GenContext::emitUnaryWithTempInstruction(const gbe::SelectionInstruction&), line 438\r\ninitUnicode(UNICODE, ASCII/ASCII)\r\nASCII -> ASCII -> ASCII\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nOptions used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=1 -D_OPENCL_COMPILER \r\nOptions used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=1 -D_OPENCL_COMPILER \r\nLocal worksize (LWS) 64, Global worksize (GWS) 8192\r\nBenchmarking: descrypt-opencl, traditional crypt(3) [DES OpenCL]... FAILED (cmp_one(0))\r\n\r\ninitUnicode(UNICODE, ASCII/ASCII)\r\nASCII -> ASCII -> ASCII\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nBenchmarking: krb5pa-md5-opencl, Kerberos 5 AS-REQ Pre-Auth etype 23 [MD4 HMAC-MD5 RC4 OpenCL]... Options used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=1 -D_OPENCL_COMPILER -DASCII -DPLAINTEXT_LENGTH=27\r\nCalculating best global worksize (GWS); max. 200ms single kernel invocation.\r\nRaw speed figures including buffer transfers:\r\nxfer: 599.920us*, init: 35.280us, crypt: 187.360us, xfer: 251.280us*\r\ngws:       256\t    238396 c/s      238396 rounds/s   1.073ms per crypt_all()!\r\nxfer: 315.920us*, init: 21.360us, crypt: 217.360us, xfer: 285.120us*\r\ngws:       512\t    609698 c/s      609698 rounds/s 839.760us per crypt_all()!\r\nxfer: 517.840us*, init: 38.720us, crypt: 420.240us, xfer: 487.600us*\r\ngws:      1024\t    699262 c/s      699262 rounds/s   1.464ms per crypt_all()+\r\nxfer: 1.057ms*, init: 76us, crypt: 852.880us, xfer: 1.022ms*\r\ngws:      2048\t    680651 c/s      680651 rounds/s   3.008ms per crypt_all()\r\nxfer: 1.868ms*, init: 152.800us, crypt: 1.574ms, xfer: 1.805ms*\r\ngws:      4096\t    758282 c/s      758282 rounds/s   5.401ms per crypt_all()+\r\nxfer: 3.460ms*, init: 257.680us, crypt: 3.044ms, xfer: 3.368ms*\r\ngws:      8192\t    808616 c/s      808616 rounds/s  10.130ms per crypt_all()+\r\nxfer: 6.591ms*, init: 538.880us, crypt: 5.843ms, xfer: 6.458ms*\r\ngws:     16384\t    843086 c/s      843086 rounds/s  19.433ms per crypt_all()+\r\nxfer: 12.555ms*, init: 1.016ms, crypt: 11.198ms, xfer: 12.309ms*\r\ngws:     32768\t    883728 c/s      883728 rounds/s  37.079ms per crypt_all()+\r\nxfer: 24.328ms*, init: 2.031ms, crypt: 21.744ms, xfer: 23.894ms*\r\ngws:     65536\t    910238 c/s      910238 rounds/s  71.998ms per crypt_all()+\r\nxfer: 1.200us, init: 4.020ms, crypt: 42.825ms, xfer: 47.095ms*\r\ngws:    131072\t   1395235 c/s     1395235 rounds/s  93.942ms per crypt_all()+\r\nxfer: 1.040us, init: 7.988ms, crypt: 85.503ms, xfer: 93.441ms*\r\ngws:    262144\t   1402330 c/s     1402330 rounds/s 186.934ms per crypt_all()\r\nxfer: 2.480us, init: 15.840ms, crypt: 168.811ms, xfer: 1.360us\r\ngws:    524288\t   2839267 c/s     2839267 rounds/s 184.656ms per crypt_all()!\r\nxfer: 1.760us, init: 32.312ms, crypt: 336.594ms (exceeds 200ms)\r\nLocal worksize (LWS) 64, global worksize (GWS) 524288\r\nFAILED (cmp_all(1))\r\n\r\ninitUnicode(UNICODE, ASCII/ASCII)\r\nASCII -> ASCII -> ASCII\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nBenchmarking: nt-opencl, NT [MD4 OpenCL]... Options used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=1 -D_OPENCL_COMPILER -D OFFSET_TABLE_SIZE=32 -D HASH_TABLE_SIZE=43 -D SHIFT64_OT_SZ=0 -D SHIFT64_HT_SZ=41 -D NUM_LOADED_HASHES=41 -D NUM_INT_KEYS=1 -D SELECT_CMP_STEPS=4 -D BITMAP_SIZE_BITS_LESS_ONE=32767 -D USE_LOCAL_BITMAPS=1 -D IS_STATIC_GPU_MASK=0 -D CONST_CACHE_SIZE=134217728 -D LOC_0=-1 -D LOC_1=-1 -D LOC_2=-1 -D LOC_3=-1\r\nSelf test GWS: 2097152, LWS: 256\r\nFAILED (cmp_all(0))\r\n\r\ninitUnicode(UNICODE, ASCII/ASCII)\r\nASCII -> ASCII -> ASCII\r\nWill run 4 OpenMP threads\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nBenchmarking: office2013-opencl, MS Office 2013 (100,000 iterations) [SHA512 OpenCL 2x AES]... (4xOMP) Options used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=1 -D_OPENCL_COMPILER -DHASH_LOOPS=100 -DUNICODE_LENGTH=96 -DV_WIDTH=2\r\nASSERTION FAILED: 0\r\n  at file /home/fd/git/beignet/backend/src/backend/gen_context.cpp, function virtual void gbe::GenContext::emitUnaryWithTempInstruction(const gbe::SelectionInstruction&), line 438\r\ninitUnicode(UNICODE, ASCII/ASCII)\r\nASCII -> ASCII -> ASCII\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nBenchmarking: PBKDF2-HMAC-SHA512-opencl, GRUB2 / OS X 10.8+, rounds=10000 [PBKDF2-SHA512 OpenCL]... Options used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=1 -D_OPENCL_COMPILER -DHASH_LOOPS=250 -DPLAINTEXT_LENGTH=110 -DMAX_SALT_SIZE=107\r\nASSERTION FAILED: 0\r\n  at file /home/fd/git/beignet/backend/src/backend/gen_context.cpp, function virtual void gbe::GenContext::emitUnaryWithTempInstruction(const gbe::SelectionInstruction&), line 438\r\ninitUnicode(UNICODE, ASCII/ASCII)\r\nASCII -> ASCII -> ASCII\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nBenchmarking: Raw-MD4-opencl [MD4 OpenCL]... Options used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=1 -D_OPENCL_COMPILER -D OFFSET_TABLE_SIZE=13 -D HASH_TABLE_SIZE=15 -D SHIFT64_OT_SZ=3 -D SHIFT64_HT_SZ=1 -D NUM_LOADED_HASHES=14 -D NUM_INT_KEYS=1 -D SELECT_CMP_STEPS=4 -D BITMAP_SIZE_BITS_LESS_ONE=32767 -D USE_LOCAL_BITMAPS=1 -D IS_STATIC_GPU_MASK=0 -D CONST_CACHE_SIZE=134217728 -D LOC_0=-1 -D LOC_1=-1 -D LOC_2=-1 -D LOC_3=-1\r\nSelf test GWS: 1048576, LWS: 256\r\nFAILED (cmp_all(0))\r\n\r\ninitUnicode(UNICODE, ASCII/ASCII)\r\nASCII -> ASCII -> ASCII\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nBenchmarking: Raw-MD5-opencl [MD5 OpenCL]... Options used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=1 -D_OPENCL_COMPILER -D OFFSET_TABLE_SIZE=11 -D HASH_TABLE_SIZE=5 -D SHIFT64_OT_SZ=5 -D SHIFT64_HT_SZ=1 -D NUM_LOADED_HASHES=5 -D NUM_INT_KEYS=1 -D SELECT_CMP_STEPS=4 -D BITMAP_SIZE_BITS_LESS_ONE=32767 -D USE_LOCAL_BITMAPS=1 -D IS_STATIC_GPU_MASK=0 -D CONST_CACHE_SIZE=134217728 -D LOC_0=-1 -D LOC_1=-1 -D LOC_2=-1 -D LOC_3=-1\r\nSelf test GWS: 1048576, LWS: 256\r\nFAILED (cmp_all(0))\r\n\r\ninitUnicode(UNICODE, ASCII/ASCII)\r\nASCII -> ASCII -> ASCII\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nBenchmarking: Raw-SHA1-opencl [SHA1 OpenCL]... Options used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=1 -D_OPENCL_COMPILER -D OFFSET_TABLE_SIZE=11 -D HASH_TABLE_SIZE=5 -D SHIFT64_OT_SZ=5 -D SHIFT64_HT_SZ=1 -D SHIFT128_OT_SZ=3 -D SHIFT128_HT_SZ=1 -D NUM_LOADED_HASHES=4 -D NUM_INT_KEYS=1 -D SELECT_CMP_STEPS=4 -D BITMAP_SIZE_BITS_LESS_ONE=32767 -D USE_LOCAL_BITMAPS=1 -D IS_STATIC_GPU_MASK=0 -D CONST_CACHE_SIZE=134217728 -D LOC_0=-1 -D LOC_1=-1 -D LOC_2=-1 -D LOC_3=-1\r\nSelf test GWS: 1048576, LWS: 256\r\nFAILED (cmp_all(0))\r\n\r\ninitUnicode(UNICODE, ASCII/ASCII)\r\nASCII -> ASCII -> ASCII\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nOptions used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=1 -D_OPENCL_COMPILER \r\nBenchmarking: Raw-SHA256-opencl [SHA256 OpenCL]... Calculating best global worksize (GWS); max. 1s total for crypt_all()\r\nRaw speed figures including buffer transfers:\r\npass xfer: 595.280us*, crypt: 404.160us, result xfer: 464.720us*, index xfer: 618.720us*\r\ngws:      1024\t    491626 c/s      491626 rounds/s   2.082ms per crypt_all()!\r\npass xfer: 827.120us*, crypt: 739.440us, result xfer: 785.360us*, index xfer: 829.840us*\r\ngws:      2048\t    643668 c/s      643668 rounds/s   3.181ms per crypt_all()+\r\npass xfer: 1.628ms*, crypt: 1.454ms, result xfer: 1.566ms*, index xfer: 1.614ms*\r\ngws:      4096\t    653920 c/s      653920 rounds/s   6.263ms per crypt_all()+\r\npass xfer: 3.012ms*, crypt: 2.780ms, result xfer: 2.892ms*, index xfer: 2.955ms*\r\ngws:      8192\t    703731 c/s      703731 rounds/s  11.640ms per crypt_all()+\r\npass xfer: 5.786ms*, crypt: 5.491ms, result xfer: 5.542ms*, index xfer: 5.627ms*\r\ngws:     16384\t    729864 c/s      729864 rounds/s  22.448ms per crypt_all()+\r\npass xfer: 11.123ms*, crypt: 10.653ms, result xfer: 10.753ms*, index xfer: 10.843ms*\r\ngws:     32768\t    755478 c/s      755478 rounds/s  43.373ms per crypt_all()+\r\npass xfer: 21.932ms*, crypt: 21.075ms, result xfer: 21.183ms*, index xfer: 21.338ms*\r\ngws:     65536\t    766222 c/s      766222 rounds/s  85.531ms per crypt_all()+\r\npass xfer: 42.630ms*, crypt: 42.106ms, result xfer: 42.217ms*, index xfer: 42.391ms*\r\ngws:    131072\t    773992 c/s      773992 rounds/s 169.345ms per crypt_all()+\r\npass xfer: 84.888ms*, crypt: 83.984ms, result xfer: 84.102ms*, index xfer: 84.428ms*\r\ngws:    262144\t    776943 c/s      776943 rounds/s 337.404ms per crypt_all()\r\npass xfer: 1.120us, crypt: 166.969ms, result xfer: 167.046ms*, index xfer: 167.826ms*\r\ngws:    524288\t   1044726 c/s     1044726 rounds/s 501.842ms per crypt_all()+\r\npass xfer: 338.504ms*, crypt: 335.524ms, result xfer: 335.601ms*, index xfer: 336.525ms*\r\ngws:   1048576\t    778940 c/s      778940 rounds/s    1.346s per crypt_all() - too slow\r\nMax local worksize 512, Local worksize (LWS) 32, global worksize (GWS) 524288\r\nFAILED (cmp_all(0))\r\n\r\ninitUnicode(UNICODE, ASCII/ASCII)\r\nASCII -> ASCII -> ASCII\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nOptions used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=1 -D_OPENCL_COMPILER \r\nASSERTION FAILED: 0\r\n  at file /home/fd/git/beignet/backend/src/backend/gen_context.cpp, function virtual void gbe::GenContext::emitUnaryWithTempInstruction(const gbe::SelectionInstruction&), line 438\r\ninitUnicode(UNICODE, ASCII/ASCII)\r\nASCII -> ASCII -> ASCII\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nBenchmarking: wpapsk-opencl, WPA/WPA2 PSK [PBKDF2-SHA1 OpenCL 4x]... Options used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=1 -D_OPENCL_COMPILER -DHASH_LOOPS=105 -DITERATIONS=4095 -DPLAINTEXT_LENGTH=63 -DV_WIDTH=4\r\nASSERTION FAILED: !(ctx->getErrCode() == OUT_OF_RANGE_IF_ENDIF && ctx->getIFENDIFFix())\r\n  at file /home/fd/git/beignet/backend/src/backend/gen_program.cpp, function virtual gbe::Kernel* gbe::GenProgram::compileKernel(const gbe::ir::Unit&, const string&, bool), line 196\r\ninitUnicode(UNICODE, ASCII/ASCII)\r\nASCII -> ASCII -> ASCII\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nOptions used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=1 -D_OPENCL_COMPILER \r\nASSERTION FAILED: 0\r\n  at file /home/fd/git/beignet/backend/src/backend/gen_context.cpp, function virtual void gbe::GenContext::emitUnaryWithTempInstruction(const gbe::SelectionInstruction&), line 438\r\n```\r\n\r\nThe common warnings\r\n```\r\ndrm_intel_gem_bo_context_exec() failed: Invalid argument\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nBeignet: self-test failed: (3, 7, 5) + (5, 7, 3) returned (3, 7, 5)\r\nSee README.md or http://www.freedesktop.org/wiki/Software/Beignet/\r\nBeignet: warning - disable atomic in L3 feature.\r\nBeignet: Warning - overriding self-test failure\r\n```\r\nwhich appeared for every single test have been filtered using `grep -v`.\r\n\r\n@magnum I wasn't sure whether I should create separate issues for every possible problem (failing format self tests etc.)\r\n\r\nMay be beignet, libdrm or the i915 driver are to blame for many of these.\r\nIf you think I should have created multiple issues, I suggest that we keep this one for the `multiple unsequenced modifications to 'p'` warning in oldoffice-opencl (and may be the failing oldoffice-opencl self test), and I'll create new issues for the other failing self tests."},"comment":{"url":"https://api.github.com/repos/magnumripper/JohnTheRipper/issues/comments/142368537","html_url":"https://github.com/magnumripper/JohnTheRipper/issues/1571#issuecomment-142368537","issue_url":"https://api.github.com/repos/magnumripper/JohnTheRipper/issues/1571","id":142368537,"user":{"login":"frank-dittrich","id":2026819,"avatar_url":"https://avatars.githubusercontent.com/u/2026819?v=3","gravatar_id":"","url":"https://api.github.com/users/frank-dittrich","html_url":"https://github.com/frank-dittrich","followers_url":"https://api.github.com/users/frank-dittrich/followers","following_url":"https://api.github.com/users/frank-dittrich/following{/other_user}","gists_url":"https://api.github.com/users/frank-dittrich/gists{/gist_id}","starred_url":"https://api.github.com/users/frank-dittrich/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/frank-dittrich/subscriptions","organizations_url":"https://api.github.com/users/frank-dittrich/orgs","repos_url":"https://api.github.com/users/frank-dittrich/repos","events_url":"https://api.github.com/users/frank-dittrich/events{/privacy}","received_events_url":"https://api.github.com/users/frank-dittrich/received_events","type":"User","site_admin":false},"created_at":"2015-09-22T18:11:03Z","updated_at":"2015-09-22T18:11:03Z","body":"With the latest beignet commit http://cgit.freedesktop.org/beignet/commit/?id=18a52ffc966027a3004b85c7c03c9416e1a84c3a the bug https://bugs.freedesktop.org/show_bug.cgi?id=91525 seems to be fixed.\r\nI don't get ASSERTION FAILED anymore.\r\n\r\nHowever, this is what happens now:\r\n```\r\n$ time ./john --test --format=office2013-opencl \r\nWill run 4 OpenMP threads\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nBenchmarking: office2013-opencl, MS Office 2013 (100,000 iterations) [SHA512 OpenCL 2x AES]... (4xOMP) \r\ntotal:1 hash_loops 100 rounds 100004 looptime 0\r\ndrm_intel_gem_bo_context_exec() failed: Input/output error\r\ndrm_intel_gem_bo_context_exec() failed: Input/output error\r\ndrm_intel_gem_bo_context_exec() failed: Input/output error\r\nOpenCL CL_INVALID_EVENT error in common-opencl.c:1357 - WaitForEvents failed\r\n\r\nreal\t0m33.325s\r\nuser\t0m17.857s\r\nsys\t0m0.074s\r\n```\r\n\r\n```\r\n$ time ./john --test --format=office2013-opencl --verbosity=5\r\ninitUnicode(UNICODE, ASCII/ASCII)\r\nASCII -> ASCII -> ASCII\r\nWill run 4 OpenMP threads\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nFailed to release test userptr object! (9) i915 kernel driver may not be sane!\r\nDevice 0: Intel(R) HD Graphics Haswell GT2 Desktop\r\nBenchmarking: office2013-opencl, MS Office 2013 (100,000 iterations) [SHA512 OpenCL 2x AES]... (4xOMP) Options used: -I ./kernels -cl-mad-enable -D__GPU__ -DDEVICE_INFO=34 -DDEV_VER_MAJOR=1 -DDEV_VER_MINOR=2 -D_OPENCL_COMPILER -DHASH_LOOPS=100 -DUNICODE_LENGTH=96 -DV_WIDTH=2\r\nCalculating best GWS for LWS=32; max. 5s total for crypt_all()\r\nRaw speed figures including buffer transfers:\r\nxfer: 5.385s*, xfer: 5.385s*, init: 1.274ms, loop: 1000x0ns,  final: 6.219ms, xfer: 6.328ms*\r\n\r\ntotal:1 hash_loops 100 rounds 100004 looptime 0\r\ngws:       128\t        23 c/s     2300092 rounds/s   10.785s per crypt_all()!\r\ndrm_intel_gem_bo_context_exec() failed: Input/output error\r\ndrm_intel_gem_bo_context_exec() failed: Input/output error\r\nOpenCL CL_OUT_OF_RESOURCES error in opencl_office2013_fmt_plug.c:328 - failed in clEnqueueNDRangeKernel\r\n (error occurred)Calculating best LWS for GWS=128\r\ndrm_intel_gem_bo_context_exec() failed: Input/output error\r\ndrm_intel_gem_bo_context_exec() failed: Input/output error\r\nOpenCL CL_OUT_OF_RESOURCES error in opencl_office2013_fmt_plug.c:320 - failed in clEnqueueNDRangeKernel\r\ndrm_intel_gem_bo_context_exec() failed: Input/output error\r\ndrm_intel_gem_bo_context_exec() failed: Input/output error\r\nOpenCL CL_OUT_OF_RESOURCES error in opencl_office2013_fmt_plug.c:320 - failed in clEnqueueNDRangeKernel\r\nOpenCL CL_INVALID_EVENT error in common-opencl.c:1357 - WaitForEvents failed\r\n\r\nreal\t0m33.834s\r\nuser\t0m17.863s\r\nsys\t0m0.079s\r\n```\r\n\r\nNot sure whether this is a bleeding-jumbo error, a beignet error, a kernel driver error, or a combination of errors.\r\nAny thoughts?\r\n\r\nI'll try to upgrade to Linux kernel version 4.3.0-0.rc1 and see whether the problem disappears."}}