 
         cuda_check(cudaMemcpy(y_gpu_1, y_gpu_2, s * sizeof(T), cudaMemcpyDeviceToDevice));
 
-        invoke_cce_loss_kernel<T, false>(s, y_gpu_1, 1, y_gpu_1, 1, y_gpu_2, numThreads, numBlocks);
+        if (Loss) {
+            invoke_cce_loss_kernel<T, true>(s, y_gpu_1, 1, y_gpu_1, 1, y_gpu_2, numThreads, numBlocks);
+        } else {
+            invoke_cce_error_kernel<T, true>(s, m, y_gpu_1, y_gpu_1, y_gpu_2, numThreads, numBlocks);
+        }
 
         s = (s + numThreads * 2 - 1) / (numThreads * 2);
     }

    sum_reduce_impl<T, blockSize>(r_output, shared_data, mySum);
}

template <typename T, bool Reduce>
void invoke_cce_loss_kernel(size_t n, const T* output, size_t incx, const T* labels, size_t incy, T* r_output, size_t numThreads, size_t numBlocks) {
    int sharedSize = (numThreads <= 32) ? 64 * sizeof(T) : numThreads * sizeof(T);

    return logf(output) * label;
}

template <class T, size_t blockSize, bool Reduce>
__global__ void cce_loss_kernel(size_t n, const T* output, size_t incx, const T* labels, size_t incy, T* r_output) {
    extern __shared__ __align__(sizeof(T)) volatile unsigned char shared_data_raw[];
