linux - CUDA - 关于 “branch” 和 “divergent branch” 的 Visual Profiler 结果的混淆 (2)-6ren

linux - CUDA - 关于 “branch” 和 “divergent branch” 的 Visual Profiler 结果的混淆 (2)

转载作者：塔克拉玛干更新时间：2023-11-03 01:01:48

我使用 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/

文章推荐： linux - shell bash 代码未运行且没有错误警告

文章推荐：安卓 : How to get SIM ID of both SIM in Dual SIM mobile

文章推荐： c++ - 将 wstring 转换为小写

cuda - CUDA/CUDA Thrust 中的多态性和派生类
这是我关于 Stack Overflow 的第一个问题，这是一个很长的问题。 tl;dr 版本是:我如何使用 thrust::device_vector如果我希望它存储不同类型的对象 DerivedC
cuda - 在 cuda 内核运行时访问 cuda 设备内存
我已使用 cudaMalloc 在设备上分配内存并将其传递给内核函数。是否可以在内核完成执行之前从主机访问该内存？最佳答案我能想到的在内核仍在执行时启动 memcpy 的唯一方法是在与内核不同的流
cuda - 我可以在没有 cuda 设备的情况下编译 cuda 程序吗
是否可以在同一节点上没有支持 CUDA 的设备的情况下编译 CUDA 程序，仅使用 NVIDIA CUDA Toolkit...？最佳答案你的问题的答案是肯定的。 nvcc编译器驱动程序与设备的物
cuda - Cuda 上的引用参数
我不知道 cuda 不支持引用参数。我的程序中有这两个函数: __global__ void ExtractDisparityKernel ( ExtractDisparity& es)
cuda - CUDA 内核中主机浮点常量的使用
我正在使用 CUDA 5.0。我注意到编译器将允许我在内核中使用主机声明的 int 常量。但是，它拒绝编译任何使用主机声明的 float 常量的内核。有谁知道这种看似差异的原因？例如，下面的代码可以
cuda - CUDA 中的协作组
自从 CUDA 9 发布以来，显然可以将不同的线程和 block 分组到同一组中，以便您可以一起管理它们。这对我来说非常有用，因为我需要启动一个包含多个 block 的内核并等待所有 block 都同
cuda - CUDA 中的三线性插值
我需要在 CUDA 中执行三线性插值。这是问题定义: 给定三个点向量:x[nx]、y[ny]、z[nz] 和一个函数值矩阵func[nx][ny][nz]，我想在 x、y 范围之间的一些随机点处找到函
cuda - CUDA 是否具有向量操作内在函数？
我认为由于 CUDA 可以执行 64 位 128 位加载/存储，因此它可能具有一些用于加/减/等的内在函数。像 float3 这样的向量类型，在像 SSE 这样更少的指令中。 CUDA 有这样的功能吗
cuda - CUDA 共享内存中的并行扫描
我有一个问题，每个线程 block (一维)必须对共享内存内的一个数组进行扫描，并执行几个其他任务。 (该数组最多有 1024 个元素。) 有没有支持这种操作的好库？我检查了 Thrust 和 Cu
cuda - CUDA 线程如何工作
我对线程的形成和执行方式有很多疑惑。首先，文档将 GPU 线程描述为轻量级线程。假设我希望将两个 100*100 矩阵相乘。如果每个元素都由不同的线程计算，则这将需要 100*100 个线程。但是，
cuda - CUDA 共享内存大小意味着什么
我正在尝试自己解决这个问题，但我不能。所以我想听听你的建议。我正在编写这样的内核代码。 VGA 是 GTX 580。 xxxx >> (... threadNum ...) (note. Shar
cuda - 让 CUDA Thrust 使用您选择的 CUDA 流
查看 CUDA Thrust 代码中的内核启动，似乎它们总是使用默认流。我可以让 Thrust 使用我选择的流吗？我在 API 中遗漏了什么吗？最佳答案我想在 Thrust 1.8 发布后更新 t
cuda - CUDA 扭曲调度是确定性的吗？
我想知道 CUDA 应用程序的扭曲调度顺序是否是确定性的。具体来说，我想知道在同一设备上使用相同输入数据多次运行同一内核时，warp 执行的顺序是否会保持不变。如果没有，是否有任何东西可以强制对扭曲
cuda - CUDA 中有多少个网格
一个 GPU 中可以有多少个 CUDA 网格？两个网格可以同时存在于 GPU 中吗？还是一台 GPU 设备只有一个网格？ Kernel1>(dst1, param1); Kernel1>(dst2,
cuda - CUDA 应用程序的计算能力是否会自动升级？
如果我编译一个计算能力较低的 CUDA 程序，例如 1.3(nvcc 标志 sm_13)，并在具有 Compute Capability 2.1 的设备上运行它，它是否会利用 Compute 2.1
cuda - CUDA 固定内存是零拷贝吗？
固定内存应该可以提高从主机到设备的传输速率(api 引用)。但是我发现我不需要为内核调用 cuMemcpyHtoD 来访问这些值，也不需要为主机调用 cuMemcpyDtoA 来读取值。我不认为这会奏
cuda - CUDA 是否会自动为您进行负载平衡？
我希望对 CUDA C 中负载平衡的最佳实践有一些一般性的建议和说明，特别是: 如果经纱中的 1 个线程比其他 31 个线程花费的时间长，它会阻止其他 31 个线程完成吗？如果是这样，多余的处理能力
cuda - CUDA 中是否有内置的交叉和点积？
CUDA 中是否有像 opencl 一样的内置交叉和点积，所以 cuda 内核可以使用它？到目前为止，我在规范中找不到任何内容。最佳答案您可以在 SDK 的 cutil_math.h 中找到这些
cuda - CUDA 中的动态共享内存
有一些与我要问的问题类似的问题，但我觉得它们都没有触及我真正要寻找的核心。我现在拥有的是一种 CUDA 方法，它需要将两个数组定义到共享内存中。现在，数组的大小由在执行开始后读入程序的变量给出。因此，
cuda - CUDA 扭曲中的线程是否在多处理器上并行执行？
经线是 32 根线。 32 个线程是否在多处理器中并行执行？如果 32 个线程没有并行执行，则扭曲中没有竞争条件。在经历了一些例子后，我有了这个疑问。最佳答案在 CUDA 编程模型中，warp

塔克拉玛干

个人简介

我是一名优秀的程序员,十分优秀！

作者热门文章

滴滴打车优惠券免费领取

全站热门文章

首页

博学

6Ren·AI

商城

linux - CUDA - 关于 “branch” 和 “divergent branch” 的 Visual Profiler 结果的混淆 (2)