如何说服 nvcc 使用 128 位宽的负载？
How to convince nvcc to use 128-bit wide loads?
我有一个内核需要对一个数组应用模板操作并将结果存储在另一个数组上。模板可以用函数表示为：
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)），其中假设我的 values 数组只包含 5 个值，而 float4 加载操作修改了数组的前 4 个值（最后一个值我还有一个额外的加载）。在那种情况下，编译器使用 LD.E.128.
我的问题是，如果目标寄存器不是本地数组中的第一个寄存器，为什么编译器不理解它可以使用 128 位宽读取。毕竟本地数组 values 只是一种编程方式，表示我需要将 6 个浮点数存储在寄存器中。在生成的 ptx 或 SASS 代码中没有像数组这样的东西。我以为我给了编译器足够的提示，让它理解 LD.E.128 是这里的正确指令。
第二个问题：如何让它在这里使用128宽负载而不需要手动编写低级代码？ （但是，如果一些 asm 指令有帮助，我愿意接受建议。）
旁注：使用 32 位加载读取输入和使用 128 位存储写入输入的决定是在生成 ptx 代码时做出的。 ptx 代码已经显示了这种多个小负载和一个大存储的模式。
我在 linux 下使用 CUDA 7.5。
根据评论中的建议，我做了一些实验。
将 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 片段中显示的相同结果（一个 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 编译。
我有一个内核需要对一个数组应用模板操作并将结果存储在另一个数组上。模板可以用函数表示为：
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)），其中假设我的 values 数组只包含 5 个值，而 float4 加载操作修改了数组的前 4 个值（最后一个值我还有一个额外的加载）。在那种情况下，编译器使用 LD.E.128.
我的问题是，如果目标寄存器不是本地数组中的第一个寄存器，为什么编译器不理解它可以使用 128 位宽读取。毕竟本地数组 values 只是一种编程方式，表示我需要将 6 个浮点数存储在寄存器中。在生成的 ptx 或 SASS 代码中没有像数组这样的东西。我以为我给了编译器足够的提示，让它理解 LD.E.128 是这里的正确指令。
第二个问题：如何让它在这里使用128宽负载而不需要手动编写低级代码？ （但是，如果一些 asm 指令有帮助，我愿意接受建议。）
旁注：使用 32 位加载读取输入和使用 128 位存储写入输入的决定是在生成 ptx 代码时做出的。 ptx 代码已经显示了这种多个小负载和一个大存储的模式。
我在 linux 下使用 CUDA 7.5。
根据评论中的建议，我做了一些实验。
将 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 片段中显示的相同结果（一个 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 编译。