clm4rm_bitwise.cpp

Go to the documentation of this file.

#include <clm4rm.h>
#include <qdebug.h>

void assertMatrixSize(clmatrix_t* C, clmatrix_t* A, clmatrix_t* B)
{
    Q_ASSERT(A->nrows==B->nrows);
    Q_ASSERT(A->nrows==C->nrows);
    Q_ASSERT(A->ncols==B->ncols);
    Q_ASSERT(A->ncols==C->ncols);
//    Q_ASSERT(A->width==B->width);
//    Q_ASSERT(A->width==C->width);
//    Q_ASSERT(A->rowstride==B->rowstride);
//    Q_ASSERT(A->rowstride==C->rowstride);
}

extern cl_kernel clm4rm_and_kernel;
extern cl_kernel clm4rm_or_kernel;
extern cl_kernel clm4rm_copy_kernel;
extern cl_kernel clm4rm_query_diagonal_kernel;


void clm4rm_or(clmatrix_t* C, clmatrix_t* A, clmatrix_t* B,
                   cl_command_queue queue, clm4rm_conditions* cond)
{
    assertMatrixSize(C,A,B);
    Q_ASSERT(clm4rm_or_kernel!=NULL);

    clm4rm_error = clSetKernelArg(clm4rm_or_kernel, 0, sizeof(cl_mem), &C->data);
    clm4rm_error = clSetKernelArg(clm4rm_or_kernel, 1, sizeof(cl_mem), &A->data);
    clm4rm_error = clSetKernelArg(clm4rm_or_kernel, 2, sizeof(cl_mem), &B->data);

    //  can't do in-place operation on Image2D
    Q_ASSERT(! IMAGE2D || (C->data!=A->data) && (C->data!=B->data));

    cl_uint work_dim = 2;
    size_t work_size[2] = { (size_t)A->nrows, (size_t)A->width };
    clm4rm_error = clEnqueueNDRangeKernel(queue, clm4rm_or_kernel,
                      work_dim, NULL, work_size, NULL,
                      pre_count(cond), pre_events(cond), push_event(cond));
    Q_ASSERT(pushed_event(cond) != NULL);
}


void clm4rm_and(clmatrix_t* C, clmatrix_t* A, clmatrix_t* B,
                    cl_command_queue queue, clm4rm_conditions* cond)
{
    assertMatrixSize(C,A,B);
    Q_ASSERT(clm4rm_and_kernel!=NULL);

    clm4rm_error = clSetKernelArg(clm4rm_and_kernel, 0, sizeof(cl_mem), &C->data);
    clm4rm_error = clSetKernelArg(clm4rm_and_kernel, 1, sizeof(cl_mem), &A->data);
    clm4rm_error = clSetKernelArg(clm4rm_and_kernel, 2, sizeof(cl_mem), &B->data);

    //  can't do in-place operation on Image2D
    Q_ASSERT(! IMAGE2D || (C->data!=A->data) && (C->data!=B->data));

    cl_uint work_dim = 2;
    size_t work_size[2] = { (size_t)A->nrows, (size_t)A->width };
    clm4rm_error = clEnqueueNDRangeKernel(queue, clm4rm_and_kernel,
            work_dim, NULL, work_size, NULL,
            pre_count(cond), pre_events(cond), push_event(cond));
    Q_ASSERT(pushed_event(cond) != NULL);
}


cl_mem clm4rm_query_diagonal(clmatrix_t* M,
                             cl_context ctx, cl_command_queue queue,
                             clm4rm_conditions* cond)
{
    Q_ASSERT(clm4rm_query_diagonal_kernel != NULL);
    
    int result = -1;
    cl_mem result_buffer = clCreateBuffer(ctx, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, 
                                        sizeof(int), &result, 
                                        &clm4rm_error);

    clm4rm_error = clSetKernelArg(clm4rm_query_diagonal_kernel, 0, sizeof(cl_mem), &M->data);
    clm4rm_error = clSetKernelArg(clm4rm_query_diagonal_kernel, 1, sizeof(int), &M->padded_rows);
    clm4rm_error = clSetKernelArg(clm4rm_query_diagonal_kernel, 2, sizeof(cl_mem), &result_buffer);

    cl_uint work_dim = 1;
    size_t work_size = MIN(M->ncols, M->nrows);

    clm4rm_error = clEnqueueNDRangeKernel(queue, clm4rm_query_diagonal_kernel,
        work_dim, NULL, &work_size, NULL,
        pre_count(cond), pre_events(cond), push_event(cond));
    Q_ASSERT(clm4rm_error == CL_SUCCESS);
    Q_ASSERT(pushed_event(cond) != NULL);
    return result_buffer;
}


int clm4rm_query_result(cl_mem result_buffer,
                        cl_command_queue queue, clm4rm_conditions* cond)
{
    int result;
    //  call is blocking
    clm4rm_error = clEnqueueReadBuffer(queue, result_buffer, CL_TRUE,
                                       0, sizeof(int), &result,
                                       pre_count(cond), pre_events(cond), NULL);
    //  was blocking: all conditions are met
    release_conditions(cond);
    clReleaseMemObject(result_buffer);
    return result;
}


void clm4rm_free(clmatrix_t* gpu_matrix)
{
    if (gpu_matrix) {
        clReleaseMemObject(gpu_matrix->data);
        allocated_size -= DATA_BYTES(gpu_matrix);
        free(gpu_matrix->local_data);
        free(gpu_matrix);
    }
}


