Skip to content

Commit ed80e9f

Browse files
committed
use CUDA_KERNEL_LOOP_TYPE
1 parent 1ff0a80 commit ed80e9f

File tree

2 files changed

+3
-4
lines changed

2 files changed

+3
-4
lines changed

paddle/phi/kernels/gpu/index_select_grad_kernel.cu

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -35,7 +35,7 @@ __global__ void index_select_grad_cuda_kernel(const T* output_grad,
3535
int64_t stride,
3636
int64_t size,
3737
int64_t delta) {
38-
CUDA_KERNEL_LOOP(idx, N) {
38+
CUDA_KERNEL_LOOP_TYPE(idx, N, int64_t) {
3939
int64_t pre_idx = idx / (stride * size);
4040
int64_t dim_idx = idx % (stride * size) / stride;
4141
IndexT src_dim_idx = index[dim_idx];
@@ -47,8 +47,7 @@ __global__ void index_select_grad_cuda_kernel(const T* output_grad,
4747

4848
template <typename T>
4949
__global__ void index_select_grad_init(T* input_grad, int64_t N) {
50-
int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;
51-
CUDA_KERNEL_LOOP(idx, N) {
50+
CUDA_KERNEL_LOOP_TYPE(idx, N, int64_t) {
5251
input_grad[idx] = 0.0;
5352
}
5453
}

paddle/phi/kernels/gpu/index_select_kernel.cu

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -32,7 +32,7 @@ __global__ void index_select_cuda_kernel(const T* input,
3232
int64_t stride,
3333
int64_t size,
3434
int64_t delta) {
35-
CUDA_KERNEL_LOOP(idx, N) {
35+
CUDA_KERNEL_LOOP_TYPE(idx, N, int64_t) {
3636
int64_t pre_idx = idx / (stride * size);
3737
int64_t dim_idx = idx % (stride * size) / stride;
3838
IndexT src_dim_idx = index[dim_idx];

0 commit comments

Comments
 (0)