如何说服 nvcc 使用 128 位宽负载？
Posted
技术标签:
【中文标题】如何说服 nvcc 使用 128 位宽负载？【英文标题】：How to convince nvcc to use 128-bit wide loads? 【发布时间】：2015-09-25 12:05:59 【问题描述】：我有一个内核需要对一个数组应用模板操作并将结果存储在另一个数组中。模板可以在函数中表示为：
float stencil(const float* data)
return *(data-1) + *(data+1);
我希望每个线程通过加载输入数组的 6 个连续值来生成输出数组的 4 个连续值。通过这样做，我将能够使用 float4 类型来加载和存储 128 字节的块。这是我的程序（您可以下载并编译它，但请首先考虑内核）：
#include<iostream>
#include<cstdlib>
#include<thrust/host_vector.h>
#include<thrust/device_vector.h>
__global__ void kernel(const float* input, float* output, int size)
int i = 4*(blockDim.x*blockIdx.x + threadIdx.x);
float values[6];
float res[4];
// Load values
values[0] = *(input+i-1);
*reinterpret_cast<float4*>(values+1) = *reinterpret_cast<const float4*>(input+i);
values[5] = *(input+i+4);
// Compute result
res[0] = values[0]+values[2];
res[1] = values[1]+values[3];
res[2] = values[2]+values[4];
res[3] = values[3]+values[5];
// Store result
*reinterpret_cast<float4*>(output+i) = *reinterpret_cast<const float4*>(res);
int main()
// Parameters
const int nBlocks = 8;
const int nThreads = 128;
const int nValues = 4 * nThreads * nBlocks;
// Allocate host and device memory
thrust::host_vector<float> input_host(nValues+64);
thrust::device_vector<float> input(nValues+64), output(nValues);
// Generate random input
srand48(42);
thrust::generate(input_host.begin(), input_host.end(), [] return drand48()+1.; );
input = input_host;
// Run kernel
kernel<<<nBlocks, nThreads>>>(thrust::raw_pointer_cast(input.data()+32), thrust::raw_pointer_cast(output.data()), nValues);
// Check output
for (int i = 0; i < nValues; ++i)
float ref = input_host[31+i] + input_host[33+i];
if (ref != output[i])
std::cout << "Error at " << i << " : " << ref << " " << output[i] << "\n";
std::cout << "Abort with errors\n";
std::exit(1);
std::cout << "Success\n";
程序完美运行。
我希望编译器为本地数组 values 的中心部分生成一条 LD.E.128 指令，并且该中心部分的寄存器是连续的（例如 R4、R5、R6、R7）；在values 的两端有两条LD.E 指令；为output 数组设置一个ST.E.128。
现实中发生的事情如下：
code for sm_21
Function : _Z6kernelPKfPfi
/*0000*/ MOV R1, c[0x1][0x100]; /* 0x2800440400005de4 */
/*0008*/ NOP; /* 0x4000000000001de4 */
/*0010*/ MOV32I R3, 0x4; /* 0x180000001000dde2 */
/*0018*/ S2R R0, SR_CTAID.X; /* 0x2c00000094001c04 */
/*0020*/ S2R R2, SR_TID.X; /* 0x2c00000084009c04 */
/*0028*/ IMAD R0, R0, c[0x0][0x8], R2; /* 0x2004400020001ca3 */
/*0030*/ SHL R6, R0, 0x2; /* 0x6000c00008019c03 */
/*0038*/ IMAD R10.CC, R6, R3, c[0x0][0x20]; /* 0x2007800080629ca3 */
/*0040*/ IMAD.HI.X R11, R6, R3, c[0x0][0x24]; /* 0x208680009062dce3 */
/*0048*/ IMAD R2.CC, R6, R3, c[0x0][0x28]; /* 0x20078000a0609ca3 */
/*0050*/ LD.E R4, [R10+0xc]; /* 0x8400000030a11c85 */
/*0058*/ IMAD.HI.X R3, R6, R3, c[0x0][0x2c]; /* 0x20868000b060dce3 */
/*0060*/ LD.E R7, [R10+0x4]; /* 0x8400000010a1dc85 */
/*0068*/ LD.E R9, [R10+-0x4]; /* 0x87fffffff0a25c85 */
/*0070*/ LD.E R5, [R10+0x8]; /* 0x8400000020a15c85 */
/*0078*/ LD.E R0, [R10+0x10]; /* 0x8400000040a01c85 */
/*0080*/ LD.E R8, [R10]; /* 0x8400000000a21c85 */
/*0088*/ FADD R6, R7, R4; /* 0x5000000010719c00 */
/*0090*/ FADD R4, R9, R7; /* 0x500000001c911c00 */
/*0098*/ FADD R7, R5, R0; /* 0x500000000051dc00 */
/*00a0*/ FADD R5, R8, R5; /* 0x5000000014815c00 */
/*00a8*/ ST.E.128 [R2], R4; /* 0x9400000000211cc5 */
/*00b0*/ EXIT; /* 0x8000000000001de7 */
................................
所有加载都是 32 位宽 (LD.E)。另一方面，正如预期的那样，只有一条存储指令ST.E.128。
我不再在这里显示整个代码，但我做了一个测试，其中模板不需要左侧的值，而只需要右侧的值（例如*data + *(data+1)），在这种情况下我的@987654336 @array 仅包含 5 个值，float4 加载操作修改了数组的前 4 个值（我仍然为最后一个值额外加载了一个值）。在这种情况下，编译器使用LD.E.128。
我的问题是，如果目标寄存器不是本地数组中的第一个寄存器，为什么编译器不明白它可以使用 128 位宽读取。毕竟本地数组values 只是一种编程方式，可以说我需要将 6 个浮点数存储在寄存器中。在生成的 ptx 或 SASS 代码中没有像数组这样的东西。我以为我给了编译器足够的提示，让它明白 LD.E.128 是正确的指令。
第二个问题：我怎样才能让它在这里使用 128 宽的负载，而不必手动编写低级代码？ （但是，如果一些 asm 说明有帮助，我愿意接受建议。）
旁注：使用 32 位加载读取输入和使用 128 位存储写入输入的决定是在生成 ptx 代码时做出的。 ptx 代码已经显示了这种多小负载和单个大存储的模式。
我在linux下使用CUDA 7.5。
根据cmets给出的建议，我做了一些实验。
将input 或output 声明为__restrict__（或同时声明两者）解决了这个问题：编译器生成了一个LD.E.128 和两个LD.E，这是我在为建筑sm_35。奇怪的是，当为sm_21 生成时，它仍然产生六个LD.E，但它产生一个ST.E.128。对我来说这听起来像是一个编译器错误，因为指令 LD.E.128 在旧架构中应该完全可用，因为它在最新架构中。
上面显示的代码使用 128 位加载，只是按照 njuffa 的建议使用了 __restrict__ 关键字的微小变化并且可以工作。我也确实听从了m.s.的建议。我复制了 pastebin sn-p 中显示的相同结果（一个 LD.E.128 + 一个 LD.E.64）。但在运行时它会崩溃并出现以下错误：
terminate called after throwing an instance of 'thrust::system::system_error'
what(): an illegal memory access was encountered
我很确定错位是导致此问题的原因。
更新：使用 cuda-memcheck 后，我确定问题是未对齐：
========= Invalid __global__ read of size 16
========= at 0x00000060 in kernel(float const *, float*, int)
========= by thread (4,0,0) in block (7,0,0)
========= Address 0xb043638bc is misaligned
【问题讨论】：
是的，您需要从其中加载的全局向量自然对齐。很明显，CUDA 假设我的输出向量是自然对齐的（确实如此）。输入向量也自然对齐。在任何情况下，编译器都可以对input 和output 采取相同的假设。
不，它们是 128 字节对齐的。 thurst 数组对齐到 128 个字节。我向前移动了 32 个浮点数，所以我传递给内核的指针是 128 字节对齐的。然后我向前移动 4*(global thread Id)，它是 4 个浮点数 = 128 位的倍数。负载必须与 128 位边界（= 4 字节）对齐，而且确实如此。还是我误解了一些信息？无论如何，这同样适用于输出数组，并且可以正常工作。
如果您连续加载前 4 个值，然后是两个单独的加载 ([i-1,i,i+1,i+2] [i+3] [i+4])，您将获得所需的 LD.E.128 甚至是 LD.E.64: pastebin.com/MLrfmD4k;但这并不能解释为什么您的代码不能按预期工作
我也解释不了；我认为编译器有more trouble generating vector loads lately（CUDA 7.5）。除了@m.s. 的建议。 （这可能是最好的 - 性能和效率方面）您还可以修改您的内核代码，使您的 i 变量中的因子不为 4，然后使用它生成必要的加载和存储。它seems to work for me.
@m.s.令人费解。也许来自可能的混叠的干扰？也许试试const float4* __restrict__ input, float4* __restrict__ output
【参考方案1】：
问题在于 nvcc 编译器无法解析内核中向量加载的基地址。这可能是一个错误，也可能只是一个不足之处。
我稍微修改了你的代码：
__global__ void kernel2(const float* input, float* output, int size)
int i = (blockDim.x*blockIdx.x + threadIdx.x);
float values[6];
float res[4];
// Load values
values[0] = *(input+(i*4)-1);
float4 test =*(reinterpret_cast<const float4*>(input)+i);
values[5] = *(input+(i*4)+4);
values[1] = test.x;
values[2] = test.y;
values[3] = test.z;
values[4] = test.w;
// Compute result
res[0] = values[0]+values[2];
res[1] = values[1]+values[3];
res[2] = values[2]+values[4];
res[3] = values[3]+values[5];
// Store result
*(reinterpret_cast<float4*>(output)+i) = *reinterpret_cast<const float4*>(res);
编译成ptx的内核代码：
.visible .entry _Z7kernel2PKfPfi(
.param .u64 _Z7kernel2PKfPfi_param_0,
.param .u64 _Z7kernel2PKfPfi_param_1,
.param .u32 _Z7kernel2PKfPfi_param_2
)
.reg .f32 %f<15>;
.reg .b32 %r<7>;
.reg .b64 %rd<10>;
ld.param.u64 %rd1, [_Z7kernel2PKfPfi_param_0];
ld.param.u64 %rd2, [_Z7kernel2PKfPfi_param_1];
mov.u32 %r1, %ntid.x;
mov.u32 %r2, %ctaid.x;
mov.u32 %r3, %tid.x;
mad.lo.s32 %r4, %r2, %r1, %r3;
shl.b32 %r5, %r4, 2;
add.s32 %r6, %r5, -1;
mul.wide.s32 %rd3, %r6, 4;
cvta.to.global.u64 %rd4, %rd1;
add.s64 %rd5, %rd4, %rd3;
ld.global.f32 %f1, [%rd5];
mul.wide.s32 %rd6, %r4, 16;
add.s64 %rd7, %rd4, %rd6;
ld.global.v4.f32 %f2, %f3, %f4, %f5, [%rd7];
ld.global.f32 %f10, [%rd5+20];
cvta.to.global.u64 %rd8, %rd2;
add.s64 %rd9, %rd8, %rd6;
add.f32 %f11, %f3, %f5;
add.f32 %f12, %f2, %f4;
add.f32 %f13, %f4, %f10;
add.f32 %f14, %f1, %f3;
st.global.v4.f32 [%rd9], %f14, %f12, %f11, %f13;
ret;
您可以清楚地看到负载地址是如何计算的（%rd6 和 %rd8）。
将内核编译为 ptx 会导致：
.visible .entry _Z6kernelPKfPfi(
.param .u64 _Z6kernelPKfPfi_param_0,
.param .u64 _Z6kernelPKfPfi_param_1,
.param .u32 _Z6kernelPKfPfi_param_2
)
.reg .f32 %f<11>;
.reg .b32 %r<6>;
.reg .b64 %rd<8>;
ld.param.u64 %rd1, [_Z6kernelPKfPfi_param_0];
ld.param.u64 %rd2, [_Z6kernelPKfPfi_param_1];
cvta.to.global.u64 %rd3, %rd2;
cvta.to.global.u64 %rd4, %rd1;
mov.u32 %r1, %ntid.x;
mov.u32 %r2, %ctaid.x;
mov.u32 %r3, %tid.x;
mad.lo.s32 %r4, %r2, %r1, %r3;
shl.b32 %r5, %r4, 2;
mul.wide.s32 %rd5, %r5, 4;
add.s64 %rd6, %rd4, %rd5;
ld.global.f32 %f1, [%rd6+-4];
ld.global.f32 %f2, [%rd6];
ld.global.f32 %f3, [%rd6+12];
ld.global.f32 %f4, [%rd6+4];
ld.global.f32 %f5, [%rd6+8];
ld.global.f32 %f6, [%rd6+16];
add.s64 %rd7, %rd3, %rd5;
add.f32 %f7, %f5, %f6;
add.f32 %f8, %f4, %f3;
add.f32 %f9, %f2, %f5;
add.f32 %f10, %f1, %f4;
st.global.v4.f32 [%rd7], %f10, %f9, %f8, %f7;
ret;
编译器只生成代码来计算一个地址 (%rd6) 并使用静态偏移量。此时编译器未能发出向量负载。为什么？老实说，我不知道，也许这里有两个优化干扰。
在 SASS 中你看到 kernel2：
.section .text._Z7kernel2PKfPfi,"ax",@progbits
.sectioninfo @"SHI_REGISTERS=18"
.align 64
.global _Z7kernel2PKfPfi
.type _Z7kernel2PKfPfi,@function
.size _Z7kernel2PKfPfi,(.L_39 - _Z7kernel2PKfPfi)
.other _Z7kernel2PKfPfi,@"STO_CUDA_ENTRY STV_DEFAULT"
_Z7kernel2PKfPfi:
.text._Z7kernel2PKfPfi:
/*0008*/ MOV R1, c[0x0][0x44];
/*0010*/ S2R R0, SR_CTAID.X;
/*0018*/ MOV R4, c[0x0][0x140];
/*0020*/ S2R R3, SR_TID.X;
/*0028*/ MOV R5, c[0x0][0x144];
/*0030*/ IMAD R3, R0, c[0x0][0x28], R3;
/*0038*/ MOV32I R8, 0x10;
/*0048*/ IMAD R16.CC, R3, 0x10, R4;
/*0050*/ ISCADD R0, R3, -0x1, 0x2;
/*0058*/ IMAD.HI.X R17, R3, 0x10, R5;
/*0060*/ IMAD R14.CC, R0, 0x4, R4;
/*0068*/ IMAD.HI.X R15, R0, 0x4, R5;
/*0070*/ LD.E.128 R4, [R16];
/*0078*/ LD.E R2, [R14];
/*0088*/ IMAD R12.CC, R3, R8, c[0x0][0x148];
/*0090*/ LD.E R0, [R14+0x14];
/*0098*/ IMAD.HI.X R13, R3, R8, c[0x0][0x14c];
/*00a0*/ FADD R9, R4, R6;
/*00a8*/ FADD R10, R5, R7;
/*00b0*/ FADD R8, R2, R5;
/*00b8*/ FADD R11, R6, R0;
/*00c8*/ ST.E.128 [R12], R8;
/*00d0*/ EXIT;
.L_1:
/*00d8*/ BRA `(.L_1);
.L_39:
这里有你的LD.E.128。
与 nvcc 版本 7.5、V7.5.17 一起编译。
【讨论】：
感谢您的回答。我的问题是，您为哪种架构编译？你能为 sm_21 编译你的kernel2 吗？我记得完全应用了您的方法（使用临时变量并将其解压缩到本地数组中），并获得了 6 个单独的负载。
我最初是为 sm_35 编译的，但对于 sm_21 我得到了相同的结果。
是的，我也是。我一定是第一次做错了什么。我也可以用 nvcc 6.5 编译并得到相同的结果。非常感谢！以上是关于如何说服 nvcc 使用 128 位宽负载？的主要内容，如果未能解决你的问题，请参考以下文章
如何将英特尔 C++ 编译器与 CUDA nvcc 一起使用？