void clm4rm_stack(clmatrix_t* C, clmatrix_t* A, clmatrix_t* B,
                  cl_command_queue queue, clm4rm_conditions* cond)
{
    Q_ASSERT(C->ncols==A->ncols);
    Q_ASSERT(C->ncols==B->ncols);
//    Q_ASSERT(C->rowstride==A->rowstride);
//    Q_ASSERT(C->rowstride==B->rowstride);
    Q_ASSERT(A->nrows+B->nrows == C->nrows);
    Q_ASSERT(DATA_BYTES(C) >= DATA_BYTES(A)+DATA_BYTES(B));

#if IMAGE2D
    size_t aorigin[3] = { 0,0,0 };
    size_t borigin[3] = { 0,A->padded_rows,0 };
    size_t aregion[3] = { A->width, A->padded_rows, 1 };
    size_t bregion[3] = { B->width, B->padded_rows, 1 };

    clm4rm_error = clEnqueueCopyImage(queue,
        A->data, C->data,
        aorigin, aorigin, aregion,
        0, NULL, &events[0]);
    clm4rm_error = clEnqueueCopyImage(queue,
        B->data, C->data,
        aorigin, borigin, bregion,
        0, NULL, &events[1]);
#else
    clm4rm_error = clEnqueueCopyBuffer(queue,
        A->data, C->data, 0, 0, DATA_BYTES(A),
        pre_count(cond),pre_events(cond),push_event(cond));
    Q_ASSERT(pushed_event(cond) != NULL);
    clm4rm_error = clEnqueueCopyBuffer(queue,
        B->data, C->data, 0, DATA_BYTES(A), DATA_BYTES(B),
        pre_count(cond),pre_events(cond),push_event(cond));
    Q_ASSERT(pushed_event(cond) != NULL);
#endif
}


void clm4rm_concat(clmatrix_t* C, clmatrix_t* A, clmatrix_t* B,
                   cl_command_queue queue, clm4rm_conditions* cond)
{
    Q_ASSERT(C->nrows==A->nrows);
    Q_ASSERT(C->nrows==B->nrows);
    Q_ASSERT(A->ncols + B->ncols == C->ncols);

    size_t Arows = A->nrows;
    size_t Acols = A->ncols;
    size_t Brows = B->nrows;
    size_t Bcols = B->ncols;
    size_t origin[3] = { 0,0,0 };
    size_t aregion[3] = { /*BYTE_WIDTH(A)*/0, Arows, 1 };

    //  Copy A -> C
#if IMAGE2D
    clm4rm_error = clEnqueueCopyImage(queue,
        A->data, C->data,
        origin, origin, aregion,
        0, NULL, &events[0]);

    if ((A->ncols % clm4rm_radix) == 0) {
        //  A is Word aligned. Let's use clEnqueueCopyBufferRect
        size_t boffset[3] = { A->ncols, 0, 1 };
        size_t bregion[3] = { B->width, B->padded_rows, 1 };
        //  Copy B -> C
        clm4rm_error = clEnqueueCopyImage(queue,
            B->data, C->data,
            origin, boffset, bregion,
            pre_count(cond),pre_events(cond),post_event(cond));
    }
    else {
        //  nasty: B needs to be shifted left by .. bits
        int shleft = clm4rm_radix - (A->ncols % clm4rm_radix);
        // TODO use a kernel ?
    }
#else
    clm4rm_error = clEnqueueCopyBufferRect (queue,
            A->data, C->data,
            origin, origin, aregion,
            /*BYTE_WIDTH(A)*/0, 0,
            /*BYTE_WIDTH(C)*/0, 0,
            pre_count(cond),pre_events(cond),push_event(cond));
    Q_ASSERT(pushed_event(cond) != NULL);

    if ((A->ncols % 8)==0) {
        //  A is Byte aligned. Let's use clEnqueueCopyBufferRect
        size_t boffset[3] = { Acols/8, 0, 0 };
        size_t bregion[3] = { /*BYTE_WIDTH(B)*/0, Brows, 1 };
        //  Copy B -> C
        clm4rm_error = clEnqueueCopyBufferRect (queue,
                B->data, C->data,
                origin, boffset, bregion,
                /*BYTE_WIDTH(B)*/0, 0,
                /*BYTE_WIDTH(C)*/0, 0,
                pre_count(cond),pre_events(cond),push_event(cond));
        Q_ASSERT(pushed_event(cond) != NULL);
    }
    else {
        /*
         *  __global unsigned int* C, unsigned int C_rowstride,
            __global unsigned int* B, unsigned int B_rowstride,
            unsigned int offset
         */
        Q_ASSERT(clm4rm_copy_kernel!=NULL);

        clm4rm_error = clSetKernelArg(clm4rm_copy_kernel, 0, sizeof(cl_mem), &C->data);
//        clm4rm_error = clSetKernelArg(clm4rm_copy_kernel, 1, sizeof(int), &C->rowstride);
        clm4rm_error = clSetKernelArg(clm4rm_copy_kernel, 2, sizeof(cl_mem), &B->data);
//        clm4rm_error = clSetKernelArg(clm4rm_copy_kernel, 3, sizeof(int), &B->rowstride);
        clm4rm_error = clSetKernelArg(clm4rm_copy_kernel, 4, sizeof(int), &A->ncols);

        cl_uint work_dim = 2;
        size_t work_size[2] = { Brows, CEILCOLS(Bcols) };
        clm4rm_error = clEnqueueNDRangeKernel(queue, clm4rm_copy_kernel,
            work_dim, NULL, work_size, NULL,
            pre_count(cond),pre_events(cond),push_event(cond));

    }
#endif
}