DynamicArray grow/shrink broken for GPU arrays

[mathijs727]

Due to a bug in Memory::realloc_impl DynamicArray cannot grow/shrink if they point to GPU memory (either GPU_Malloc or GPU_Managed).
Grow/shrinking calls the Memory::realloc_impl function which checks whether the allocation contains CPU memory or GPU accessible memory.
In case of the latter cuda_memcpy_impl is called with cudaMemcpyDefault as cudaMemcpyKind.

if (is_gpu_type(oldAlloc.type)) {
        ptr = malloc_impl(newSize, oldAlloc.name, oldAlloc.type);
        if (ptr)
            cuda_memcpy_impl(static_cast<uint8*>(ptr), static_cast<uint8*>(oldPtr), oldAlloc.size, cudaMemcpyDefault);
        free_impl(oldPtr);
    } else {
        ...
    }

However, cuda_memcpy_impl does not support cudaMemcpyDefault and silently fails, causing the DynamicArray to contain junk values (whatever malloc returned):

void Memory::cuda_memcpy_impl(uint8* dst, const uint8* src, uint64 size, cudaMemcpyKind memcpyKind)
{
    const auto BlockCopy = [&]() {
        const double Start = Utils::seconds();
        CUDA_CHECKED_CALL cudaMemcpy(dst, src, size, memcpyKind);
        const double End = Utils::seconds();

        return size / double(1u << 30) / (End - Start);
    };

    if (memcpyKind == cudaMemcpyDeviceToDevice) {
        PROFILE_SCOPEF("Memcpy HtH %fMB", size / double(1u << 20));
        [[maybe_unused]] const double Bandwidth = BlockCopy();
        ZONE_METADATA("%fGB/s", Bandwidth);
    } else if (memcpyKind == cudaMemcpyDeviceToHost) {
        PROFILE_SCOPEF("Memcpy DtH %fMB", size / double(1u << 20));
        [[maybe_unused]] const double Bandwidth = BlockCopy();
        ZONE_METADATA("%fGB/s", Bandwidth);
    } else if (memcpyKind == cudaMemcpyHostToDevice) {
        PROFILE_SCOPEF("Memcpy HtD %fMB", size / double(1u << 20));
        [[maybe_unused]] const double Bandwidth = BlockCopy();
        ZONE_METADATA("%fGB/s", Bandwidth);
    }
}

This could be fixed by adding an extra else if statement to cuda_memcpy_impl to handle the cudaMemcpyDefault case.