diff --git a/example/gpt2/main.cc b/example/gpt2/main.cc
index 60c0c908..3fdc9cf8 100644
--- a/example/gpt2/main.cc
+++ b/example/gpt2/main.cc
@@ -479,10 +479,10 @@ void Train(const nn::parallel::Rank &rank) {
 
                 LOG(INFO) << "Rank " << rank.GlobalRank() << ": finish loss forward";
 
-                auto loss_cpu = loss->To(Device());
-                lossf += static_cast<const float *>(loss_cpu.DataPtr())[0];
                 LOG(INFO) << "Rank " << rank.GlobalRank() << ": start backward";
                 loss->Backward();
+                auto loss_cpu = loss->To(Device());
+                lossf += static_cast<const float *>(loss_cpu.DataPtr())[0];
                 LOG(INFO) << "Rank " << rank.GlobalRank() << ": finish backward";
             }
 
diff --git a/example/llama3/main.cc b/example/llama3/main.cc
index 302e0808..76801f28 100644
--- a/example/llama3/main.cc
+++ b/example/llama3/main.cc
@@ -456,10 +456,10 @@ void Train(const nn::parallel::Rank &rank) {
 
                 LOG(INFO) << "Rank " << rank.GlobalRank() << ": finish loss forward";
 
-                auto loss_cpu = loss->To(Device());
-                lossf += static_cast<const float *>(loss_cpu.DataPtr())[0];
                 LOG(INFO) << "Rank " << rank.GlobalRank() << ": start backward";
                 loss->Backward();
+                auto loss_cpu = loss->To(Device());
+                lossf += static_cast<const float *>(loss_cpu.DataPtr())[0];
                 LOG(INFO) << "Rank " << rank.GlobalRank() << ": finish backward";
             }
 
diff --git a/example/mnist/main.cc b/example/mnist/main.cc
index e62257d7..726a873d 100644
--- a/example/mnist/main.cc
+++ b/example/mnist/main.cc
@@ -70,6 +70,8 @@ int main(int argc, char *argv[]) {
             optimizer.ZeroGrad();
 
             auto loss = loss_fn.Forward({outputs[0], new_label});
+            loss[0]->Backward();
+
             auto loss_cpu = loss[0]->To(cpu_device);
             float current_loss = static_cast<float *>(loss_cpu.DataPtr())[0];
             total_loss += current_loss;
@@ -79,7 +81,6 @@ int main(int argc, char *argv[]) {
                            << " loss: " << current_loss;
             }
 
-            loss[0]->Backward();
             optimizer.Step();
             train_idx += 1;
         }
