Hello, As far as I understand, this is somehow connected with pycuda::complex having default constructor. When you write (in reduction kernel): __shared__ out_type sdata[BLOCK_SIZE]; and "out_type" has the default constructor, it is called from _every_ thread for _every_ value of the array. The reduction kernel starts like: __shared__ out_type sdata[BLOCK_SIZE]; ... collecting values ... sdata[tid] = <some value> which means that there is a race between the default constructor and actual initialization. I am not sure why nvcc does not complain about this (bug?), but the following things do help: 1. Using "extern __shared__ out_type sdata[]" and set the size of shared memory when preparing the kernel. or 2. Putting "__syncthreads()" between default initialization and actual initialization (not very good, since it leaves all those calls to default constructor, but still removes the symptom). I googled a bit, but could not find any actual rules about default constructors and shared memory. Best regards, Bogdan On Fri, Jan 27, 2012 at 4:19 AM, Andreas Kloeckner &lt;lists(a)informa.tiker.net&gt; wrote: