2

我按照教科书“专业CUDA C编程”中的翘曲发散示例（以下代码）。

__global__ void math_kernel1(float *c) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float a, b;
    a = b = 0.f;
    if (tid % 2 == 0) {
        a = 100.0f;
    } else {
        b = 200.0f;
    }
    c[tid] = a + b;
}

__global__ void math_kernel2(float *c) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float a, b;
    a = b = 0.f;
    if ((tid / warpSize) % 2 == 0) {
        a = 100.0f;
    } else {
        b = 200.0f;
    }
    c[tid] = a + b;
}

__global__ void math_kernel3(float *c) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float a, b;
    bool ipred = (tid % 2 == 0);
    a = b = 0.f;
    if (ipred) {
        a = 100.0f;
    }
    if (!ipred) {
        b = 200.0f;
    }

    c[tid] = a + b;
}

显然（并且写在教科书上），math_kernel2应该有最好的分支效率，math_kernel1遵循并math_kernel3有最差的结果。然而，nvprof报告给我的结果与教科书相矛盾。我使用 CUDA 8.0 在 GTX 1080 Ti 上对这些内核进行了基准测试（我还添加了编译器标志-g -G以nvcc禁用优化），它报告了以下分支效率：

  1. 数学内核1 83.33%
  2. 数学内核2 100.00%
  3. math_kernel3 100.00%（预计小于math_kernel1，教科书上是71.43%）
4

0 回答 0