1

我有以下内核来获得一堆向量的大小：

__global__ void norm_v1(double *in, double *out, int n)
{
    const uint i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < n)
    {
        double x = in[3*i], y = in[3*i+1], z = in[3*i+2];
        out[i] = sqrt(x*x + y*y + z*z);
    }
}

然而，由于in它[x0,y0,z0,...,xn,yn,zn]在分析器中表现不佳，表明全局负载效率为 32%。将数据重新包装为[x0, x1, ..., xn, y0, y1, ..., yn, z0, z1, ..., zn]极大地改善了事情（偏移量为x，y，并z相应地改变）。运行时间减少，效率高达 100%。

但是，这种包装对我的应用程序来说根本不实用。因此，我希望研究共享内存的使用。我的想法是让块中的每个线程blockDim.x从全局内存中复制三个值（分开）——产生合并访问。在最大值的假设下，blockDim.x = 256我想出了：

#define BLOCKDIM 256

__global__ void norm_v2(double *in, double *out, int n)
{
    __shared__ double invec[3*BLOCKDIM];

    const uint i = blockIdx.x * blockDim.x + threadIdx.x;

    invec[0*BLOCKDIM + threadIdx.x] = in[0*BLOCKDIM+i];
    invec[1*BLOCKDIM + threadIdx.x] = in[1*BLOCKDIM+i];
    invec[2*BLOCKDIM + threadIdx.x] = in[2*BLOCKDIM+i];
    __syncthreads();

    if (i < n)
    {
        double x = invec[3*threadIdx.x];
        double y = invec[3*threadIdx.x+1];
        double z = invec[3*threadIdx.x+2];

        out[i] = sqrt(x*x + y*y + z*z);
    }
}

但是，这显然是有缺陷的，需要提前n % blockDim.x != 0知道最大值，并且在使用. 我应该如何最好地解决这个问题？blockDimout[i > 255]n = 1024

4

1 回答 1

1

我认为这可以解决out[i > 255]问题：

__shared__ double shIn[3*BLOCKDIM];

const uint blockStart = blockIdx.x * blockDim.x;

invec[0*blockDim.x+threadIdx.x] = in[ blockStart*3 + 0*blockDim.x + threadIdx.x];
invec[1*blockDim.x+threadIdx.x] = in[ blockStart*3 + 1*blockDim.x + threadIdx.x];
invec[2*blockDim.x+threadIdx.x] = in[ blockStart*3 + 2*blockDim.x + threadIdx.x];
__syncthreads();

double x = shIn[3*threadIdx.x];
double y = shIn[3*threadIdx.x+1];
double z = shIn[3*threadIdx.x+2];

out[blockStart+threadIdx.x] = sqrt(x*x + y*y + z*z);

至于n % blockDim.x != 0我建议用 0 填充输入/输出数组以匹配要求。

如果您不喜欢BLOCKDIM宏 - 探索使用extern __shared__ shArr[]然后将第三个参数传递给内核配置：

norm_v2<<<gridSize,blockSize,dynShMem>>>(...)

这dynShMem是动态共享内存使用量（以字节为单位）。这是一个额外的共享内存池，其大小在运行时指定，所有extern __shared__变量都将最初分配到其中。

你用的是什么GPU？Fermi 或 Kepler可能会通过 L1 缓存帮助您的原始代码。

如果您不想填充您的in数组，或者您最终在其他地方做类似的技巧，您可能需要考虑实现一个 device-side memcopy，如下所示：

template <typename T>
void memCopy(T* destination, T* source, size_t numElements) {
    //assuming sizeof(T) is a multiple of sizeof(int)
    //assuming one-dimentional kernel (only threadIdx.x and blockDim.x matters) 
    size_t totalSize = numElements*sizeof(T)/sizeof(int);
    int* intDest = (int*)destination;
    int* intSrc = (int*)source;
    for (size_t i = threadIdx.x; i < totalSize; i += blockDim.x) {
        intDest[i] = intSrc[i];
    }
    __syncthreads();
}

它基本上将任何数组视为int-s 数组，并将数据从一个位置复制到另一个位置。您可能希望将基础int类型替换为double-s 或者long long int仅使用 64 位类型。

然后您可以将复制行替换为：

memCopy(invec, in+blockStart*3, min(blockDim.x, n-blockStart));
于 2012-11-06T04:19:56.720 回答