saymrwulf/onnxruntime
1
0
Fork 0
mirror of https://github.com/saymrwulf/onnxruntime.git synced 2026-07-30 20:18:08 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

update roialign cuda impl to onnx opset16 (#12036)

Browse source 
* roialign opset16

* fix

* fix
This commit is contained in:
George Wu 2022-06-29 17:32:59 -07:00 committed by GitHub
parent c8cd36da01
commit 102d01b206
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
3 changed files with 8 additions and 3 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

1
onnxruntime/core/providers/cuda/object_detection/roialign.cc
View file

@ -59,6 +59,7 @@ Status RoiAlign<T>::ComputeInternal(OpKernelContext* context) const {
num_roi_cols,
reinterpret_cast<typename ToCudaType<T>::MappedType*>(Y.template MutableData<T>()),
this->mode_ == RoiAlignMode::avg,
this->half_pixel_,
batch_indices_ptr->template Data<int64_t>());
}

9
onnxruntime/core/providers/cuda/object_detection/roialign_impl.cu
View file

@ -94,6 +94,7 @@ __global__ void RoIAlignForward(
int64_t roi_cols,
T* top_data,
const bool is_mode_avg,
const bool half_pixel,
const int64_t* batch_indices_ptr) {
for (size_t index = blockIdx.x * blockDim.x + threadIdx.x; index < nthreads; index += blockDim.x * gridDim.x) {
// (n, c, ph, pw) is an element in the pooled output
@ -106,9 +107,8 @@ __global__ void RoIAlignForward(
const T* offset_bottom_rois = bottom_rois + n * roi_cols;
const auto roi_batch_ind = batch_indices_ptr[n];
bool continuous_coordinate = false;
// Do not using rounding; this implementation detail is critical
T roi_offset = continuous_coordinate ? T(0.5) : T(0);
T roi_offset = half_pixel ? T(0.5) : T(0);
T roi_start_w = offset_bottom_rois[0] * spatial_scale - roi_offset;
T roi_start_h = offset_bottom_rois[1] * spatial_scale - roi_offset;
T roi_end_w = offset_bottom_rois[2] * spatial_scale - roi_offset;
@ -116,7 +116,7 @@ __global__ void RoIAlignForward(
T roi_width = roi_end_w - roi_start_w;
T roi_height = roi_end_h - roi_start_h;
if (!continuous_coordinate) { // backward compatiblity
if (!half_pixel) { // backward compatiblity
// Force malformed ROIs to be 1x1
roi_width = max(roi_width, (T)1.);
roi_height = max(roi_height, (T)1.);
@ -188,6 +188,7 @@ void RoiAlignImpl(
int64_t roi_cols,
T* top_data,
const bool is_mode_avg,
const bool half_pixel,
const int64_t* batch_indices_ptr) {
int blocksPerGrid = (int)(ceil(static_cast<float>(nthreads) / GridDim::maxThreadsPerBlock));
RoIAlignForward<T><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0, stream>>>(
@ -204,6 +205,7 @@ void RoiAlignImpl(
roi_cols,
top_data,
is_mode_avg,
half_pixel,
batch_indices_ptr);
}
@ -223,6 +225,7 @@ void RoiAlignImpl(
int64_t roi_cols, \
T* top_data, \
const bool is_mode_avg, \
const bool half_pixel, \
const int64_t* batch_indices_ptr);
SPECIALIZED_IMPL(float)

1
onnxruntime/core/providers/cuda/object_detection/roialign_impl.h
View file

@ -24,6 +24,7 @@ void RoiAlignImpl(
int64_t roi_cols,
T* top_data,
const bool is_mode_avg,
const bool half_pixel,
const int64_t* batch_indices_ptr);
} // namespace cuda