- iOS/Objective-C 元类和类别
- objective-c - -1001 错误,当 NSURLSession 通过 httpproxy 和/etc/hosts
- java - 使用网络类获取 url 地址
- ios - 推送通知中不播放声音
我使用 NVIDIA Visual Profiler 来分析我的代码。测试内核是:
//////////////////////////////////////////////////////////////// Group 1
static __global__ void gpu_test_divergency_0(float *a, float *b)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < 0)
{
a[tid] = tid;
}
else
{
b[tid] = tid;
}
}
static __global__ void gpu_test_divergency_1(float *a, float *b)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid == 0)
{
a[tid] = tid;
}
else
{
b[tid] = tid;
}
}
static __global__ void gpu_test_divergency_2(float *a, float *b)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid >= 0)
{
a[tid] = tid;
}
else
{
b[tid] = tid;
}
}
static __global__ void gpu_test_divergency_3(float *a, float *b)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid > 0)
{
a[tid] = tid;
}
else
{
b[tid] = tid;
}
}
//////////////////////////////////////////////////////////////// Group 2
static __global__ void gpu_test_divergency_4(float *a, float *b)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < 0)
{
a[tid] = tid + 1;
}
else
{
b[tid] = tid + 2;
}
}
static __global__ void gpu_test_divergency_5(float *a, float *b)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid == 0)
{
a[tid] = tid + 1;
}
else
{
b[tid] = tid + 2;
}
}
static __global__ void gpu_test_divergency_6(float *a, float *b)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid >= 0)
{
a[tid] = tid + 1;
}
else
{
b[tid] = tid + 2;
}
}
static __global__ void gpu_test_divergency_7(float *a, float *b)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid > 0)
{
a[tid] = tid + 1;
}
else
{
b[tid] = tid + 2;
}
}
//////////////////////////////////////////////////////////////// Group 3
static __global__ void gpu_test_divergency_8(float *a, float *b)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < 0)
{
a[tid] = tid + 1.0;
}
else
{
b[tid] = tid + 2.0;
}
}
static __global__ void gpu_test_divergency_9(float *a, float *b)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid == 0)
{
a[tid] = tid + 1.0;
}
else
{
b[tid] = tid + 2.0;
}
}
static __global__ void gpu_test_divergency_10(float *a, float *b)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid >= 0)
{
a[tid] = tid + 1.0;
}
else
{
b[tid] = tid + 2.0;
}
}
static __global__ void gpu_test_divergency_11(float *a, float *b)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid > 0)
{
a[tid] = tid + 1.0;
}
else
{
b[tid] = tid + 2.0;
}
}
当我使用 <<< 1, 32 >>> 启动测试内核时,我从分析器中得到了这样的结果:
gpu_test_divergency_0 : Branch Efficiency = 100% branch = 1 divergent branch = 0
gpu_test_divergency_1 : Branch Efficiency = 100% branch = 1 divergent branch = 0
gpu_test_divergency_2 : Branch Efficiency = 100% branch = 1 divergent branch = 0
gpu_test_divergency_3 : Branch Efficiency = 100% branch = 1 divergent branch = 0
gpu_test_divergency_4 : Branch Efficiency = 100% branch = 3 divergent branch = 0
gpu_test_divergency_5 : Branch Efficiency = 100% branch = 3 divergent branch = 0
gpu_test_divergency_6 : Branch Efficiency = 100% branch = 2 divergent branch = 0
gpu_test_divergency_7 : Branch Efficiency = 100% branch = 3 divergent branch = 0
gpu_test_divergency_8 : Branch Efficiency = 100% branch = 3 divergent branch = 0
gpu_test_divergency_9 : Branch Efficiency = 75% branch = 4 divergent branch = 1
gpu_test_divergency_10 : Branch Efficiency = 100% branch = 2 divergent branch = 0
gpu_test_divergency_11 : Branch Efficiency = 75% branch = 4 divergent branch = 1
当我使用 <<< 1, 64 >>> 启动测试内核时,我从分析器中得到了这样的结果:
gpu_test_divergency_0 : Branch Efficiency = 100% branch = 2 divergent branch = 0
gpu_test_divergency_1 : Branch Efficiency = 100% branch = 2 divergent branch = 0
gpu_test_divergency_2 : Branch Efficiency = 100% branch = 2 divergent branch = 0
gpu_test_divergency_3 : Branch Efficiency = 100% branch = 2 divergent branch = 0
gpu_test_divergency_4 : Branch Efficiency = 100% branch = 6 divergent branch = 0
gpu_test_divergency_5 : Branch Efficiency = 100% branch = 6 divergent branch = 0
gpu_test_divergency_6 : Branch Efficiency = 100% branch = 4 divergent branch = 0
gpu_test_divergency_7 : Branch Efficiency = 100% branch = 5 divergent branch = 0
gpu_test_divergency_8 : Branch Efficiency = 100% branch = 6 divergent branch = 0
gpu_test_divergency_9 : Branch Efficiency = 85.7% branch = 7 divergent branch = 1
gpu_test_divergency_10 : Branch Efficiency = 100% branch = 4 divergent branch = 0
gpu_test_divergency_11 : Branch Efficiency = 83.3% branch = 6 divergent branch = 1
我在 Linux 上使用 CUDA Capability 2.0 和 NVIDIA Visual Profiler v4.2 的“GeForce GTX 570”。根据文件:
“分支”-“执行内核的线程采用的分支数。如果 warp 中至少有一个线程采用分支,则此计数器将递增 1。”
“发散分支” - “一个 warp 中发散分支的数量。如果 warp 中至少有一个线程通过数据依赖项发散(即遵循不同的执行路径),则此计数器将递增 1条件分支。”
但我对结果真的很困惑。为什么每个测试组的“分支”数量不同?为什么只有第三个测试组似乎有正确的“发散分支”?
@JackOLantern:我在 Release模式下编译。我按照你的方法拆了它。 “gpu_test_divergency_4”的结果和你的完全一样,但是“gpu_test_divergency_0”的结果不同:
Function : _Z21gpu_test_divergency_0PfS_
/*0000*/ /*0x00005de428004404*/ MOV R1, c [0x1] [0x100];
/*0008*/ /*0x94001c042c000000*/ S2R R0, SR_CTAid_X;
/*0010*/ /*0x84009c042c000000*/ S2R R2, SR_Tid_X;
/*0018*/ /*0x20009ca320044000*/ IMAD R2, R0, c [0x0] [0x8], R2;
/*0020*/ /*0xfc21dc23188e0000*/ ISETP.LT.AND P0, pt, R2, RZ, pt;
/*0028*/ /*0x0920de0418000000*/ I2F.F32.S32 R3, R2;
/*0030*/ /*0x9020204340004000*/ @!P0 ISCADD R0, R2, c [0x0] [0x24], 0x2;
/*0038*/ /*0x8020804340004000*/ @P0 ISCADD R2, R2, c [0x0] [0x20], 0x2;
/*0040*/ /*0x0000e08590000000*/ @!P0 ST [R0], R3;
/*0048*/ /*0x0020c08590000000*/ @P0 ST [R2], R3;
/*0050*/ /*0x00001de780000000*/ EXIT;
我想,正如您所说,转换指令(在本例中为 I2F)不会添加额外的分支。
但我看不到这些反汇编代码与 Profiler 结果之间的关系。我从另一篇文章(https://devtalk.nvidia.com/default/topic/463316/branch-divergent-branches/)了解到,发散分支是根据 SM 上的实际线程(warp)运行情况计算的。所以我估计我们不能仅仅根据这些反汇编代码来推导出每次实际运行的分支发散。我对吗?
最佳答案
跟进 - 使用 VOTE Intrinsics 检查线程分歧
我认为检查 warps 内线程分歧的最佳方法是使用投票内在函数,尤其是 __ballot
和 __popc
内在函数。关于 __ballot
和 __popc
的很好的解释可以在 Shane Cook,CUDA Programming,Morgan Kaufmann 的书中找到。
__ballot
的原型(prototype)如下
unsigned int __ballot(int predicate);
如果谓词非零,__ballot
返回一个设置了第 N
位的值,其中 N
是 threadIdx.x
.
另一方面,__popc
返回使用 32
位参数设置的位数。
因此,通过联合使用 __ballot
、__popc
和 atomicAdd
,可以检查 warp 是否发散。
为此,我设置了如下代码
#include <cuda.h>
#include <stdio.h>
#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
__device__ unsigned int __ballot_non_atom(int predicate)
{
if (predicate != 0) return (1 << (threadIdx.x % 32));
else return 0;
}
__global__ void gpu_test_divergency_0(unsigned int* d_ballot, int Num_Warps_per_Block)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
const unsigned int warp_num = threadIdx.x >> 5;
atomicAdd(&d_ballot[warp_num+blockIdx.x*Num_Warps_per_Block],__popc(__ballot_non_atom(tid > 2)));
// atomicAdd(&d_ballot[warp_num+blockIdx.x*Num_Warps_per_Block],__popc(__ballot(tid > 2)));
}
#include <conio.h>
int main(int argc, char *argv[])
{
unsigned int Num_Threads_per_Block = 64;
unsigned int Num_Blocks_per_Grid = 1;
unsigned int Num_Warps_per_Block = Num_Threads_per_Block/32;
unsigned int Num_Warps_per_Grid = (Num_Threads_per_Block*Num_Blocks_per_Grid)/32;
unsigned int* h_ballot = (unsigned int*)malloc(Num_Warps_per_Grid*sizeof(unsigned int));
unsigned int* d_ballot; cudaMalloc((void**)&d_ballot, Num_Warps_per_Grid*sizeof(unsigned int));
for (int i=0; i<Num_Warps_per_Grid; i++) h_ballot[i] = 0;
cudaMemcpy(d_ballot, h_ballot, Num_Warps_per_Grid*sizeof(unsigned int), cudaMemcpyHostToDevice);
gpu_test_divergency_0<<<Num_Blocks_per_Grid,Num_Threads_per_Block>>>(d_ballot,Num_Warps_per_Block);
cudaMemcpy(h_ballot, d_ballot, Num_Warps_per_Grid*sizeof(unsigned int), cudaMemcpyDeviceToHost);
for (int i=0; i<Num_Warps_per_Grid; i++) {
if ((h_ballot[i] == 0)||(h_ballot[i] == 32)) std::cout << "Warp " << i << " IS NOT divergent- Predicate true for " << h_ballot[i] << " threads\n";
else std::cout << "Warp " << i << " IS divergent - Predicate true for " << h_ballot[i] << " threads\n";
}
getch();
return EXIT_SUCCESS;
}
请注意,我现在正在计算能力为 1.2 的卡上运行代码,因此在上面的示例中,我使用的是 __ballot_non_atom
,它是 的非固有等效项__ballot
,因为 __ballot
仅适用于 >= 2.0 的计算能力。换句话说,如果你有一张计算能力>=2.0的卡,请取消注释内核函数中使用__ballot
的指令。
使用上面的代码,您可以通过简单地更改内核函数中的相关谓词来使用上面的所有内核函数。
上一个答案
我在 release 模式下为计算能力 2.0
编译了你的代码,我使用 -keep
来保留中间文件和 cuobjdump
实用程序来生成两个内核的反汇编,即:
static __global__ void gpu_test_divergency_0(float *a, float *b)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < 0) a[tid] = tid;
else b[tid] = tid;
}
和
static __global__ void gpu_test_divergency_4(float *a, float *b)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < 0) a[tid] = tid + 1;
else b[tid] = tid + 2;
}
结果如下
gpu_test_divergency_0
/*0000*/ MOV R1, c[0x1][0x100]; /* 0x2800440400005de4 */
/*0008*/ S2R R0, SR_CTAID.X; /* 0x2c00000094001c04 */
/*0010*/ S2R R2, SR_TID.X; /* 0x2c00000084009c04 */
/*0018*/ IMAD R2, R0, c[0x0][0x8], R2; /* 0x2004400020009ca3 */
/*0020*/ ISETP.LT.AND P0, PT, R2, RZ, PT; /* 0x188e0000fc21dc23 */
/*0028*/ I2F.F32.S32 R0, R2; /* 0x1800000009201e04 */
/*0030*/ @!P0 ISCADD R3, R2, c[0x0][0x24], 0x2; /* 0x400040009020e043 */
/*0038*/ @P0 ISCADD R2, R2, c[0x0][0x20], 0x2; /* 0x4000400080208043 */
/*0040*/ @!P0 ST [R3], R0; /* 0x9000000000302085 */
/*0048*/ @P0 ST [R2], R0; /* 0x9000000000200085 */
/*0050*/ EXIT ; /* 0x8000000000001de7 */
和
gpu_test_divergency_4
/*0000*/ MOV R1, c[0x1][0x100]; /* 0x2800440400005de4 */
/*0008*/ S2R R0, SR_CTAID.X; /* 0x2c00000094001c04 */ R0 = BlockIdx.x
/*0010*/ S2R R2, SR_TID.X; /* 0x2c00000084009c04 */ R2 = ThreadIdx.x
/*0018*/ IMAD R0, R0, c[0x0][0x8], R2; /* 0x2004400020001ca3 */ R0 = R0 * c + R2
/*0020*/ ISETP.LT.AND P0, PT, R0, RZ, PT; /* 0x188e0000fc01dc23 */ If statement
/*0028*/ @P0 BRA.U 0x58; /* 0x40000000a00081e7 */ Branch 1 - Jump to 0x58
/*0030*/ @!P0 IADD R2, R0, 0x2; /* 0x4800c0000800a003 */ Branch 2 - R2 = R0 + 2
/*0038*/ @!P0 ISCADD R0, R0, c[0x0][0x24], 0x2; /* 0x4000400090002043 */ Branch 2 - Calculate gmem address
/*0040*/ @!P0 I2F.F32.S32 R2, R2; /* 0x180000000920a204 */ Branch 2 - R2 = R2 after int to float cast
/*0048*/ @!P0 ST [R0], R2; /* 0x900000000000a085 */ Branch 2 - gmem store
/*0050*/ @!P0 BRA.U 0x78; /* 0x400000008000a1e7 */ Branch 2 - Jump to 0x78 (exit)
/*0058*/ @P0 IADD R2, R0, 0x1; /* 0x4800c00004008003 */ Branch 1 - R2 = R0 + 1
/*0060*/ @P0 ISCADD R0, R0, c[0x0][0x20], 0x2; /* 0x4000400080000043 */ Branch 1 - Calculate gmem address
/*0068*/ @P0 I2F.F32.S32 R2, R2; /* 0x1800000009208204 */ Branch 1 - R2 = R2 after int to float cast
/*0070*/ @P0 ST [R0], R2; /* 0x9000000000008085 */ Branch 1 - gmem store
/*0078*/ EXIT ; /* 0x8000000000001de7 */
从上面的反汇编中,我希望你的分支发散测试的结果是相同的。
您是在 Debug模式还是 Release模式下编译?
关于linux - CUDA - 关于 “branch” 和 “divergent branch” 的 Visual Profiler 结果的混淆 (2),我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/19334589/
我试图让 Branch.io 在 Android 上工作,但我遇到了: myapplication.MainActivity cannot be cast to android.app.Applica
当我执行 git branch 时,我知道我在分支 v0.2 上。 git branch v0.1 * v0.2 但是当我执行 git push 时它说“你当前分支的上游分支与你当前分支的名称不
在使用 Git GUI 检查远程分支 releases/rel_5.4.1 之后,当我尝试 push 时看到了这个意外的错误消息: fatal: The upstream branch of your
SO 上有一个相关问题处理如何更改 push 命令的参数以避免出现此消息: fatal: The upstream branch of your current branch does not mat
arc feature [branch-name] 和 git branch [branch-name] 有什么区别? 他们似乎都创建了一个新分支。 最佳答案 arc feature [branch-
FIFO、LIFO 和LC Branch and Bound 有什么区别? 最佳答案 Branch & Bound 通过使用估计边界来限制可能解决方案的数量来发现完整搜索空间内的分支。不同的类型(FI
有人知道这两个切换和跟踪远程分支的命令之间的区别吗? git checkout -b branch origin/branch git checkout --track origin/branch 我
关于 git-svn 工作流程有很多问题,但我一直无法弄清楚这一点: This section of the svn book谈到 SVN 的一个常见做法:创建一个分支,并在主干更新时不断合并主干中的
我正在构建一个控制 git 存储库的 PHP 应用程序。我有一个执行命令 git status 的同步函数,虽然没有返回“你的分支是最新的”,但可能会提前采取必要的行动,比如远程或本地分支。 我还构建
是否可以使用 branch.io 创建自定义链接,例如 https://example.app.link/fzmLEhobLD所以我可以用我的自定义 10 位参数(如 amitpp8888)控制 fz
我从 github 克隆了一个分支,它的名字是 dev。我已经开始使用它, pull 和推送代码更改并确保我的本地存储库与远程存储库保持同步。我要开始实现一个新功能,因此创建了一个新分支,如下所示:
我们有一个发布模型,为简单起见,我们假设每月 1 次。所以,我们通常会去: Jan -> trunk trunk -> Feb trunk trunk et
使用 Branch.io HTTP API 创建的链接不会在 Branch 门户中显示为快速链接。快速链接很方便,因为它们在一个 View 中显示“点击”、“打开”等内容 用于创建链接的 API:li
我创建了一个分支,当我第一次从源代码合并到分支时,出现了一大堆旧的变更集,它说没有合并,但它们在分支之前就存在,我确认它们在那里。 例子: 假设当 Source 中有 9 个变更集时,我从 Sourc
这是关于我为什么这样做 不是 收到错误“致命:当前分支 A 没有上游分支”。 我删除了 远程 分公司 一个 使用命令 git push origin :A . 然后我切换到本地 分支 A 使用命令 g
我正在使用 clover 插件来检查我的 java 代码测试覆盖率。 我为所有行编写了单元测试。当我点击红线时,它显示“true分支执行了2次,分支执行了0次”。这是什么意思?我该如何解决这个问题?
很确定我误解了 git。 我的目标 我在 github 上有一个带有“master”分支的私有(private)存储库。 我还想有一个生产分支,我会将所有更改从 master 推送到该分支。 然后我想
我将一个相当老的主题分支重新定位到 master 上。由于在 rebase 期间有很多冲突,我想将旧主题分支与 rebased 分支进行比较,以确保我没有意外删除或搞砸主题上的任何更改。我最接近的是比
我正在尝试将我的一个项目推送到 github,但我一直收到此错误: peeplesoft@jane3:~/846156 (master) $ git push fatal: The current b
Jenkins Git 插件根据我的引用规范在控制台输出中生成了以下命令 下面两个命令有什么区别?他们的输出看起来没什么不同。我在下面给出了他们的输出: 命令 1: git fetch --no-ta
我是一名优秀的程序员,十分优秀!