Ejemplo n.º 1
0
void weightedMeanAllTest(af::dim4 dims)
{
    typedef typename meanOutType<T>::type outType;

    if (noDoubleTests<T>()) return;
    if (noDoubleTests<outType>()) return;
    if (noDoubleTests<wtsType>()) return;

    using af::array;
    using af::mean;

    std::srand(std::time(0));

    vector<T> data(dims.elements());
    vector<wtsType> wts(dims.elements());
    std::generate(data.begin(), data.end(), random<T>);
    std::generate(wts.begin(), wts.end(), random<wtsType>);

    outType wtdSum = outType(0);
    wtsType wtsSum = wtsType(0);

    for(int i = 0; i < (int)data.size(); i++) {
        wtdSum = wtdSum + data[i]*wts[i];
        wtsSum = wtsSum + wts[i];
    }

    outType gold = wtdSum / wtsSum;

    array a(dims, &(data.front()));
    array w(dims, &(wts.front()));
    outType output = mean<outType>(a, w);

    ASSERT_NEAR(::real(output), ::real(gold), 1.0e-2);
    ASSERT_NEAR(::imag(output), ::imag(gold), 1.0e-2);
}
Ejemplo n.º 2
0
    Array<T>::Array(af::dim4 dims, const T * const in_data, bool is_device) :
        info(getActiveDeviceId(), dims, af::dim4(0,0,0,0), calcStrides(dims), (af_dtype)dtype_traits<T>::af_type),
        data((is_device ? (T *)in_data : memAlloc<T>(dims.elements())), memFree<T>),
        data_dims(dims),
        node(), offset(0), ready(true), owner(true)
    {
#if __cplusplus > 199711L
        static_assert(std::is_standard_layout<Array<T>>::value, "Array<T> must be a standard layout type");
        static_assert(offsetof(Array<T>, info) == 0, "Array<T>::info must be the first member variable of Array<T>");
#endif
        if (!is_device) {
            CUDA_CHECK(cudaMemcpy(data.get(), in_data, dims.elements() * sizeof(T), cudaMemcpyHostToDevice));
        }
    }
Ejemplo n.º 3
0
    static Array<T> diff(const Array<T> &in, const int dim)
    {
        const af::dim4 iDims = in.dims();
        af::dim4 oDims = iDims;
        oDims[dim] -= (isDiff2 + 1);

        if(iDims.elements() == 0 || oDims.elements() == 0) {
            throw std::runtime_error("Elements are 0");
        }

        Array<T> out = createEmptyArray<T>(oDims);

        switch (dim) {

            case (0):    kernel::diff<T, 0, isDiff2>(out, in, in.ndims());
                         break;

            case (1):    kernel::diff<T, 1, isDiff2>(out, in, in.ndims());
                         break;

            case (2):    kernel::diff<T, 2, isDiff2>(out, in, in.ndims());
                         break;

            case (3):    kernel::diff<T, 3, isDiff2>(out, in, in.ndims());
                         break;
        }

        return out;
    }
Ejemplo n.º 4
0
af_err af_tile(af_array *out, const af_array in, const af::dim4 &tileDims)
{
    try {
        ArrayInfo info = getInfo(in);
        af_dtype type = info.getType();

        if(info.ndims() == 0) {
            return af_retain_array(out, in);
        }
        DIM_ASSERT(1, info.dims().elements() > 0);
        DIM_ASSERT(2, tileDims.elements() > 0);

        af_array output;

        switch(type) {
            case f32: output = tile<float  >(in, tileDims);  break;
            case c32: output = tile<cfloat >(in, tileDims);  break;
            case f64: output = tile<double >(in, tileDims);  break;
            case c64: output = tile<cdouble>(in, tileDims);  break;
            case b8:  output = tile<char   >(in, tileDims);  break;
            case s32: output = tile<int    >(in, tileDims);  break;
            case u32: output = tile<uint   >(in, tileDims);  break;
            case s64: output = tile<intl   >(in, tileDims);  break;
            case u64: output = tile<uintl  >(in, tileDims);  break;
            case s16: output = tile<short  >(in, tileDims);  break;
            case u16: output = tile<ushort >(in, tileDims);  break;
            case u8:  output = tile<uchar  >(in, tileDims);  break;
            default:  TYPE_ERROR(1, type);
        }
        std::swap(*out,output);
    }
    CATCHALL;

    return AF_SUCCESS;
}
Ejemplo n.º 5
0
    Array<T> resize(const Array<T> &in, const dim_type odim0, const dim_type odim1,
                    const af_interp_type method)
    {
        if ((std::is_same<T, double>::value || std::is_same<T, cdouble>::value) &&
            !isDoubleSupported(getActiveDeviceId())) {
            OPENCL_NOT_SUPPORTED();
        }
        const af::dim4 iDims = in.dims();
        af::dim4 oDims(odim0, odim1, iDims[2], iDims[3]);

        if(iDims.elements() == 0 || oDims.elements() == 0) {
            throw std::runtime_error("Elements is 0");
        }

        Array<T> out = createEmptyArray<T>(oDims);

        switch(method) {
            case AF_INTERP_NEAREST:
                kernel::resize<T, AF_INTERP_NEAREST> (out, in);
                break;
            case AF_INTERP_BILINEAR:
                kernel::resize<T, AF_INTERP_BILINEAR>(out, in);
                break;
            default:
                break;
        }
        return out;
    }
