Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

add index initialization in the block loop for index_sample kernel when dealing with a input tensor whose shape is larger than block_dim * grid_dim #39736

Merged
merged 7 commits into from
Feb 20, 2022
Merged

add index initialization in the block loop for index_sample kernel when dealing with a input tensor whose shape is larger than block_dim * grid_dim #39736

Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
9861e2f
add block and grid loop for index_sample kernel to deal with a large-…
FlyingQianMM Dec 3, 2021
bc32edd
Merge branch 'develop' of https://github.com/PaddlePaddle/Paddle into…
FlyingQianMM Jan 11, 2022
f375970
fix code format
FlyingQianMM Jan 12, 2022
10ecd53
limit grid dim
FlyingQianMM Jan 20, 2022
fe2d3c4
fix the omissive initialization of index_i in the second cycle for in…
FlyingQianMM Feb 18, 2022
85fa09e
fix conflicts
FlyingQianMM Feb 18, 2022
6cfe566
fix conflicts
FlyingQianMM Feb 18, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 2 additions & 0 deletions paddle/fluid/operators/index_sample_op.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,6 +44,7 @@ __global__ void IndexSampleForward(const IndexT* index, const T* in_data,
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;
for (; index_i < index_length; index_i += blockDim.x * gridDim.x) {
unsigned int index_idx = index_j * index_length + index_i;
unsigned int in_idx = index_j * input_length + index_i;
Expand All @@ -62,6 +63,7 @@ __global__ void IndexSampleGrad(const IndexT* index, T* in_grad,
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;
for (; index_i < index_length; index_i += blockDim.x * gridDim.x) {
unsigned int index_idx = index_j * index_length + index_i;
unsigned int in_idx = index_j * input_length + index_i;
Expand Down