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

Support broadcast elementwise operators with int64 index type #45741

Merged
merged 16 commits into from
Sep 16, 2022
Merged

Support broadcast elementwise operators with int64 index type #45741

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
16 commits
Select commit Hold shift + click to select a range
55c1543
support int64 non-broadcast
sneaxiy Sep 5, 2022
b7b0100
support broadcast case for int64 index
sneaxiy Sep 7, 2022
c86679a
fix bug
sneaxiy Sep 7, 2022
ed3cc0a
merge upstream/develop to solve conflict
sneaxiy Sep 7, 2022
b1bc21f
merge upstream/develop
sneaxiy Sep 7, 2022
727229d
support more Arity
sneaxiy Sep 8, 2022
c081861
remove some codes
sneaxiy Sep 8, 2022
4535d32
upgrade patchelf to v0.15.0 to pass CI build
sneaxiy Sep 8, 2022
d0d3981
fix bug
sneaxiy Sep 8, 2022
914642d
fix patchelf installation
sneaxiy Sep 9, 2022
6490000
Merge branch 'develop' of https://github.com/PaddlePaddle/Paddle into…
sneaxiy Sep 13, 2022
1f018cb
add debug flags
sneaxiy Sep 13, 2022
0358f91
remove useless codes
sneaxiy Sep 13, 2022
86c83b4
fix viterbi_decode and set_value op uts
sneaxiy Sep 13, 2022
5cf1854
remove always enable int64
sneaxiy Sep 14, 2022
b7419c0
Merge branch 'develop' of https://github.com/PaddlePaddle/Paddle into…
sneaxiy Sep 15, 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
11 changes: 6 additions & 5 deletions paddle/fluid/platform/device/gpu/cuda/cuda_helper.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -70,11 +70,12 @@ namespace platform {
*
*/

#define CUDA_KERNEL_LOOP_TYPE(i, num, index_type) \
int64_t __index__ = \
static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x; \
for (index_type i = __index__; __index__ < (num); \
__index__ += blockDim.x * gridDim.x, i = __index__)
#define CUDA_KERNEL_LOOP_TYPE(i, num, index_type) \
int64_t __index__ = \
static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x; \
int64_t __stride__ = static_cast<int64_t>(blockDim.x) * gridDim.x; \
for (index_type i = __index__; __index__ < (num); \
__index__ += __stride__, i = __index__)

class CublasHandleHolder {
public:
Expand Down
3 changes: 2 additions & 1 deletion paddle/fluid/platform/device/gpu/rocm/rocm_helper.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -70,8 +70,9 @@ namespace platform {
#define CUDA_KERNEL_LOOP_TYPE(i, num, index_type) \
int64_t __index__ = \
static_cast<int64_t>(hipBlockIdx_x) * hipBlockDim_x + hipThreadIdx_x; \
int64_t __stride__ = static_cast<int64_t>(hipBlockDim_x) * hipGridDim_x; \
for (index_type i = __index__; __index__ < (num); \
__index__ += hipBlockDim_x * hipGridDim_x, i = __index__)
__index__ += __stride__, i = __index__)

class CublasHandleHolder {
public:
Expand Down
11 changes: 6 additions & 5 deletions paddle/phi/backends/gpu/cuda/cuda_helper.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -62,11 +62,12 @@ namespace gpu {
*
*/

#define CUDA_KERNEL_LOOP_TYPE(i, num, index_type) \
int64_t __index__ = \
static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x; \
for (index_type i = __index__; __index__ < (num); \
__index__ += blockDim.x * gridDim.x, i = __index__)
#define CUDA_KERNEL_LOOP_TYPE(i, num, index_type) \
int64_t __index__ = \
static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x; \
int64_t __stride__ = static_cast<int64_t>(blockDim.x) * gridDim.x; \
for (index_type i = __index__; __index__ < (num); \
__index__ += __stride__, i = __index__)

} // namespace gpu
} // namespace backends
Expand Down
3 changes: 2 additions & 1 deletion paddle/phi/backends/gpu/rocm/rocm_helper.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -65,8 +65,9 @@ namespace gpu {
#define CUDA_KERNEL_LOOP_TYPE(i, num, index_type) \
int64_t __index__ = \
static_cast<int64_t>(hipBlockIdx_x) * hipBlockDim_x + hipThreadIdx_x; \
int64_t __stride__ = static_cast<int64_t>(hipBlockDim_x) * hipGridDim_x; \
for (index_type i = __index__; __index__ < (num); \
__index__ += hipBlockDim_x * hipGridDim_x, i = __index__)
__index__ += __stride__, i = __index__)

} // namespace gpu
} // namespace backends
Expand Down
Loading