PyOpenCL August 2013

pyopencl@tiker.net

9 participants
14 discussions

Invalid command Queue when using big data sets on nVidia

by Justin Heinermann

Dear all, we are trying to implement a K nearest neighbor search on GPUs with PyOpenCL. The goal of the algorithm is: For a given target point, find the nearest points from a given set (training data). The distance between two points is computed by the squared euclidean distance. One of our implementations is a brute force approach, which aims at processing big data sets in parallel, e.g. 1 million training data and some millions of targets (test data). For every target point one kernel instance is created which finds the k nearest points out of the training points. Our problem is the following. Everything works fine for small data sets and the results are as expected on both GPU (GeForce GTX 650 with nVidia Driver 313.09.) and CPU(Intel Core i5-3450 with AMD APP SDK) running Ubuntu 12.10, PyOpenCL 2013.1-py2.7-linux-x86_64. But if we increase the size of the data sets, the GPU version crashes with the following error: > File "brutegpu.py", line 65, in query > cl.enqueue_copy(self.queue, d_min, self.d_min_buf).wait() > File "/usr/local/lib/python2.7/dist-packages/ > pyopencl-2013.1-py2.7-linux-x86_64.egg/pyopencl/__init__.py", > line 935, in enqueue_copy > return _cl._enqueue_read_buffer(queue, src, dest, **kwargs) > pyopencl.LogicError: clEnqueueReadBuffer failed: invalid command queue The CPU-Version still works fine with 1 million training points and 1 million of test points. Attached you can find the corresponding source code as working minimal example, which consists of on Host-Python-File and one OpenCL-Kernel-File. We would highly apprecriate any help - maybe we made a mistake which is already known to you. So the big question for us is: Why is it working on CPU and why isn't it working on the GPU? Are there nVidia-specific pitfalls for such big data sets? The compiler says: > ptxas info : Compiling entry function 'find_knn' for 'sm_30' > ptxas info : Function properties for find_knn > 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads > ptxas info : Used 17 registers, 336 bytes cmem[0], 4 bytes cmem[3] Or are there any rules for using a kernel for big data sets such as setting the work group sizes or maximum memory usage? The error message "invalid command queue" is confusing and I wasn't able to find any helpful information (except that oftentimes "invalid command queue" means segfault, but i could not find any wrong array adress yet.) Maybe one of you could have a look at our code and finds some stupid mistake. We would be very grateful for every hint. Best regards, Justin Heinermann, University Oldenburg

5 years, 11 months

Segmentation fault in pyopencl.image_from_array

by Jerome Kieffer

