diff --git a/include/infinicore/ops.hpp b/include/infinicore/ops.hpp
index 5e93e1457..1d794ee26 100644
--- a/include/infinicore/ops.hpp
+++ b/include/infinicore/ops.hpp
@@ -90,6 +90,7 @@
 #include "ops/swap.hpp"
 #include "ops/swiglu.hpp"
 #include "ops/tanh.hpp"
+#include "ops/timestep_embedding.hpp"
 #include "ops/topksoftmax.hpp"
 #include "ops/vocab_parallel_embedding.hpp"
 #include "ops/zeros.hpp"
diff --git a/include/infinicore/ops/timestep_embedding.hpp b/include/infinicore/ops/timestep_embedding.hpp
new file mode 100644
index 000000000..debd6c518
--- /dev/null
+++ b/include/infinicore/ops/timestep_embedding.hpp
@@ -0,0 +1,19 @@
+#pragma once
+
+#include "../device.hpp"
+#include "../graph/graph.hpp"
+#include "common/op.hpp"
+
+namespace infinicore::op {
+
+INFINICORE_GRAPH_OP_CLASS(TimestepEmbedding, Tensor, const Tensor &, float);
+
+Tensor timestep_embedding(const Tensor &timestep,
+                          size_t embedding_dim = 256,
+                          float max_period = 10000.0f);
+
+void timestep_embedding_(Tensor output,
+                         const Tensor &timestep,
+                         float max_period = 10000.0f);
+
+} // namespace infinicore::op
diff --git a/include/infiniop.h b/include/infiniop.h
index 9f632e27f..d26e5767f 100644
--- a/include/infiniop.h
+++ b/include/infiniop.h
@@ -155,6 +155,7 @@
 #include "infiniop/ops/tan.h"
 #include "infiniop/ops/tanh.h"
 #include "infiniop/ops/tanhshrink.h"
+#include "infiniop/ops/timestep_embedding.h"
 #include "infiniop/ops/topk.h"
 #include "infiniop/ops/topkrouter.h"
 #include "infiniop/ops/topksoftmax.h"
diff --git a/include/infiniop/ops/timestep_embedding.h b/include/infiniop/ops/timestep_embedding.h
new file mode 100644
index 000000000..7403c3f65
--- /dev/null
+++ b/include/infiniop/ops/timestep_embedding.h
@@ -0,0 +1,24 @@
+#ifndef __INFINIOP_TIMESTEP_EMBEDDING_API_H__
+#define __INFINIOP_TIMESTEP_EMBEDDING_API_H__
+
+#include "../operator_descriptor.h"
+
+typedef struct InfiniopDescriptor *infiniopTimestepEmbeddingDescriptor_t;
+
+__INFINI_C __export infiniStatus_t infiniopCreateTimestepEmbeddingDescriptor(
+    infiniopHandle_t handle,
+    infiniopTimestepEmbeddingDescriptor_t *desc_ptr,
+    infiniopTensorDescriptor_t output_desc,
+    infiniopTensorDescriptor_t timestep_desc);
+
+__INFINI_C __export infiniStatus_t infiniopTimestepEmbedding(
+    infiniopTimestepEmbeddingDescriptor_t desc,
+    void *output,
+    const void *timestep,
+    float max_period,
+    void *stream);
+
+__INFINI_C __export infiniStatus_t infiniopDestroyTimestepEmbeddingDescriptor(
+    infiniopTimestepEmbeddingDescriptor_t desc);
+
+#endif
diff --git a/python/infinicore/__init__.py b/python/infinicore/__init__.py
index 612db614c..aa1f2d440 100644
--- a/python/infinicore/__init__.py
+++ b/python/infinicore/__init__.py
@@ -165,6 +165,7 @@
 from infinicore.ops.scatter import scatter
 from infinicore.ops.sinh import sinh
 from infinicore.ops.situ_and_mul import situ_and_mul
+from infinicore.ops.timestep_embedding import timestep_embedding
 from infinicore.ops.squeeze import squeeze
 from infinicore.ops.sum import sum
 from infinicore.ops.swap import swap
@@ -280,6 +281,7 @@
     "mul",
     "mul_scalar",
     "situ_and_mul",
+    "timestep_embedding",
     "diff",
     "digamma",
     "dist",
