0 
spmv_csr_scalar_kernel(const int num_rows,
                       const int * ptr,
                       const int * indices,
                       const float * data,
                       const float * x,
                       float * y )
{
    int row = get_global_id(0);
    if(row < num_rows)
    {
        float dot = 0;
        int row_start = ptr[row];
        int row_end = ptr[row+1];
        for (int jj = row_start; jj < row_end; jj++)
        {  
            dot += data[jj] * x[indices[jj]];
        }
        y[row] += dot;
    }
}

上記は、CSR 形式のスパース行列を列ベクトルで乗算するための Open Cl コードです。これは、for ループごとに 1 つのグローバル作業項目を使用します。最小のものでも変更すると多くの問題が発生します。助けてください。これは私のプロジェクトの一部です。これを並列にしましたが、もっと並列にしたいです。できれば助けてください.plzzzz

単一の作業項目は、row_start から row_end までの for ループを実行します。この行または for ループをさらに 2 つの部分に分割し、それぞれが単一の作業項目によって実行されるようにします。どうすればそれを達成できますか?

これは私が思いついたものですが、間違った出力を返します。

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

    int thread_id=get_global_id(0);
    int warp_id=thread_id/2;
    int lane=(thread_id)&1;
    int row=warp_id;

    if(row<4)
    {
        int row_start = rowptr[row];
        int row_end = rowptr[row+1];
        vals[thread_id]=0;

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

        vals[thread_id]+=vals[thread_id+1];

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

1 に答える 1

0

これを実現するには、ワーク グループの数とグループ ID を使用する必要があります。

spmv_csr_scalar_kernel(const int num_rows, const int * ptr, const int * indices, const float * data, const float * x, float * y) {

    __local float[64] dot; //local memory for dot product. assumes you want to use up to 64-item workgroups.
    int localId = get_local_id(0); //offest value for jj loop
    int localSize = get_local_size(0); //this is the number of work items computing the jj loop
    int nGroups = get_num_groups(0); //number of workgroups executing this kernel

    for (int row = get_group_id(0); row < num_rows; row += nGroups) {
        dot[localId] = 0;
        int row_start = ptr[row] + localId;
        int row_end = ptr[row + 1];
        for (int jj = row_start; jj < row_end; jj++) {
            dot[localId] += data[jj] * x[indices[jj]];
        }
        barrier(CLK_LOCAL_MEM_FENCE);

        //here's where work item #0 in the group sums the values stored in local memory
        if (localId == 0) {
            for (int i = 1; i < localSize; i++) {
                dot[0] += dot[i];
            }
            y[row] += dot[0]; //only a single write to the global memory location
            barrier(CLK_LOCAL_MEM_FENCE);
        }
    }

このカーネルは、すべての作業を完了するために必要な数のワークグループを使用します。(つまり、グループは複数の行で作業できます) これは、'if(row < num_rows)' を 'for(int row...' ループに変換することによって行われます。アイテムは、多くの小さな作業グループよりも優れている場合があります。

ローカル メモリ サイズのハード コード値として 64 を使用しました。上記のカーネルは、最大 64 のワークグループ サイズで問題なく動作します。今日の多くのデバイスでは 64 が最適です。他のハードコードされた値を試すか、__local パラメータを渡してローカル メモリを割り当てることができます。

于 2012-10-21T15:29:12.263 に答える