Dear Python/OpenCL community, I am pretty new (py)opencl and encountered a problem, maybe it a lack of understanding of openCL, but I found strange python seg-faults: test program: #!/usr/bin/python import numpy, pyopencl ctx = pyopencl.create_some_context() data=numpy.random.random((1024,1024)).astype(numpy.float32) img = pyopencl.image_from_array(ctx, ary=data, mode="r", norm_int=False, num_channels=1) print img System: debian sid: pyopencl2012.1 (the same code works with debian stable and v2011.2) Here is the backtrace obtained with GDB: 0x0000000000000000 in ?? () (gdb) bt #0 0x0000000000000000 in ?? () #1 0x00007ffff340c253 in pyopencl::create_image_from_desc(pyopencl::context const&, unsigned long, _cl_image_format const&, _cl_image_desc&, boost::python::api::object) () from /usr/lib/python2.7/dist-packages/pyopencl/_cl.so #2 0x00007ffff342de36 in _object* boost::python::detail::invoke<boost::python::detail::install_holder<pyopencl::image*>, pyopencl::image* (*)(pyopencl::context const&, unsigned long, _cl_image_format const&, _cl_image_desc&, boost::python::api::object), boost::python::arg_from_python<pyopencl::context const&>, boost::python::arg_from_python<unsigned long>, boost::python::arg_from_python<_cl_image_format const&>, boost::python::arg_from_python<_cl_image_desc&>, boost::python::arg_from_python<boost::python::api::object> >(boost::python::detail::invoke_tag_<false, false>, boost::python::detail::install_holder<pyopencl::image*> const&, pyopencl::image* (*&)(pyopencl::context const&, unsigned long, _cl_image_format const&, _cl_image_desc&, boost::python::api::object), boost::python::arg_from_python<pyopencl::context const&>&, boost::python::arg_from_python<unsigned long>&, boost::python::arg_from_python<_cl_image_format const&>&, boost::python::arg_from_python<_cl_image_desc&>&, boost::python::arg_from_python<boost::python::api::object>&) () from /usr/lib/python2.7/dist-packages/pyopencl/_cl.so #3 0x00007ffff342e06f in boost::python::detail::caller_arity<5u>::impl<pyopencl::image* (*)(pyopencl::context const&, unsigned long, _cl_image_format const&, _cl_image_desc&, boost::python::api::object), boost::python::detail::constructor_policy<boost::python::default_call_policies>, boost::mpl::vector6<pyopencl::image*, pyopencl::context const&, unsigned long, _cl_image_format const&, _cl_image_desc&, boost::python::api::object> >::operator()(_object*, _object*) () from /usr/lib/python2.7/dist-packages/pyopencl/_cl.so #4 0x00007ffff311715b in boost::python::objects::function::call(_object*, _object*) const () from /usr/lib/libboost_python-py27.so.1.49.0 #5 0x00007ffff3117378 in ?? () from /usr/lib/libboost_python-py27.so.1.49.0 #6 0x00007ffff3120593 in boost::python::detail::exception_handler::operator()(boost::function0<void> const&) const () from /usr/lib/libboost_python-py27.so.1.49.0 #7 0x00007ffff3445983 in boost::detail::function::function_obj_invoker2<boost::_bi::bind_t<bool, boost::python::detail::translate_exception<pyopencl::error, void (*)(pyopencl::error const&)>, boost::_bi::list3<boost::arg<1>, boost::arg<2>, boost::_bi::value<void (*)(pyopencl::error const&)> > >, bool, boost::python::detail::exception_handler const&, boost::function0<void> const&>::invoke(boost::detail::function::function_buffer&, boost::python::detail::exception_handler const&, boost::function0<void> const&) () from /usr/lib/python2.7/dist-packages/pyopencl/_cl.so #8 0x00007ffff3120373 in boost::python::handle_exception_impl(boost::function0<void>) () from /usr/lib/libboost_python-py27.so.1.49.0 #9 0x00007ffff3115635 in ?? () from /usr/lib/libboost_python-py27.so.1.49.0 Thanks for your help. If you are not able to reproduce this bug, I should mention it to debian. Cheers, -- Jérôme Kieffer Data analysis unit - ESRF

6 years, 5 months

Re: [PyOpenCL] Windows 8 and PyOpenCL missing pytools

by Andreas Kloeckner

"=?UTF-8?Q?marcus.desto?=" <marcus.desto(a)o2.pl> writes: > Hello Andreas, > > thanks for your mail. > > First of all I would like to say that I had subscribed to you pyopencl mailing list few days before I mailed you first time. But somehow the mailing list did not work. > > However. > > Dnia 11 lipca 2013 23:27 Andreas Kloeckner <lists(a)informa.tiker.net> napisał(a): > >> >> Wrote a howto for you, since I agree that that area of PyOpenCL is >> somewhat underdocumented: >> >> http://documen.tician.de/pyopencl/howto.html > > Thank you a lot. It helped me a lot. > > Now, sticking to the subject, I have installed python 2.7 and 3.3 on my testing windows 8 system using external sources for python addons (ifd.uci.edu/~gohlke/pythonlibs). > > Installation was fine, but as I tried to import pyopencl the python shell complained, there is no module named pytools (ImportError). It occured in pyopencl\__init__.py, line 679 and 448) > > Are there any solutions available solving that problem? Download pytools from here: https://pypi.python.org/pypi/pytools unpack, cd pytools-2013.5.3, python setup.py install. HTH, Andreas

6 years, 11 months

clmathlibraries

by Sam

Hello, I was wondering if anyone has looked at the clblas and clfft libraries that AMD recently open sourced? Would it be possible to have these inside pyopencl or should they be a seperate wrapper on top of it like pyfft? Also are they vendor agnostic? The reason I ask is I suck at opencl but; I am a okay at python and it would be awesome to be able to do cl.math.matrix.multiply(A,B) and have it work. pyopencl is awesome but is still difficult to do work like train a neural net on it. Currently most python gpu neural net libraries are written in/for cuda (theano, cudamat, etc.) this is probably because of the lack of a good BLAS for opencl. So is clBlas a candidate for a python wrapper? github.com/clMathLibraries

6 years, 11 months

