 
     template <typename T>
     __global__ static void copyBuffers(Nd4jLong* destination, void const* source, Nd4jLong bufferLength) {
-        const auto tid = blockIdx.x * gridDim.x + threadIdx.x;
+        const auto tid = blockIdx.x * blockDim.x + threadIdx.x;
         const auto step = gridDim.x * blockDim.x;
         for (int t = tid; t < bufferLength; t += step) {
-            destination[t] = reinterpret_cast<T const*>(source)[t];
+            destination[t] = static_cast<Nd4jLong>(reinterpret_cast<T const*>(source)[t]);
         }
     }
 

            if (err != 0)
                throw nd4j::cuda_exception::build("Cannot deallocate memory for labels long buffer", err);
        }
        manager.synchronize();
    }

    void confusionFunctor(nd4j::LaunchContext * context, NDArray* labels, NDArray* predictions, NDArray* weights, NDArray* output) {
        auto xType = output->dataType(); // weights can be null

        BUILD_SINGLE_SELECTOR(xType, _confusionFunctor, (context, labels, predictions, weights, output), NUMERIC_TYPES);
    }

    BUILD_SINGLE_TEMPLATE(template void _confusionFunctor, (nd4j::LaunchContext * context, NDArray* labels, NDArray* predictions, NDArray* weights, NDArray* output);, NUMERIC_TYPES);

}
}
}
        }
        __syncthreads();

        const auto tid = blockIdx.x * blockDim.x + threadIdx.x;
        const auto step = gridDim.x * blockDim.x;
        for (int t = tid; t < bufferLength; t += step) {
            auto label = labelsBuffer[t]; //->e<Nd4jLong>(j);
            auto pred = predictionBuffer[t]; //->e<Nd4jLong>(j);
            auto tZ = z + tadOffsets[label];
            T val = (weightsBuffer == nullptr ? (T)1.0f : w[t]);

            auto idx = shape::getIndexOffset(pred, tadShape, arrLen);
            tZ[idx] = val;
        }
    }

    template <typename X, typename Z>
    void _confusionFunctor(nd4j::LaunchContext * context, NDArray* labels, NDArray* predictions, NDArray* weights, NDArray* output) {
        auto stream = context->getCudaStream();

        auto pack = nd4j::ConstantTadHelper::getInstance()->tadForDimensions(output->shapeInfo(), 1);

        PointersManager manager(context, "helpers::confusion");

