 
     // Run the first reduction on GPU
 
-    invoke_cce_loss_kernel<T, false>(n, output, incx, labels, incy, y_gpu_2, numThreads, numBlocks);
+    if (Loss) {
+        invoke_cce_loss_kernel<T, false>(n, output, incx, labels, incy, y_gpu_2, numThreads, numBlocks);
+    } else {
+        invoke_cce_error_kernel<T, false>(n, m, output, labels, y_gpu_2, numThreads, numBlocks);
+    }
 
     size_t s = numBlocks;
 

    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[];
