#include "EmbeddingExecution.hpp" #include "core/MusaBackend.hpp" namespace MNN { namespace MUSA { template __global__ void EmbeddingKernel(const T* embedding, const int* indices, T* output, int numIndices, int embeddingDim) { int index = blockIdx.x * blockDim.x + threadIdx.x; int totalSize = numIndices * embeddingDim; if (index < totalSize) { int dim = index % embeddingDim; int idx = index / embeddingDim; int embeddingIdx = indices[idx]; output[index] = embedding[embeddingIdx * embeddingDim + dim]; } } EmbeddingExecution::EmbeddingExecution(const std::vector& inputs, const MNN::Op* op, Backend* backend) : Execution(inputs, {}, backend) { mBackend = static_cast(backend); } ErrorCode EmbeddingExecution::onResize(const std::vector& inputs, const std::vector& outputs) { auto embedding = inputs[0]; auto indices = inputs[1]; auto output = outputs[0]; mNumIndices = 1; for (int i = 0; i < indices->dimensions(); i++) { mNumIndices *= indices->length(i); } mEmbeddingDim = embedding->length(1); int threads = 256; int totalSize = mNumIndices * mEmbeddingDim; int blocks = (totalSize + threads - 1) / threads; mDim3Grid = {blocks, 1, 1}; mDim3Block = {threads, 1, 1}; return NO_ERROR; } ErrorCode EmbeddingExecution::onExecute(const std::vector& inputs, const std::vector& outputs) { auto embedding = inputs[0]; auto indices = inputs[1]; auto output = outputs[0]; auto embeddingPtr = embedding->host(); auto indicesPtr = indices->host(); auto outputPtr = output->host(); EmbeddingKernel<<>>( embeddingPtr, indicesPtr, outputPtr, mNumIndices, mEmbeddingDim ); musaError_t err = musaGetLastError(); if (err != musaSuccess) { return COMPUTE_NO_SUPPORT; } return NO_ERROR; } class EmbeddingCreator : public Creator { public: virtual Execution* onCreate(const std::vector& inputs, const MNN::Op* op, Backend* backend) const override { return new EmbeddingExecution(inputs, op, backend); } }; MNNCreatorRegister gEmbeddingRegistration(OpType_Embedding); } // namespace MUSA } // namespace MNN