CUDA L2传输开销

4
我有一个内核需要使用atomicMin测试点的渲染。在理想情况下的内存布局中,测试设置有大量点。有两个缓冲区,一个用于256x uint32的簇的uint32
namespace Point
{
struct PackedBitfield
{
    glm::uint32_t x : 6;
    glm::uint32_t y : 6;
    glm::uint32_t z : 6;
    glm::uint32_t nx : 4;
    glm::uint32_t ny : 4;
    glm::uint32_t nz : 4;
    glm::uint32_t unused : 2;
};

union __align__(4) Packed
{
    glm::uint32_t bits;
    PackedBitfield field;
};

struct ClusterPositionBitfield
{
    glm::uint32_t x : 10;
    glm::uint32_t y : 10;
    glm::uint32_t z : 10;
    glm::uint32_t w : 2;
};

union ClusterPosition
{
    glm::uint32_t bits;
    ClusterPositionBitfield field;
};
}

//
// launch with blockSize=(256, 1, 1) and grid=(numberOfClusters, 1, 1)
//
extern "C" __global__ void pointsRenderKernel(mat4 u_mvp,
                    ivec2 u_resolution,
                    uint64_t* rasterBuffer,
                    Point::Packed* points, 
                    Point::ClusterPosition* clusterPosition)
{
// extract and compute world position
const Point::ClusterPosition cPosition(clusterPosition[blockIdx.x]);
const Point::Packed point(points[blockIdx.x*256 + threadIdx.x]);

...use points and write to buffer...

}

生成的SASS代码如下所示:

enter image description here

查看内存分析器输出结果:从Point::Packed*缓冲区读取的L2传输开销为3.0。这是为什么呢?内存应该是完美对齐和连续的。此外,为什么会自动生成LDG(compute_50, sm_50)?我不需要这个缓存。

2
对于只读数据,LDG 使用通常最有效的加载路径,因此 CUDA 工具链更倾向于使用它。 - njuffa
对于我们这些没有那个分析器的人(我猜它是Windows版本?),你能提一下开销使用的单位吗? - einpoklum
@einpoklum - 开销是获取数据所需的内存事务的倍数。Linux 中也存在分析器(除了基于 Eclipse 的我相信)。 - FHoenig
@njuffa- 我的问题不是编译器生成什么样的指令。 - FHoenig
我们看不到整个代码,只有 ...use points and write to buffer...,以及汇编的前10.8行。您只展示了您认为相关的部分,显然这并没有帮助。请提供一个最小化完整可验证实例 - Jakub Klinkovský
1个回答

0
L2传输开销的工具提示中,它表示它测量了“L1和L2之间实际传输的字节数除以L1中请求的每个字节的数量”,并且它还表示“数值越低越好”。
在我的情况下,读取Point::Packed的L2传输开销为1.0

enter image description here

测试代码

namespace Point
{
    struct PackedBitfield
    {
        uint32_t x : 6;
        uint32_t y : 6;
        uint32_t z : 6;
        uint32_t nx : 4;
        uint32_t ny : 4;
        uint32_t nz : 4;
        uint32_t unused : 2;
    };

    union __align__(4) Packed
    {
        uint32_t bits;
        PackedBitfield field;
    };

    struct ClusterPositionBitfield
    {
        uint32_t x : 10;
        uint32_t y : 10;
        uint32_t z : 10;
        uint32_t w : 2;
    };

    union ClusterPosition
    {
        uint32_t bits;
        ClusterPositionBitfield field;
    };
}

__global__ void pointsRenderKernel(Point::Packed* points, Point::ClusterPosition* clusterPosition)
{
    int t_id = blockIdx.x * blockDim.x + threadIdx.x;

    clusterPosition[blockIdx.x + blockDim.x] = clusterPosition[blockIdx.x];
    points[t_id + blockDim.x * gridDim.x] = points[t_id];
}

void main()
{
    int blockSize = 256;
    int numberOfClusters = 256;

    std::cout << sizeof(Point::Packed) << std::endl;
    std::cout << sizeof(Point::ClusterPosition) << std::endl;

    Point::Packed *d_points;
    cudaMalloc(&d_points, sizeof(Point::Packed) * numberOfClusters * blockSize * 2);

    Point::ClusterPosition *d_clusterPositions;
    cudaMalloc(&d_points, sizeof(Point::ClusterPosition) * numberOfClusters * 2);

    pointsRenderKernel<<<numberOfClusters, blockSize>>>(d_points, d_clusterPositions);
}

更新

我之前在使用最新的驱动程序时,遇到了一些Nsight的问题。我将驱动程序降级到与默认CUDA 8.0.61安装程序(从这里下载)一起提供的版本,并解决了该问题。安装程序附带的版本是376.51。在Windows 10 64位和Visual Studio 2015上进行测试,Nsight版本为5.2,我的显卡是cc6.1。

这是我的完整编译器命令:

nvcc.exe -gencode=arch=compute_61,code=\"sm_61,compute_61\" --use-local-env --cl-version 2015 -Xcompiler "/wd 4819" -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 14.0\VC\bin\x86_amd64" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\include" -lineinfo --keep-dir x64\Release -maxrregcount=0 --machine 64 --compile -cudart static -DWIN32 -DWIN64 -DNDEBUG -D_CONSOLE -D_MBCS -Xcompiler "/EHsc /W3 /nologo /O2 /FS /Zi /MD " -o x64\Release\kernel.cu.obj kernel.cu

更新2

当我使用 sm_50,compute_50 选项编译时,得到了相同的结果:1.0 用于 L2 传输开销。


哇,谢谢你抽出时间来尝试这个。我想为了让问题和答案都被接受,我们都需要发布使用的确切编译器标志。 - FHoenig
@FHoenig 我之前也遇到过一些Nsight的问题,降级NVIDIA图形驱动程序解决了这个问题,正如答案中所提到的。 - nglee
使用计算能力为5.0 / SM 5.0,你能得到相同的结果吗? - FHoenig
@FHoenig 是的,L2传输开销的值为1.0,同时编译时使用compute_50、sm_50。 - nglee
虽然并不是一个直接的答案,但我认为这是一个在旧版本的Nsight中的测量误差。 - FHoenig

网页内容由stack overflow 提供, 点击上面的
可以查看英文原文,
原文链接