PyOpenCL on PyPy

by Matti Picus

Hi. I am interested in getting PyOpenCL to work with PyPy, an implementation of cpython with a JITwww.pypy.org . Has there been any discussion or thought about doing this? PyPy has a basic implementation of numpy called numpypy that I contribute to, and it has a rudimentary numpy-compatible c interface available as an external module at https://bitbucket.org/antocuni/numpypy_c The PyPy team has a cpython-compatible replacement for ctypes called cffi, that is jit-friendly on PyPy and no slower than ctypes on cpython. So it seems like all the pieces exist to start, is anyone else interested in getting the work done? Or are there blocking issues I do not understand? Matti

6 years, 11 months

pyOpenCL advice needed

by Pongsak Suvanpong

Dear All I have a code a Multivariable Linear Regression in python using numpy show below. I am trying to learn pyOpenCL by writing the same algorithm in OpenCL(code shown below), some how no matter I tried(even using very large samples 200 features and 5000 samples), the openCL code always runs slower than serial code using numpy. I wonder if anyone could point out what have I down wrong. My setup is Mac OSX 10.8.4, video card NVIDIA GTX650, python3, numpy171 and pyopen check out from github. """ Multivariable Linear regression by psksvp(a)gmail.com based on Andrew Ng lecture (https://class.coursera.org/ml-003/lecture/index) """ import numpy as np """ Multivariables Linear Regression, Currently there is no bias term. class MultivariablesLinearRegressionSerial(object): def __init__(self, numberOfFeatures): self.parameterVector = np.array(np.random.random(numberOfFeatures)) def recall(self, inputVector): return np.dot(inputVector, self.parameterVector) def error(self, inputVector, target): return self.recall(inputVector) - target def learn(self, sampleMatrix, targetVector, learningRate=0.05, convergeAtError=0.001): #if(len(sampleMatrix) != 2) #error checking mSamples = sampleMatrix.shape[0] ## number of samples #nFeatures = sampleMatrix.shape[1] ## number of features howmuch = (learningRate * (1.0 / float(mSamples))) converge = False mse = 0.0 count = 0 while not converge: output = np.dot(sampleMatrix, self.parameterVector) #print(sampleMatrix) #print(self.parameterVector) #print(output) error = output - targetVector gradient = np.dot(error, sampleMatrix) self.parameterVector = self.parameterVector - (gradient * howmuch) error = np.dot(sampleMatrix, self.parameterVector) - targetVector mse = np.dot(error, error) print(mse) if(mse <= convergeAtError): converge = True else: count = count + 1 return count ####################################################### import numpy as np #import pyopencl as ocl import pyopencl.array as oclArray from psksvp import MathOCL ################################################################################### #### openCL ################################################################################## class MultivariablesLinearRegressionOpenCL(object): def __init__(self, numberOfFeatures, openCLContext, queue): self.parameterVector = np.array(np.random.random(numberOfFeatures), dtype=np.float32) #self.parameterVector = np.zeros(shape=numberOfFeatures, dtype=np.float32) self.openCLContext = openCLContext self.queue = queue def recall(self, inputVector): return np.dot(inputVector, self.parameterVector) def error(self, inputVector, target): return self.recall(inputVector) - target def learn(self, sampleMatrix, targetVector, learningRate=0.05, convergeAtError=0.001): #if(len(sampleMatrix) != 2) #error checking mSamples = sampleMatrix.shape[0] ## number of samples nFeatures = sampleMatrix.shape[1] ## number of features howMuch = (learningRate * (1.0 / float(mSamples))) converge = False #mse = 0.0 deviceParameterVector = oclArray.to_device(self.queue, self.parameterVector.T) deviceTargetVector = oclArray.to_device(self.queue, targetVector) deviceSampleMatrix = oclArray.to_device(self.queue, sampleMatrix) deviceGradient = oclArray.zeros(self.queue, shape=nFeatures, dtype=np.float32) deviceError = oclArray.zeros(self.queue, shape=mSamples, dtype=np.float32) deviceOutput = oclArray.zeros(self.queue, shape=mSamples, dtype=np.float32) #deviceConvergeAtError = oclArray.to_device(self.queue, np.array(convergeAtError)) outputCalculator = MathOCL.MatrixMultipleKernel(AWidth=nFeatures, AHeight=mSamples, BWidth=1, BHeight=nFeatures, OpenCLContext=self.openCLContext, queue=self.queue) gradientCalculator = MathOCL.MatrixMultipleKernel(AWidth=mSamples, AHeight=1, BWidth=nFeatures, BHeight=mSamples, OpenCLContext=self.openCLContext, queue=self.queue) vectorToVectorSubtractor = MathOCL.VectorToVectorOperatorKernel("-", OpenCLContext=self.openCLContext, queue=self.queue) from pyopencl.elementwise import ElementwiseKernel parameterUpdateKernel = ElementwiseKernel(context = self.openCLContext, arguments = "float howMuch, float *gradient, float* parameters", operation = "parameters[i] = parameters[i] - (gradient[i] * howMuch)", name="parameterUpdated") count = 0 while not converge: #output = np.dot(sampleMatrix, self.parameterVector) outputCalculator.multiply(deviceSampleMatrix.data, deviceParameterVector.data, deviceOutput.data) #error = output - targetVector vectorToVectorSubtractor.do(deviceOutput.data, deviceTargetVector.data, deviceError.data, deviceError.shape) #gradient = np.dot(error, sampleMatrix) gradientCalculator.multiply(deviceError.data, deviceSampleMatrix.data, deviceGradient.data) #self.parameterVector = self.parameterVector - (gradient * howMuch) parameterUpdateKernel(howMuch, deviceGradient, deviceParameterVector) #error = np.dot(sampleMatrix, self.parameterVector) - targetVector #output = np.dot(sampleMatrix, self.parameterVector) #outputCalculator.multiply(deviceSampleMatrix.data, deviceParameterVector.data, deviceOutput.data) #error = output - targetVector #vectorToVectorSubtractor.do(deviceOutput.data, deviceTargetVector.data, deviceError.data, deviceError.shape) mse = oclArray.dot(deviceError, deviceError) print(mse) if(mse <= convergeAtError): self.parameterVector = deviceParameterVector.get(self.queue) converge = True else: count = count + 1 return count /////////// """ Math routines for OpenCL. by psksvp(a)gmail.com Matrix Multiplication is based on example provided by pyopencl package. """ from __future__ import division matrixMultiplyKernelCode = """ // Thread block size #define BLOCK_SIZE %(block_size)d // Matrix dimensions // (chosen as multiples of the thread block size for simplicity) #define WA %(w_a)d // Matrix A width #define HA %(h_a)d // Matrix A height #define WB %(w_b)d // Matrix B width #define HB WA // Matrix B height #define WC WB // Matrix C width #define HC HA // Matrix C height /* * Copyright 1993-2009 NVIDIA Corporation. All rights reserved. * * NVIDIA Corporation and its licensors retain all intellectual property and * proprietary rights in and to this software and related documentation. * Any use, reproduction, disclosure, or distribution of this software * and related documentation without an express license agreement from * NVIDIA Corporation is strictly prohibited. * * Please refer to the applicable NVIDIA end user license agreement (EULA) * associated with this source code for terms and conditions that govern * your use of this NVIDIA software. * */ /* Matrix multiplication: C = A * B. * Device code. */ #define AS(j, i) As[i + j * BLOCK_SIZE] #define BS(j, i) Bs[i + j * BLOCK_SIZE] //////////////////////////////////////////////////////////////////////////////// //! Matrix multiplication on the device: C = A * B //! WA is A's width and WB is B's width //////////////////////////////////////////////////////////////////////////////// __kernel __attribute__((reqd_work_group_size(BLOCK_SIZE,BLOCK_SIZE,1))) void matrixMultiple( __global float* C, __global float* A, __global float* B) { __local float As[BLOCK_SIZE*BLOCK_SIZE]; __local float Bs[BLOCK_SIZE*BLOCK_SIZE]; // Block index int bx = get_group_id(0); int by = get_group_id(1); // Thread index int tx = get_local_id(0); int ty = get_local_id(1); // Index of the first sub-matrix of A processed by the block int aBegin = WA * BLOCK_SIZE * by; // Index of the last sub-matrix of A processed by the block int aEnd = aBegin + WA - 1; // Step size used to iterate through the sub-matrices of A int aStep = BLOCK_SIZE; // Index of the first sub-matrix of B processed by the block int bBegin = BLOCK_SIZE * bx; // Step size used to iterate through the sub-matrices of B int bStep = BLOCK_SIZE * WB; // Csub is used to store the element of the block sub-matrix // that is computed by the thread float Csub = 0.0f; // Loop over all the sub-matrices of A and B // required to compute the block sub-matrix for (int a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep) { // Load the matrices from device memory // to shared memory; each thread loads // one element of each matrix AS(ty, tx) = A[a + WA * ty + tx]; BS(ty, tx) = B[b + WB * ty + tx]; // Synchronize to make sure the matrices are loaded barrier(CLK_LOCAL_MEM_FENCE); // Multiply the two matrices together; // each thread computes one element // of the block sub-matrix for (int k = 0; k < BLOCK_SIZE; ++k) Csub += AS(ty, k) * BS(k, tx); // Synchronize to make sure that the preceding // computation is done before loading two new // sub-matrices of A and B in the next iteration barrier(CLK_LOCAL_MEM_FENCE); } // Write the block sub-matrix to device memory; // each thread writes one element C[get_global_id(1) * get_global_size(0) + get_global_id(0)] = Csub; } """ ####################################### vectorToVectorKernelCode = """ __kernel void vectorToVectorOp(__global const float *a, __global const float *b, __global float *c) { int gid = get_global_id(0); c[gid] = a[gid] %(OPERATOR)s b[gid]; } """ import pyopencl as ocl import pyopencl.array as oclArray def makeMatrixMultipleKernel(AWidth, AHeight, BWidth, BHeight, OpenCLContext, queue, blockSize=1): kernelParams = {"block_size": blockSize, "w_a":AWidth, "h_a":AHeight, "w_b":BWidth} if "NVIDIA" in queue.device.vendor: options = "-cl-mad-enable -cl-fast-relaxed-math" else: options = "" kernel = ocl.Program(OpenCLContext, matrixMultiplyKernelCode % kernelParams,).build(options=options) return kernel ## A * B = C ##c_width = b_width ##c_height = a_height class MatrixMultipleKernel(object): def __init__(self, AWidth, AHeight, BWidth, BHeight, OpenCLContext, queue, blockSize=1): self.resultDim = (BWidth, AHeight) self.blockSize = blockSize self.queue = queue self.kernel = makeMatrixMultipleKernel(AWidth, AHeight, BWidth, BHeight, OpenCLContext, queue, blockSize) def multiply(self, deviceMatrix1, deviceMatrix2, deviceMatrixResult): self.kernel.matrixMultiple(self.queue, self.resultDim , (self.blockSize, self.blockSize), deviceMatrixResult, deviceMatrix1, deviceMatrix2).wait() class VectorToVectorOperatorKernel(object): def __init__(self, operator, OpenCLContext, queue): self.OpenCLContext = OpenCLContext self.queue = queue self.kernel = ocl.Program(OpenCLContext, vectorToVectorKernelCode % {"OPERATOR": operator}).build() def do(self, deviceVector1, deviceVector2, deviceResult, shape): self.kernel.vectorToVectorOp(self.queue, shape, None, deviceVector1, deviceVector2, deviceResult).wait() class VectorToVectorSubtractKernel(VectorToVectorOperatorKernel): def __init__(self, OpenCLContext, queue): VectorToVectorOperatorKernel.__init__(self, "-", OpenCLContext, queue) ####################################################### def test(): import numpy matrix1 = numpy.array([[ 7, 3], [ 2, 5], [ 6, 8], [ 9, 0]], dtype=numpy.float32) matrix2 = numpy.array([[ 7, 4, 9], [ 8, 1, 5]], dtype=numpy.float32) ctx = ocl.create_some_context() queue = ocl.CommandQueue(ctx) matrix3 = numpy.array([[0, 0, 0], [0, 0, 0], [0, 0, 0], [0, 0, 0]], dtype=numpy.float32) dA = oclArray.to_device(queue, matrix1) dB = oclArray.to_device(queue, matrix2) dC = oclArray.to_device(queue, matrix3) #deviceA = ocl.Buffer(ctx, ocl.mem_flags.READ_ONLY | ocl.mem_flags.COPY_HOST_PTR, hostbuf=matrix1) #deviceB = ocl.Buffer(ctx, ocl.mem_flags.READ_ONLY | ocl.mem_flags.COPY_HOST_PTR, hostbuf=matrix2) #deviceC = ocl.Buffer(ctx, ocl.mem_flags.WRITE_ONLY, size=matrix3.nbytes) print(numpy.dot(matrix1, matrix2)) mK = MatrixMultipleKernel(AWidth=matrix1.shape[1], AHeight=matrix1.shape[0], BWidth=matrix2.shape[1], BHeight=matrix2.shape[0], OpenCLContext=ctx, queue=queue) mK.multiply(dA.data, dB.data, dC.data) ocl.enqueue_copy(queue, matrix3, dC.data) print(matrix3) if __name__ == "__main__": test()