Ejemplo n.º 6
0
    Array<T>::Array(af::dim4 dims, const T * const in_data, bool is_device, bool copy_device) :
        info(getActiveDeviceId(), dims, 0, calcStrides(dims), (af_dtype)dtype_traits<T>::af_type),
        data(((is_device & !copy_device) ? const_cast<T*>(in_data) : memAlloc<T>(dims.elements()).release()), memFree<T>),
        data_dims(dims),
        node(bufferNodePtr<T>()), ready(true), owner(true)
    {
#if __cplusplus > 199711L
        static_assert(std::is_standard_layout<Array<T>>::value, "Array<T> must be a standard layout type");
        static_assert(offsetof(Array<T>, info) == 0, "Array<T>::info must be the first member variable of Array<T>");
#endif
        if (!is_device) {
            CUDA_CHECK(cudaMemcpyAsync(data.get(), in_data, dims.elements() * sizeof(T),
                                       cudaMemcpyHostToDevice, cuda::getActiveStream()));
            CUDA_CHECK(cudaStreamSynchronize(cuda::getActiveStream()));
        } else if (copy_device) {
            CUDA_CHECK(cudaMemcpyAsync(data.get(), in_data, dims.elements() * sizeof(T),
                                       cudaMemcpyDeviceToDevice, cuda::getActiveStream()));
            CUDA_CHECK(cudaStreamSynchronize(cuda::getActiveStream()));
        }
    }
Ejemplo n.º 7
0
void testCPPMean(T const_value, af::dim4 dims)
{
    typedef typename meanOutType<T>::type outType;
    if (noDoubleTests<T>()) return;

    using af::array;
    using af::mean;

    vector<T> hundred(dims.elements(), const_value);

    outType gold = outType(0);
    //for(auto i:hundred) gold += i;
    for(int i = 0; i < (int)hundred.size(); i++) {
        gold += hundred[i];
    }
    gold /= dims.elements();

    array a(dims, &(hundred.front()));
    outType output = mean<outType>(a);

    ASSERT_NEAR(::real(output), ::real(gold), 1.0e-3);
    ASSERT_NEAR(::imag(output), ::imag(gold), 1.0e-3);
}
Ejemplo n.º 8
0
    Array<T> tile(const Array<T> &in, const af::dim4 &tileDims)
    {
        const af::dim4 iDims = in.dims();
        af::dim4 oDims = iDims;
        oDims *= tileDims;

        if(iDims.elements() == 0 || oDims.elements() == 0) {
            throw std::runtime_error("Elements are 0");
        }

        Array<T> out = createEmptyArray<T>(oDims);

        T* outPtr = out.get();
        const T* inPtr = in.get();

        const af::dim4 ist = in.strides();
        const af::dim4 ost = out.strides();

        for(dim_t ow = 0; ow < oDims[3]; ow++) {
            const dim_t iw = ow % iDims[3];
            const dim_t iW = iw * ist[3];
            const dim_t oW = ow * ost[3];
            for(dim_t oz = 0; oz < oDims[2]; oz++) {
                const dim_t iz = oz % iDims[2];
                const dim_t iZW = iW + iz * ist[2];
                const dim_t oZW = oW + oz * ost[2];
                for(dim_t oy = 0; oy < oDims[1]; oy++) {
                    const dim_t iy = oy % iDims[1];
                    const dim_t iYZW = iZW + iy * ist[1];
                    const dim_t oYZW = oZW + oy * ost[1];
                    for(dim_t ox = 0; ox < oDims[0]; ox++) {
                        const dim_t ix = ox % iDims[0];
                        const dim_t iMem = iYZW + ix;
                        const dim_t oMem = oYZW + ox;
                        outPtr[oMem] = inPtr[iMem];
                    }
                }
            }
        }

        return out;
    }
Ejemplo n.º 9
0
void testCPPVar(T const_value, af::dim4 dims)
{
    typedef typename varOutType<T>::type outType;
    if (noDoubleTests<T>()) return;
    if (noDoubleTests<outType>()) return;

    using af::array;
    using af::var;

    vector<T> hundred(dims.elements(), const_value);

    outType gold = outType(0);

    array a(dims, &(hundred.front()));
    outType output = var<outType>(a, false);

    ASSERT_NEAR(::real(output), ::real(gold), 1.0e-3);
    ASSERT_NEAR(::imag(output), ::imag(gold), 1.0e-3);

    output = var<outType>(a, true);

    ASSERT_NEAR(::real(output), ::real(gold), 1.0e-3);
    ASSERT_NEAR(::imag(output), ::imag(gold), 1.0e-3);

    gold = outType(2.5);
    outType tmp[] = { outType(0), outType(1), outType(2), outType(3),
        outType(4) };
    array b(5, tmp);
    output = var<outType>(b, false);

    ASSERT_NEAR(::real(output), ::real(gold), 1.0e-3);
    ASSERT_NEAR(::imag(output), ::imag(gold), 1.0e-3);

    gold = outType(2);
    output = var<outType>(b, true);

    ASSERT_NEAR(::real(output), ::real(gold), 1.0e-3);
    ASSERT_NEAR(::imag(output), ::imag(gold), 1.0e-3);
}
Ejemplo n.º 10
0
 size_t elements() const             { return dim_size.elements();   }
Ejemplo n.º 11
0
 Array<T>::Array(af::dim4 dims) :
     info(getActiveDeviceId(), dims, 0, calcStrides(dims), (af_dtype)dtype_traits<T>::af_type),
     data((dims.elements() ? memAlloc<T>(dims.elements()).release() : nullptr), memFree<T>), data_dims(dims),
     node(bufferNodePtr<T>()), ready(true), owner(true)
 {}
Ejemplo n.º 12
0
 Array<T>::Array(af::dim4 dims) :
     info(getActiveDeviceId(), dims, af::dim4(0,0,0,0), calcStrides(dims), (af_dtype)dtype_traits<T>::af_type),
     data(memAlloc<T>(dims.elements()), memFree<T>), data_dims(dims),
     node(), offset(0), ready(true), owner(true)
 {}