<html> <head> <base href="https://bugs.freedesktop.org/"> </head> <body> <div> <a class="bz_bug_link bz_status_NEW " title="NEW - Make Theano OpenCL support work on Clover and RadeonSI" href="https://bugs.freedesktop.org/show_bug.cgi?id=100105#c6">Comment # 6</a> on <a class="bz_bug_link bz_status_NEW " title="NEW - Make Theano OpenCL support work on Clover and RadeonSI" href="https://bugs.freedesktop.org/show_bug.cgi?id=100105">bug 100105</a> from <a class="email" href="mailto:ben@besd.de" title="ben@besd.de <ben@besd.de>"> ben@besd.de</a> <pre>Seems the error is still there: CLOVER_DEVICE_VERSION_OVERRIDE=1.2 CLOVER_DEVICE_CLC_VERSION_OVERRIDE=1.2 DEVICE="opencl0:0" python3 -c "import pygpu;pygpu.test()" fails with: pygpu is installed in /usr/local/lib/python3.6/dist-packages/pygpu-0.7.6+20.g9cec614-py3.6-linux-x86_64.egg/pygpu NumPy version 1.16.3 NumPy relaxed strides checking option: True NumPy is installed in /home/nano/.local/lib/python3.6/site-packages/numpy Python version 3.6.7 (default, Oct 22 2018, 11:32:17) [GCC 8.2.0] nose version 1.3.7 *** Testing for Radeon RX 560 Series (POLARIS11, DRM 3.30.0, 5.1.0-rc5+, LLVM 8.0.0) mpi4py found: True ................................................. ======================================================== AN INTERNAL KERNEL BUILD ERROR OCCURRED! device name = Radeon RX 560 Series (POLARIS11, DRM 3.30.0, 5.1.0-rc5+, LLVM 8.0.0) error = -43 memory pattern = Register accumulation based swap, computing kernel generator Subproblem dimensions: dims[0].itemY = 32, dims[0].itemX = 32, dims[0].y = 32, dims[0].x = 32, dims[0].bwidth = 64; ; dims[1].itemY = 4, dims[1].itemX = 4, dims[1].y = 4, dims[1].x = 4, dims[1].bwidth = 8; ; Parallelism granularity: pgran->wgDim = 1, pgran->wgSize[0] = 64, pgran->wgSize[1] = 1, pgran->wfSize = 64 Kernel extra flags: 369130144 Source: #ifdef DOUBLE_PRECISION #ifdef cl_khr_fp64 #pragma OPENCL EXTENSION cl_khr_fp64 : enable #else #pragma OPENCL EXTENSION cl_amd_fp64 : enable #endif #endif __kernel void Sdot_kernel( __global float *_X, __global float *_Y, __global float *scratchBuff, uint N, uint offx, int incx, uint offy, int incy, int doConj ) { __global float *X = _X + offx; __global float *Y = _Y + offy; float dotP = (float) 0.0; if ( incx < 0 ) { X = X + (N - 1) * abs(incx); } if ( incy < 0 ) { Y = Y + (N - 1) * abs(incy); } int gOffset; for( gOffset=(get_global_id(0) * 4); (gOffset + 4 - 1)<N; gOffset+=( get_global_size(0) * 4 ) ) { float4 vReg1, vReg2, res; #ifdef INCX_NONUNITY vReg1 = (float4)( (X + (gOffset*incx))[0 + ( incx * 0)], (X + (gOffset*incx))[0 + ( incx * 1)], (X + (gOffset*incx))[0 + ( incx * 2)], (X + (gOffset*incx))[0 + ( incx * 3)]); #else vReg1 = vload4( 0, (__global float *) (X + gOffset) ); #endif #ifdef INCY_NONUNITY vReg2 = (float4)( (Y + (gOffset*incy))[0 + ( incy * 0)], (Y + (gOffset*incy))[0 + ( incy * 1)], (Y + (gOffset*incy))[0 + ( incy * 2)], (Y + (gOffset*incy))[0 + ( incy * 3)]); #else vReg2 = vload4( 0, (__global float *) (Y + gOffset) ); #endif ; res = vReg1 * vReg2 ; dotP += res .S0 + res .S1 + res .S2 + res .S3; ; // Add-up elements in the vector to give a scalar } // Loop for the last thread to handle the tail part of the vector // Using the same gOffset used above for( ; gOffset<N; gOffset++ ) { float sReg1, sReg2, res; sReg1 = X[gOffset * incx]; sReg2 = Y[gOffset * incy]; ; res = sReg1 * sReg2 ; dotP = dotP + res ; } // Note: this has to be called outside any if-conditions- because REDUCTION uses barrier // dotP of work-item 0 will have the final reduced item of the work-group __local float bixzI [ 64 ]; uint yBrfY = get_local_id(0); bixzI [ yBrfY ] = dotP ; barrier(CLK_LOCAL_MEM_FENCE); if( yBrfY < 32 ) { bixzI [ yBrfY ] = bixzI [ yBrfY ] + bixzI [ yBrfY + 32 ]; } barrier(CLK_LOCAL_MEM_FENCE); if( yBrfY < 16 ) { bixzI [ yBrfY ] = bixzI [ yBrfY ] + bixzI [ yBrfY + 16 ]; } barrier(CLK_LOCAL_MEM_FENCE); if( yBrfY < 8 ) { bixzI [ yBrfY ] = bixzI [ yBrfY ] + bixzI [ yBrfY + 8 ]; } barrier(CLK_LOCAL_MEM_FENCE); if( yBrfY < 4 ) { bixzI [ yBrfY ] = bixzI [ yBrfY ] + bixzI [ yBrfY + 4 ]; } barrier(CLK_LOCAL_MEM_FENCE); if( yBrfY < 2 ) { bixzI [ yBrfY ] = bixzI [ yBrfY ] + bixzI [ yBrfY + 2 ]; } barrier(CLK_LOCAL_MEM_FENCE); if( yBrfY == 0 ) { dotP = bixzI [0] + bixzI [1]; } if( (get_local_id(0)) == 0 ) { scratchBuff[ get_group_id(0) ] = dotP; } } -------------------------------------------------------- Build log: ======================================================== [nano2:28210] *** Process received signal *** [nano2:28210] Signal: Segmentation fault (11) [nano2:28210] Signal code: Address not mapped (1) [nano2:28210] Failing at address: (nil) [nano2:28210] [ 0] /lib/x86_64-linux-gnu/libc.so.6(+0x3ef20)[0x7fbff3a90f20] [nano2:28210] [ 1] /usr/lib/x86_64-linux-gnu/libclBLAS.so(makeKernelCached+0x2a0)[0x7fbf9eaefcf0] [nano2:28210] [ 2] /usr/lib/x86_64-linux-gnu/libclBLAS.so(makeSolutionSeq+0x101b)[0x7fbf9eaf445b] [nano2:28210] [ 3] /usr/lib/x86_64-linux-gnu/libclBLAS.so(doDot+0x2b2)[0x7fbf9ead7c52] [nano2:28210] [ 4] /usr/lib/x86_64-linux-gnu/libclBLAS.so(clblasSdot+0x98)[0x7fbf9ead7da8] [nano2:28210] [ 5] /home/nano/.local/lib/libgpuarray.so.3(+0x32529)[0x7fbff23aa529] [nano2:28210] [ 6] /home/nano/.local/lib/libgpuarray.so.3(GpuArray_rdot+0x393)[0x7fbff23879f3] [nano2:28210] [ 7] /usr/local/lib/python3.6/dist-packages/pygpu-0.7.6+20.g9cec614-py3.6-linux-x86_64.egg/pygpu/blas.cpython-36m-x86_64-linux-gnu.so(+0x6032)[0x7fbf9c0fa032] [nano2:28210] [ 8] /usr/local/lib/python3.6/dist-packages/pygpu-0.7.6+20.g9cec614-py3.6-linux-x86_64.egg/pygpu/blas.cpython-36m-x86_64-linux-gnu.so(+0x67ba)[0x7fbf9c0fa7ba] [nano2:28210] [ 9] python3[0x5030d5] [nano2:28210] [10] python3(_PyEval_EvalFrameDefault+0x1231)[0x507641] [nano2:28210] [11] python3[0x504c28] [nano2:28210] [12] python3[0x58650d] [nano2:28210] [13] python3(PyObject_Call+0x3e)[0x59ebbe] [nano2:28210] [14] python3(_PyEval_EvalFrameDefault+0x1807)[0x507c17] [nano2:28210] [15] python3[0x504c28] [nano2:28210] [16] python3[0x58644b] [nano2:28210] [17] python3(PyObject_Call+0x3e)[0x59ebbe] [nano2:28210] [18] python3(_PyEval_EvalFrameDefault+0x1807)[0x507c17] [nano2:28210] [19] python3[0x502209] [nano2:28210] [20] python3[0x502f3d] [nano2:28210] [21] python3(_PyEval_EvalFrameDefault+0x449)[0x506859] [nano2:28210] [22] python3[0x504c28] [nano2:28210] [23] python3(_PyFunction_FastCallDict+0x2de)[0x501b2e] [nano2:28210] [24] python3[0x591461] [nano2:28210] [25] python3(PyObject_Call+0x3e)[0x59ebbe] [nano2:28210] [26] python3(_PyEval_EvalFrameDefault+0x1807)[0x507c17] [nano2:28210] [27] python3[0x504c28] [nano2:28210] [28] python3(_PyFunction_FastCallDict+0x2de)[0x501b2e] [nano2:28210] [29] python3[0x591461] [nano2:28210] *** End of error message *** Speicherzugriffsfehler (Speicherabzug geschrieben)</pre> </div> <hr> You are receiving this mail because: <ul> <li>You are the assignee for the bug.</li> <li>You are the QA Contact for the bug.</li> </ul> </body> </html>