我有一个内核，需要对一个数组应用模板操作并将结果存储在另一个数组上。模板可以用函数表示为:
float stencil(const float* data)
{
return *(data-1) + *(data+1);
}
我希望每个线程通过加载输入数组的 6 个连续值来生成输出数组的 4 个连续值。通过这样做，我将能够使用 float4 类型来加载和存储 128 字节的 block 。这是我的程序(您可以下载并编译它，但请首先考虑内核):
#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)；对于值的两端都有两个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 ))，在这种情况下，我的 values 数组仅包含 5 个值，而 float4 加载操作会修改数组的前 4 个值(我还有一个额外的值)加载最后一个值)。在这种情况下，编译器使用 LD.E.128。
我的问题是，如果目标寄存器不是本地数组中的第一个寄存器，为什么编译器不明白它可以使用 128 位宽读取。毕竟，本地数组 values 只是一种编程方式，表示我需要将 6 个 float 存储在寄存器中。在生成的 ptx 或 SASS 代码中不存在像数组这样的东西。我认为我给了编译器足够的提示，让它理解 LD.E.128 这里是正确的指令。
第二个问题:如何让它在此处使用128宽的加载，而无需手动编写低级代码？ (但是，如果一些汇编指令有帮助，我愿意接受建议。)
旁注:在生成 ptx 代码时决定使用 32 位加载来读取输入，使用 128 位存储来写入输入。 ptx 代码已经显示了这种多个小负载和单个大存储的模式。
我在linux下使用CUDA 7.5。
根据评论中给出的建议，我做了一些实验。
将input或output声明为__restrict__(或两者)可以解决问题:编译器生成一个LD.E。 128 和两个 LD.E，这是我在为架构 sm_35 生成代码时想要实现的目标。奇怪的是，在生成 sm_21 时，它仍然会生成 6 个 LD.E，但会生成 1 个 ST.E.128。对我来说，这听起来像是一个编译器错误，因为指令 LD.E.128 在旧架构中应该完全可用，就像在最新架构中一样。
上面提供的代码使用 128 位加载，只是按照 njuffa 的建议使用 __restrict__ 关键字进行了小改动，并且可以正常工作。我也遵循了m.s.的建议。我重现了 Pastebin 代码段中显示的相同结果(一个 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
最佳答案
问题是 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 编译。
关于vector - 如何说服 nvcc 使用 128 位宽负载？，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/32781846/