There was a lot of discussion about how to select #blocks and blockSize, but I still can't see anything. Many of my problems address this question: How are CUDA Blocks / Warps / Threads mapped to CUDA kernels? (To simplify the discussion, enough memory perThread and perBlock. This is not a problem).
kernelA<<<nBlocks, nThreads>>>(varA,constB, nThreadsTotal);
1) To keep SM as busy as possible, I have to set nThreadsto multiplicity warpSize. Truth?
2) SM can only execute one core at a time. These are all SM HWcores that run only kernelA. (Not some HWcores run kernelA, while others run kernelB.) Therefore, if I have only one thread to run, I "waste" other HWcores. Truth?
3) If the problems with the warp scheduler work in units warpSize(32 threads) and each SM has 32 HWcores, then the SM will be fully used. What happens when an SM has 48 HWcores? How can I save all 48 cores to the fullest when the scheduler releases work in pieces of 32? (If the previous paragraph is correct, would it be better if the scheduler issued the work in units of size HWcore?)
4) It appears that warp-scheduler queues 2 tasks at a time. Thus, when the current executable kernel is stopped or blocked, the second core is swapped. (This is not clear, but I guess there are more than two cores in line). Is it correct?
5) HW 512 (nThreadsMax), , 512 . ( , mem .) , , 512- , . SM. ?
5a) , , , , nBlocks? , nBlocks, ? ( .) nBlocks, trial-n-err.