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 循环使用一个全局工作项。任何人都可以帮助我在每个 for 循环中使用两个工作项。我是 open cl 和如果我修改即使是最小的东西也会遇到很多问题。请帮助我。这是我项目的一部分。我让它平行但我想让它更平行。如果可以的话请帮助我。plzzzz

单个工作项执行从 row_start 到 row_end 的 for 循环。我希望将此行或 for 循环进一步分为两部分，每个部分由单个工作项执行。我该如何继续完成呢？

这是我能想到的，但它返回了错误的输出。

__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 回答