clm4rm_bitwise.cl

Go to the documentation of this file.

typedef unsigned int gpuword;

#if IMAGE2D
//
// Matrix stored in texture memory
//
# define read_only_global     __read_only image2d_t
# define write_only_global    __write_only image2d_t
// Note: column-major format
// a matrix colum is actually a row (y-coordinate) in Image2D
// a matrix row is actually a column (x-coordinate) in Image2D
// Pixel contains only one (red) component
# define read(M,row,col)      read_imageui(M,(int2)(row,col)).x
# define write(M,row,col,x)   write_imageui(M,(int2)(row,col),(uint4)(x,0,0,0))
#else
//
// Matrix stored in __global memory
//
# define read_only_global     __global gpuword*
# define write_only_global    __global gpuword*
# define read(M,row,col)      M[(col)*M ## _nrows + row]
# define write(M,row,col,x)   M[(col)*M ## _nrows + row]=x
#endif


__kernel void clm4rm_and(
    write_only_global C,
    read_only_global A,
    read_only_global B)
{
    const int A_nrows = get_global_size(0);
#define B_nrows A_nrows
#define C_nrows A_nrows

    const int row = get_global_id(0);
    const int col = get_global_id(1);

    gpuword a = read(A,row,col);
    gpuword b = read(B,row,col);

    write(C,row,col, a & b);
}

__kernel void clm4rm_or(
    write_only_global C,
    read_only_global A,
    read_only_global B)
{
    const int A_nrows = get_global_size(0);

    const int row = get_global_id(0);
    const int col = get_global_id(1);

    gpuword a = read(A,row,col);
    gpuword b = read(B,row,col);

    write(C,row,col, a | b);
}

__kernel void clm4rm_copy(
    __global unsigned int* C, int C_rowstride,
    __global unsigned int* B, int B_rowstride,
    int offset )
{
    const int i = get_global_id(0);
    const int j = get_global_id(1);

    unsigned int word_offset = offset/32;
    unsigned int bit_offset = offset%32;

    C = C + i*C_rowstride + word_offset;
    B = B + i*B_rowstride;

    unsigned int cj;
    if (j==0)
        cj = (C[0] & ((1<<bit_offset)-1));
    else
        cj = (B[j-1] >> (32-bit_offset));
    cj |= (B[j] << bit_offset);
    C[j] = cj;
}


#define WRITE_ATOMIC 1

__kernel void clm4rm_query_diagonal(
    read_only_global M,
    int M_nrows,
#if WRITE_ATOMIC
    volatile
#endif
    __global int* result)
{
    const int i = get_global_id(0);

    // query M[i][i]
    gpuword word = read(M, i, i/32);
    word >>= i%32;

#if WRITE_ATOMIC
    if ((word & 1) && (*result==-1))
        atomic_xchg (result, i);
    //  If atomic fails, someone else was faster. No matter.
#else
    if ((word & 1) && (*result==-1)) 
        *result = i;
#endif
}