Refine for_range #8152
Refine for_range #8152
Conversation
paddle/platform/for_range.h
Outdated
| if (block_size < 1024) { | ||
| int size = 1; | ||
| while (size < block_size) size <<= 1; | ||
| block_size = size; |
There was a problem hiding this comment.
- The warp size is 32 in Nvidia GPU, if limit_ is less than 1024, make the block size be divided by 32 by the following calculation:
if (block_size < 1024) {
block_size = ((blokc_size + 31) >> 5) << 5;
}
- The following function from line 44 to line 48 also needs to be updated.
template <typename Function>
__global__ static void ForRangeElemwiseOpGridIsOne(Function func) {
size_t idx = static_cast<size_t>(threadIdx.x);
func(idx);
}
There was a problem hiding this comment.
I am not sure this optimization is necessary.
If the number of thread cannot be divided by wrap size, there will be inactive threads. However, even we make the number of thread can be divided by wrap size, there will be inactive threads, too. Because all threads run the same code, there will be some thread not to fulfil the if statement and wait for other threads.
There was a problem hiding this comment.
@qingqing01 If so, ForRangeElemwiseOpGridIsOne should be removed.
There was a problem hiding this comment.
@reyoung It seems reasonable, but I have not seen any document directing uses to set block_size in this way.
I don't do benchmark, so I don't tell you which is better.
… feature/refine_for_range
|
I do not think this PR is useful, considering that it is not saving any computation and SMs in CUDA. |
fix #7081
Don't merge. Need to further discuss and analyze.