1

これは、圧縮された列形式で疎行列を乗算するための私のコードです

__kernel void mykernel(__global int* colvector,
                       __global int* val,
                       __global int* result,
                       __global int* index,
                       __global int* rowptr,
                       __global int* sync )
{   
    __local int vals[1000];
    for(int i=0;i<4;i++)
    {
        result[i]=0;
    }
    barrier(CLK_LOCAL_MEM_FENCE);
    barrier(CLK_GLOBAL_MEM_FENCE);

    const int items_per_row=32;//total threads working in a row

    const int thread_id=get_global_id(0)+get_local_id(0);//total threads in the program

    const int warpid = thread_id/items_per_row;//warp id is actual row

    int lane=thread_id&(items_per_row-1);//thread id within the warp

    int row = warpid;

    if(row<4)
    {
        int sum = 0;

        int row_start = rowptr[row];
        int row_end = rowptr[row+1];

        vals[get_global_id(0)]=0;
        barrier(CLK_LOCAL_MEM_FENCE);
        barrier(CLK_GLOBAL_MEM_FENCE);


        for (int i = row_start+lane; i<row_end; i+=items_per_row)
        {
            vals[get_local_id(0)]+=val[i]*colvector[index[i]];
        }

        barrier(CLK_LOCAL_MEM_FENCE);
        barrier(CLK_GLOBAL_MEM_FENCE);

        if (lane < 16 ) vals[get_local_id(0)] += vals[get_local_id(0) + 16];

        if (lane < 8 ) vals[get_local_id(0)] += vals[get_local_id(0) + 8];

        if (lane < 4 ) vals[get_local_id(0)] += vals[get_local_id(0) +4];

        if (lane < 2 ) vals[get_local_id(0)] += vals[get_local_id(0) + 2];

        if (lane < 1 ) vals[get_local_id(0)] += vals[get_local_id(0) + 1];

        barrier(CLK_LOCAL_MEM_FENCE);
        barrier(CLK_GLOBAL_MEM_FENCE);

        if(lane==0)
        {
            result[row] += vals[get_local_id(0)];
        }
    }
}

上記の OpenCL コードは、以下の CUDA コードから変換されたものです。

spmv_csr_vector_kernel(const int num_rows,
                       const int * ptr,
                       const int * indices,
                       const float * data,
                       const float * x,
                       float * y )
{
    __shared__ float vals[];

    int thread_id = blockDim.x * blockIdx.x + threadIdx.x; // global thread index

    int warp_id = thread_id / 32; // global warp index

    int lane = thread_id & (32 - 1); // thread index within the warp

    // one warp per row

    int row = warp_id;

    if (row < num_rows)
    {
        int row_start = ptr[row];

        int row_end = ptr[row+1];

        // compute running sum per thread

        vals[threadIdx.x] = 0;

        for(int jj = row_start + lane; jj < row_end; jj += 32)
        {
            vals[threadIdx.x] += data[jj] * x[indices[jj]];
        }
        // parallel reduction in shared memory

        if (lane < 16) vals[threadIdx.x] += vals[threadIdx.x + 16];

        if (lane < 8) vals[threadIdx.x] += vals[threadIdx.x + 8];

        if (lane < 4) vals[threadIdx.x] += vals[threadIdx.x + 4];

        if (lane < 2) vals[threadIdx.x] += vals[threadIdx.x + 2];

        if (lane < 1) vals[threadIdx.x] += vals[threadIdx.x + 1];

        // first thread writes the result

        if (lane == 0)
        {
            y[row] += vals[threadIdx.x];
        }
    }
}

CUDA コードは正しいのですが、OpenCL カーネルが正しい出力を返しません。私は今一週間試していますが、解決策はありません。私が犯している間違いを誰か知っていますか？

4

2 に答える 2

2

少なくとも1つの間違いを見ることができます。thread_id は各コードで同じではありません。blockDim.x * blockIdx.x + CUDA の threadIdx.x == OpenCL の get_global_id(0)、get_global_id(0)+get_local_id(0) ではありません。また、get_local_id(0) == threadIdx.x

于 2012-10-21T08:51:14.400 に答える
0

白鳥を使用してみてください。これは、問題を理解するのに役立つ場合があります。

あなたはそれについての記事をここで見つけることができます。

于 2013-08-24T15:47:30.727 に答える