Skip to content

Migrate std::floor to sycl::floor in device kernels #3721

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

Open
orrangetabby17 wants to merge 1 commit into
base: main
Choose a base branch
from
Open

Migrate std::floor to sycl::floor in device kernels #3721

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
4 changes: 2 additions & 2 deletions src/ATen/native/xpu/UpSample.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -185,15 +185,15 @@ static inline accscalar_t area_pixel_compute_source_index(
struct NearestIndexOp {
int operator()(const float scale, int dst_index, int input_size) const {
const int src_index =
min(static_cast<int>(floorf((dst_index)*scale)), input_size - 1);
min(static_cast<int>(sycl::floor((dst_index)*scale)), input_size - 1);
return src_index;
}
};

struct NearestExactIndexOp {
int operator()(const float scale, int dst_index, int input_size) const {
const int src_index = min(
static_cast<int>(floorf((dst_index + static_cast<float>(0.5)) * scale)),
static_cast<int>(sycl::floor((dst_index + static_cast<float>(0.5)) * scale)),
input_size - 1);
return src_index;
}
Expand Down
2 changes: 1 addition & 1 deletion src/ATen/native/xpu/sycl/BinaryDivFloorKernel.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -53,7 +53,7 @@ struct DivFloorWithScalarFunctor {

scalar_t floordiv;
if (div != 0) {
floordiv = std::floor(div);
floordiv = sycl::floor(div);
if (div - floordiv > scalar_t(0.5)) {
floordiv += scalar_t(1.0);
}
Expand Down
11 changes: 7 additions & 4 deletions src/ATen/native/xpu/sycl/DeformConv2dKernels.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -78,6 +78,7 @@
DISABLE_RETURN_TYPE_WARNING_BEGIN
// clang-format on
#include <ATen/ceil_div.h>
#include <ATen/OpMathType.h>
#include <ATen/native/xpu/sycl/Atomics.h>
#include <ATen/native/xpu/sycl/DistributionTemplates.h>
#include <ATen/native/xpu/sycl/KernelUtils.h>
Expand All @@ -101,8 +102,9 @@ scalar_t bilinear_interpolate(
return 0;
}

index_t h_low = std::floor(h);
index_t w_low = std::floor(w);
using opmath_t = at::opmath_type<scalar_t>;
Comment thread
weishi-deng marked this conversation as resolved.
index_t h_low = sycl::floor(static_cast<opmath_t>(h));
index_t w_low = sycl::floor(static_cast<opmath_t>(w));
index_t h_high = h_low + 1;
index_t w_high = w_low + 1;

Expand Down Expand Up @@ -609,8 +611,9 @@ scalar_t get_coordinate_weight(
scalar_t y,
scalar_t x,
bool is_y_direction) {
index_t y_l = std::floor(y);
index_t x_l = std::floor(x);
using opmath_t = at::opmath_type<scalar_t>;
Comment thread
weishi-deng marked this conversation as resolved.
index_t y_l = sycl::floor(static_cast<opmath_t>(y));
index_t x_l = sycl::floor(static_cast<opmath_t>(x));
index_t y_h = y_l + 1;
index_t x_h = x_l + 1;

Expand Down
30 changes: 16 additions & 14 deletions src/ATen/native/xpu/sycl/GridSampler.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -56,8 +56,8 @@ struct GridSampler2dKernelFunctor {

if (interpolation_mode_ == GridSamplerInterpolation::Bilinear) {
// get NE, NW, SE, SW pixel values from (x, y)
index_t ix_nw = static_cast<index_t>(std::floor(ix));
index_t iy_nw = static_cast<index_t>(std::floor(iy));
index_t ix_nw = static_cast<index_t>(sycl::floor(ix));
index_t iy_nw = static_cast<index_t>(sycl::floor(iy));
index_t ix_ne = ix_nw + 1;
index_t iy_ne = iy_nw;
index_t ix_sw = ix_nw;
Expand Down Expand Up @@ -113,8 +113,8 @@ struct GridSampler2dKernelFunctor {
ix = grid_sampler_unnormalize(x, inp_W_, align_corners_);
iy = grid_sampler_unnormalize(y, inp_H_, align_corners_);

opmath_t ix_nw = std::floor(ix);
opmath_t iy_nw = std::floor(iy);
opmath_t ix_nw = sycl::floor(ix);
opmath_t iy_nw = sycl::floor(iy);

const opmath_t tx = ix - ix_nw;
const opmath_t ty = iy - iy_nw;
Expand Down Expand Up @@ -370,6 +370,7 @@ Tensor grid_sampler_2d_kernel(

template <typename scalar_t, typename index_t>
struct GridSampler2dBackwardKernelFunctor {
using opmath_t = at::opmath_type<scalar_t>;
void operator()(sycl::nd_item<1> item) const {
auto index = item.get_global_linear_id();
if (index >= nthreads_)
Expand All @@ -392,8 +393,8 @@ struct GridSampler2dBackwardKernelFunctor {

if (interpolation_mode_ == GridSamplerInterpolation::Bilinear) {
// get NE, NW, SE, SW pixel values from (x, y)
index_t ix_nw = static_cast<index_t>(std::floor(ix));
index_t iy_nw = static_cast<index_t>(std::floor(iy));
index_t ix_nw = static_cast<index_t>(sycl::floor(static_cast<opmath_t>(ix)));
index_t iy_nw = static_cast<index_t>(sycl::floor(static_cast<opmath_t>(iy)));
index_t ix_ne = ix_nw + 1;
index_t iy_ne = iy_nw;
index_t ix_sw = ix_nw;
Expand Down Expand Up @@ -530,8 +531,8 @@ struct GridSampler2dBackwardKernelFunctor {
iy = grid_sampler_unnormalize_set_grad(
y, inp_H_, align_corners_, &giy_mult);

scalar_t ix_nw = std::floor(ix);
scalar_t iy_nw = std::floor(iy);
scalar_t ix_nw = static_cast<scalar_t>(sycl::floor(static_cast<opmath_t>(ix)));
scalar_t iy_nw = static_cast<scalar_t>(sycl::floor(static_cast<opmath_t>(iy)));

const scalar_t tx = ix - ix_nw;
const scalar_t ty = iy - iy_nw;
Expand Down Expand Up @@ -882,9 +883,9 @@ struct GridSampler3dKernelFunctor {
// get corner pixel values from (x, y, z)
// for 4d, we used north-east-south-west
// for 5d, we add top-bottom
index_t ix_tnw = static_cast<index_t>(std::floor(ix));
index_t iy_tnw = static_cast<index_t>(std::floor(iy));
index_t iz_tnw = static_cast<index_t>(std::floor(iz));
index_t ix_tnw = static_cast<index_t>(sycl::floor(ix));
index_t iy_tnw = static_cast<index_t>(sycl::floor(iy));
index_t iz_tnw = static_cast<index_t>(sycl::floor(iz));

index_t ix_tne = ix_tnw + 1;
index_t iy_tne = iy_tnw;
Expand Down Expand Up @@ -1227,6 +1228,7 @@ Tensor grid_sampler_3d_kernel(

template <typename scalar_t, typename index_t>
struct GridSampler3dBackwardKernelFunctor {
using opmath_t = at::opmath_type<scalar_t>;
void operator()(sycl::nd_item<1> item) const {
auto index = item.get_global_linear_id();
if (index >= nthreads_)
Expand Down Expand Up @@ -1257,9 +1259,9 @@ struct GridSampler3dBackwardKernelFunctor {
// get corner pixel values from (x, y, z)
// for 4d, we used north-east-south-west
// for 5d, we add top-bottom
index_t ix_tnw = static_cast<index_t>(std::floor(ix));
index_t iy_tnw = static_cast<index_t>(std::floor(iy));
index_t iz_tnw = static_cast<index_t>(std::floor(iz));
index_t ix_tnw = static_cast<index_t>(sycl::floor(static_cast<opmath_t>(ix)));
index_t iy_tnw = static_cast<index_t>(sycl::floor(static_cast<opmath_t>(iy)));
index_t iz_tnw = static_cast<index_t>(sycl::floor(static_cast<opmath_t>(iz)));

index_t ix_tne = ix_tnw + 1;
index_t iy_tne = iy_tnw;
Expand Down
2 changes: 1 addition & 1 deletion src/ATen/native/xpu/sycl/Philox4x32.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -569,7 +569,7 @@ static inline unsigned int rand_poisson_gammainc(
while (true) {
y = rand_uniform(state);
x = pgammaincinv(lambda, y);
x = floorf(x);
x = sycl::floor(x);
z = rand_uniform(state);
v = (pgammainc(lambda, x + 1.0f) - pgammainc(lambda, x)) * 1.3f;
z = z * v;
Expand Down
13 changes: 9 additions & 4 deletions src/ATen/native/xpu/sycl/PsRoiPoolKernels.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,6 +17,7 @@
DISABLE_RETURN_TYPE_WARNING_BEGIN
// clang-format on
#include <ATen/ceil_div.h>
#include <ATen/OpMathType.h>
#include <ATen/native/xpu/sycl/Atomics.h>
#include <ATen/native/xpu/sycl/KernelUtils.h>
#include <comm/SYCLContext.h>
Expand Down Expand Up @@ -52,11 +53,13 @@ struct PsRoiPoolForwardKernel {
T bin_size_h =
static_cast<T>(roi_height) / static_cast<T>(pooled_height_);
T bin_size_w = static_cast<T>(roi_width) / static_cast<T>(pooled_width_);

using opmath_t = at::opmath_type<T>;

int hstart =
static_cast<int>(std::floor(static_cast<T>(ph) * bin_size_h));
static_cast<int>(sycl::floor(static_cast<opmath_t>(static_cast<T>(ph) * bin_size_h)));
int wstart =
static_cast<int>(std::floor(static_cast<T>(pw) * bin_size_w));
static_cast<int>(sycl::floor(static_cast<opmath_t>(static_cast<T>(pw) * bin_size_w)));
int hend =
static_cast<int>(std::ceil(static_cast<T>(ph + 1) * bin_size_h));
int wend =
Expand Down Expand Up @@ -204,11 +207,13 @@ struct PsRoiPoolBackwardKernel {
T bin_size_h =
static_cast<T>(roi_height) / static_cast<T>(pooled_height_);
T bin_size_w = static_cast<T>(roi_width) / static_cast<T>(pooled_width_);

using opmath_t = at::opmath_type<T>;

int hstart =
static_cast<int>(std::floor(static_cast<T>(ph) * bin_size_h));
static_cast<int>(sycl::floor(static_cast<opmath_t>(static_cast<T>(ph) * bin_size_h)));
int wstart =
static_cast<int>(std::floor(static_cast<T>(pw) * bin_size_w));
static_cast<int>(sycl::floor(static_cast<opmath_t>(static_cast<T>(pw) * bin_size_w)));
int hend =
static_cast<int>(std::ceil(static_cast<T>(ph + 1) * bin_size_h));
int wend =
Expand Down
7 changes: 5 additions & 2 deletions src/ATen/native/xpu/sycl/RoiPoolKernels.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,6 +17,7 @@
DISABLE_RETURN_TYPE_WARNING_BEGIN
// clang-format on
#include <ATen/ceil_div.h>
#include <ATen/OpMathType.h>
#include <ATen/native/xpu/sycl/Atomics.h>
#include <ATen/native/xpu/sycl/KernelUtils.h>
#include <comm/SYCLContext.h>
Expand Down Expand Up @@ -48,11 +49,13 @@ struct RoiPoolForwardKernel {
T bin_size_h =
static_cast<T>(roi_height) / static_cast<T>(pooled_height_);
T bin_size_w = static_cast<T>(roi_width) / static_cast<T>(pooled_width_);

using opmath_t = at::opmath_type<T>;

int hstart =
static_cast<int>(std::floor(static_cast<T>(ph) * bin_size_h));
static_cast<int>(sycl::floor(static_cast<opmath_t>(static_cast<T>(ph) * bin_size_h)));
int wstart =
static_cast<int>(std::floor(static_cast<T>(pw) * bin_size_w));
static_cast<int>(sycl::floor(static_cast<opmath_t>(static_cast<T>(pw) * bin_size_w)));
int hend =
static_cast<int>(std::ceil(static_cast<T>(ph + 1) * bin_size_h));
int wend =
Expand Down
4 changes: 3 additions & 1 deletion src/ATen/native/xpu/sycl/UnaryFractionKernels.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,6 +9,7 @@
*/

#include <ATen/Dispatch.h>
#include <ATen/OpMathType.h>
#include <ATen/native/TensorIterator.h>
#include <c10/core/ScalarType.h>

Expand Down Expand Up @@ -172,7 +173,8 @@ void round_decimals_kernel(TensorIteratorBase& iter, int64_t decimals) {
template <typename scalar_t>
struct FloorFunctor {
scalar_t operator()(scalar_t a) const {
return std::floor(a);
using opmath_t = at::opmath_type<scalar_t>;
return sycl::floor(static_cast<opmath_t>(a));
}
};

Expand Down
8 changes: 4 additions & 4 deletions src/ATen/native/xpu/sycl/UpSampleBicubic2dKernels.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -53,12 +53,12 @@ struct UpsampleBicubic2dKernelFunctor {
// Interpolation kernel
accscalar_t real_x = area_pixel_compute_source_index(
width_scale_, output_x, align_corners_, /*cubic=*/true);
int in_x = floorf(real_x);
int in_x = sycl::floor(real_x);
accscalar_t t_x = real_x - in_x;

accscalar_t real_y = area_pixel_compute_source_index(
height_scale_, output_y, align_corners_, /*cubic=*/true);
int in_y = floorf(real_y);
int in_y = sycl::floor(real_y);
accscalar_t t_y = real_y - in_y;
for (int n = 0; n < nbatch; n++) {
for (int c = 0; c < channels; c++) {
Expand Down Expand Up @@ -223,12 +223,12 @@ struct UpsampleBicubic2dBackwardKernelFunctor {

accscalar_t real_x = area_pixel_compute_source_index(
width_scale_, output_x, align_corners_, /*cubic=*/true);
int input_x = floorf(real_x);
int input_x = sycl::floor(real_x);
accscalar_t t_x = real_x - input_x;

accscalar_t real_y = area_pixel_compute_source_index(
height_scale_, output_y, align_corners_, /*cubic=*/true);
int input_y = floorf(real_y);
int input_y = sycl::floor(real_y);
accscalar_t t_y = real_y - input_y;

accscalar_t x_coeffs[4];
Expand Down
Loading