wordpress 简介主题,seo优化是啥,html5 metro风格网站模板,wordpress可视化不显示Bug描述
在测量如下一个简单的核函数的执行时间的时候#xff0c;发现测量的时间和循环的次数完全无关#xff0c;觉得很奇怪#xff0c;因为循环的次数已经很大了#xff0c;不管我再怎么提升循环次数#xff0c;这么大的计算量#xff0c;不可能保持时间的恒定。
__g…Bug描述
在测量如下一个简单的核函数的执行时间的时候发现测量的时间和循环的次数完全无关觉得很奇怪因为循环的次数已经很大了不管我再怎么提升循环次数这么大的计算量不可能保持时间的恒定。
__global__ void setRowReadRow(int * out)
{unsigned int idxthreadIdx.y*blockDim.xthreadIdx.x;for(unsigned int l00; l065536; l0)for(unsigned int l10; l165536; l1)for(unsigned int l20; l265536; l2)for(unsigned int l30; l365536; l3)for(unsigned int m0; m65536; m){out[idx] m ;}
}于是去查看该Kernel的PTX代码发现该函数主体只有一条ret指令用于函数返回没有任何计算过程:
.visible .entry setRowReadRow(int*)(.param .u64 setRowReadRow(int*)_param_0
)
{ret;}
这就解释得通为什么执行时间不变了于是尝试调小循环次数只保留变量m这一层嵌套此时PTX代码如下
.visible .entry setRowReadRow(int*)(.param .u64 setRowReadRow(int*)_param_0
)
{ld.param.u64 %rd1, [setRowReadRow(int*)_param_0];cvta.to.global.u64 %rd2, %rd1;mov.u32 %r1, %tid.y;mov.u32 %r2, %ntid.x;mov.u32 %r3, %tid.x;mad.lo.s32 %r4, %r1, %r2, %r3;mul.wide.u32 %rd3, %r4, 4;add.s64 %rd4, %rd2, %rd3;ld.global.u32 %r5, [%rd4];add.s32 %r6, %r5, 2147450880;st.global.u32 [%rd4], %r6;ret;}
这里不解释每条指令的具体含义了可以用GPT等大模型帮忙翻译一下重点解释这两条指令 add.s32 %r6, %r5, 2147450880;st.global.u32 [%rd4], %r6;%r5保存的是out[idx]的原始值%rd4保存的是out[idx]在内存中的地址所以这两条指令的意思就是out[idx]加上2147450880的值再存回去。
因为这部分代码只保留了m变量所在的那一层循环分析可得Kernel函数得到的结果就是把out[idx]的值再加上0123…65535)2147450880。
很显然编译器帮我们做了优化把65536次循环加法变成了一次加法指令再加上英伟达官方论坛的解答可以大致推测出循环次数过多导致PTX代码只有一条ret指令的原因是编译器在做优化时把循环的加法拿出去计算导致了溢出了所以产生了不可预期的错误。
但是测试的时候发现把加法改成乘法后不会产生ret错误分析ptx是因为对于乘法没有做这方面的优化老老实实按照循环嵌套写的PTX代码所以此时虽然out[idx]的计算会出现溢出但是并不影响程序的运行。加法由于编译器会对循环优化所以出现PTX的异常。
