diff --git a/CMakeLists.txt b/CMakeLists.txt
index 709bc30f2..451d5ae1a 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -16,6 +16,7 @@ set(CMAKE_CXX_EXTENSIONS OFF)
 set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
 
 add_compile_options(-Wunused -Wunused-function)
+add_compile_options(-D_USE_MATH_DEFINES)
 
 # ------------------------------------------------------------------------------
 # GoogleTest (submodule)
diff --git a/example/mnist/main.cc b/example/mnist/main.cc
index 7744e0947..8afa4f5a4 100644
--- a/example/mnist/main.cc
+++ b/example/mnist/main.cc
@@ -17,6 +17,9 @@
 #include "example/mnist/dataset.h"
 #include "example/mnist/net.h"
 
+#include "infini_train/include/dispatcher.h"
+#include "infini_train/include/device.h"
+
 DEFINE_string(dataset, "", "mnist dataset path");
 DEFINE_int32(bs, 64, "batch size");
 DEFINE_int32(num_epoch, 1, "num epochs");
@@ -36,6 +39,7 @@ constexpr char kDeviceCUDA[] = "cuda";
 DEFINE_validator(device,
                  [](const char *, const std::string &value) { return value == kDeviceCPU || value == kDeviceCUDA; });
 
+
 int main(int argc, char *argv[]) {
     gflags::ParseCommandLineFlags(&argc, &argv, true);
     google::InitGoogleLogging(argv[0]);
@@ -47,14 +51,15 @@ int main(int argc, char *argv[]) {
     auto test_dataset = std::make_shared<MNISTDataset>(FLAGS_dataset, false);
     DataLoader test_dataloader(test_dataset, FLAGS_bs);
 
-    auto network = MNIST();
+    auto network = std::make_shared<MNIST>();
     Device device = FLAGS_device == kDeviceCPU ? Device() : Device(Device::DeviceType::kCUDA, 0);
     Device cpu_device = Device();
-    network.To(device);
+    network->To(device);
+    network->To(cpu_device);
 
     auto loss_fn = nn::CrossEntropyLoss();
     loss_fn.To(device);
-    auto optimizer = optimizers::SGD(network.Parameters(), FLAGS_lr);
+    auto optimizer = optimizers::SGD(network->Parameters(), FLAGS_lr);
 
     for (int epoch = 0; epoch < FLAGS_num_epoch; ++epoch) {
         int train_idx = 0;
@@ -66,12 +71,33 @@ int main(int argc, char *argv[]) {
             auto new_image = std::make_shared<Tensor>(image->To(device));
             auto new_label = std::make_shared<Tensor>(label->To(device));
 
-            auto outputs = network.Forward({new_image});
+            auto outputs = network->Forward({new_image});
             optimizer.ZeroGrad();
 
             auto loss = loss_fn.Forward({outputs[0], new_label});
             loss[0]->Backward();
 
+            // 临时加：检查梯度
+            auto params = network->Parameters();
+            float total_grad = 0.0f;
+            for (auto &param : params) {
+                if (param->grad()) {
+                    float *grad_ptr = static_cast<float *>(param->grad()->DataPtr());
+                    for (int i = 0; i < param->grad()->NumElements(); ++i) {
+                        total_grad += std::abs(grad_ptr[i]);
+                    }
+                }
+            }
+            LOG(ERROR) << "total grad sum: " << total_grad;
+
+            // 临时加：检查参数更新前的值
+            float *first_param = static_cast<float *>(params[0]->DataPtr());
+            float before = first_param[0];
+
+            // 临时加：检查参数更新后的值
+            float after = first_param[0];
+            LOG(ERROR) << "param before: " << before << ", after: " << after;
+
             // Defer the loss D2H copy until after backward; reading it earlier would synchronize CUDA
             // between forward and backward.
             auto loss_cpu = loss[0]->To(cpu_device);
@@ -104,7 +130,7 @@ int main(int argc, char *argv[]) {
         auto new_label = std::make_shared<Tensor>(label->To(device));
 
         auto label_cpu = label->To(cpu_device);
-        auto outputs = network.Forward({new_image});
+        auto outputs = network->Forward({new_image});
         auto output_cpu = outputs[0]->To(cpu_device);
         auto loss = loss_fn.Forward({outputs[0], new_label});
         auto loss_cpu = loss[0]->To(cpu_device);
diff --git a/example/mnist/net.cc b/example/mnist/net.cc
index 501fee7ef..7a81b2246 100644
--- a/example/mnist/net.cc
+++ b/example/mnist/net.cc
@@ -8,6 +8,7 @@
 
 #include "infini_train/include/nn/modules/activations.h"
 #include "infini_train/include/nn/modules/container.h"
+#include "infini_train/include/nn/modules/conv2d.h"
 #include "infini_train/include/nn/modules/linear.h"
 #include "infini_train/include/nn/modules/module.h"
 #include "infini_train/include/tensor.h"
@@ -16,16 +17,21 @@ namespace nn = infini_train::nn;
 
 MNIST::MNIST() {
     std::vector<std::shared_ptr<nn::Module>> layers;
-    layers.push_back(std::make_shared<nn::Linear>(784, 30));
+    layers.push_back(std::make_shared<nn::Conv2d>(1, 16, 3));
+    layers.push_back(std::make_shared<nn::Sigmoid>());
+    layers.push_back(std::make_shared<nn::Conv2d>(16, 32, 3));
     layers.push_back(std::make_shared<nn::Sigmoid>());
     modules_["sequential"] = std::make_shared<nn::Sequential>(std::move(layers));
-    modules_["linear2"] = std::make_shared<nn::Linear>(30, 10);
+    modules_["linear2"] = std::make_shared<nn::Linear>(32 * 24 * 24, 10);
 }
 
 std::vector<std::shared_ptr<infini_train::Tensor>>
 MNIST::Forward(const std::vector<std::shared_ptr<infini_train::Tensor>> &x) {
     CHECK_EQ(x.size(), 1);
-    auto x1 = (*modules_["sequential"])(x);
-    auto x2 = (*modules_["linear2"])(x1);
-    return x2;
-}
+    auto reshaped = x[0]->View({x[0]->Dims()[0], 1, 28, 28});
+    std::vector<std::shared_ptr<infini_train::Tensor>> x_reshaped = {reshaped};
+    auto x1 = (*modules_["sequential"])(x_reshaped);
+    auto x2 = x1[0]->View({x1[0]->Dims()[0], 32 * 24 * 24})->Contiguous();
+    auto x3 = (*modules_["linear2"])({x2});
+    return x3;
+}
\ No newline at end of file
diff --git a/infini_train/include/autograd/conv2d.h b/infini_train/include/autograd/conv2d.h
new file mode 100644
index 000000000..e5e28a291
--- /dev/null
+++ b/infini_train/include/autograd/conv2d.h
@@ -0,0 +1,35 @@
+#pragma once
+
+#include <cstdint>
+#include <memory>
+#include <vector>
+
+#include "infini_train/include/autograd/function.h"
+
+namespace infini_train {
+class Tensor;
+}
+
+namespace infini_train::autograd {
+
+class Conv2d : public Function {
+public:
+    static constexpr char kType[] = "Conv2dFunction";
+
+    Conv2d(int64_t stride, int64_t padding) : Function(kType), stride_(stride), padding_(padding) {}
+
+    std::vector<std::shared_ptr<Tensor>> Forward(const std::vector<std::shared_ptr<Tensor>> &input_tensors) override;
+    void SetupContext(const std::vector<std::shared_ptr<Tensor>> &input_tensors,
+                      const std::vector<std::shared_ptr<Tensor>> &output_tensors) override;
+    std::vector<std::shared_ptr<Tensor>> Backward(const std::vector<std::shared_ptr<Tensor>> &grad_outputs) override;
+
+private:
+    int64_t stride_ = 1;
+    int64_t padding_ = 0;
+    bool bias_ = false;
+    int64_t in_channels_ = 0;
+    int64_t out_channels_ = 0;
+    int64_t kernel_size_ = 0;
+    std::vector<int64_t> input_dims_;
+};
+} // namespace infini_train::autograd
diff --git a/infini_train/include/nn/modules/conv2d.h b/infini_train/include/nn/modules/conv2d.h
new file mode 100644
index 000000000..bf246867f
--- /dev/null
+++ b/infini_train/include/nn/modules/conv2d.h
@@ -0,0 +1,36 @@
+#pragma once
+
+#include <cstdint>
+#include <memory>
+#include <vector>
+
+#include "infini_train/include/device.h"
+#include "infini_train/include/nn/modules/module.h"
+
+namespace infini_train {
+class Tensor;
+class Device;
+}
+
+namespace infini_train::nn {
+class Conv2d : public CloneableModule<Conv2d> {
+public:
+    static constexpr char kType[] = "Conv2d";
+    static constexpr char kParamWeightName[] = "weight";
+    static constexpr char kParamBiasName[] = "bias";
+
+    Conv2d(int64_t in_channels, int64_t out_channels, int64_t kernel_size,
+           int64_t stride = 1, int64_t padding = 0, bool bias = true, Device device = Device());
+
+    std::vector<std::shared_ptr<Tensor>> Forward(const std::vector<std::shared_ptr<Tensor>> &input_tensors) override;
+
+private:
+    void ResetParameters();
+    int64_t in_channels_ = 0;
+    int64_t out_channels_ = 0;
+    int64_t kernel_size_ = 0;
+    int64_t stride_ = 1;
+    int64_t padding_ = 0;
+    bool bias_ = true;
+};
+} // namespace infini_train::nn
\ No newline at end of file
diff --git a/infini_train/include/utils/string_utils.h b/infini_train/include/utils/string_utils.h
index 070bb992c..87e2ec644 100644
--- a/infini_train/include/utils/string_utils.h
+++ b/infini_train/include/utils/string_utils.h
@@ -2,6 +2,7 @@
 
 #include <string>
 #include <vector>
+#include <cstdint>
 
 namespace infini_train::utils {
 std::string DimsToString(const std::vector<int64_t> &dims);
diff --git a/infini_train/src/autograd/conv2d.cc b/infini_train/src/autograd/conv2d.cc
new file mode 100644
index 000000000..b4dca2375
--- /dev/null
+++ b/infini_train/src/autograd/conv2d.cc
@@ -0,0 +1,84 @@
+#include "infini_train/include/autograd/conv2d.h"
+
+#include "glog/logging.h"
+
+#include "infini_train/include/dispatcher.h"
+#include "infini_train/include/tensor.h"
+
+namespace infini_train::autograd {
+std::vector<std::shared_ptr<Tensor>> Conv2d::Forward(const std::vector<std::shared_ptr<Tensor>> &input_tensors) {
+    CHECK_GE(input_tensors.size(), 2);
+    const auto &input = input_tensors[0];
+    const auto &weight = input_tensors[1];
+    std::shared_ptr<Tensor> bias = nullptr;
+    if (input_tensors.size() == 3) {
+        bias = input_tensors[2];
+    }
+
+    auto device = input->GetDevice().type();
+    return {Dispatcher::Instance().Call<std::shared_ptr<Tensor>>(
+        {device, "Conv2dForward"}, input, weight, bias, stride_, padding_)};
+}
+
+void Conv2d::SetupContext(const std::vector<std::shared_ptr<Tensor>> &input_tensors,
+                          const std::vector<std::shared_ptr<Tensor>> &) {
+    const auto &input = input_tensors[0];
+    const auto &weight = input_tensors[1];
+    bool need_input = ctx_.needs_input_grad().size() > 0 && ctx_.needs_input_grad()[0];
+    bool need_weight = ctx_.needs_input_grad().size() > 1 && ctx_.needs_input_grad()[1];
+
+    ctx_.SaveForBackward({need_weight ? input : nullptr, need_input ? weight : nullptr});
+
+    input_dims_ = input->Dims();
+    LOG(ERROR) << "SetupContext input_dims_ size: " << input_dims_.size();
+    in_channels_ = weight->Dims()[1];
+    out_channels_ = weight->Dims()[0];
+    kernel_size_ = weight->Dims()[2];
+    bias_ = input_tensors.size() == 3;
+}
+
+std::vector<std::shared_ptr<Tensor>> Conv2d::Backward(const std::vector<std::shared_ptr<Tensor>> &grad_outputs) {
+    auto saved_tensors = ctx_.GetSavedTensors();
+    LOG(ERROR) << "Backward input_dims_ size: " << input_dims_.size();
+    CHECK_EQ(saved_tensors.size(), 2);
+    const auto &input = saved_tensors[0];
+    const auto &weight = saved_tensors[1];
+    LOG(ERROR) << "Backward weight ptr: " << weight.get();
+    if (weight) {
+        LOG(ERROR) << "Backward weight dims size: " << weight->Dims().size();
+    } else {
+        LOG(ERROR) << "Backward weight is null!";
+    }
+    const auto &grad_output = grad_outputs[0];
+
+    CHECK(!ctx_.needs_input_grad().empty()) << "needs_input_grad not populated in Conv2d::Backward";
+    bool need_grad_input = ctx_.needs_input_grad()[0];
+    bool need_grad_weight = ctx_.needs_input_grad().size() > 1 && ctx_.needs_input_grad()[1];
+    bool need_grad_bias = bias_ && ctx_.needs_input_grad().size() > 2 && ctx_.needs_input_grad()[2];
+
+    auto device = grad_output->GetDevice().type();
+
+    std::shared_ptr<Tensor> grad_input = nullptr;
+    std::shared_ptr<Tensor> grad_weight = nullptr;
+    std::shared_ptr<Tensor> grad_bias = nullptr;
+
+    if (need_grad_input) {
+        grad_input = Dispatcher::Instance().Call<std::shared_ptr<Tensor>>(
+            {device, "Conv2dBackwardInput"}, weight, grad_output, input_dims_, stride_, padding_);
+    }
+    if (need_grad_weight && weight) {
+    grad_weight = Dispatcher::Instance().Call<std::shared_ptr<Tensor>>(
+        {device, "Conv2dBackwardWeight"}, input, grad_output, weight->Dims(), stride_, padding_);
+    }
+    if (need_grad_bias) {
+        grad_bias = Dispatcher::Instance().Call<std::shared_ptr<Tensor>>(
+            {device, "Conv2dBackwardBias"}, grad_output, out_channels_);
+    }
+
+    if (bias_) {
+        return {grad_input, grad_weight, grad_bias};
+    } else {
+        return {grad_input, grad_weight};
+    }
+}
+} // namespace infini_train::autograd
\ No newline at end of file
diff --git a/infini_train/src/kernels/cpu/conv2d.cc b/infini_train/src/kernels/cpu/conv2d.cc
new file mode 100644
index 000000000..4163db2fd
--- /dev/null
+++ b/infini_train/src/kernels/cpu/conv2d.cc
@@ -0,0 +1,146 @@
+#include <cstdint>
+#include <memory>
+#include <vector>
+
+#include "glog/logging.h"
+
+#include "infini_train/include/dispatcher.h"
+#include "infini_train/include/tensor.h"
+
+namespace infini_train::kernels::cpu {
+
+std::shared_ptr<Tensor> Conv2dForward(const std::shared_ptr<Tensor> &input,
+                                      const std::shared_ptr<Tensor> &weight,
+                                      const std::shared_ptr<Tensor> &bias,
+                                      int64_t stride, int64_t padding) {
+    const auto &input_dims = input->Dims();
+    const auto &weight_dims = weight->Dims();
+    int64_t N = input_dims[0], C_in = input_dims[1], H = input_dims[2], W = input_dims[3];
+    int64_t C_out = weight_dims[0], K = weight_dims[2];
+    int64_t H_out = (H + 2 * padding - K) / stride + 1;
+    int64_t W_out = (W + 2 * padding - K) / stride + 1;
+
+    auto output = std::make_shared<Tensor>(std::vector<int64_t>{N, C_out, H_out, W_out}, DataType::kFLOAT32);
+
+    const float *input_ptr = static_cast<const float *>(input->DataPtr());
+    const float *weight_ptr = static_cast<const float *>(weight->DataPtr());
+    const float *bias_ptr = bias ? static_cast<const float *>(bias->DataPtr()) : nullptr;
+    float *output_ptr = static_cast<float *>(output->DataPtr());
+
+    for (int64_t n = 0; n < N; ++n)
+        for (int64_t co = 0; co < C_out; ++co)
+            for (int64_t oh = 0; oh < H_out; ++oh)
+                for (int64_t ow = 0; ow < W_out; ++ow) {
+                    float sum = bias_ptr ? bias_ptr[co] : 0.0f;
+                    for (int64_t ci = 0; ci < C_in; ++ci)
+                        for (int64_t kh = 0; kh < K; ++kh)
+                            for (int64_t kw = 0; kw < K; ++kw) {
+                                int64_t ih = oh * stride + kh - padding;
+                                int64_t iw = ow * stride + kw - padding;
+                                if (ih < 0 || ih >= H || iw < 0 || iw >= W) continue;
+                                sum += input_ptr[((n * C_in + ci) * H + ih) * W + iw]
+                                     * weight_ptr[((co * C_in + ci) * K + kh) * K + kw];
+                            }
+                    output_ptr[((n * C_out + co) * H_out + oh) * W_out + ow] = sum;
+                }
+    return output;
+}
+
+std::shared_ptr<Tensor> Conv2dBackwardInput(const std::shared_ptr<Tensor> &weight,
+                                            const std::shared_ptr<Tensor> &grad_output,
+                                            const std::vector<int64_t> &input_dims,
+                                            int64_t stride, int64_t padding) {
+    int64_t N = input_dims[0], C_in = input_dims[1], H = input_dims[2], W = input_dims[3];
+    const auto &weight_dims = weight->Dims();
+    int64_t C_out = weight_dims[0], K = weight_dims[2];
+    int64_t H_out = (H + 2 * padding - K) / stride + 1;
+    int64_t W_out = (W + 2 * padding - K) / stride + 1;
+
+    auto grad_input = std::make_shared<Tensor>(input_dims, DataType::kFLOAT32);
+    float *grad_input_ptr = static_cast<float *>(grad_input->DataPtr());
+    const float *grad_output_ptr = static_cast<const float *>(grad_output->DataPtr());
+    const float *weight_ptr = static_cast<const float *>(weight->DataPtr());
+
+    for (int64_t i = 0; i < N * C_in * H * W; ++i) grad_input_ptr[i] = 0.0f;
+
+    for (int64_t n = 0; n < N; ++n)
+        for (int64_t co = 0; co < C_out; ++co)
+            for (int64_t oh = 0; oh < H_out; ++oh)
+                for (int64_t ow = 0; ow < W_out; ++ow) {
+                    float g = grad_output_ptr[((n * C_out + co) * H_out + oh) * W_out + ow];
+                    for (int64_t ci = 0; ci < C_in; ++ci)
+                        for (int64_t kh = 0; kh < K; ++kh)
+                            for (int64_t kw = 0; kw < K; ++kw) {
+                                int64_t ih = oh * stride + kh - padding;
+                                int64_t iw = ow * stride + kw - padding;
+                                if (ih < 0 || ih >= H || iw < 0 || iw >= W) continue;
+                                grad_input_ptr[((n * C_in + ci) * H + ih) * W + iw]
+                                    += g * weight_ptr[((co * C_in + ci) * K + kh) * K + kw];
+                            }
+                }
+    return grad_input;
+}
+
+std::shared_ptr<Tensor> Conv2dBackwardWeight(const std::shared_ptr<Tensor> &input,
+                                             const std::shared_ptr<Tensor> &grad_output,
+                                             const std::vector<int64_t> &weight_dims,
+                                             int64_t stride, int64_t padding) {
+    int64_t N = input->Dims()[0], C_in = input->Dims()[1], H = input->Dims()[2], W = input->Dims()[3];
+    int64_t C_out = weight_dims[0], K = weight_dims[2];
+    int64_t H_out = (H + 2 * padding - K) / stride + 1;
+    int64_t W_out = (W + 2 * padding - K) / stride + 1;
+
+    auto grad_weight = std::make_shared<Tensor>(weight_dims, DataType::kFLOAT32);
+    float *grad_weight_ptr = static_cast<float *>(grad_weight->DataPtr());
+    const float *input_ptr = static_cast<const float *>(input->DataPtr());
+    const float *grad_output_ptr = static_cast<const float *>(grad_output->DataPtr());
+
+    for (int64_t i = 0; i < C_out * C_in * K * K; ++i) grad_weight_ptr[i] = 0.0f;
+
+    for (int64_t n = 0; n < N; ++n)
+        for (int64_t co = 0; co < C_out; ++co)
+            for (int64_t oh = 0; oh < H_out; ++oh)
+                for (int64_t ow = 0; ow < W_out; ++ow) {
+                    float g = grad_output_ptr[((n * C_out + co) * H_out + oh) * W_out + ow];
+                    for (int64_t ci = 0; ci < C_in; ++ci)
+                        for (int64_t kh = 0; kh < K; ++kh)
+                            for (int64_t kw = 0; kw < K; ++kw) {
+                                int64_t ih = oh * stride + kh - padding;
+                                int64_t iw = ow * stride + kw - padding;
+                                if (ih < 0 || ih >= H || iw < 0 || iw >= W) continue;
+                                grad_weight_ptr[((co * C_in + ci) * K + kh) * K + kw]
+                                    += g * input_ptr[((n * C_in + ci) * H + ih) * W + iw];
+                            }
+                }
+    return grad_weight;
+}
+
+std::shared_ptr<Tensor> Conv2dBackwardBias(const std::shared_ptr<Tensor> &grad_output, int64_t out_channels) {
+    const auto &grad_dims = grad_output->Dims();
+    int64_t N = grad_dims[0], H_out = grad_dims[2], W_out = grad_dims[3];
+    auto grad_bias = std::make_shared<Tensor>(std::vector<int64_t>{out_channels}, DataType::kFLOAT32);
+    float *grad_bias_ptr = static_cast<float *>(grad_bias->DataPtr());
+    const float *grad_output_ptr = static_cast<const float *>(grad_output->DataPtr());
+
+    for (int64_t co = 0; co < out_channels; ++co) {
+        float sum = 0.0f;
+        for (int64_t n = 0; n < N; ++n)
+            for (int64_t oh = 0; oh < H_out; ++oh)
+                for (int64_t ow = 0; ow < W_out; ++ow)
+                    sum += grad_output_ptr[((n * out_channels + co) * H_out + oh) * W_out + ow];
+        grad_bias_ptr[co] = sum;
+    }
+    return grad_bias;
+}
+
+} // namespace infini_train::kernels::cpu
+
+#define REGISTER_CPU_CONV2D_KERNEL(kernel_name) \
+    REGISTER_KERNEL(infini_train::Device::DeviceType::kCPU, kernel_name, infini_train::kernels::cpu::kernel_name)
+
+REGISTER_CPU_CONV2D_KERNEL(Conv2dForward)
+REGISTER_CPU_CONV2D_KERNEL(Conv2dBackwardInput)
+REGISTER_CPU_CONV2D_KERNEL(Conv2dBackwardWeight)
+REGISTER_CPU_CONV2D_KERNEL(Conv2dBackwardBias)
+
+#undef REGISTER_CPU_CONV2D_KERNEL
\ No newline at end of file
diff --git a/infini_train/src/kernels/cuda/conv2d.cu b/infini_train/src/kernels/cuda/conv2d.cu
new file mode 100644
index 000000000..2956a7f1a
--- /dev/null
+++ b/infini_train/src/kernels/cuda/conv2d.cu
@@ -0,0 +1,220 @@
+#include <cuda_runtime.h>
+
+#include <cstdint>
+#include <memory>
+#include <vector>
+
+#include "glog/logging.h"
+
+#include "infini_train/include/dispatcher.h"
+#include "infini_train/include/tensor.h"
+
+namespace infini_train::kernels::cuda {
+
+__global__ void Conv2dForwardKernel(const float *input, const float *weight, const float *bias,
+                                    float *output, int64_t N, int64_t C_in, int64_t H, int64_t W,
+                                    int64_t C_out, int64_t K, int64_t H_out, int64_t W_out,
+                                    int64_t stride, int64_t padding) {
+    int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;
+    int64_t total = N * C_out * H_out * W_out;
+    if (idx >= total) return;
+
+    int64_t ow = idx % W_out;
+    int64_t oh = (idx / W_out) % H_out;
+    int64_t co = (idx / (W_out * H_out)) % C_out;
+    int64_t n = idx / (W_out * H_out * C_out);
+
+    float sum = bias ? bias[co] : 0.0f;
+    for (int64_t ci = 0; ci < C_in; ++ci) {
+        for (int64_t kh = 0; kh < K; ++kh) {
+            for (int64_t kw = 0; kw < K; ++kw) {
+                int64_t ih = oh * stride + kh - padding;
+                int64_t iw = ow * stride + kw - padding;
+                if (ih < 0 || ih >= H || iw < 0 || iw >= W) continue;
+                sum += input[((n * C_in + ci) * H + ih) * W + iw]
+                     * weight[((co * C_in + ci) * K + kh) * K + kw];
+            }
+        }
+    }
+    output[idx] = sum;
+}
+
+std::shared_ptr<Tensor> Conv2dForward(const std::shared_ptr<Tensor> &input,
+                                      const std::shared_ptr<Tensor> &weight,
+                                      const std::shared_ptr<Tensor> &bias,
+                                      int64_t stride, int64_t padding) {
+    const auto &input_dims = input->Dims();
+    const auto &weight_dims = weight->Dims();
+    int64_t N = input_dims[0], C_in = input_dims[1], H = input_dims[2], W = input_dims[3];
+    int64_t C_out = weight_dims[0], K = weight_dims[2];
+    int64_t H_out = (H + 2 * padding - K) / stride + 1;
+    int64_t W_out = (W + 2 * padding - K) / stride + 1;
+
+    auto output = std::make_shared<Tensor>(std::vector<int64_t>{N, C_out, H_out, W_out}, DataType::kFLOAT32, input->GetDevice());
+
+    int64_t total = N * C_out * H_out * W_out;
+    int threads = 256;
+    int blocks = (total + threads - 1) / threads;
+
+    Conv2dForwardKernel<<<blocks, threads>>>(
+        static_cast<const float *>(input->DataPtr()),
+        static_cast<const float *>(weight->DataPtr()),
+        bias ? static_cast<const float *>(bias->DataPtr()) : nullptr,
+        static_cast<float *>(output->DataPtr()),
+        N, C_in, H, W, C_out, K, H_out, W_out, stride, padding);
+
+    return output;
+}
+
+__global__ void Conv2dBackwardInputKernel(const float *weight, const float *grad_output,
+                                          float *grad_input, int64_t N, int64_t C_in, int64_t H,
+                                          int64_t W, int64_t C_out, int64_t K, int64_t H_out,
+                                          int64_t W_out, int64_t stride, int64_t padding) {
+    int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;
+    int64_t total = N * C_in * H * W;
+    if (idx >= total) return;
+
+    int64_t iw = idx % W;
+    int64_t ih = (idx / W) % H;
+    int64_t ci = (idx / (W * H)) % C_in;
+    int64_t n = idx / (W * H * C_in);
+
+    float sum = 0.0f;
+    for (int64_t co = 0; co < C_out; ++co) {
+        for (int64_t kh = 0; kh < K; ++kh) {
+            for (int64_t kw = 0; kw < K; ++kw) {
+                int64_t oh = (ih + padding - kh);
+                int64_t ow = (iw + padding - kw);
+                if (oh % stride != 0 || ow % stride != 0) continue;
+                oh /= stride;
+                ow /= stride;
+                if (oh < 0 || oh >= H_out || ow < 0 || ow >= W_out) continue;
+                sum += grad_output[((n * C_out + co) * H_out + oh) * W_out + ow]
+                     * weight[((co * C_in + ci) * K + kh) * K + kw];
+            }
+        }
+    }
+    grad_input[idx] = sum;
+}
+
+std::shared_ptr<Tensor> Conv2dBackwardInput(const std::shared_ptr<Tensor> &weight,
+                                            const std::shared_ptr<Tensor> &grad_output,
+                                            const std::vector<int64_t> &input_dims,
+                                            int64_t stride, int64_t padding) {
+    int64_t N = input_dims[0], C_in = input_dims[1], H = input_dims[2], W = input_dims[3];
+    const auto &weight_dims = weight->Dims();
+    int64_t C_out = weight_dims[0], K = weight_dims[2];
+    int64_t H_out = (H + 2 * padding - K) / stride + 1;
+    int64_t W_out = (W + 2 * padding - K) / stride + 1;
+
+    auto grad_input = std::make_shared<Tensor>(input_dims, DataType::kFLOAT32, weight->GetDevice());
+
+    int64_t total = N * C_in * H * W;
+    int threads = 256;
+    int blocks = (total + threads - 1) / threads;
+
+    Conv2dBackwardInputKernel<<<blocks, threads>>>(
+        static_cast<const float *>(weight->DataPtr()),
+        static_cast<const float *>(grad_output->DataPtr()),
+        static_cast<float *>(grad_input->DataPtr()),
+        N, C_in, H, W, C_out, K, H_out, W_out, stride, padding);
+
+    return grad_input;
+}
+
+__global__ void Conv2dBackwardWeightKernel(const float *input, const float *grad_output,
+                                           float *grad_weight, int64_t N, int64_t C_in, int64_t H,
+                                           int64_t W, int64_t C_out, int64_t K, int64_t H_out,
+                                           int64_t W_out, int64_t stride, int64_t padding) {
+    int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;
+    int64_t total = C_out * C_in * K * K;
+    if (idx >= total) return;
+
+    int64_t kw = idx % K;
+    int64_t kh = (idx / K) % K;
+    int64_t ci = (idx / (K * K)) % C_in;
+    int64_t co = idx / (K * K * C_in);
+
+    float sum = 0.0f;
+    for (int64_t n = 0; n < N; ++n) {
+        for (int64_t oh = 0; oh < H_out; ++oh) {
+            for (int64_t ow = 0; ow < W_out; ++ow) {
+                int64_t ih = oh * stride + kh - padding;
+                int64_t iw = ow * stride + kw - padding;
+                if (ih < 0 || ih >= H || iw < 0 || iw >= W) continue;
+                sum += grad_output[((n * C_out + co) * H_out + oh) * W_out + ow]
+                     * input[((n * C_in + ci) * H + ih) * W + iw];
+            }
+        }
+    }
+    grad_weight[idx] = sum;
+}
+
+std::shared_ptr<Tensor> Conv2dBackwardWeight(const std::shared_ptr<Tensor> &input,
+                                             const std::shared_ptr<Tensor> &grad_output,
+                                             const std::vector<int64_t> &weight_dims,
+                                             int64_t stride, int64_t padding) {
+    int64_t N = input->Dims()[0], C_in = input->Dims()[1], H = input->Dims()[2], W = input->Dims()[3];
+    int64_t C_out = weight_dims[0], K = weight_dims[2];
+    int64_t H_out = (H + 2 * padding - K) / stride + 1;
+    int64_t W_out = (W + 2 * padding - K) / stride + 1;
+
+    auto grad_weight = std::make_shared<Tensor>(weight_dims, DataType::kFLOAT32, input->GetDevice());
+
+    int64_t total = C_out * C_in * K * K;
+    int threads = 256;
+    int blocks = (total + threads - 1) / threads;
+
+    Conv2dBackwardWeightKernel<<<blocks, threads>>>(
+        static_cast<const float *>(input->DataPtr()),
+        static_cast<const float *>(grad_output->DataPtr()),
+        static_cast<float *>(grad_weight->DataPtr()),
+        N, C_in, H, W, C_out, K, H_out, W_out, stride, padding);
+
+    return grad_weight;
+}
+
+__global__ void Conv2dBackwardBiasKernel(const float *grad_output, float *grad_bias,
+                                         int64_t N, int64_t C_out, int64_t H_out, int64_t W_out) {
+    int64_t co = blockIdx.x * blockDim.x + threadIdx.x;
+    if (co >= C_out) return;
+
+    float sum = 0.0f;
+    for (int64_t n = 0; n < N; ++n) {
+        for (int64_t oh = 0; oh < H_out; ++oh) {
+            for (int64_t ow = 0; ow < W_out; ++ow) {
+                sum += grad_output[((n * C_out + co) * H_out + oh) * W_out + ow];
+            }
+        }
+    }
+    grad_bias[co] = sum;
+}
+
+std::shared_ptr<Tensor> Conv2dBackwardBias(const std::shared_ptr<Tensor> &grad_output, int64_t out_channels) {
+    const auto &grad_dims = grad_output->Dims();
+    int64_t N = grad_dims[0], H_out = grad_dims[2], W_out = grad_dims[3];
+
+    auto grad_bias = std::make_shared<Tensor>(std::vector<int64_t>{out_channels}, DataType::kFLOAT32, grad_output->GetDevice());
+
+    int threads = 256;
+    int blocks = (out_channels + threads - 1) / threads;
+
+    Conv2dBackwardBiasKernel<<<blocks, threads>>>(
+        static_cast<const float *>(grad_output->DataPtr()),
+        static_cast<float *>(grad_bias->DataPtr()),
+        N, out_channels, H_out, W_out);
+
+    return grad_bias;
+}
+
+} // namespace infini_train::kernels::cuda
+
+#define REGISTER_CUDA_CONV2D_KERNEL(kernel_name) \
+    REGISTER_KERNEL(infini_train::Device::DeviceType::kCUDA, kernel_name, infini_train::kernels::cuda::kernel_name)
+
+REGISTER_CUDA_CONV2D_KERNEL(Conv2dForward)
+REGISTER_CUDA_CONV2D_KERNEL(Conv2dBackwardInput)
+REGISTER_CUDA_CONV2D_KERNEL(Conv2dBackwardWeight)
+REGISTER_CUDA_CONV2D_KERNEL(Conv2dBackwardBias)
+
+#undef REGISTER_CUDA_CONV2D_KERNEL
\ No newline at end of file
diff --git a/infini_train/src/nn/modules/conv2d.cc b/infini_train/src/nn/modules/conv2d.cc
new file mode 100644
index 000000000..f8228ed99
--- /dev/null
+++ b/infini_train/src/nn/modules/conv2d.cc
@@ -0,0 +1,47 @@
+#include "infini_train/include/nn/modules/conv2d.h"
+
+#include <cmath>
+#include <memory>
+#include <vector>
+
+#include "infini_train/include/autograd/conv2d.h"
+#include "infini_train/include/device.h"
+#include "infini_train/include/nn/init.h"
+#include "infini_train/include/tensor.h"
+
+namespace infini_train::nn {
+Conv2d::Conv2d(int64_t in_channels, int64_t out_channels, int64_t kernel_size,
+               int64_t stride, int64_t padding, bool bias, Device device)
+    : CloneableModule(kType), in_channels_(in_channels), out_channels_(out_channels),
+      kernel_size_(kernel_size), stride_(stride), padding_(padding), bias_(bias) {
+    device_ = device;
+
+    parameters_[kParamWeightName] = std::make_shared<Tensor>(
+        std::vector<int64_t>{out_channels, in_channels, kernel_size, kernel_size},
+        DataType::kFLOAT32, device_)->RequiresGrad();
+
+    if (bias) {
+        parameters_[kParamBiasName] = std::make_shared<Tensor>(
+            std::vector<int64_t>{out_channels}, DataType::kFLOAT32, device_)->RequiresGrad();
+    }
+    ResetParameters();
+}
+
+std::vector<std::shared_ptr<Tensor>> Conv2d::Forward(const std::vector<std::shared_ptr<Tensor>> &input_tensors) {
+    auto op = std::make_shared<autograd::Conv2d>(stride_, padding_);
+    if (bias_) {
+        return op->Apply({input_tensors[0], parameters_[kParamWeightName], parameters_[kParamBiasName]});
+    } else {
+        return op->Apply({input_tensors[0], parameters_[kParamWeightName]});
+    }
+}
+
+void Conv2d::ResetParameters() {
+    init::KaimingUniform(parameters_[kParamWeightName], sqrt(5.0f));
+    if (bias_) {
+        const auto [fan_in, _] = init::CalculateFanInAndFanOut(parameters_[kParamWeightName]);
+        const float bound = fan_in > 0 ? 1.0f / sqrt(fan_in) : 0.0f;
+        init::Uniform(parameters_[kParamBiasName], -bound, bound);
+    }
+}
+} // namespace infini_train::nn
\ No newline at end of file
diff --git a/infini_train/src/optimizer.cc b/infini_train/src/optimizer.cc
index 39b999c77..1a26b8748 100644
--- a/infini_train/src/optimizer.cc
+++ b/infini_train/src/optimizer.cc
@@ -54,15 +54,18 @@ SGD::SGD(const std::vector<std::shared_ptr<Tensor>> &params, float learning_rate
 SGD::SGD(const NamedParameterList &named_params, float learning_rate) : Optimizer(named_params, learning_rate) {}
 
 void SGD::Step() {
+    LOG(ERROR) << "SGD Step: params_.size() = " << params_.size();
     for (auto param : params_) {
         if (!param->grad()) {
-            LOG(INFO) << "Skipping param with null grad.";
+            LOG(ERROR) << "SGD Step: param has null grad, skipping";
             continue;
         }
         auto device = param->GetDevice();
         core::DeviceGuard guard(device);
         auto kernel = Dispatcher::Instance().GetKernel({device.type(), "AccumulateGrad"});
+        LOG(ERROR) << "SGD Step: calling kernel, device type = " << static_cast<int>(device.type());
         kernel.Call<void>(param->grad(), -learning_rate_, param);
+        LOG(ERROR) << "SGD Step: kernel call done";
     }
 }
 
diff --git a/infini_train/src/profiler.cc b/infini_train/src/profiler.cc
index d53cc351a..0d0e8ebed 100644
--- a/infini_train/src/profiler.cc
+++ b/infini_train/src/profiler.cc
@@ -8,6 +8,7 @@
 #include "glog/logging.h"
 
 #include "infini_train/include/core/runtime/device_guard.h"
+#include <iomanip>
 #include "infini_train/include/device.h"
 
 namespace infini_train {
diff --git a/infini_train/src/tensor.cc b/infini_train/src/tensor.cc
index 4e61e221b..32ff6a193 100644
--- a/infini_train/src/tensor.cc
+++ b/infini_train/src/tensor.cc
@@ -24,6 +24,7 @@
 #include "infini_train/include/device.h"
 #include "infini_train/include/dispatcher.h"
 #include "infini_train/include/nn/init.h"
+#include <iomanip>
 
 namespace infini_train {
 TensorBuffer::TensorBuffer(Device device, size_t size) : device_(device), size_(size) {
diff --git a/infini_train/src/utils/precision_check_config.cc b/infini_train/src/utils/precision_check_config.cc
index 7e1b7176c..1cae01b58 100644
--- a/infini_train/src/utils/precision_check_config.cc
+++ b/infini_train/src/utils/precision_check_config.cc
@@ -62,7 +62,11 @@ void PrecisionCheckEnv::Init(const PrecisionCheckConfig &config) {
         auto now = std::chrono::system_clock::now();
         auto time_t = std::chrono::system_clock::to_time_t(now);
         std::tm tm;
-        localtime_r(&time_t, &tm);
+        #ifdef _WIN32
+            localtime_s(&tm, &time_t);
+        #else
+            localtime_r(&time_t, &tm);
+        #endif
         char buf[32];
         std::strftime(buf, sizeof(buf), "%Y%m%d_%H%M%S", &tm);
 
diff --git a/infini_train/src/utils/string_utils.cc b/infini_train/src/utils/string_utils.cc
index e3d9c4ce7..be9ece3ec 100644
--- a/infini_train/src/utils/string_utils.cc
+++ b/infini_train/src/utils/string_utils.cc
@@ -1,6 +1,7 @@
 #include "infini_train/include/utils/string_utils.h"
 
 #include <sstream>
+#include <cstdint>
 
 namespace infini_train::utils {
 std::string DimsToString(const std::vector<int64_t> &dims) {
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
index 5b9838781..407ffce08 100644
--- a/tests/CMakeLists.txt
+++ b/tests/CMakeLists.txt
@@ -41,6 +41,8 @@ add_subdirectory(transformer)
 # Checkpoint tests
 add_subdirectory(checkpoint)
 
+# Conv2d tests
+add_subdirectory(conv2d)
 # Suite declaration above is device-independent. Instantiate the built-in
 # devices after every test subdirectory has registered its sources.
 _infini_train_instantiate_test_suites(
diff --git a/tests/conv2d/CMakeLists.txt b/tests/conv2d/CMakeLists.txt
new file mode 100644
index 000000000..c3c88f0c8
--- /dev/null
+++ b/tests/conv2d/CMakeLists.txt
@@ -0,0 +1,3 @@
+add_executable(conv2d_test conv2d_test.cc)
+target_link_libraries(conv2d_test PRIVATE infini_train gtest gtest_main)
+add_test(NAME conv2d_test COMMAND conv2d_test)
\ No newline at end of file
diff --git a/tests/conv2d/conv2d_test.cc b/tests/conv2d/conv2d_test.cc
new file mode 100644
index 000000000..22adc5755
--- /dev/null
+++ b/tests/conv2d/conv2d_test.cc
@@ -0,0 +1,27 @@
+#include <gtest/gtest.h>
+#include <memory>
+#include <vector>
+
+#include "infini_train/include/tensor.h"
+#include "infini_train/include/dispatcher.h"
+
+using namespace infini_train;
+
+TEST(Conv2dTest, ForwardBasic) {
+    auto input = std::make_shared<Tensor>(std::vector<int64_t>{1, 1, 5, 5}, DataType::kFLOAT32);
+    auto weight = std::make_shared<Tensor>(std::vector<int64_t>{1, 1, 3, 3}, DataType::kFLOAT32);
+
+    float *in_ptr = static_cast<float *>(input->DataPtr());
+    for (int i = 0; i < 25; ++i) in_ptr[i] = static_cast<float>(i);
+
+    float *w_ptr = static_cast<float *>(weight->DataPtr());
+    for (int i = 0; i < 9; ++i) w_ptr[i] = 1.0f;
+
+    auto output = Dispatcher::Instance().Call<std::shared_ptr<Tensor>>(
+        {Device::DeviceType::kCPU, "Conv2dForward"}, input, weight, nullptr, 1, 0);
+
+    float *out_ptr = static_cast<float *>(output->DataPtr());
+    EXPECT_NEAR(out_ptr[0], 54.0f, 1e-5);
+    EXPECT_NEAR(out_ptr[1], 63.0f, 1e-5);
+    EXPECT_NEAR(out_ptr[4], 108.0f, 1e-5);
+}
\ No newline at end of file