add block and grid loop for index_sample kernel to deal with a large-shape tensor #37816
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
Thanks for your contribution! |
|
Sorry to inform you that 9861e2f's CIs have passed for more than 7 days. To prevent PR conflicts, you need to re-run all CIs manually. |
thisjiang
reviewed
Jan 19, 2022
| unsigned int index_i = blockDim.x * blockIdx.x + threadIdx.x; | ||
| unsigned int index_j = blockDim.y * blockIdx.y + threadIdx.y; | ||
| for (; index_j < batch_size; index_j += blockDim.y * gridDim.y) { | ||
| index_i = blockDim.x * blockIdx.x + threadIdx.x; |
There was a problem hiding this comment.
这个确定不是冗余的么😂完全没必要重新计算一遍吧？
| unsigned int index_j = blockDim.y * blockIdx.y + threadIdx.y; | ||
|
|
||
| for (; index_j < batch_size; index_j += blockDim.y * gridDim.y) { | ||
| index_i = blockDim.x * blockIdx.x + threadIdx.x; |
| @@ -153,9 +166,16 @@ class IndexSampleGradKernel<platform::CUDADeviceContext, T> | |||
| auto block_height = | |||
| platform::RoundToPowerOfTwo(index_length * batch_size) / block_width; | |||
| dim3 block_dim(block_width, block_height); | |||
| unsigned int threads = 512; | |||
There was a problem hiding this comment.
重复代码，可以提取出来：
void CheckLaunchParamValid(const framework::ExecutionContext& ctx, dim3* block_dim, dim3* grid_dim) {
unsigned int threads = 512;
block_dim->x = block_dim->x < threads ? block_dim->x : threads;
block_dim->y = block_dim->y < threads ? block_dim->y : threads;
dim3 max_grid_dim =
ctx.template device_context<platform::CUDADeviceContext>()
.GetCUDAMaxGridDimSize();
grid_dim->x = grid_dim->x < max_grid_dim.x ? grid_dim->x : max_grid_dim.x;
grid_dim->y = grid_dim->y < max_grid_dim.y ? grid_dim->y : max_grid_dim.y;
}然后调用
CheckLaunchParamValid(ctx, &block_dim, &grid_dim);而非重复写两次。
There was a problem hiding this comment.
定义了函数MIN检查block dim，函数LimitGridDim检查grid dim。感谢～
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
PR types
Bug fixes
PR changes
OPs
Describe
When the length of input tensor is large than block_dim * grid_dim，the index_sample kernel would not deal with the exceeding part. So we add block and grid loop in the kernel.