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

【问题标题】：Only half of the shared memory array is assigned只分配了一半的共享内存数组
【发布时间】：2015-03-29 07:24:26
【问题描述】：

当我使用 Nsight 时，我看到只分配了一半的共享内存数组s_f[sidx] = 5;

__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);
    }



}

这是我观察到的

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

【问题讨论】：

什么时候你观察到的？ 32是经线。您确定您不只是查看部分执行的结果吗？
@talonmies 在第一次命中并执行断点时观察。断点设置在 s_f[sidx] = 5;

标签： cuda shared-memory nsight

【解决方案1】：

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

【讨论】：

如何使用 Nsight 查看所有经线？当我使用 Nsight 观看时，观看窗口并没有说明它正在监控哪个扭曲。谢谢。
我现在无法访问 Nsight，但 here is a good documentation about your issue
您可能还想研究如何在 nsight VSE 中使用条件断点，以确保您要查看的特定扭曲与您命中的断点相对应。

【解决方案2】：

我想知道为什么只有 32 个 5？

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

s_f中不应该有六十四个5吗？

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

如何使用 Nsight 查看所有经线？

您可以通过将其放入监视字段轻松查看所有线程的值：

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 是第二个经纱中第一个线程的索引。有关块和网格尺寸的其他组合，请参阅this useful cheatsheet。

【讨论】：