Unkown Issue with input sequence size of FFT in OpenCL

Question

everyone, I recently tried to implement Fast Fourier Transformation in OpenCL, and I tried some code provided in Book "OpenCL in Action". The following code intends to initialize input sequence( or discrete signal ) by perform 4-compoent-wise FFT after bit reversing :

fft.cl:

__kernel void fft_init( __global float2* g_data, __global float2* l_data, uint points_per_group, uint size, int dir )

{

uint g_addr, l_addr, points_per_item;
points_per_item = points_per_group / get_local_size( 0 );

l_addr = get_local_id( 0 ) * points_per_item;

g_addr = get_group_id( 0 ) * points_per_group + l_addr;


uint4 index;
uint mask_left, mask_right, shift_pos;
uint4 br;

float2 x1, x2, x3, x4;
float2 sum12, diff12, sum34, diff34;

for (int i = 0; i < points_per_item; i += 4)
{
    index = (uint4){ g_addr, g_addr + 1, g_addr + 2, g_addr + 3 };

    mask_left = size / 2;
    mask_right = 1;

    shift_pos = log2( ( float ) size ) - 1;

    while( shift_pos > 1 )
    {
        br = ( index << shift_pos ) & mask_left;
        br |= ( index >> shift_pos ) & mask_right;

        mask_left >>= 1;
        mask_right <<= 1;
        shift_pos -= 2;
    }

    x1 = g_data[ br.s0 ];
    x2 = g_data[ br.s1 ];
    x3 = g_data[ br.s2 ];
    x4 = g_data[ br.s3 ];

    sum12 = x1 + x2;
    diff12 = x1 - x2;
    sum34 = x3 + x4;
    diff34 = ( float2 ){
        x3.s1 - x4.s1,
        x4.s0 - x3.s1
    };

    l_data[ l_addr ] = sum12 + sum34;
    l_data[ l_addr + 1 ] = diff12 + diff34;
    l_data[ l_addr + 2 ] = sum12 - sum34;
    l_data[ l_addr + 3 ] = diff12 - diff34;

    g_addr += 4;
    l_addr += 4;
}

}

I tried a very short sequence of 4 points ( 4 complex numbers ), when I set the parameter "points_per_group" and "size" to 4, the GPU ran out of resource; while I set these two parameter to 8, which is incorrect apparently with the input sequence, it ran through without crash; I also tried another sequence of size 8, with "points_per_group" and "size" set to 8, and it ran through well too. Here is part of the host side code; I wrote in Python, using pyopencl:

def setInputs( self ):

    self.numPoints = 8

    self.points = ( 1.0, 1.0 ) * self.numPoints

    self.pointsArray = numpy.array( self.points, dtype = numpy.float32 )

    self.resultArray = numpy.array( ( -1.0, -1.0 ) * self.numPoints, dtype = numpy.float32 )


    self.pointsBuffer = cl.Buffer( self.context,
                                   cl.mem_flags.READ_WRITE | cl.mem_flags.COPY_HOST_PTR,
                                   hostbuf = self.pointsArray
                                   )

    self.resultBuffer = cl.Buffer( self.context,
                                   cl.mem_flags.READ_WRITE | cl.mem_flags.COPY_HOST_PTR,
                                   hostbuf = self.resultArray
                                   )


def runKernel( self ):

    globalWorkSize = ( 1, )

    localWorkSize = ( 1, )

    print 'Input points : ', self.pointsArray

    event = self.program.fft_init( self.commandQueues[ 0 ],
                                   globalWorkSize,
                                   localWorkSize,
                                   self.pointsBuffer,
                                   self.resultBuffer,
                                   numpy.uint32( 8 ),
                                   numpy.uint32( 8 ),
                                   numpy.int32( 1 )
                                   )

    event.wait()

    print 'Time Consumption : ', ( event.profile.end - event.profile.start ) * 1e-9, ' seconds'

    cl.enqueue_copy( self.commandQueues[ 0 ],
                     self.resultArray,
                     self.resultBuffer
                     ).wait()

    print 'Result : ', self.resultArray

Any idea about this weird issue? :-o