diff --git a/infini_train/src/kernels/cuda/concat.cu b/infini_train/src/kernels/cuda/concat.cu
index a7fa7490..16b5da65 100644
--- a/infini_train/src/kernels/cuda/concat.cu
+++ b/infini_train/src/kernels/cuda/concat.cu
@@ -186,7 +186,7 @@ std::vector<std::shared_ptr<Tensor>> ConcatBackward(const std::shared_ptr<Tensor
     grads.reserve(input_dims_list.size());
     for (const auto &dvec : input_dims_list) {
         auto t = std::make_shared<Tensor>(dvec, dtype, device);
-        t->Fill(0.0);
+        // ConcatBackwardKernel maps every grad_output element to exactly one grad tensor element; no Fill is needed.
         grads.push_back(t);
     }
 
diff --git a/infini_train/src/kernels/cuda/cross_entropy.cu b/infini_train/src/kernels/cuda/cross_entropy.cu
index 8ba2785b..96f987fe 100644
--- a/infini_train/src/kernels/cuda/cross_entropy.cu
+++ b/infini_train/src/kernels/cuda/cross_entropy.cu
@@ -204,7 +204,7 @@ std::shared_ptr<Tensor> CrossEntropyBackward(const std::shared_ptr<Tensor> &inpu
                                  DataTypeList<INFINI_ALL_FLOATING_TYPES>>(
         {target->Dtype(), input_casted->Dtype()},
         [=]<typename Ttarget, typename Tinput>() {
-            grad_input->Fill(0.0);
+            // One sample block writes all of its num_classes gradient elements; no Fill is needed.
             const Tinput *output_grad_ptr = static_cast<const Tinput *>(grad_output->DataPtr());
             const Ttarget *target_ptr = static_cast<const Ttarget *>(target->DataPtr());
             const Tinput *input_ptr = static_cast<const Tinput *>(input_casted->DataPtr());
diff --git a/infini_train/src/kernels/cuda/layernorm.cu b/infini_train/src/kernels/cuda/layernorm.cu
index 12f3c08d..05d43c37 100644
--- a/infini_train/src/kernels/cuda/layernorm.cu
+++ b/infini_train/src/kernels/cuda/layernorm.cu
@@ -89,8 +89,7 @@ LayerNormForward(const std::shared_ptr<Tensor> &input, const std::shared_ptr<Ten
     core::cuda::DispatchCudaFunc<INFINI_ALL_FLOATING_TYPES>(
         dtype,
         [=]<typename T>() {
-            mean->Fill(0.0);
-            rstd->Fill(0.0);
+            // Each token block writes its mean and rstd exactly once; no Fill is needed.
             LayerNormForwardKernel<BLOCK_SIZE><<<num_blocks, threads_per_block, 0, cuda_stream>>>(
                 static_cast<const T *>(input->DataPtr()), static_cast<const T *>(weight->DataPtr()),
                 static_cast<const T *>(bias->DataPtr()), static_cast<float *>(mean->DataPtr()),
@@ -183,7 +182,7 @@ LayerNormBackward(const std::shared_ptr<Tensor> &input, const std::shared_ptr<Te
     core::cuda::DispatchCudaFunc<INFINI_ALL_FLOATING_TYPES>(
         dtype,
         [=]<typename T>() {
-            grad_input->Fill(0.0);
+            // Each token block writes its complete grad_input slice; no Fill is needed.
             grad_weight->Fill(0.0);
             grad_bias->Fill(0.0);
             LayerNormBackwardKernel<BLOCK_SIZE><<<num_blocks, threads_per_block, 0, cuda_stream>>>(
diff --git a/infini_train/src/kernels/cuda/linear.cu b/infini_train/src/kernels/cuda/linear.cu
index 2725053e..1b4c1819 100644
--- a/infini_train/src/kernels/cuda/linear.cu
+++ b/infini_train/src/kernels/cuda/linear.cu
@@ -65,6 +65,7 @@ std::shared_ptr<Tensor> LinearForward(const std::shared_ptr<Tensor> &input, cons
                                  infini_train::core::GetDeviceGuardImpl(device.type())->GetStream(device))
                                  ->cuda_stream();
 
+    const float beta = bias ? 1.0f : 0.0f;
     if (bias) {
         CHECK_EQ(bias->Dims().size(), 1);
         CHECK_EQ(bias->Dims()[0], out_features);
@@ -78,9 +79,8 @@ std::shared_ptr<Tensor> LinearForward(const std::shared_ptr<Tensor> &input, cons
                     static_cast<T *>(output->DataPtr()), static_cast<const T *>(bias->DataPtr()), bs, out_features);
             },
             "CUDA LinearForward");
-    } else {
-        output->Fill(0.0);
     }
+    // In the no-bias path, beta=0 makes cuBLAS fully overwrite output; no Fill is needed.
 
     // When bs==1 and fp32, use cublasSgemv (more efficient than GEMM for matrix-vector).
     // cublasSgemv does not support bf16, so bf16 falls through to Gemm.
@@ -94,7 +94,7 @@ std::shared_ptr<Tensor> LinearForward(const std::shared_ptr<Tensor> &input, cons
                               .x = static_cast<const float *>(input->DataPtr()),
                               .y = static_cast<float *>(output->DataPtr()),
                               .alpha = 1.0f,
-                              .beta = 1.0f, // output already initialized with bias or zero above
+                              .beta = beta,
                           });
     } else {
         // cuBLAS is colmun-major
@@ -125,7 +125,7 @@ std::shared_ptr<Tensor> LinearForward(const std::shared_ptr<Tensor> &input, cons
                 .C = output->DataPtr(),
                 .ldc = static_cast<int>(out_features),
                 .alpha = 1.0f,
-                .beta = 1.0f, // bias already written into output; beta=1 accumulates
+                .beta = beta,
                 .batch_count = 1,
                 .input_dtype = dtype,
                 .output_dtype = dtype,
diff --git a/infini_train/src/kernels/cuda/reduction.cu b/infini_train/src/kernels/cuda/reduction.cu
index c56470e3..b9fb57e0 100644
--- a/infini_train/src/kernels/cuda/reduction.cu
+++ b/infini_train/src/kernels/cuda/reduction.cu
@@ -181,7 +181,7 @@ std::shared_ptr<Tensor> ReduceOpBackward(const std::shared_ptr<Tensor> &grad_out
     core::cuda::DispatchCudaFunc<INFINI_ALL_FLOATING_TYPES>(
         dtype,
         [=]<typename T>() {
-            grad_input->Fill(0.0);
+            // The backward kernel assigns every grad_input element on all reduction branches; no Fill is needed.
             GenericReduceBackwardKernel<<<num_blocks, threads_per_block, 0, cuda_stream>>>(
                 static_cast<T *>(grad_input->DataPtr()), static_cast<const T *>(grad_output->DataPtr()),
                 input ? static_cast<const T *>(input->DataPtr()) : nullptr,
diff --git a/infini_train/src/kernels/cuda/slice.cu b/infini_train/src/kernels/cuda/slice.cu
index d030d73a..f71b4be6 100644
--- a/infini_train/src/kernels/cuda/slice.cu
+++ b/infini_train/src/kernels/cuda/slice.cu
@@ -48,8 +48,7 @@ std::shared_ptr<Tensor> SliceForward(const std::shared_ptr<Tensor> &input, const
 
     auto dtype = input->Dtype();
     auto new_tensor = std::make_shared<Tensor>(new_dims, dtype, input->GetDevice());
-    // NOTE(zbl): must initialize with 0
-    new_tensor->Fill(0.0);
+    // SliceForwardKernel writes every output index in [0, total_elements); no Fill is needed.
 
     std::vector<int64_t> src_strides(dims.size(), 0), dst_strides(new_dims.size(), 0);
     int64_t stride = 1;
diff --git a/infini_train/src/kernels/cuda/softmax.cu b/infini_train/src/kernels/cuda/softmax.cu
index a7f3b612..ce709d38 100644
--- a/infini_train/src/kernels/cuda/softmax.cu
+++ b/infini_train/src/kernels/cuda/softmax.cu
@@ -198,7 +198,7 @@ std::shared_ptr<Tensor> SoftmaxBackward(const std::shared_ptr<Tensor> &grad_outp
     CHECK(dim >= 0 && dim < output->Dims().size());
 
     auto grad_input = std::make_shared<Tensor>(output_dims, promoted_type, output->GetDevice());
-    grad_input->Fill(0.0);
+    // For non-empty tensors, the grid covers every outer/axis/inner index exactly once; no Fill is needed.
 
     switch (promoted_type) {
         DISPATCH_CASE(WRAP(LaunchBackward<256, float>(grad_input, grad_output_promoted, output_promoted, dim);),
diff --git a/infini_train/src/kernels/cuda/split.cu b/infini_train/src/kernels/cuda/split.cu
index bda0dd70..b52513a2 100644
--- a/infini_train/src/kernels/cuda/split.cu
+++ b/infini_train/src/kernels/cuda/split.cu
@@ -114,6 +114,7 @@ std::shared_ptr<Tensor> LaunchSplitBackward(const std::vector<int64_t> &input_di
     const auto &grad = grad_outputs[0];
     auto dtype = grad->Dtype();
     auto grad_input = std::make_shared<Tensor>(input_dims, dtype, grad->GetDevice());
+    // Keep initialization: defensive early returns in SplitBackwardKernel can leave elements unwritten.
     grad_input->Fill(0.0);
 
     int64_t N = std::accumulate(input_dims.begin(), input_dims.begin() + dim, 1, std::multiplies<int64_t>());
diff --git a/infini_train/src/kernels/cuda/stack.cu b/infini_train/src/kernels/cuda/stack.cu
index 841940ea..2ffe7795 100644
--- a/infini_train/src/kernels/cuda/stack.cu
+++ b/infini_train/src/kernels/cuda/stack.cu
@@ -113,7 +113,7 @@ std::vector<std::shared_ptr<Tensor>> StackBackward(const std::vector<int64_t> &i
     std::vector<std::shared_ptr<Tensor>> grads;
     for (int i = 0; i < num_inputs; ++i) {
         auto t = std::make_shared<Tensor>(base_dims, dtype, grad_output->GetDevice());
-        t->Fill(0.0);
+        // StackBackwardKernel writes every element of every grad tensor exactly once; no Fill is needed.
         grads.push_back(t);
     }
 
diff --git a/infini_train/src/nn/parallel/pp/pipeline_schedule.cc b/infini_train/src/nn/parallel/pp/pipeline_schedule.cc
index 090b7b15..ff6df532 100644
--- a/infini_train/src/nn/parallel/pp/pipeline_schedule.cc
+++ b/infini_train/src/nn/parallel/pp/pipeline_schedule.cc
@@ -255,9 +255,8 @@ float PipelineSchedule::StepMicroBatches(const std::vector<std::shared_ptr<Tenso
                         {activations[task.local_chunk_idx][mb][0], std::make_shared<Tensor>(target_on_device)})[0];
                     loss = loss / n;
                 }
-                total_loss += static_cast<const float *>(loss->To(Device()).DataPtr())[0];
-
                 loss->Backward();
+                total_loss += static_cast<const float *>(loss->To(Device()).DataPtr())[0];
             } else {
                 auto out_tensor = activations[task.local_chunk_idx][mb][0];
 
diff --git a/tests/autograd/test_autograd_linear_forward.cc b/tests/autograd/test_autograd_linear_forward.cc
index 89d7c537..33ba9ea1 100644
--- a/tests/autograd/test_autograd_linear_forward.cc
+++ b/tests/autograd/test_autograd_linear_forward.cc
@@ -3,6 +3,7 @@
 #include "gtest/gtest.h"
 
 #include "infini_train/include/autograd/linear.h"
+#include "infini_train/include/core/runtime/device_guard.h"
 #include "infini_train/include/nn/parallel/global.h"
 #include "infini_train/include/tensor.h"
 
@@ -12,17 +13,30 @@ using namespace infini_train;
 
 class AutogradLinearForwardTest : public infini_train::test::InfiniTrainTest {};
 
+namespace {
+std::shared_ptr<Tensor> CopyToCPU(const std::shared_ptr<Tensor> &tensor) {
+    auto cpu = std::make_shared<Tensor>(tensor->Dims(), tensor->Dtype(), Device());
+    cpu->CopyFrom(tensor);
+    core::GetDeviceGuardImpl(tensor->GetDevice().type())->SynchronizeDevice(tensor->GetDevice());
+    return cpu;
+}
+} // namespace
+
 TEST_P(AutogradLinearForwardTest, LinearForward) {
     auto input = std::make_shared<Tensor>(std::vector<int64_t>{2, 3}, DataType::kFLOAT32, GetDevice(), true);
     input->Fill(1.0f);
     auto weight = std::make_shared<Tensor>(std::vector<int64_t>{4, 3}, DataType::kFLOAT32, GetDevice(), true);
     weight->Fill(1.0f);
     auto bias = std::make_shared<Tensor>(std::vector<int64_t>{4}, DataType::kFLOAT32, GetDevice(), true);
-    bias->Fill(0.0f);
+    bias->Fill(2.0f);
     auto linear_fn = std::make_shared<autograd::Linear>();
     auto result = linear_fn->Apply({input, weight, bias});
     EXPECT_EQ(result.size(), 1);
     EXPECT_EQ(result[0]->Dims(), (std::vector<int64_t>{2, 4}));
+
+    auto output = CopyToCPU(result[0]);
+    const float *actual = static_cast<const float *>(output->DataPtr());
+    for (int64_t i = 0; i < output->NumElements(); ++i) { EXPECT_FLOAT_EQ(actual[i], 5.0f); }
 }
 
 TEST_P(AutogradLinearForwardTest, LinearNoBias) {
@@ -34,6 +48,10 @@ TEST_P(AutogradLinearForwardTest, LinearNoBias) {
     auto result = linear_fn->Apply({input, weight});
     EXPECT_EQ(result.size(), 1);
     EXPECT_EQ(result[0]->Dims(), (std::vector<int64_t>{2, 4}));
+
+    auto output = CopyToCPU(result[0]);
+    const float *actual = static_cast<const float *>(output->DataPtr());
+    for (int64_t i = 0; i < output->NumElements(); ++i) { EXPECT_FLOAT_EQ(actual[i], 3.0f); }
 }
 
 TEST_P(AutogradLinearForwardTest, LinearBatch) {
diff --git a/tests/autograd/test_autograd_loss.cc b/tests/autograd/test_autograd_loss.cc
new file mode 100644
index 00000000..d55cf7ac
--- /dev/null
+++ b/tests/autograd/test_autograd_loss.cc
@@ -0,0 +1,63 @@
+#include <cmath>
+#include <vector>
+
+#include "gtest/gtest.h"
+
+#include "infini_train/include/autograd/loss.h"
+#include "infini_train/include/core/runtime/device_guard.h"
+#include "infini_train/include/nn/parallel/global.h"
+#include "infini_train/include/tensor.h"
+
+#include "tests/common/test_utils.h"
+
+using namespace infini_train;
+
+class AutogradLossTest : public infini_train::test::InfiniTrainTest {};
+
+namespace {
+std::shared_ptr<Tensor> CopyToCPU(const std::shared_ptr<Tensor> &tensor) {
+    auto cpu = std::make_shared<Tensor>(tensor->Dims(), tensor->Dtype(), Device());
+    cpu->CopyFrom(tensor);
+    core::GetDeviceGuardImpl(tensor->GetDevice().type())->SynchronizeDevice(tensor->GetDevice());
+    return cpu;
+}
+} // namespace
+
+TEST_P(AutogradLossTest, CrossEntropyForwardAndBackwardValues) {
+    std::vector<float> logits_values{1.0f, 2.0f, 3.0f, 2.0f, 1.0f, 0.0f};
+    auto logits
+        = std::make_shared<Tensor>(logits_values.data(), std::vector<int64_t>{2, 3}, DataType::kFLOAT32, GetDevice());
+    auto target = std::make_shared<Tensor>(std::vector<int64_t>{2}, DataType::kINT64, GetDevice());
+    target->Fill(0);
+
+    auto cross_entropy = std::make_shared<autograd::CrossEntropy>();
+    auto result = cross_entropy->Apply({logits, target});
+    ASSERT_EQ(result.size(), 1);
+    ASSERT_TRUE(result[0]->Dims().empty());
+
+    const float row0_sum = std::exp(1.0f) + std::exp(2.0f) + std::exp(3.0f);
+    const float row1_sum = std::exp(2.0f) + std::exp(1.0f) + std::exp(0.0f);
+    const float expected_loss = 0.5f * ((std::log(row0_sum) - 1.0f) + (std::log(row1_sum) - 2.0f));
+    auto loss_cpu = CopyToCPU(result[0]);
+    EXPECT_NEAR(static_cast<const float *>(loss_cpu->DataPtr())[0], expected_loss, 1e-5f);
+
+    auto grad_output = std::make_shared<Tensor>(std::vector<int64_t>{}, DataType::kFLOAT32, GetDevice());
+    grad_output->Fill(1.0f);
+    auto grad_inputs = cross_entropy->Backward({grad_output});
+    ASSERT_EQ(grad_inputs.size(), 2);
+    ASSERT_NE(grad_inputs[0], nullptr);
+    EXPECT_EQ(grad_inputs[1], nullptr);
+
+    auto grad_cpu = CopyToCPU(grad_inputs[0]);
+    const float *actual_grad = static_cast<const float *>(grad_cpu->DataPtr());
+    for (int row = 0; row < 2; ++row) {
+        const float sum = row == 0 ? row0_sum : row1_sum;
+        for (int col = 0; col < 3; ++col) {
+            const float probability = std::exp(logits_values[row * 3 + col]) / sum;
+            const float expected_grad = 0.5f * (probability - (col == 0 ? 1.0f : 0.0f));
+            EXPECT_NEAR(actual_grad[row * 3 + col], expected_grad, 1e-5f);
+        }
+    }
+}
+
+INFINI_TRAIN_REGISTER_TEST(AutogradLossTest);
diff --git a/tests/autograd/test_autograd_normalization_backward.cc b/tests/autograd/test_autograd_normalization_backward.cc
index d3926ce8..463860ea 100644
--- a/tests/autograd/test_autograd_normalization_backward.cc
+++ b/tests/autograd/test_autograd_normalization_backward.cc
@@ -3,6 +3,7 @@
 #include "gtest/gtest.h"
 
 #include "infini_train/include/autograd/normalization.h"
+#include "infini_train/include/core/runtime/device_guard.h"
 #include "infini_train/include/nn/parallel/global.h"
 #include "infini_train/include/tensor.h"
 
@@ -12,6 +13,17 @@ using namespace infini_train;
 
 class AutogradNormalizationBackwardTest : public infini_train::test::InfiniTrainTest {};
 
+namespace {
+void ExpectValues(const std::shared_ptr<Tensor> &tensor, float expected) {
+    auto cpu = std::make_shared<Tensor>(tensor->Dims(), tensor->Dtype(), Device());
+    cpu->CopyFrom(tensor);
+    core::GetDeviceGuardImpl(tensor->GetDevice().type())->SynchronizeDevice(tensor->GetDevice());
+
+    const float *actual = static_cast<const float *>(cpu->DataPtr());
+    for (int64_t i = 0; i < cpu->NumElements(); ++i) { EXPECT_NEAR(actual[i], expected, 1e-5f); }
+}
+} // namespace
+
 TEST_P(AutogradNormalizationBackwardTest, LayerNormBackward) {
     auto a = std::make_shared<Tensor>(std::vector<int64_t>{2, 3, 4}, DataType::kFLOAT32, GetDevice(), true);
     a->Fill(1.0f);
@@ -25,6 +37,9 @@ TEST_P(AutogradNormalizationBackwardTest, LayerNormBackward) {
     grad->Fill(1.0f);
     auto grad_inputs = layernorm_fn->Backward({grad});
     EXPECT_EQ(grad_inputs.size(), 3);
+    ExpectValues(grad_inputs[0], 0.0f);
+    ExpectValues(grad_inputs[1], 0.0f);
+    ExpectValues(grad_inputs[2], 6.0f);
 }
 
 TEST_P(AutogradNormalizationBackwardTest, LayerNormBackwardZeroBias) {
diff --git a/tests/autograd/test_autograd_reduction_backward.cc b/tests/autograd/test_autograd_reduction_backward.cc
index 04709dd2..871116fc 100644
--- a/tests/autograd/test_autograd_reduction_backward.cc
+++ b/tests/autograd/test_autograd_reduction_backward.cc
@@ -3,6 +3,7 @@
 #include "gtest/gtest.h"
 
 #include "infini_train/include/autograd/reduction.h"
+#include "infini_train/include/core/runtime/device_guard.h"
 #include "infini_train/include/nn/parallel/global.h"
 #include "infini_train/include/tensor.h"
 
@@ -12,6 +13,18 @@ using namespace infini_train;
 
 class AutogradReductionBackwardTest : public infini_train::test::InfiniTrainTest {};
 
+namespace {
+void ExpectValues(const std::shared_ptr<Tensor> &tensor, const std::vector<float> &expected) {
+    auto cpu = std::make_shared<Tensor>(tensor->Dims(), tensor->Dtype(), Device());
+    cpu->CopyFrom(tensor);
+    core::GetDeviceGuardImpl(tensor->GetDevice().type())->SynchronizeDevice(tensor->GetDevice());
+
+    ASSERT_EQ(cpu->NumElements(), expected.size());
+    const float *actual = static_cast<const float *>(cpu->DataPtr());
+    for (size_t i = 0; i < expected.size(); ++i) { EXPECT_FLOAT_EQ(actual[i], expected[i]); }
+}
+} // namespace
+
 TEST_P(AutogradReductionBackwardTest, SumBackward) {
     auto a = std::make_shared<Tensor>(std::vector<int64_t>{2, 3}, DataType::kFLOAT32, GetDevice(), true);
     a->Fill(1.0f);
@@ -21,6 +34,7 @@ TEST_P(AutogradReductionBackwardTest, SumBackward) {
     grad->Fill(1.0f);
     auto grad_inputs = sum_fn->Backward({grad});
     EXPECT_EQ(grad_inputs.size(), 1);
+    ExpectValues(grad_inputs[0], {1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f});
 }
 
 TEST_P(AutogradReductionBackwardTest, MeanBackward) {
@@ -32,28 +46,31 @@ TEST_P(AutogradReductionBackwardTest, MeanBackward) {
     grad->Fill(1.0f);
     auto grad_inputs = mean_fn->Backward({grad});
     EXPECT_EQ(grad_inputs.size(), 1);
+    ExpectValues(grad_inputs[0], {1.0f / 3.0f, 1.0f / 3.0f, 1.0f / 3.0f, 1.0f / 3.0f, 1.0f / 3.0f, 1.0f / 3.0f});
 }
 
 TEST_P(AutogradReductionBackwardTest, MaxBackward) {
-    auto a = std::make_shared<Tensor>(std::vector<int64_t>{2, 3}, DataType::kFLOAT32, GetDevice(), true);
-    a->Fill(1.0f);
+    std::vector<float> values{1.0f, 3.0f, 2.0f, 4.0f, 0.0f, -1.0f};
+    auto a = std::make_shared<Tensor>(values.data(), std::vector<int64_t>{2, 3}, DataType::kFLOAT32, GetDevice());
     auto max_fn = std::make_shared<autograd::Max>(1, false);
     auto result = max_fn->Apply({a});
     auto grad = std::make_shared<Tensor>(std::vector<int64_t>{2}, DataType::kFLOAT32, GetDevice(), true);
     grad->Fill(1.0f);
     auto grad_inputs = max_fn->Backward({grad});
     EXPECT_EQ(grad_inputs.size(), 1);
+    ExpectValues(grad_inputs[0], {0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f});
 }
 
 TEST_P(AutogradReductionBackwardTest, MinBackward) {
-    auto a = std::make_shared<Tensor>(std::vector<int64_t>{2, 3}, DataType::kFLOAT32, GetDevice(), true);
-    a->Fill(1.0f);
+    std::vector<float> values{1.0f, 3.0f, 2.0f, 4.0f, 0.0f, -1.0f};
+    auto a = std::make_shared<Tensor>(values.data(), std::vector<int64_t>{2, 3}, DataType::kFLOAT32, GetDevice());
     auto min_fn = std::make_shared<autograd::Min>(1, false);
     auto result = min_fn->Apply({a});
     auto grad = std::make_shared<Tensor>(std::vector<int64_t>{2}, DataType::kFLOAT32, GetDevice(), true);
     grad->Fill(1.0f);
     auto grad_inputs = min_fn->Backward({grad});
     EXPECT_EQ(grad_inputs.size(), 1);
+    ExpectValues(grad_inputs[0], {1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f});
 }
 
 TEST_P(AutogradReductionBackwardTest, SumBackwardKeepDim) {
@@ -65,6 +82,7 @@ TEST_P(AutogradReductionBackwardTest, SumBackwardKeepDim) {
     grad->Fill(1.0f);
     auto grad_inputs = sum_fn->Backward({grad});
     EXPECT_EQ(grad_inputs.size(), 1);
+    ExpectValues(grad_inputs[0], {1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f});
 }
 
 TEST_P(AutogradReductionBackwardTest, MeanBackwardKeepDim) {
@@ -76,6 +94,7 @@ TEST_P(AutogradReductionBackwardTest, MeanBackwardKeepDim) {
     grad->Fill(1.0f);
     auto grad_inputs = mean_fn->Backward({grad});
     EXPECT_EQ(grad_inputs.size(), 1);
+    ExpectValues(grad_inputs[0], {1.0f / 3.0f, 1.0f / 3.0f, 1.0f / 3.0f, 1.0f / 3.0f, 1.0f / 3.0f, 1.0f / 3.0f});
 }
 
 INFINI_TRAIN_REGISTER_TEST(AutogradReductionBackwardTest);
diff --git a/tests/autograd/test_autograd_softmax_backward.cc b/tests/autograd/test_autograd_softmax_backward.cc
index fb18f00e..73516a5a 100644
--- a/tests/autograd/test_autograd_softmax_backward.cc
+++ b/tests/autograd/test_autograd_softmax_backward.cc
@@ -3,6 +3,7 @@
 #include "gtest/gtest.h"
 
 #include "infini_train/include/autograd/softmax.h"
+#include "infini_train/include/core/runtime/device_guard.h"
 #include "infini_train/include/nn/parallel/global.h"
 #include "infini_train/include/tensor.h"
 
@@ -12,15 +13,29 @@ using namespace infini_train;
 
 class AutogradSoftmaxBackwardTest : public infini_train::test::InfiniTrainTest {};
 
+namespace {
+void ExpectValues(const std::shared_ptr<Tensor> &tensor, const std::vector<float> &expected) {
+    auto cpu = std::make_shared<Tensor>(tensor->Dims(), tensor->Dtype(), Device());
+    cpu->CopyFrom(tensor);
+    core::GetDeviceGuardImpl(tensor->GetDevice().type())->SynchronizeDevice(tensor->GetDevice());
+
+    ASSERT_EQ(cpu->NumElements(), expected.size());
+    const float *actual = static_cast<const float *>(cpu->DataPtr());
+    for (size_t i = 0; i < expected.size(); ++i) { EXPECT_NEAR(actual[i], expected[i], 1e-6f); }
+}
+} // namespace
+
 TEST_P(AutogradSoftmaxBackwardTest, SoftmaxBackward) {
     auto a = std::make_shared<Tensor>(std::vector<int64_t>{2, 3}, DataType::kFLOAT32, GetDevice(), true);
     a->Fill(1.0f);
     auto softmax_fn = std::make_shared<autograd::Softmax>(1);
     auto result = softmax_fn->Apply({a});
-    auto grad = std::make_shared<Tensor>(std::vector<int64_t>{2, 3}, DataType::kFLOAT32, GetDevice(), true);
-    grad->Fill(1.0f);
+    std::vector<float> grad_values{1.0f, 2.0f, 4.0f, 4.0f, 2.0f, 1.0f};
+    auto grad
+        = std::make_shared<Tensor>(grad_values.data(), std::vector<int64_t>{2, 3}, DataType::kFLOAT32, GetDevice());
     auto grad_inputs = softmax_fn->Backward({grad});
     EXPECT_EQ(grad_inputs.size(), 1);
+    ExpectValues(grad_inputs[0], {-4.0f / 9.0f, -1.0f / 9.0f, 5.0f / 9.0f, 5.0f / 9.0f, -1.0f / 9.0f, -4.0f / 9.0f});
 }
 
 TEST_P(AutogradSoftmaxBackwardTest, SoftmaxBackwardDim0) {
@@ -28,10 +43,13 @@ TEST_P(AutogradSoftmaxBackwardTest, SoftmaxBackwardDim0) {
     a->Fill(1.0f);
     auto softmax_fn = std::make_shared<autograd::Softmax>(0);
     auto result = softmax_fn->Apply({a});
-    auto grad = std::make_shared<Tensor>(std::vector<int64_t>{4, 3}, DataType::kFLOAT32, GetDevice(), true);
-    grad->Fill(1.0f);
+    std::vector<float> grad_values{1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f, 9.0f, 10.0f, 11.0f, 12.0f};
+    auto grad
+        = std::make_shared<Tensor>(grad_values.data(), std::vector<int64_t>{4, 3}, DataType::kFLOAT32, GetDevice());
     auto grad_inputs = softmax_fn->Backward({grad});
     EXPECT_EQ(grad_inputs.size(), 1);
+    ExpectValues(grad_inputs[0], {-1.125f, -1.125f, -1.125f, -0.375f, -0.375f, -0.375f, 0.375f, 0.375f, 0.375f, 1.125f,
+                                  1.125f, 1.125f});
 }
 
 INFINI_TRAIN_REGISTER_TEST(AutogradSoftmaxBackwardTest);
diff --git a/tests/autograd/test_autograd_transform_backward.cc b/tests/autograd/test_autograd_transform_backward.cc
index cdb3b367..cdde7a37 100644
--- a/tests/autograd/test_autograd_transform_backward.cc
+++ b/tests/autograd/test_autograd_transform_backward.cc
@@ -1,8 +1,10 @@
+#include <numeric>
 #include <vector>
 
 #include "gtest/gtest.h"
 
 #include "infini_train/include/autograd/transform.h"
+#include "infini_train/include/core/runtime/device_guard.h"
 #include "infini_train/include/nn/parallel/global.h"
 #include "infini_train/include/tensor.h"
 
@@ -12,6 +14,22 @@ using namespace infini_train;
 
 class AutogradTransformBackwardTest : public infini_train::test::InfiniTrainTest {};
 
+namespace {
+std::shared_ptr<Tensor> CopyToCPU(const std::shared_ptr<Tensor> &tensor) {
+    auto cpu = std::make_shared<Tensor>(tensor->Dims(), tensor->Dtype(), Device());
+    cpu->CopyFrom(tensor);
+    core::GetDeviceGuardImpl(tensor->GetDevice().type())->SynchronizeDevice(tensor->GetDevice());
+    return cpu;
+}
+
+void ExpectValues(const std::shared_ptr<Tensor> &tensor, const std::vector<float> &expected) {
+    auto cpu = CopyToCPU(tensor);
+    ASSERT_EQ(cpu->NumElements(), expected.size());
+    const float *actual = static_cast<const float *>(cpu->DataPtr());
+    for (size_t i = 0; i < expected.size(); ++i) { EXPECT_FLOAT_EQ(actual[i], expected[i]); }
+}
+} // namespace
+
 TEST_P(AutogradTransformBackwardTest, TransposeBackward) {
     auto a = std::make_shared<Tensor>(std::vector<int64_t>{2, 3}, DataType::kFLOAT32, GetDevice(), true);
     a->Fill(1.0f);
@@ -23,4 +41,56 @@ TEST_P(AutogradTransformBackwardTest, TransposeBackward) {
     EXPECT_EQ(grad_inputs.size(), 1);
 }
 
+TEST_P(AutogradTransformBackwardTest, SplitBackwardValues) {
+    auto input = std::make_shared<Tensor>(std::vector<int64_t>{2, 5}, DataType::kFLOAT32, GetDevice(), true);
+    auto split = std::make_shared<autograd::Split>(2, 1);
+    auto outputs = split->Apply({input});
+    ASSERT_EQ(outputs.size(), 3);
+
+    std::vector<std::shared_ptr<Tensor>> grad_outputs;
+    for (size_t i = 0; i < outputs.size(); ++i) {
+        auto grad = std::make_shared<Tensor>(outputs[i]->Dims(), DataType::kFLOAT32, GetDevice());
+        grad->Fill(static_cast<float>(i + 1));
+        grad_outputs.push_back(grad);
+    }
+
+    auto grad_inputs = split->Backward(grad_outputs);
+    ASSERT_EQ(grad_inputs.size(), 1);
+    ExpectValues(grad_inputs[0], {1.0f, 1.0f, 2.0f, 2.0f, 3.0f, 1.0f, 1.0f, 2.0f, 2.0f, 3.0f});
+}
+
+TEST_P(AutogradTransformBackwardTest, StackBackwardValues) {
+    auto a = std::make_shared<Tensor>(std::vector<int64_t>{2, 3}, DataType::kFLOAT32, GetDevice(), true);
+    auto b = std::make_shared<Tensor>(std::vector<int64_t>{2, 3}, DataType::kFLOAT32, GetDevice(), true);
+    auto stack = std::make_shared<autograd::Stack>(1);
+    auto outputs = stack->Apply({a, b});
+    ASSERT_EQ(outputs.size(), 1);
+
+    std::vector<float> grad_values(12);
+    std::iota(grad_values.begin(), grad_values.end(), 0.0f);
+    auto grad_output
+        = std::make_shared<Tensor>(grad_values.data(), outputs[0]->Dims(), DataType::kFLOAT32, GetDevice());
+    auto grad_inputs = stack->Backward({grad_output});
+    ASSERT_EQ(grad_inputs.size(), 2);
+    ExpectValues(grad_inputs[0], {0.0f, 1.0f, 2.0f, 6.0f, 7.0f, 8.0f});
+    ExpectValues(grad_inputs[1], {3.0f, 4.0f, 5.0f, 9.0f, 10.0f, 11.0f});
+}
+
+TEST_P(AutogradTransformBackwardTest, ConcatBackwardValues) {
+    auto a = std::make_shared<Tensor>(std::vector<int64_t>{2, 2}, DataType::kFLOAT32, GetDevice(), true);
+    auto b = std::make_shared<Tensor>(std::vector<int64_t>{2, 1}, DataType::kFLOAT32, GetDevice(), true);
+    auto concat = std::make_shared<autograd::Concat>(1);
+    auto outputs = concat->Apply({a, b});
+    ASSERT_EQ(outputs.size(), 1);
+
+    std::vector<float> grad_values(6);
+    std::iota(grad_values.begin(), grad_values.end(), 0.0f);
+    auto grad_output
+        = std::make_shared<Tensor>(grad_values.data(), outputs[0]->Dims(), DataType::kFLOAT32, GetDevice());
+    auto grad_inputs = concat->Backward({grad_output});
+    ASSERT_EQ(grad_inputs.size(), 2);
+    ExpectValues(grad_inputs[0], {0.0f, 1.0f, 3.0f, 4.0f});
+    ExpectValues(grad_inputs[1], {2.0f, 5.0f});
+}
+
 INFINI_TRAIN_REGISTER_TEST(AutogradTransformBackwardTest);
diff --git a/tests/autograd/test_autograd_transform_forward.cc b/tests/autograd/test_autograd_transform_forward.cc
index 680d8b0d..055f2dba 100644
--- a/tests/autograd/test_autograd_transform_forward.cc
+++ b/tests/autograd/test_autograd_transform_forward.cc
@@ -1,8 +1,10 @@
+#include <numeric>
 #include <vector>
 
 #include "gtest/gtest.h"
 
 #include "infini_train/include/autograd/transform.h"
+#include "infini_train/include/core/runtime/device_guard.h"
 #include "infini_train/include/nn/parallel/global.h"
 #include "infini_train/include/tensor.h"
 
@@ -12,6 +14,18 @@ using namespace infini_train;
 
 class AutogradTransformForwardTest : public infini_train::test::InfiniTrainTest {};
 
+namespace {
+void ExpectValues(const std::shared_ptr<Tensor> &tensor, const std::vector<float> &expected) {
+    auto cpu = std::make_shared<Tensor>(tensor->Dims(), tensor->Dtype(), Device());
+    cpu->CopyFrom(tensor);
+    core::GetDeviceGuardImpl(tensor->GetDevice().type())->SynchronizeDevice(tensor->GetDevice());
+
+    ASSERT_EQ(cpu->NumElements(), expected.size());
+    const float *actual = static_cast<const float *>(cpu->DataPtr());
+    for (size_t i = 0; i < expected.size(); ++i) { EXPECT_FLOAT_EQ(actual[i], expected[i]); }
+}
+} // namespace
+
 TEST_P(AutogradTransformForwardTest, TransposeForward) {
     auto a = std::make_shared<Tensor>(std::vector<int64_t>{2, 3}, DataType::kFLOAT32, GetDevice(), true);
     a->Fill(1.0f);
@@ -22,12 +36,14 @@ TEST_P(AutogradTransformForwardTest, TransposeForward) {
 }
 
 TEST_P(AutogradTransformForwardTest, SliceForward) {
-    auto a = std::make_shared<Tensor>(std::vector<int64_t>{4, 4}, DataType::kFLOAT32, GetDevice(), true);
-    a->Fill(1.0f);
+    std::vector<float> values(16);
+    std::iota(values.begin(), values.end(), 0.0f);
+    auto a = std::make_shared<Tensor>(values.data(), std::vector<int64_t>{4, 4}, DataType::kFLOAT32, GetDevice());
     auto slice_fn = std::make_shared<autograd::Slice>(std::vector<int64_t>{1, 1}, std::vector<int64_t>{3, 3},
                                                       std::vector<int64_t>{1, 1});
     auto result = slice_fn->Apply({a});
     EXPECT_EQ(result.size(), 1);
+    ExpectValues(result[0], {5.0f, 6.0f, 9.0f, 10.0f});
 }
 
 TEST_P(AutogradTransformForwardTest, SplitForward) {