Cuda - Confusion About The Visual Profiler Results of "Branch" and "Divergent Branch" (2)

Can CUDA branch divergence help me in this case?

Does this mean that if the other threads in a warp diverge, I can work around this lack of resources, and still get decent performance, as long as the double-precision values that I want can be broadcast to all threads in a warp?

No, you can't.

An instruction issue always occurs at the warp level, even in a warp-diverged scenario. Since it is issued at the warp level, it will require/use/schedule enough execution resources for the warp, even for inactive threads.

Therefore a computation done on only one thread will still use the same resources/scheduling slot as a computation done on all 32 threads in the warp.

For example, a floating point multiply will require 32 instances of usage of a floating point ALU. The exact scheduling of this will vary based on the specific GPU, but you cannot reduce the 32 instance usage to a lower number through warp divergence or any other mechanism.

Based on a question in the comments, here's a worked example on CUDA 7.5, Fedora 20, GT640 (GK208 - has 1/24 ratio of DP to SP units):

$ cat t1241.cu
#include <stdio.h>
#include <time.h>
#include <sys/time.h>
#define USECPSEC 1000000ULL

unsigned long long dtime_usec(unsigned long long start){

  timeval tv;
  gettimeofday(&tv, 0);
  return ((tv.tv_sec*USECPSEC)+tv.tv_usec)-start;
}

const int nTPB = 32;
const int nBLK = 1;
const int rows = 1048576;
const int nSD = 128;

typedef double mytype;
template <bool use_warp>
__global__ void mpy_k(const mytype * in, mytype * out){
  __shared__ mytype sdata[nTPB*nSD];
  int idx = threadIdx.x + blockDim.x*blockIdx.x;
  mytype accum = in[idx];
#pragma unroll 128
  for (int i = 0; i < rows; i++)
    if (use_warp)
      accum += accum*sdata[threadIdx.x+(i&(nSD-1))*nTPB];
    else
      if (threadIdx.x == 0)
        accum += accum*sdata[threadIdx.x+(i&(nSD-1))*nTPB];
  out[idx] = accum;
}

int main(){
  mytype *din, *dout;
  cudaMalloc(&din, nTPB*nBLK*rows*sizeof(mytype));
  cudaMalloc(&dout, nTPB*nBLK*sizeof(mytype));
  cudaMemset(din, 0, nTPB*nBLK*rows*sizeof(mytype));
  cudaMemset(dout, 0, nTPB*nBLK*sizeof(mytype));
  mpy_k<true><<<nBLK, nTPB>>>(din, dout); // warm-up
  cudaDeviceSynchronize();
  unsigned long long dt = dtime_usec(0);
  mpy_k<true><<<nBLK, nTPB>>>(din, dout);
  cudaDeviceSynchronize();
  dt = dtime_usec(dt);
  printf("full warp elapsed time: %f\n", dt/(float)USECPSEC);
  mpy_k<false><<<nBLK, nTPB>>>(din, dout); //warm up
  cudaDeviceSynchronize();
  dt = dtime_usec(0);
  mpy_k<false><<<nBLK, nTPB>>>(din, dout);
  cudaDeviceSynchronize();
  dt = dtime_usec(dt);
  printf("one thread elapsed time: %f\n", dt/(float)USECPSEC);
  cudaError_t res = cudaGetLastError();
  if (res != cudaSuccess) printf("CUDA runtime failure %s\n", cudaGetErrorString(res));
  return 0;
}

$ nvcc -arch=sm_35 -o t1241 t1241.cu
$ CUDA_VISIBLE_DEVICES="1" ./t1241
full warp elapsed time: 0.034346
one thread elapsed time: 0.049174
$

It is not faster to use just one thread in the warp for a floating-point multiply

ballot' behavior on inactive lanes

From the documentation:

For each of these warp vote operations, the result excludes threads that are inactive (e.g., due to warp divergence). Inactive threads are represented by 0 bits in the value returned by __ballot() and are not considered in the reductions performed by __all() and __any().

CUDA atomics causes branch divergence

To some degree atomic implementation in CUDA may vary by GPU architecture. But specifically for the GTX 480 (a Fermi-class GPU), __shared__ memory atomics are implemented not as a single machine instruction, but in fact by a sequence of machine (SASS) instructions that form a loop.

This loop is essentially contending for a lock. When the lock is acquired by a particular thread, that thread will then complete the requested memory operation atomically on the identified shared memory cell, and then release the lock.

The process of looping to acquire the lock necessarily involves branch divergence. The possibility for branch divergence in this case is not evident from the C/C++ source code, but will be evident if you inspect the SASS code.

Global atomics are generally implemented as a single (ATOM or RED) SASS instruction. However global atomics may still involve serialization of access if executed by multiple threads in the warp. I wouldn't normally think of this as a case of "divergence" but I'm not entirely sure how the profiler would report it. If you ran an experiment that involved only global atomics, I think it would become clear.

It's possible that the reported divergence in your case is entirely due to the shared memory divergence (which is expected) as described above.

CUDA: Divergent warps penalty specifics

I found an interesting doc on this issue: pdf

From what I understand is that control flow statements (including break) define sync points for threads. In your case it would be at
i += BLOCKS*THREADS;
So lane X leaves for loop and waits for the other threads to reach abovementioned line.

---

Related Topics