对于我的项目，我以两种不同的方式为一些函数生成PTX指令。第一种方法使用CUDA C实现函数并使用nvcc编译它们，使用
我在这里遇到的问题是一些函数在第二种情况下表现更差。其他函数的性能相对更好或者差不多。
现在我想知道为什么有些函数的性能会出现如此大的差异（以及为什么其他函数没有），但使用nsight进行分析还没有给我任何好的想法。
到目前为止，我找到的唯一区别是寄存器使用情况。在生成的ptx代码中，我可以看到以下内容：
使用nvcc编译
据我所知，这表示使用的虚拟寄存器的数量和类型，但显然，第二种情况下这是不正确的。通过使用nsight进行分析，我可以看到每个线程实际使用的寄存器数量在第一种情况下为8，在第二种情况下为31。当然，这可能是第二种情况下代码运行较慢的原因之一，但问题在于，我从LLVM IR编译成PTX使用NVPTX的所有函数都存在这个问题。它们都有396个使用的虚拟寄存器，并且nsight报告对于所有这些函数，每个线程使用了31个寄存器，即使有些函数的性能几乎与第一种情况完全相同。
这个寄存器是否是我的速度变慢的问题？为什么不影响所有函数？如果不是，可能会导致速度变慢的原因是什么？您能给出任何提示，我应该朝哪个方向查找？
谢谢！
（使用的LLVM版本为3.3）
编辑：我注意到的另一个区别是停顿原因：
NVCC：
图1 Stall reasons for nvcc compiled code NVPTX：
图2 Stall reasons for nvptx compiled code 显然，“其他”原因的相对增加。也许这可以解释问题？
编辑：添加了PTX源代码
此处显示的函数将数据从全局内存复制到共享内存。然后，每个线程将自己的元素和前一个元素与数组中的最后一个元素进行比较。如果比较结果为正，则将索引写入输出数组。
1）使用NVPTX将LLVM IR编译成PTX
nvcc -ptx <file>.cu -o <file>.ptx。另一种方法使用不同的语言编写代码，从中生成LLVM IR，并使用NVPTX后端将其编译成ptx。我在这里遇到的问题是一些函数在第二种情况下表现更差。其他函数的性能相对更好或者差不多。
现在我想知道为什么有些函数的性能会出现如此大的差异（以及为什么其他函数没有），但使用nsight进行分析还没有给我任何好的想法。
到目前为止，我找到的唯一区别是寄存器使用情况。在生成的ptx代码中，我可以看到以下内容：
使用nvcc编译
.reg .u32 %r<8>;
.reg .u64 %rd<17>;
.reg .f32 %f<8>;
.reg .pred %p<5>;
使用nvptx编译
.reg .pred %p<396>;
.reg .s16 %rc<396>;
.reg .s16 %rs<396>;
.reg .s32 %r<396>;
.reg .s64 %rl<396>;
.reg .f64 %fl<396>;
据我所知，这表示使用的虚拟寄存器的数量和类型，但显然，第二种情况下这是不正确的。通过使用nsight进行分析，我可以看到每个线程实际使用的寄存器数量在第一种情况下为8，在第二种情况下为31。当然，这可能是第二种情况下代码运行较慢的原因之一，但问题在于，我从LLVM IR编译成PTX使用NVPTX的所有函数都存在这个问题。它们都有396个使用的虚拟寄存器，并且nsight报告对于所有这些函数，每个线程使用了31个寄存器，即使有些函数的性能几乎与第一种情况完全相同。
这个寄存器是否是我的速度变慢的问题？为什么不影响所有函数？如果不是，可能会导致速度变慢的原因是什么？您能给出任何提示，我应该朝哪个方向查找？
谢谢！
（使用的LLVM版本为3.3）
编辑：我注意到的另一个区别是停顿原因：
NVCC：
图1 Stall reasons for nvcc compiled code NVPTX：
图2 Stall reasons for nvptx compiled code 显然，“其他”原因的相对增加。也许这可以解释问题？
编辑：添加了PTX源代码
此处显示的函数将数据从全局内存复制到共享内存。然后，每个线程将自己的元素和前一个元素与数组中的最后一个元素进行比较。如果比较结果为正，则将索引写入输出数组。
1）使用NVPTX将LLVM IR编译成PTX
// .globl julia_cuda_find_weighted_median18585
.entry julia_cuda_find_weighted_median18585(
.param .u64 .ptr .global .align 4 julia_cuda_find_weighted_median18585_param_0,
.param .u64 .ptr .global .align 4 julia_cuda_find_weighted_median18585_param_1
) // @julia_cuda_find_weighted_median18585
{
.reg .pred %p<396>;
.reg .s16 %rc<396>;
.reg .s16 %rs<396>;
.reg .s32 %r<396>;
.reg .s64 %rl<396>;
.reg .f64 %fl<396>;
// BB#0: // %top
mov.u32 %r0, %tid.x;
cvt.s64.s32 %rl4, %r0;
mov.u32 %r0, %ctaid.x;
cvt.s64.s32 %rl0, %r0;
mov.u32 %r1, %ntid.x;
cvt.s64.s32 %rl3, %r1;
mad.lo.s64 %rl5, %rl3, %rl0, %rl4;
setp.gt.s64 %p0, %rl5, -1;
@%p0 bra BB9_2;
bra.uni BB9_1;
BB9_2: // %idxend
ld.param.u64 %rl6, [julia_cuda_find_weighted_median18585_param_0];
ld.param.u64 %rl2, [julia_cuda_find_weighted_median18585_param_1];
add.s64 %rl1, %rl4, 1;
mov.u64 %rl7, shmem1;
cvta.shared.u64 %rl7, %rl7;
shl.b64 %rl5, %rl5, 2;
add.s64 %rl5, %rl6, %rl5;
ld.global.f32 %f0, [%rl5];
cvta.to.shared.u64 %rl5, %rl7;
shl.b64 %rl6, %rl4, 2;
add.s64 %rl4, %rl5, %rl6;
st.shared.f32 [%rl4], %f0;
bar.sync 0;
setp.lt.s64 %p0, %rl1, 2;
@%p0 bra BB9_8;
// BB#3: // %if
shl.b64 %rl3, %rl3, 2;
add.s64 %rl3, %rl3, %rl5;
ld.shared.f32 %f0, [%rl3+-4];
cvt.f64.f32 %fl0, %f0;
mul.f64 %fl0, %fl0, 0d3FE0000000000000;
add.s64 %rl3, %rl6, %rl5;
ld.shared.f32 %f0, [%rl3+-4];
cvt.f64.f32 %fl1, %f0;
setp.geu.f64 %p0, %fl1, %fl0;
@%p0 bra BB9_8;
// BB#4: // %L2
ld.shared.f32 %f0, [%rl4];
cvt.f64.f32 %fl1, %f0;
setp.gtu.f64 %p0, %fl0, %fl1;
@%p0 bra BB9_8;
// BB#5: // %if3
setp.gt.s32 %p0, %r0, -1;
@%p0 bra BB9_7;
bra.uni BB9_6;
BB9_7: // %idxend5
shl.b64 %rl0, %rl0, 2;
add.s64 %rl0, %rl2, %rl0;
st.global.u32 [%rl0], %rl1;
BB9_8: // %L6
ret;
BB9_1: // %oob
mov.u64 %rl0, cu_oob;
// Callseq Start 26
{
.reg .b32 temp_param_reg;
// <end>}
.param .b64 param0;
st.param.b64 [param0+0], %rl0;
.param .b64 param1;
st.param.b64 [param1+0], %rl0;
.param .b32 retval0;
call.uni (retval0),
vprintf,
(
param0,
param1
);
ld.param.b32 %r0, [retval0+0];
//{
}// Callseq End 26
ret;
BB9_6: // %oob4
mov.u64 %rl0, cu_oob;
// Callseq Start 27
{
.reg .b32 temp_param_reg;
// <end>}
.param .b64 param0;
st.param.b64 [param0+0], %rl0;
.param .b64 param1;
st.param.b64 [param1+0], %rl0;
.param .b32 retval0;
call.uni (retval0),
vprintf,
(
param0,
param1
);
ld.param.b32 %r0, [retval0+0];
//{
}// Callseq End 27
ret;
}
2) 使用nvcc编译的CUDA C代码转换为PTX格式
.entry findWeightedMedian_kernel (
.param .u64 __cudaparm_findWeightedMedian_kernel_input,
.param .u64 __cudaparm_findWeightedMedian_kernel_prescan,
.param .u64 __cudaparm_findWeightedMedian_kernel_output)
{
.reg .u32 %r<8>;
.reg .u64 %rd<17>;
.reg .f32 %f<8>;
.reg .pred %p<5>;
.loc 4 93 0
$LDWbegin_findWeightedMedian_kernel:
mov.u64 %rd1, temp;
.loc 4 103 0
cvt.s32.u16 %r1, %tid.y;
cvt.s64.s32 %rd2, %r1;
mul.wide.s32 %rd3, %r1, 4;
add.u64 %rd4, %rd1, %rd3;
cvt.s32.u16 %r2, %ntid.y;
cvt.s32.u16 %r3, %ctaid.x;
ld.param.u64 %rd5, [__cudaparm_findWeightedMedian_kernel_prescan];
mul.lo.s32 %r4, %r2, %r3;
add.s32 %r5, %r1, %r4;
cvt.s64.s32 %rd6, %r5;
mul.wide.s32 %rd7, %r5, 4;
add.u64 %rd8, %rd5, %rd7;
ld.global.f32 %f1, [%rd8+0];
st.shared.f32 [%rd4+0], %f1;
.loc 4 104 0
bar.sync 0;
mov.u32 %r6, 0;
setp.le.s32 %p1, %r1, %r6;
@%p1 bra $Lt_1_3074;
.loc 4 107 0
cvt.s64.s32 %rd9, %r2;
mul.wide.s32 %rd10, %r2, 4;
add.u64 %rd11, %rd1, %rd10;
ld.shared.f32 %f2, [%rd11+-4];
mov.f32 %f3, 0f3f000000; // 0.5
mul.f32 %f4, %f2, %f3;
ld.shared.f32 %f5, [%rd4+-4];
setp.lt.f32 %p2, %f5, %f4;
@!%p2 bra $Lt_1_3074;
ld.shared.f32 %f6, [%rd4+0];
setp.ge.f32 %p3, %f6, %f4;
@!%p3 bra $Lt_1_3074;
.loc 4 109 0
ld.param.u64 %rd12, [__cudaparm_findWeightedMedian_kernel_output];
cvt.s64.s32 %rd13, %r3;
mul.wide.s32 %rd14, %r3, 4;
add.u64 %rd15, %rd12, %rd14;
st.global.s32 [%rd15+0], %r1;
$Lt_1_3074:
$L_1_2050:
$Lt_1_2562:
.loc 4 111 0
exit;
$LDWend_findWeightedMedian_kernel:
} // findWeightedMedian_kernel