diff --git a/python/infinicore/ops/timestep_embedding.py b/python/infinicore/ops/timestep_embedding.py
new file mode 100644
index 000000000..146d12f6d
--- /dev/null
+++ b/python/infinicore/ops/timestep_embedding.py
@@ -0,0 +1,21 @@
+from infinicore.lib import _infinicore
+from infinicore.tensor import Tensor
+
+
+def timestep_embedding(timestep, embedding_dim=256, max_period=10000.0, *, out=None):
+    max_period = float(max_period)
+    if out is None:
+        return Tensor(
+            _infinicore.timestep_embedding(
+                timestep._underlying,
+                int(embedding_dim),
+                max_period,
+            )
+        )
+
+    _infinicore.timestep_embedding_(
+        out._underlying,
+        timestep._underlying,
+        max_period,
+    )
+    return out
diff --git a/src/infinicore/ops/timestep_embedding/timestep_embedding.cc b/src/infinicore/ops/timestep_embedding/timestep_embedding.cc
new file mode 100644
index 000000000..308e6a3a9
--- /dev/null
+++ b/src/infinicore/ops/timestep_embedding/timestep_embedding.cc
@@ -0,0 +1,57 @@
+#include "infinicore/ops/timestep_embedding.hpp"
+
+#include "../../utils.hpp"
+
+#include <stdexcept>
+
+namespace infinicore::op {
+
+INFINICORE_GRAPH_OP_DISPATCHERS_IMPL(TimestepEmbedding);
+
+TimestepEmbedding::TimestepEmbedding(Tensor output,
+                                     const Tensor &timestep,
+                                     float max_period) {
+    INFINICORE_ASSERT_TENSORS_SAME_DEVICE(output, timestep);
+    INFINICORE_GRAPH_OP_DISPATCH(
+        output->device().getType(), output, timestep, max_period);
+}
+
+void TimestepEmbedding::execute(Tensor output,
+                                const Tensor &timestep,
+                                float max_period) {
+    INFINICORE_GRAPH_OP_RECORD_OR_RUN(
+        TimestepEmbedding, output, timestep, max_period);
+}
+
+Tensor timestep_embedding(const Tensor &timestep,
+                          size_t embedding_dim,
+                          float max_period) {
+    auto output = Tensor::empty(
+        {timestep->numel(), embedding_dim},
+        DataType::F32,
+        timestep->device());
+    timestep_embedding_(output, timestep, max_period);
+    return output;
+}
+
+void timestep_embedding_(Tensor output,
+                         const Tensor &timestep,
+                         float max_period) {
+    if (timestep->ndim() != 1) {
+        throw std::runtime_error("timestep_embedding expects timestep shape [N]");
+    }
+    if (output->ndim() != 2 || output->size(0) != timestep->size(0)
+        || output->size(1) == 0 || output->size(1) % 2 != 0) {
+        throw std::runtime_error(
+            "timestep_embedding expects output shape [N, even embedding_dim]");
+    }
+    if (output->dtype() != DataType::F32) {
+        throw std::runtime_error("timestep_embedding output must be float32");
+    }
+    if (max_period <= 0.0f) {
+        throw std::runtime_error("timestep_embedding max_period must be positive");
+    }
+    TimestepEmbedding::execute(output, timestep, max_period);
+}
+
+} // namespace infinicore::op
diff --git a/src/infinicore/ops/timestep_embedding/timestep_embedding_infiniop.cc b/src/infinicore/ops/timestep_embedding/timestep_embedding_infiniop.cc
new file mode 100644
index 000000000..4c98e87b6
--- /dev/null
+++ b/src/infinicore/ops/timestep_embedding/timestep_embedding_infiniop.cc
@@ -0,0 +1,46 @@
+#include "../infiniop_impl.hpp"
+#include "infinicore/ops/timestep_embedding.hpp"
+
+namespace infinicore::op::timestep_embedding_impl::infiniop {
+
+INFINIOP_CACHABLE_DESCRIPTOR(Descriptor, TimestepEmbedding, 100);
+
+struct PlannedMeta {
+    std::shared_ptr<Descriptor> descriptor;
+    graph::GraphTensor output, timestep;
+    float max_period;
+};
+
+void *plan(Tensor output, const Tensor &timestep, float max_period) {
+    size_t seed = hash_combine(output, timestep);
+
+    INFINIOP_CACHABLE_DESCRIPTOR_GET_OR_CREATE(
+        Descriptor, descriptor, TimestepEmbedding,
+        seed, output->desc(), timestep->desc());
+
+    return new PlannedMeta{
+        descriptor,
+        graph::GraphTensor(output),
+        graph::GraphTensor(timestep),
+        max_period};
+}
+
+void run(void *planned_meta) {
+    auto planned = reinterpret_cast<PlannedMeta *>(planned_meta);
+    INFINICORE_CHECK_ERROR(infiniopTimestepEmbedding(
+        planned->descriptor->desc,
+        planned->output->data(),
+        planned->timestep->data(),
+        planned->max_period,
+        context::getStream()));
+}
+
+void cleanup(void **planned_meta_ptr) {
+    delete *reinterpret_cast<PlannedMeta **>(planned_meta_ptr);
+    *planned_meta_ptr = nullptr;
+}
+
+INFINICORE_GRAPH_OP_REGISTER_ALLDEVICE(
+    TimestepEmbedding, &plan, &run, &cleanup);
+
+} // namespace infinicore::op::timestep_embedding_impl::infiniop
diff --git a/src/infinicore/pybind11/ops.hpp b/src/infinicore/pybind11/ops.hpp
index e78adf275..198385e04 100644
--- a/src/infinicore/pybind11/ops.hpp
+++ b/src/infinicore/pybind11/ops.hpp
@@ -135,6 +135,7 @@
 #include "ops/tan.hpp"
 #include "ops/tanh.hpp"
 #include "ops/tanhshrink.hpp"
