只分配了共享内存数组的一半

Question

当我在 s_f[sidx] = 5;

之后使用 Nsight 时，我看到只有一半的共享内存数组被分配了

__global__ void BackProjectPixel(double* val,   
                                    double* projection,
                                    double* focalPtPos,
                                    double* pxlPos,
                                    double* pxlGrid,
                                    double* detPos, 
                                    double *detGridPos,
                                    unsigned int nN,
                                    unsigned int nS,
                                    double perModDetAngle,
                                    double perModSpaceAngle,
                                    double perModAngle)                 
{
    const double fx = focalPtPos[0];
    const double fy = focalPtPos[1];

    //extern __shared__ double s_f[64]; // 

    __shared__ double s_f[64]; // 

    unsigned int i = (blockIdx.x * blockDim.x) + threadIdx.x;
    unsigned int j = (blockIdx.y * blockDim.y) + threadIdx.y;
    unsigned int idx = j*nN + i;

    unsigned int sidx = threadIdx.y * blockDim.x + threadIdx.x;

    unsigned int threadsPerSharedMem = 64;

    if (sidx < threadsPerSharedMem)
    {
        s_f[sidx] = 5;
    }

    __syncthreads();

    //double * angle;
    //

    if (sidx < threadsPerSharedMem)
    {

        s_f[idx] = TriPointAngle(detGridPos[0], detGridPos[1],fx, fy, pxlPos[idx*2], pxlPos[idx*2+1], nN);
    }



}

这是我观察到的

我想知道为什么只有三十二个5？ s_f中不应该有六十四个5吗？谢谢

Answer 1

线程在线程组（通常为 32 个）中执行，也称为线程束。 Warps 按顺序对线程进行分组。在您的情况下，一个 warp 将获得线程 0-31，另一个 32-63。在您的调试上下文中，您可能只看到包含线程 0-31 的 warp 的结果。

Answer 2

I am wondering why there are only thirty-two 5?

有 32 个五，因为正如 mete 所说，内核仅由大小为 32 的线程组同时执行，在 CUDA 术语中称为扭曲。

Shouldn't there be sixty-four 5 in s_f?

将在同步屏障之后是 64 个五，即 __syncthreads()。因此，如果您将断点放在 __syncthreads() 调用之后的第一条指令上，您将看到所有的五。那是因为到那时来自一个块的所有扭曲将完成 __syncthreads().

之前的所有代码的执行

How can I see all warps with Nsight?

您可以通过将其放入 watchfield 来轻松查看所有线程的值：

s_f[sidx]

虽然 sidx 值可能由于优化而变得未定义，所以我最好注意以下值：

s_f[((blockIdx.y * blockDim.y) + threadIdx.y) * nN + (blockIdx.x * blockDim.x) + threadIdx.x]

事实上，如果您想研究特定扭曲的值，那么正如 Robert Crovella 指出的那样，您应该使用条件断点。如果你想在第二个扭曲内打破，那么在二维块的二维网格（我假设你正在使用）的情况下，这样的事情应该有效：

((blockIdx.x + blockIdx.y * gridDim.x) * (blockDim.x * blockDim.y) + (threadIdx.y * blockDim.x) + threadIdx.x) == 32

因为 32 是第二个 warp 中第一个线程的索引。对于块和网格维度的其他组合，请参阅 this useful cheatsheet。

只分配了共享内存数组的一半

Only half of the shared memory array is assigned

cuda

shared-memory

nsight