#include "InterpExecution.hpp" #include "core/MusaBackend.hpp" namespace MNN { namespace MUSA { template __global__ void InterpNearestKernel(const T* src, T* dst, int inBatch, int inChannels, int inHeight, int inWidth, int outHeight, int outWidth, float heightScale, float widthScale) { int index = blockIdx.x * blockDim.x + threadIdx.x; int totalSize = inBatch * inChannels * outHeight * outWidth; if (index < totalSize) { int tmp = index; int w = tmp % outWidth; tmp /= outWidth; int h = tmp % outHeight; tmp /= outHeight; int c = tmp % inChannels; int b = tmp / inChannels; int inX = __float2int_rd(w * widthScale); int inY = __float2int_rd(h * heightScale); inX = min(max(inX, 0), inWidth - 1); inY = min(max(inY, 0), inHeight - 1); int inIndex = ((b * inChannels + c) * inHeight + inY) * inWidth + inX; dst[index] = src[inIndex]; } } template __global__ void InterpBilinearKernel(const T* src, T* dst, int inBatch, int inChannels, int inHeight, int inWidth, int outHeight, int outWidth, float heightScale, float widthScale) { int index = blockIdx.x * blockDim.x + threadIdx.x; int totalSize = inBatch * inChannels * outHeight * outWidth; if (index < totalSize) { int tmp = index; int w = tmp % outWidth; tmp /= outWidth; int h = tmp % outHeight; tmp /= outHeight; int c = tmp % inChannels; int b = tmp / inChannels; float inX = (w + 0.5f) * widthScale - 0.5f; float inY = (h + 0.5f) * heightScale - 0.5f; int x0 = __float2int_rd(inX); int y0 = __float2int_rd(inY); int x1 = x0 + 1; int y1 = y0 + 1; x0 = max(0, x0); y0 = max(0, y0); x1 = min(x1, inWidth - 1); y1 = min(y1, inHeight - 1); float dx = inX - x0; float dy = inY - y0; int idx00 = ((b * inChannels + c) * inHeight + y0) * inWidth + x0; int idx01 = ((b * inChannels + c) * inHeight + y0) * inWidth + x1; int idx10 = ((b * inChannels + c) * inHeight + y1) * inWidth + x0; int idx11 = ((b * inChannels + c) * inHeight + y1) * inWidth + x1; float v00 = src[idx00]; float v01 = src[idx01]; float v10 = src[idx10]; float v11 = src[idx11]; float v0 = v00 * (1.0f - dx) + v01 * dx; float v1 = v10 * (1.0f - dx) + v11 * dx; dst[index] = v0 * (1.0f - dy) + v1 * dy; } } InterpExecution::InterpExecution(const std::vector& inputs, const MNN::Op* op, Backend* backend) : Execution(inputs, {}, backend) { mBackend = static_cast(backend); mOp = op->main_as_Interp(); } ErrorCode InterpExecution::onResize(const std::vector& inputs, const std::vector& outputs) { auto input = inputs[0]; auto output = outputs[0]; mInBatch = input->batch(); mInChannels = input->channel(); mInHeight = input->height(); mInWidth = input->width(); mOutHeight = output->height(); mOutWidth = output->width(); mHeightScale = static_cast(mInHeight) / mOutHeight; mWidthScale = static_cast(mInWidth) / mOutWidth; int threads = 256; int blocks = (mInBatch * mInChannels * mOutHeight * mOutWidth + threads - 1) / threads; mDim3Grid = {blocks, 1, 1}; mDim3Block = {threads, 1, 1}; return NO_ERROR; } ErrorCode InterpExecution::onExecute(const std::vector& inputs, const std::vector& outputs) { auto input = inputs[0]; auto output = outputs[0]; auto inputPtr = input->host(); auto outputPtr = output->host(); int totalSize = mInBatch * mInChannels * mOutHeight * mOutWidth; if (mOp->resizeType() == 1) { // NEAREST InterpNearestKernel<<>>( inputPtr, outputPtr, mInBatch, mInChannels, mInHeight, mInWidth, mOutHeight, mOutWidth, mHeightScale, mWidthScale ); } else { // BILINEAR InterpBilinearKernel<<>>( inputPtr, outputPtr, mInBatch, mInChannels, mInHeight, mInWidth, mOutHeight, mOutWidth, mHeightScale, mWidthScale ); } musaError_t err = musaGetLastError(); if (err != musaSuccess) { return COMPUTE_NO_SUPPORT; } return NO_ERROR; } class InterpCreator : public Creator { public: virtual Execution* onCreate(const std::vector& inputs, const MNN::Op* op, Backend* backend) const override { return new InterpExecution(inputs, op, backend); } }; MNNCreatorRegister gInterpRegistration(OpType_Interp); } // namespace MUSA } // namespace MNN