+#include "ops/timestep_embedding.hpp"
 #include "ops/topk.hpp"
 #include "ops/topksoftmax.hpp"
 #include "ops/triplet_margin_loss.hpp"
@@ -318,6 +319,7 @@ inline void bind(py::module &m) {
     bind_linear_w8a8i8(m);
     bind_silu_and_mul(m);
     bind_situ_and_mul(m);
+    bind_timestep_embedding(m);
     bind_sum(m);
     bind_var_mean(m);
     bind_var(m);
diff --git a/src/infinicore/pybind11/ops/timestep_embedding.hpp b/src/infinicore/pybind11/ops/timestep_embedding.hpp
new file mode 100644
index 000000000..da533d044
--- /dev/null
+++ b/src/infinicore/pybind11/ops/timestep_embedding.hpp
@@ -0,0 +1,27 @@
+#pragma once
+
+#include <pybind11/pybind11.h>
+
+#include "infinicore/ops/timestep_embedding.hpp"
+
+namespace py = pybind11;
+
+namespace infinicore::ops {
+
+inline void bind_timestep_embedding(py::module &m) {
+    m.def("timestep_embedding",
+          &op::timestep_embedding,
+          py::arg("timestep"),
+          py::arg("embedding_dim") = 256,
+          py::arg("max_period") = 10000.0f,
+          R"doc(Build sinusoidal timestep embeddings on the active device.)doc");
+
+    m.def("timestep_embedding_",
+          &op::timestep_embedding_,
+          py::arg("output"),
+          py::arg("timestep"),
+          py::arg("max_period") = 10000.0f,
+          R"doc(Build sinusoidal timestep embeddings into output.)doc");
+}
+
+} // namespace infinicore::ops
diff --git a/src/infiniop/ops/timestep_embedding/nvidia/timestep_embedding_nvidia.cu b/src/infiniop/ops/timestep_embedding/nvidia/timestep_embedding_nvidia.cu
new file mode 100644
index 000000000..9766aec58
--- /dev/null
+++ b/src/infiniop/ops/timestep_embedding/nvidia/timestep_embedding_nvidia.cu
@@ -0,0 +1,130 @@
+#include "timestep_embedding_nvidia.cuh"
+
+#include "../../../devices/nvidia/nvidia_common.cuh"
+#include "../../../devices/nvidia/nvidia_kernel_common.cuh"
+#include "../../../tensor.h"
+
+#include <cmath>
+#include <cuda_runtime.h>
+
+namespace {
+
+template <typename T>
+INFINIOP_CUDA_KERNEL timestepEmbeddingKernel(
+    float *__restrict__ output,
+    const T *__restrict__ timestep,
+    size_t num_timesteps,
+    size_t embedding_dim,
+    float log_max_period) {
+    const size_t index = blockIdx.x * blockDim.x + threadIdx.x;
+    const size_t numel = num_timesteps * embedding_dim;
+    if (index >= numel) {
+        return;
+    }
+
+    const size_t half_dim = embedding_dim / 2;
+    const size_t timestep_index = index / embedding_dim;
+    const size_t output_dim = index % embedding_dim;
+    const size_t frequency_index = output_dim % half_dim;
+    const float frequency = expf(
+        -log_max_period * static_cast<float>(frequency_index)
+        / static_cast<float>(half_dim));
+    const float angle = static_cast<float>(timestep[timestep_index]) * frequency;
+    output[index] = output_dim < half_dim ? cosf(angle) : sinf(angle);
+}
+
+} // namespace
+
+namespace op::timestep_embedding::nvidia {
+
+struct Descriptor::Opaque {
+    std::shared_ptr<device::nvidia::Handle::Internal> internal;
+};
+
+Descriptor::~Descriptor() {
+    delete _opaque;
+}
+
+infiniStatus_t Descriptor::create(
+    infiniopHandle_t handle,
+    Descriptor **desc_ptr,
+    infiniopTensorDescriptor_t output_desc,
+    infiniopTensorDescriptor_t timestep_desc) {
+    CHECK_OR_RETURN(timestep_desc->shape().size() == 1,
+                    INFINI_STATUS_BAD_TENSOR_SHAPE);
+    CHECK_OR_RETURN(output_desc->shape().size() == 2,
+                    INFINI_STATUS_BAD_TENSOR_SHAPE);
+    CHECK_OR_RETURN(output_desc->shape()[0] == timestep_desc->shape()[0]
+                        && output_desc->shape()[1] > 0
+                        && output_desc->shape()[1] % 2 == 0,
+                    INFINI_STATUS_BAD_TENSOR_SHAPE);
+    CHECK_OR_RETURN(output_desc->dtype() == INFINI_DTYPE_F32,
+                    INFINI_STATUS_BAD_TENSOR_DTYPE);
+    CHECK_DTYPE(timestep_desc->dtype(),
+                INFINI_DTYPE_F16, INFINI_DTYPE_BF16, INFINI_DTYPE_F32);
+    CHECK_OR_RETURN(output_desc->isContiguous() && timestep_desc->isContiguous(),
+                    INFINI_STATUS_BAD_TENSOR_STRIDES);
+
+    *desc_ptr = new Descriptor(
+        timestep_desc->shape()[0],
+        output_desc->shape()[1],
+        timestep_desc->dtype(),
+        new Opaque{reinterpret_cast<device::nvidia::Handle *>(handle)->internal()},
+        handle->device,
+        handle->device_id);
+    return INFINI_STATUS_SUCCESS;
+}
+
+infiniStatus_t Descriptor::calculate(
+    void *output,
+    const void *timestep,
+    float max_period,
+    void *stream) const {
+    if (max_period <= 0.0f) {
+        return INFINI_STATUS_BAD_PARAM;
+    }
+    const size_t numel = _num_timesteps * _embedding_dim;
+    if (numel == 0) {
+        return INFINI_STATUS_SUCCESS;
+    }
+
+    constexpr size_t block_size = 256;
+    const size_t grid_size = (numel + block_size - 1) / block_size;
+    const float log_max_period = std::log(max_period);
+    auto cuda_stream = reinterpret_cast<cudaStream_t>(stream);
+
+    switch (_input_dtype) {
+    case INFINI_DTYPE_F16:
+        timestepEmbeddingKernel<half><<<grid_size, block_size, 0, cuda_stream>>>(
+            reinterpret_cast<float *>(output),
+            reinterpret_cast<const half *>(timestep),
+            _num_timesteps,
+            _embedding_dim,
+            log_max_period);
+        break;
+    case INFINI_DTYPE_BF16:
+        timestepEmbeddingKernel<cuda_bfloat16><<<grid_size, block_size, 0, cuda_stream>>>(
+            reinterpret_cast<float *>(output),
+            reinterpret_cast<const cuda_bfloat16 *>(timestep),
+            _num_timesteps,
+            _embedding_dim,
+            log_max_period);
+        break;
+    case INFINI_DTYPE_F32:
+        timestepEmbeddingKernel<float><<<grid_size, block_size, 0, cuda_stream>>>(
+            reinterpret_cast<float *>(output),
+            reinterpret_cast<const float *>(timestep),
+            _num_timesteps,
+            _embedding_dim,
+            log_max_period);
+        break;
+    default:
+        return INFINI_STATUS_BAD_TENSOR_DTYPE;
+    }
+
+    return cudaGetLastError() == cudaSuccess
+             ? INFINI_STATUS_SUCCESS
+             : INFINI_STATUS_INTERNAL_ERROR;
+}
+
+} // namespace op::timestep_embedding::nvidia
diff --git a/src/infiniop/ops/timestep_embedding/nvidia/timestep_embedding_nvidia.cuh b/src/infiniop/ops/timestep_embedding/nvidia/timestep_embedding_nvidia.cuh
new file mode 100644
index 000000000..f84c8e275
--- /dev/null
+++ b/src/infiniop/ops/timestep_embedding/nvidia/timestep_embedding_nvidia.cuh
@@ -0,0 +1,8 @@
+#ifndef __TIMESTEP_EMBEDDING_NVIDIA_CUH__
+#define __TIMESTEP_EMBEDDING_NVIDIA_CUH__
+
+#include "../timestep_embedding.h"
+
+TIMESTEP_EMBEDDING_DESCRIPTOR(nvidia)
+
+#endif
diff --git a/src/infiniop/ops/timestep_embedding/operator.cc b/src/infiniop/ops/timestep_embedding/operator.cc
new file mode 100644
index 000000000..ba2b133a6
--- /dev/null
+++ b/src/infiniop/ops/timestep_embedding/operator.cc
@@ -0,0 +1,86 @@
+#include "../../handle.h"
+#include "../../operator.h"
+#include "infiniop/ops/timestep_embedding.h"
+
+#if defined(ENABLE_NVIDIA_API) || defined(ENABLE_HYGON_API)
+#include "nvidia/timestep_embedding_nvidia.cuh"
+#endif
+
+__INFINI_C infiniStatus_t infiniopCreateTimestepEmbeddingDescriptor(
+    infiniopHandle_t handle,
+    infiniopTimestepEmbeddingDescriptor_t *desc_ptr,
+    infiniopTensorDescriptor_t output_desc,
+    infiniopTensorDescriptor_t timestep_desc) {
+
+#define CREATE(CASE, NAMESPACE)                                                   \
+    case CASE:                                                                    \
+        return op::timestep_embedding::NAMESPACE::Descriptor::create(             \
+            handle,                                                               \
+            reinterpret_cast<op::timestep_embedding::NAMESPACE::Descriptor **>(   \
+                desc_ptr),                                                        \
+            output_desc,                                                          \
+            timestep_desc)
+
+    switch (handle->device) {
+#ifdef ENABLE_NVIDIA_API
+        CREATE(INFINI_DEVICE_NVIDIA, nvidia);
+#endif
+#ifdef ENABLE_HYGON_API
+        CREATE(INFINI_DEVICE_HYGON, nvidia);
+#endif
+    default:
+        return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
+    }
+
+#undef CREATE
+}
+
+__INFINI_C infiniStatus_t infiniopTimestepEmbedding(
+    infiniopTimestepEmbeddingDescriptor_t desc,
+    void *output,
+    const void *timestep,
+    float max_period,
+    void *stream) {
+
+#define CALCULATE(CASE, NAMESPACE)                                              \
+    case CASE:                                                                  \
+        return reinterpret_cast<                                               \
+                   const op::timestep_embedding::NAMESPACE::Descriptor *>(desc) \
+            ->calculate(output, timestep, max_period, stream)
+
+    switch (desc->device_type) {
+#ifdef ENABLE_NVIDIA_API
+        CALCULATE(INFINI_DEVICE_NVIDIA, nvidia);
+#endif
+#ifdef ENABLE_HYGON_API
+        CALCULATE(INFINI_DEVICE_HYGON, nvidia);
+#endif
+    default:
+        return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
+    }
+
+#undef CALCULATE
+}
+
+__INFINI_C infiniStatus_t infiniopDestroyTimestepEmbeddingDescriptor(
+    infiniopTimestepEmbeddingDescriptor_t desc) {
+
+#define DESTROY(CASE, NAMESPACE)                                               \
+    case CASE:                                                                 \
+        delete reinterpret_cast<                                               \
+            const op::timestep_embedding::NAMESPACE::Descriptor *>(desc);      \
+        return INFINI_STATUS_SUCCESS
+
+    switch (desc->device_type) {
+#ifdef ENABLE_NVIDIA_API
+        DESTROY(INFINI_DEVICE_NVIDIA, nvidia);
+#endif
+#ifdef ENABLE_HYGON_API
+        DESTROY(INFINI_DEVICE_HYGON, nvidia);
+#endif
+    default:
+        return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
+    }
+
+#undef DESTROY
+}
diff --git a/src/infiniop/ops/timestep_embedding/timestep_embedding.h b/src/infiniop/ops/timestep_embedding/timestep_embedding.h
new file mode 100644
index 000000000..8711dbd64
--- /dev/null
+++ b/src/infiniop/ops/timestep_embedding/timestep_embedding.h
@@ -0,0 +1,47 @@
+#ifndef __TIMESTEP_EMBEDDING_H__
+#define __TIMESTEP_EMBEDDING_H__
+
+#include "../../../utils.h"
+#include "../../operator.h"
+
+#define TIMESTEP_EMBEDDING_DESCRIPTOR(NAMESPACE)          \
+                                                          \
+    namespace op::timestep_embedding::NAMESPACE {         \
+    class Descriptor final : public InfiniopDescriptor {  \
+        struct Opaque;                                    \
+        Opaque *_opaque;                                  \
+        size_t _num_timesteps;                            \
+        size_t _embedding_dim;                            \
+        infiniDtype_t _input_dtype;                       \
+                                                          \
+        Descriptor(                                       \
+            size_t num_timesteps,                         \
+            size_t embedding_dim,                         \
+            infiniDtype_t input_dtype,                    \
+            Opaque *opaque,                               \
+            infiniDevice_t device_type,                   \
+            int device_id)                                \
+            : InfiniopDescriptor{device_type, device_id}, \
+              _opaque(opaque),                            \
+              _num_timesteps(num_timesteps),              \
+              _embedding_dim(embedding_dim),              \
+              _input_dtype(input_dtype) {}                \
+                                                          \
+    public:                                               \
+        ~Descriptor();                                    \
+                                                          \
+        static infiniStatus_t create(                     \
+            infiniopHandle_t handle,                      \
+            Descriptor **desc_ptr,                        \
+            infiniopTensorDescriptor_t output_desc,       \
+            infiniopTensorDescriptor_t timestep_desc);    \
+                                                          \
+        infiniStatus_t calculate(                         \
+            void *output,                                 \
+            const void *timestep,                         \
+            float max_period,                             \
+            void *stream) const;                          \
+    };                                                    \
+    }
+
+#endif
diff --git a/test/infinicore/ops/timestep_embedding.py b/test/infinicore/ops/timestep_embedding.py
new file mode 100644
index 000000000..a8dee3cce
--- /dev/null
+++ b/test/infinicore/ops/timestep_embedding.py
@@ -0,0 +1,81 @@
+import os
+import sys
+
+sys.path.insert(0, os.path.join(os.path.dirname(__file__), ".."))
+
+import torch
+from framework import BaseOperatorTest, GenericTestRunner, TensorSpec, TestCase
+
+import infinicore
+
+
+_TIMESTEP_SHAPES = [(1,), (4,), (17,)]
+_EMBEDDING_DIMS = [16, 256]
+_INPUT_DTYPES = [infinicore.float16, infinicore.bfloat16, infinicore.float32]
+
+
+class OpTest(BaseOperatorTest):
+    def __init__(self):
+        super().__init__("TimestepEmbedding")
+
+    def get_test_cases(self):
+        cases = []
+        for shape in _TIMESTEP_SHAPES:
+            for embedding_dim in _EMBEDDING_DIMS:
+                for dtype in _INPUT_DTYPES:
+                    cases.append(
+                        TestCase(
+                            inputs=[TensorSpec.from_tensor(shape, None, dtype)],
+                            kwargs={
+                                "embedding_dim": embedding_dim,
+                                "max_period": 10000.0,
+                            },
+                            output_spec=None,
+                            comparison_target=None,
+                            tolerance={"atol": 2e-5, "rtol": 2e-5},
+                            description="TimestepEmbedding - OUT_OF_PLACE",
+                        )
+                    )
+        return cases
+
+    def torch_operator(
+        self,
+        timestep,
+        embedding_dim=256,
+        max_period=10000.0,
+        out=None,
+    ):
+        half_dim = embedding_dim // 2
+        exponent = -torch.log(torch.tensor(max_period)) * torch.arange(
+            half_dim,
+            dtype=torch.float32,
+            device=timestep.device,
+        ) / half_dim
+        angles = timestep.float().unsqueeze(1) * exponent.exp().unsqueeze(0)
+        result = torch.cat((angles.cos(), angles.sin()), dim=1)
+        if out is not None:
+            out.copy_(result)
+            return out
+        return result
+
+    def infinicore_operator(
+        self,
+        timestep,
+        embedding_dim=256,
+        max_period=10000.0,
+        out=None,
+    ):
+        return infinicore.timestep_embedding(
+            timestep,
+            embedding_dim,
+            max_period,
+            out=out,
+        )
+
+
+def main():
+    GenericTestRunner(OpTest).run_and_exit()
+
+
+if __name__ == "__main__":
+    main()