     return logf(output) * label;
 }
 
-/*template <typename T>*/
-/*__global__ void cce_loss_kernel(size_t n, const T* output, size_t incx, const T* labels, size_t incy, T* y) {*/
-    /*auto index  = threadIdx.x + blockIdx.x * blockDim.x;*/
-    /*auto stride = blockDim.x * gridDim.x;*/
-
-    /*for (; index < n; index += stride) {*/
-        /*y[incy * index] = cce_loss(output[incx * index], labels[incx * index]);*/
-    /*}*/
-/*}*/
-
 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[];

    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);


        cuda_check(cudaMemcpy(y_gpu_1, y_gpu_2, s * sizeof(T), cudaMemcpyDeviceToDevice));

        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);
    }
