Should I unify two similar kernels with an ‘if’ statement, risking performance loss?

by

You have a third alternative, which is to use C++ templating and make the variable which is used in the if/switch statement a template parameter. Instantiate each version of the kernel you need, and then you have multiple kernels doing different things with no branch divergence or conditional evaluation to worry about, because the compiler will optimize away the dead code and the branching with it.

Perhaps something like this:

template<int action>
__global__ void kernel()
{
    switch(action) {
       case 1:
       // First code
       break;

       case 2:
       // Second code
       break;
    }
}

template void kernel<1>();
template void kernel<2>();

More Related Contents:

  1. Why is NVIDIA Pascal GPUs slow on running CUDA Kernels when using cudaMallocManaged
  2. What is option -O3 for g++ and nvcc?
  3. Why does C++ code for testing the Collatz conjecture run faster than hand-written assembly?
  4. Meaning of acronym SSO in the context of std::string
  5. Position of least significant bit that is set
  6. Modular arithmetics and NTT (finite field DFT) optimizations
  7. Is “for(;;)” faster than “while (true)”? If not, why do people use it?
  8. Deoptimizing a program for the pipeline in Intel Sandybridge-family CPUs
  9. Fastest way to get IPv4 address from string
  10. Different floating point result with optimization enabled – compiler bug?
  11. __builtin_prefetch, How much does it read?
  12. GCC C++ “Hello World” program -> .exe is 500kb big when compiled on Windows. How can I reduce its size?
  13. C/C++: Optimization of pointers to string constants
  14. How to read file content into istringstream?
  15. What kinds of optimizations does ‘volatile’ prevent in C++?
  16. Trouble launching CUDA kernels from static initialization code
  17. Cuda kernel returning vectors
  18. Beyond Stack Sampling: C++ Profilers
  19. Fastest implementation of sine, cosine and square root in C++ (doesn’t need to be much accurate)
  20. Get sum of values stored in __m256d with SSE/AVX
  21. how to optimize matrix multiplication (matmul) code to run fast on a single processor core
  22. Is a logical right shift by a power of 2 faster in AVR?
  23. Branchless code that maps zero, negative, and positive to 0, 1, 2
  24. Optimizing member variable order in C++
  25. Which is faster/preferred: memset or for loop to zero out an array of doubles?
  26. How can adding code to a loop make it faster?
  27. Fastest way to produce a mask with n ones starting at position i
  28. Compiling code containing dynamic parallelism fails
  29. VS: unexpected optimization behavior with _BitScanReverse64 intrinsic
  30. how to use gpu::Stream in OpenCV?

Leave a Comment