Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
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: 1 addition & 1 deletion include/core/detail/casting.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -61,7 +61,7 @@ __device__ __host__ T ScalarSaturateCast(U v) {
// Any integral signed -> Any integral unsigned, big -> small
return v <= static_cast<U>(std::numeric_limits<T>::min())
? std::numeric_limits<T>::min()
: (v >= static_cast<U>(std::numeric_limits<T>::max()) ? std::numeric_limits<T>::max
: (v >= static_cast<U>(std::numeric_limits<T>::max()) ? std::numeric_limits<T>::max()
: static_cast<T>(v));
} else {
// All other cases fall into this
Expand Down
6 changes: 6 additions & 0 deletions include/core/image_format.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -114,6 +114,12 @@ constexpr ImageFormat FMT_RGB32(eDataType::DATA_TYPE_U32, 3, eSwizzle::XYZW);
// Single plane with interleaved RGBA 32-bit channel.
constexpr ImageFormat FMT_RGBA32(eDataType::DATA_TYPE_U32, 4, eSwizzle::XYZW);

// Single plane with interleaved RGB signed 32-bit channel.
constexpr ImageFormat FMT_RGBs32(eDataType::DATA_TYPE_S32, 3, eSwizzle::XYZW);

// Single plane with interleaved RGBA signed 32-bit channel.
constexpr ImageFormat FMT_RGBAs32(eDataType::DATA_TYPE_S32, 4, eSwizzle::XYZW);

// Single plane with interleaved RGB float32 channel.
constexpr ImageFormat FMT_RGBf32(eDataType::DATA_TYPE_F32, 3, eSwizzle::XYZW);

Expand Down
47 changes: 47 additions & 0 deletions include/kernels/device/convert_to_device.hpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,47 @@
/**
Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.

Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:

The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.

THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/

#pragma once

#include <hip/hip_runtime.h>
#include "core/detail/casting.hpp"
#include <core/detail/type_traits.hpp>
#include <core/wrappers/image_wrapper.hpp>

namespace Kernels {
namespace Device {
template <typename SrcWrapper, typename DstWrapper, typename DT_AB>
__global__ void convert_to(SrcWrapper input, DstWrapper output, DT_AB alpha, DT_AB beta) {
using namespace roccv::detail; // For RangeCast, NumElements, etc.
using dst_type = typename DstWrapper::ValueType;

const int x = threadIdx.x + blockIdx.x * blockDim.x;
const int y = threadIdx.y + blockIdx.y * blockDim.y;
const int batch = blockIdx.z;

if (x >= output.width() || y >= output.height() || batch >= output.batches()) return;

output.at(batch, y, x, 0) = SaturateCast<dst_type>(alpha * (input.at(batch, y, x, 0)) + beta);

}
} // namespace Device
} // namespace Kernels
47 changes: 47 additions & 0 deletions include/kernels/host/convert_to_host.hpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,47 @@
/**
Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.

Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:

The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.

THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/

#pragma once

#include <hip/hip_runtime.h>
#include "core/detail/casting.hpp"
#include <core/detail/type_traits.hpp>
#include <core/wrappers/image_wrapper.hpp>

namespace Kernels {
namespace Host {
template <typename SrcWrapper, typename DstWrapper, typename DT_AB>
void convert_to(SrcWrapper input, DstWrapper output, DT_AB alpha, DT_AB beta) {
using namespace roccv::detail; // For RangeCast, NumElements, etc.
using dst_type = typename DstWrapper::ValueType;

for (int batch = 0; batch < output.batches(); batch++) {
#pragma omp parallel for
for (int y = 0; y < output.height(); y++) {
for (int x = 0; x < output.width(); x++) {
output.at(batch, y, x, 0) = SaturateCast<dst_type>(alpha * (input.at(batch, y, x, 0)) + beta);
}
}
}
}
} // namespace Host
} // namespace Kernels
74 changes: 74 additions & 0 deletions include/op_convert_to.hpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,74 @@
/**
Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.

Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:

The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.

THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#pragma once
#include <hip/hip_runtime.h>
#include <operator_types.h>

#include <i_operator.hpp>

#include "core/tensor.hpp"

namespace roccv {
/**
* @brief Class for managing the Warp Perspective operator.
*
*/
class ConvertTo final : public IOperator {
public:
/**
* @brief Construct a new Op Convert To object. The object can be used
* to convert the datatype of an image.
* outputs(x,y) = SaturateCast<out_Datatype>(alpha * inputs(x,y) + beta)
*
* Limitations:
*
* Input:
* Supported TensorLayout(s): [HWC, NHWC]
* Channels: [1, 2, 3, 4]
* Supported DataType(s): [U8, S8, U16, S16, S32, F32, F64]
*
* Output:
* Supported TensorLayout(s): [HWC, NHWC]
* Channels: [1, 2, 3, 4]
* Supported DataType(s): [U8, S8, U16, S16, S32, F32, F64]
*
* Input/Output dependency:
*
* Property | Input == Output
* -------------- | -------------
* TensorLayout | Yes
* DataType | No
* Channels | Yes
* Width | Yes
* Height | Yes
*
* @param[in] stream The HIP stream to run this operator on.
* @param[in] input Input tensor with image data.
* @param[out] output Output tensor for storing modified image data.
* @param[in] alpha Scalar for output data. (Default: 1.0)
* @param[in] beta Offset for the data. (Default: 0.0)
* @param[in] device The device to run this operator on. (Default: GPU)
*/
void operator()(hipStream_t stream, const roccv::Tensor &input, const roccv::Tensor &output,
const double alpha = 1.0, const double beta = 0.0, const eDeviceType device = eDeviceType::GPU) const;
};
} // namespace roccv
3 changes: 2 additions & 1 deletion include/roccv_operators.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -39,4 +39,5 @@ THE SOFTWARE.
#include "op_rotate.hpp"
#include "op_thresholding.hpp"
#include "op_warp_affine.hpp"
#include "op_warp_perspective.hpp"
#include "op_warp_perspective.hpp"
#include "op_convert_to.hpp"
40 changes: 40 additions & 0 deletions python/include/operators/py_op_convert_to.hpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,40 @@
/**
Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.

Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:

The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.

THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/

#pragma once

#include <operator_types.h>
#include <pybind11/pybind11.h>

#include "py_stream.hpp"
#include "py_tensor.hpp"

namespace py = pybind11;

class PyOpConvertTo {
public:
static void Export(py::module& m);
static PyTensor Execute(PyTensor& input, eDataType dtype, double alpha, double beta,
std::optional<std::reference_wrapper<PyStream>> stream, eDeviceType device);
static void ExecuteInto(PyTensor& output, PyTensor& input, double alpha, double beta,
std::optional<std::reference_wrapper<PyStream>> stream, eDeviceType device);
};
2 changes: 2 additions & 0 deletions python/src/main.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -41,6 +41,7 @@ THE SOFTWARE.
#include "operators/py_op_thresholding.hpp"
#include "operators/py_op_warp_affine.hpp"
#include "operators/py_op_warp_perspective.hpp"
#include "operators/py_op_convert_to.hpp"
#include "py_enums.hpp"
#include "py_exception.hpp"
#include "py_stream.hpp"
Expand Down Expand Up @@ -76,4 +77,5 @@ PYBIND11_MODULE(rocpycv, m) {
PyOpCopyMakeBorder::Export(m);
PyOpCenterCrop::Export(m);
PyOpHistogram::Export(m);
PyOpConvertTo::Export(m);
}
85 changes: 85 additions & 0 deletions python/src/operators/py_op_convert_to.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,85 @@
/**
Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.

Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:

The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.

THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/

#include "operators/py_op_convert_to.hpp"

#include <op_convert_to.hpp>

PyTensor PyOpConvertTo::Execute(PyTensor& input, eDataType dtype, double alpha, double beta,
std::optional<std::reference_wrapper<PyStream>> stream, eDeviceType device) {
hipStream_t hipStream = stream.has_value() ? stream.value().get().getStream() : nullptr;
auto inputTensor = input.getTensor();
auto outputTensor = std::make_shared<roccv::Tensor>(inputTensor->shape(), roccv::DataType(dtype), device);

roccv::ConvertTo op;
op(hipStream, *inputTensor, *outputTensor, alpha, beta, device);
return PyTensor(outputTensor);
}

void PyOpConvertTo::ExecuteInto(PyTensor& output, PyTensor& input, double alpha, double beta,
std::optional<std::reference_wrapper<PyStream>> stream, eDeviceType device) {
hipStream_t hipStream = stream.has_value() ? stream.value().get().getStream() : nullptr;
roccv::ConvertTo op;
op(hipStream, *input.getTensor(), *output.getTensor(), alpha, beta, device);
}

void PyOpConvertTo::Export(py::module& m) {
using namespace py::literals;
m.def("convert_to", &PyOpConvertTo::Execute, "src"_a, "dtype"_a, "alpha"_a = 1.0, "beta"_a = 0.0,
"stream"_a = nullptr, "device"_a = eDeviceType::GPU, R"pbdoc(

Executes the Convert To operation on the given HIP stream.

See also:
Refer to the rocCV C++ API reference for more information on this operation.

Args:
src (rocpycv.Tensor): Input tensor containing one or more images.
dtype (eDataType): Datatype of the output tensor.
alpha (double, optional): Scalar for output data. Defaults to 1.0.
beta (double, optional): Offset for the data. Defaults to 0.0.
stream (rocpycv.Stream, optional): HIP stream to run this operation on.
device (rocpycv.Device, optional): The device to run this operation on. Defaults to GPU.

Returns:
rocpycv.Tensor: The output tensor.
)pbdoc");
m.def("convert_to_into", &PyOpConvertTo::ExecuteInto, "dst"_a, "src"_a, "alpha"_a = 1.0, "beta"_a = 0.0,
"stream"_a = nullptr, "device"_a = eDeviceType::GPU, R"pbdoc(

Executes the Convert To operation on the given HIP stream.

See also:
Refer to the rocCV C++ API reference for more information on this operation.

Args:
dst (rocpycv.Tensor): The output tensor with gamma correction applied.
src (rocpycv.Tensor): Input tensor containing one or more images.
alpha (double, optional): Scalar for output data. Defaults to 1.0.
beta (double, optional): Offset for the data. Defaults to 0.0.
stream (rocpycv.Stream, optional): HIP stream to run this operation on.
device (rocpycv.Device, optional): The device to run this operation on. Defaults to GPU.

Returns:
None
)pbdoc");
}
Loading