6 years, 11 months

Re: [PyOpenCL] demo_mandelbrot.py: poor serial algorithm

by Andreas Kloeckner

"CRV§ADER//KY" <crusaderky(a)gmail.com> writes: > Hi, > in pyopencl 2013.1, I've been looking into examples/demo_mandelbrot.py. > > The serial implementation of the algorithm is a brain-off unloop of the > numpy implementation, which really does poor pure python no justice at all. > > > The original implementation: > > def calc_fractal_serial(q, maxiter): > # calculate z using numpy > # this routine unrolls calc_fractal_numpy as an intermediate > # step to the creation of calc_fractal_opencl > # it runs slower than calc_fractal_numpy > z = np.zeros(q.shape, np.complex64) > output = np.resize(np.array(0,), q.shape) > for i in range(len(q)): > for iter in range(maxiter): > z[i] = z[i]*z[i] + q[i] > if abs(z[i]) > 2.0: > q[i] = 0+0j > z[i] = 0+0j > output[i] = iter > return output > > > This doesn't take into account that > 1) it uses an iterative approach, yet it doesn't take the opportunity to > break out of the inner for cycle when it can! > 2) python is NOT optimized for non-native data types such as np.complex64. > > My version: > > def calc_fractal_serial(q, maxiter): > # calculate z using pure python on a numpy array > # note that, unlike the other two implementations, > # the number of iterations per point is NOT constant > z = np.zeros(q.shape, complex) > output = np.resize(np.array(0,), q.shape) > for i in range(len(q)): > for iter in range(maxiter): > z[i] = z[i]*z[i] + q[i] > if abs(z[i]) > 2.0: > output[i] = iter > break > return output > > > Time to process a 256x256 image: > Original implementation: 12.56s > After introducing the break in the loop: 4.44s > After changing the data type from np.complex64 to complex: 1.54s Thanks for the improvement--changed in git. Andreas

