<commit_before>/*******************************************************
 * Copyright (c) 2014, ArrayFire
 * All rights reserved.
 *
 * This file is distributed under 3-clause BSD license.
 * The complete license agreement can be obtained at:
 * http://arrayfire.com/licenses/BSD-3-Clause
 ********************************************************/

#include <cuda_runtime_api.h>
#include <af/array.h>
#include <af/defines.h>
#include <Array.hpp>
#include <copy.hpp>
#include <kernel/memcopy.hpp>
#include <err_cuda.hpp>
#include <math.hpp>

namespace cuda
{

    template<typename T>
    void copyData(T *data, const Array<T> &A)
    {
        // FIXME: Merge this with copyArray
        A.eval();

        Array<T> out = A;
        const T *ptr = NULL;

        if (A.isOwner() || // No offsets, No strides
            A.ndims() == 1 // Simple offset, no strides.
            ) {

            //A.get() gets data with offsets
            ptr = A.get();
        } else {
            //FIXME: Think about implementing eval
            out = copyArray(A);
            ptr = out.get();
        }

        CUDA_CHECK(cudaMemcpy(data, ptr,
                              A.elements() * sizeof(T),
                              cudaMemcpyDeviceToHost));

        return;
    }

    template<typename T>
    Array<T> copyArray(const Array<T> &A)
    {
        Array<T> out = createEmptyArray<T>(A.dims());

        if (A.isLinear()) {
            CUDA_CHECK(cudaMemcpyAsync(out.get(), A.get(),
                                       A.elements() * sizeof(T),
                                       cudaMemcpyDeviceToDevice,
                                       cuda::getStream(cuda::getActiveDeviceId())));
        } else {
            // FIXME: Seems to fail when using Param<T>
            kernel::memcopy(out.get(), out.strides().get(), A.get(), A.dims().get(),
                            A.strides().get(), (uint)A.ndims());
        }
        return out;
    }

    template<typename inType, typename outType>
    Array<outType> padArray(Array<inType> const &in, dim4 const &dims, outType default_value, double factor)
    {
        ARG_ASSERT(1, (in.ndims() == dims.ndims()));
        Array<outType> ret = createEmptyArray<outType>(dims);
        kernel::copy<inType, outType>(ret, in, in.ndims(), default_value, factor);
        CUDA_CHECK(cudaDeviceSynchronize());
        return ret;
    }

    template<typename T>
    void multiply_inplace(Array<T> &in, double val)
    {
        kernel::copy<T, T>(in, in, in.ndims(), scalar<T>(0), val);
    }

    template<typename inType, typename outType>
    struct copyWrapper {
        void operator()(Array<outType> &out, Array<inType> const &in)
        {
            kernel::copy<inType, outType>(out, in, in.ndims(), scalar<outType>(0), 1);
        }
    };<commit_msg>Removing unneeded cudaDeviceSynchronize()<commit_after>/*******************************************************
 * Copyright (c) 2014, ArrayFire
 * All rights reserved.
 *
 * This file is distributed under 3-clause BSD license.
 * The complete license agreement can be obtained at:
 * http://arrayfire.com/licenses/BSD-3-Clause
 ********************************************************/

#include <cuda_runtime_api.h>
#include <af/array.h>
#include <af/defines.h>
#include <Array.hpp>
#include <copy.hpp>
#include <kernel/memcopy.hpp>
#include <err_cuda.hpp>
#include <math.hpp>

namespace cuda
{

    template<typename T>
    void copyData(T *data, const Array<T> &A)
    {
        // FIXME: Merge this with copyArray
        A.eval();

        Array<T> out = A;
        const T *ptr = NULL;

        if (A.isOwner() || // No offsets, No strides
            A.ndims() == 1 // Simple offset, no strides.
            ) {

            //A.get() gets data with offsets
            ptr = A.get();
        } else {
            //FIXME: Think about implementing eval
            out = copyArray(A);
            ptr = out.get();
        }

        CUDA_CHECK(cudaMemcpy(data, ptr,
                              A.elements() * sizeof(T),
                              cudaMemcpyDeviceToHost));

        return;
    }

    template<typename T>
    Array<T> copyArray(const Array<T> &A)
    {
        Array<T> out = createEmptyArray<T>(A.dims());

        if (A.isLinear()) {
            CUDA_CHECK(cudaMemcpyAsync(out.get(), A.get(),
                                       A.elements() * sizeof(T),
                                       cudaMemcpyDeviceToDevice,
                                       cuda::getStream(cuda::getActiveDeviceId())));
        } else {
            // FIXME: Seems to fail when using Param<T>
            kernel::memcopy(out.get(), out.strides().get(), A.get(), A.dims().get(),
                            A.strides().get(), (uint)A.ndims());
        }
        return out;
    }

    template<typename inType, typename outType>
    Array<outType> padArray(Array<inType> const &in, dim4 const &dims, outType default_value, double factor)
    {
        ARG_ASSERT(1, (in.ndims() == dims.ndims()));
        Array<outType> ret = createEmptyArray<outType>(dims);
        kernel::copy<inType, outType>(ret, in, in.ndims(), default_value, factor);
        return ret;
    }

    template<typename T>
    void multiply_inplace(Array<T> &in, double val)
    {
        kernel::copy<T, T>(in, in, in.ndims(), scalar<T>(0), val);
    }

    template<typename inType, typename outType>
    struct copyWrapper {
        void operator()(Array<outType> &out, Array<inType> const &in)
        {
            kernel::copy<inType, outType>(out, in, in.ndims(), scalar<outType>(0), 1);
        }
    };
{'fraction_non_alphanumeric': 0.17351841964762413, 'fraction_numerical': 0.004271222637479978, 'mean_word_length': 2.0777656078860898, 'pattern_counts': {'":': 0, '<': 67, '<?xml version=': 0, '>': 67, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 6, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '7024888', 'n_tokens_mistral': 1735, 'n_tokens_neox': 1676, 'n_words': 387}