- html - 出于某种原因,IE8 对我的 Sass 文件中继承的 html5 CSS 不友好?
- JMeter 在响应断言中使用 span 标签的问题
- html - 在 :hover and :active? 上具有不同效果的 CSS 动画
- html - 相对于居中的 html 内容固定的 CSS 重复背景?
在 Nsight Visual Studio 中,我们将有一个图表来呈现“已采取”、“未采取”和“分歧”分支的统计信息。我对“不采取”和“分歧”之间的区别感到困惑。例如
kernel()
{
if(tid % 32 != 31)
{...}
else
{...}
}
在我看来,当 tid %31 == 31
在扭曲中时,就会发生分歧,但什么是“未采取”?
最佳答案
来自 Nsight Visual Studio 版用户指南:
Not Taken / Taken Total: number of executed branch instructions with a uniform control flow decision; that is all active threads of a warp either take or not take the branch.
Diverged: Total number of executed branch instruction for which the conditional resulted in different outcomes across the threads of the warp. All code paths with at least one participating thread get executed sequentially. Lower numbers are better, however, check the Flow Control Efficiency to understand the impact of control flow on the device utilization.
现在,让我们考虑以下简单的代码,这可能是您当前在测试中考虑的内容:
#include<thrust\device_vector.h>
__global__ void test_divergence(int* d_output) {
int tid = threadIdx.x;
if(tid % 32 != 31)
d_output[tid] = tid;
else
d_output[tid] = 30000;
}
void main() {
const int N = 32;
thrust::device_vector<int> d_vec(N,0);
test_divergence<<<2,32>>>(thrust::raw_pointer_cast(d_vec.data()));
}
下面报告了 Nsight 生成的分支统计图表。正如您所看到的,Taken 等于 100%
,因为所有线程都会遇到 if
语句。令人惊讶的结果是您没有分歧。这可以通过查看内核函数的反汇编代码(针对 2.1 的计算能力进行编译)来解释:
MOV R1, c[0x1][0x100];
S2R R0, SR_TID.X;
SHR R2, R0, 0x1f;
IMAD.U32.U32.HI R2, R2, 0x20, R0;
LOP.AND R2, R2, -0x20;
ISUB R2, R0, R2;
ISETP.EQ.AND P0, PT, R2, 0x1f, PT;
ISCADD R2, R0, c[0x0][0x20], 0x2;
SEL R0, R0, 0x7530, !P0;
ST [R2], R0;
EXIT;
正如您所看到的,编译器能够优化反汇编代码,以便不存在分支,除了由于 EXIT
指令导致的统一分支,如所指出的由 Greg Smith 在下面的评论中提出。
编辑:遵循格雷格·史密斯评论的更复杂的示例
我现在正在考虑以下更复杂的示例
/**************************/
/* TEST DIVERGENCE KERNEL */
/**************************/
__global__ void testDivergence(float *a, float *b)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < 16) a[tid] = tid + 1;
else b[tid] = tid + 2;
}
/********/
/* MAIN */
/********/
void main() {
const int N = 64;
float* d_a; cudaMalloc((void**)&d_a,N*sizeof(float));
float* d_b; cudaMalloc((void**)&d_b,N*sizeof(float));
testDivergence<<<2,32>>>(d_a, d_b);
}
这是分支统计图表
这是反汇编代码
MOV R1, c[0x1][0x100];
S2R R0, SR_CTAID.X; R0 = blockIdx.x
S2R R2, SR_TID.X; R0 = threadIdx.x
IMAD R0, R0, c[0x0][0x8], R2; R0 = threadIdx.x + blockIdx.x * blockDim.x
ISETP.LT.AND P0, PT, R0, 0x10, PT; Checks if R0 < 16 and puts the result in predicate register P0
/*0028*/ @P0 BRA.U 0x58; If P0 = true, jumps to line 58
@!P0 IADD R2, R0, 0x2; If P0 = false, R2 = R0 + 2
@!P0 ISCADD R0, R0, c[0x0][0x24], 0x2; If P0 = false, calculates address to store b[tid] in global memory
@!P0 I2F.F32.S32 R2, R2; "
@!P0 ST [R0], R2; "
/*0050*/ @!P0 BRA.U 0x78; If P0 = false, jumps to line 78
/*0058*/ @P0 IADD R2, R0, 0x1; R2 = R0 + 1
@P0 ISCADD R0, R0, c[0x0][0x20], 0x2;
@P0 I2F.F32.S32 R2, R2;
@P0 ST [R0], R2;
/*0078*/ EXIT;
可以看出,现在反汇编代码中有两条BRA
指令。从上图中,每个扭曲都会遇到 3
分支(一个用于 EXIT
和两个 BRA
)。两个 warp 都有 1
taken 分支,因为所有线程一致地碰到 EXIT
指令。第一个经纱有 2
未采用 分支,因为两个 BRA
路径在经 yarn 程中并未统一遵循。第二个 warp 有 1
not take 分支和 1
taken 分支,因为所有 warp 线程都统一遵循其中一个分支两个BRA
。我想说,diverged* 再次等于零,因为两个分支中的指令完全相同,尽管在不同的操作数上执行。
关于cuda - CUDA中分支的概念(采取、不采取、发散),我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/24693800/
我的类有一个 foo 方法和一个 main 方法,其中有一些变量和一个 print 语句。 public static boolean foo(int x, boolean b) { if (
我正在尝试对每几列取行平均值。这是一个示例数据集。 d = {'2000-01': range(0,10), '2000-02': range(10,20), '2000-03': range(10,
在 Nsight Visual Studio 中,我们将有一个图表来呈现“已采取”、“未采取”和“分歧”分支的统计信息。我对“不采取”和“分歧”之间的区别感到困惑。例如 kernel() { if
在 Nsight Visual Studio 中,我们将有一个图表来呈现“已采取”、“未采取”和“分歧”分支的统计信息。我对“不采取”和“分歧”之间的区别感到困惑。例如 kernel() { if
int main() { long int i,t,n,q[500],d[500],s[500],res[500]={0},j,h; scanf("%ld",&t); whil
我在 Linux 上使用 racket v6.5 repl 并尝试运行流教程中的 take 函数示例 https://docs.racket-lang.org/functional-data-stru
tl;博士无法在 ggpairs 中获得独立的图例(描述整个情节的常用颜色)令我满意。 对不起,长度。 我正在尝试使用 GGally::ggpairs 绘制(下三角形)对图(用于绘制各种绘图矩阵的扩展
几个月前我问过this question 。我想添加一个具有不同背景的相同 div。我想知道为什么 jQuery 在第二个 div 中不起作用?我发现仅当我单击第二个 div 中的小图像时,图像才会在
引用Performing a right join in django ,当我尝试类似的方法时(字段略有不同): class Student: user = ForeignKey(User)
所以我使用带有 Action Sheet 样式的 UIAlertController 来显示两个选项,一个用于取消操作,另一个用于删除数据。按钮工作正常,删除按钮工作,操作表关闭。我的问题是,在后台从
我有一个列表,其中每个单元格都是一个可放置的对象,可以接受某个类的可拖动对象。该表的边框是可见的,但我不希望固定大小的单元格着色且可见,这对我来说很难看。当我拖动一个可拖动对象与一个单元格相交时,该单
我有一个 RDD,它是通过读取一个大小约为 117MB 的本地文本文件形成的。 scala> rdd res87: org.apache.spark.rdd.RDD[String] = MapPart
如果我们有 n 级台阶并且我们可以一次上 1 或 2 级台阶,则台阶数和攀登台阶的方式之间存在斐波那契关系。当且仅当我们不认为 2+1 和 1+2 不同。 但是,情况不再如此,我们还必须添加第三个选项
var query = from ch in Client.wcf.context.CashHeading where ch.Id_customer == customern//cc.Id
我是一名优秀的程序员,十分优秀!