6 years, 11 months

pyopencl matrix-multiply.py problem

by Pongsak Suvanpong

Dear all I am trying to run the example program matrix-multiply.py in examples directory from pyopencl, but got the below error pyopencl.LogicError: clEnqueueNDRangeKernel failed: invalid work group size reading this message http://lists.tiker.net/pipermail/pyopencl/2011-August/000867.html seems to be similar problem, but not sure how to fix it. my setup is python3.3 numpy 1.7.x osx 10.8.4, system reports OpenCL 1.2, pyopencl from clone from github psksvp

6 years, 11 months

Re: [PyOpenCL] help with reductionKernel

by Pongsak Suvanpong

Hi Bogdan numpy.version.version reports 1.8.0.dev-Unknow. I am running on python3.3 and pyopencl checked out from github. thanks for any help psksvp On 15/08/2013, at 12:08 AM, Bogdan Opanchuk <mantihor(a)gmail.com> wrote: > Hi Pongsak, > > What version of numpy do you have? This code works normally for me on > the same configuration + numpy 1.7.1. > > On Wed, Aug 14, 2013 at 11:42 PM, Pongsak Suvanpong <psksvp(a)gmail.com> wrote: >> Dear all >> >> I am trying to run the following from example in documentation but got error show below, I wonder what have I done wong >> >> I am running pyopencl on python3.3 osx 10.8.4 >> >> thanks for any help >> >> import numpy >> import pyopencl as cl >> import pyopencl.array as clarray >> from pyopencl.reduction import ReductionKernel >> >> ctx = cl.create_some_context() >> queue = cl.CommandQueue(ctx) >> krnl = ReductionKernel(ctx, numpy.float32, neutral="0", >> reduce_expr="a+b", map_expr="x[i]*y[i]", >> arguments="__global float *x, __global float *y") >> x = clarray.arange(queue, 400, dtype=numpy.float32) >> y = clarray.arange(queue, 400, dtype=numpy.float32) >> m = krnl(x, y).get() >> >> >> >> --------------------- >> /Users/psksvp/Local/Library/Python.framework/Versions/3.3/lib/python3.3/site-packages/pyopencl-2013.2-py3.3-macosx-10.8-x86_64.egg/pyopencl/__init__.py:61: CompilerWarning: Non-empty compiler output encountered. Set the environment variable PYOPENCL_COMPILER_OUTPUT=1 to see more. >> "to see more.", CompilerWarning) >> Traceback (most recent call last): >> File "p.py", line 13, in <module> >> m = krnl(x, y).get() >> File "/Users/psksvp/Local/Library/Python.framework/Versions/3.3/lib/python3.3/site-packages/pyopencl-2013.2-py3.3-macosx-10.8-x86_64.egg/pyopencl/array.py", line 617, in get >> ary = _as_strided(ary, strides=self.strides) >> File "/Users/psksvp/Local/Library/Python.framework/Versions/3.3/lib/python3.3/site-packages/pyopencl-2013.2-py3.3-macosx-10.8-x86_64.egg/pyopencl/compyte/array.py", line 170, in as_strided >> return np.asarray(_DummyArray(interface, base=x)) >> File "/Users/psksvp/Local/Library/Python.framework/Versions/3.3/lib/python3.3/site-packages/numpy/core/numeric.py", line 325, in asarray >> return array(a, dtype, copy=False, order=order) >> ValueError: __array_interface__ shape must be at least size 1 >> _______________________________________________ >> PyOpenCL mailing list >> PyOpenCL(a)tiker.net >> http://lists.tiker.net/listinfo/pyopencl

6 years, 12 months

Re: [PyOpenCL] Good performance from a hacked-together GEMM?

by James Bergstra

Hi Bogdan, thanks for the plug. Part of writing was to find a home for ongoing development. I'll check out Reikna. On Wed, Aug 14, 2013 at 10:17 AM, Bogdan Opanchuk <mantihor(a)gmail.com>wrote: > Hi James, > > My pet project, Reikna (http://reikna.publicfields.net/) has a simple > implementation of matrix multiplication (mostly lifted from the > corresponding example in nVidia SDK). More specifically, see > https://github.com/Manticore/reikna/blob/develop/reikna/matrixmul.mako > > Now I understand that it is a shameless self-advertisement, but one of > the main ideas of Reikna is to provide a framework for implementing > GPGPU algorithms. I will appreciate it if you had a look at Reikna and > consider the idea of implementing your GEMM with it. > > Best regards, > Bogdan > > On Wed, Aug 14, 2013 at 11:58 PM, James Bergstra > <james.bergstra(a)gmail.com> wrote: > > Hi list, sort of writing a show-and-tell post here. > > > > Yesterday I hacked together a GEMM code generator using mako and PyOpenCL > > and by including some of the tips from [1] I made a *CPU* GEMM > > implementation that gets up around 35 GFLOP/s in single precision on my > core > > i7-3770 (using the intel OCL compiler). On another machine with an Intel > > i7-960 and the AMD compiler it gets around 8.5 GFLOP/s with the same > > settings that I used on the 3770. Proper auto-tuning would presumably > raise > > these rates to some extent. > > > > The code I used is here: > > > https://github.com/jaberg/python-benchmarks/blob/gemm_ocl/gemm/gemm_pyopenc… > > > > Basically I'm writing because I looked around a little and didn't find > > PyOpenCL source code for GEMM. If I failed to find a more mature project, > > could someone let me know? Also, if anyone has tips to make it work > better, > > patches welcome. The kernel doesn't prefetch anything, I'm slightly > > suspicious of the clunkiness associated with "transposing" the little > block > > of A that's stored in vector registers, and I haven't yet plugged it into > > any kind of profiler. > > > > More obviously, the `vectorized_text` code generator has some bugs in it, > > but seems to work with some combinations of block values (e.g. 4, 4, 4 > and > > 8, 8, 8). > > > > [1] > > Kazuya Matsumoto, Naohito Nakasato, Stanislav G. Sedukhin > > Implementing a Code Generator for Fast Matrix Multiplication in OpenCL on > > the GPU > > ftp://ftp.u-aizu.ac.jp/u-aizu/doc/Tech-Report/2012/2012-002.pdf > > > > _______________________________________________ > > PyOpenCL mailing list > > PyOpenCL(a)tiker.net > > http://lists.tiker.net/listinfo/pyopencl > > >

6 years, 12 months

2020

2019

2018

2017

2016

2015

2014

2013

2012

2011

2010

2009

PyOpenCL August 2013