Fix dpcpp memory issue and capture exception in raw_free

[yhmtsai]

This PR limits the direct memcpy between different backend or device, add free after kernel test and capture exception in raw_free

the memcpy issue between backend on the same device is on https://gitlab.com/ginkgo-project/ginkgo-public-ci/-/jobs/1413646244
the raw_free isssue is on https://gitlab.com/ginkgo-project/ginkgo-public-ci/-/jobs/1415200944

the codeplay_host_task does not work as expectation yet.
need to investigate furthermore and use the new api.
Thus, we stay wait workaround to ensure free is after kernel

[upsj]

LGTM!

[upsj]

We need to think about other ways to approach this at some point, since killing the process may not be desirable for application users.

[yhmtsai]

Probably, I am not sure whether the application allows the error situation is happened or not.

[upsj]

Is this necessary? The case we want to catch is the one without synchronization, isn't it?

[yhmtsai]

Not sure honestly, I do not know whether queue wait the all kernel finished when queue destroy.
If it is like cuda, just terminate. It is possible that we only have allocation and free but no execution.
I use curly brackets to contain the Array allocation to make sure the free is called before the synchronize.
Yes, it catches the free is not after the kernel finishes

[upsj]

When the array gets freed, we wait on the queue, that isn't sufficient?

[yhmtsai]

Yes, it is fixed version. when the free does not contain synchronization, we will probably not wait the kernel.

[upsj]

Can you remove this line to make sure the program can actually fail due to use-after-free issues?

[yhmtsai]

it's shown in the https://gitlab.com/ginkgo-project/ginkgo-public-ci/-/jobs/1415200944
and the code contains sync https://gitlab.com/ginkgo-project/ginkgo-public-ci/-/commit/f9da28bc5b8f6847b4d587b0283b2cc7f7e2f60a
I mean fixed version is the we wait on the queue when we call free.
(the following description is from cuda thinking)
Without sync before program end

1. when free does not contain wait:
   submit job (not run) -> free -> program end (the job may not run, so we may does not face the free after use)
2. when free does contain wait: (fixed version)
   submit job (not run) -> run job (due to free contain wait) -> free -> program end
   this case does not need sync in the end

[upsj]

I see, thanks! I completely missed that the arrays have their own scope, so this is fine to me

[codecov]

Codecov Report

Merging #832 (883dd5c) into develop (26ee87e) will decrease coverage by 0.04%.
The diff coverage is 0.00%.

@@             Coverage Diff             @@
##           develop     #832      +/-   ##
===========================================
- Coverage    94.29%   94.25%   -0.05%     
===========================================
  Files          408      408              
  Lines        32742    32756      +14     
===========================================
- Hits         30875    30874       -1     
- Misses        1867     1882      +15     

Impacted Files	Coverage Δ
core/test/base/executor.cpp	87.60% <0.00%> (-4.58%)	⬇️
include/ginkgo/core/base/executor.hpp	79.47% <0.00%> (-0.85%)	⬇️
include/ginkgo/core/base/math.hpp	100.00% <ø> (ø)
core/base/extended_float.hpp	91.26% <0.00%> (-0.98%)	⬇️

---

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update 26ee87e...883dd5c. Read the comment docs.

[yhmtsai]

After more test, the combination of device gives different result:

1. opencl:gpu vs opencl:cpu -> use opencl:gpu queue can get the opencl:cpu memory but reverse case can not. (opencl:cpu is like normal host memory)
2. opencl:gpu vs opencl:gpu -> no matter whether two queue is mapped to the same device, they can not transfer data to each other
3. level_zero:gpu vs opencl:cpu -> similar to 1.
4. level_zero:gpu vs level_zero:gpu -> if they are on the same device, they can transfer the memory to each other and use the memory from another one.
5. opencl:cpu vs opencl:cpu -> can use the memory on different queue and transfer memory

host option does not work in filter_selector (ref), so I need to check it.

I only find one way to get the backend, which using dynamic_cast (with backend type) to type provided by level_zero.
queue has get_backend() returns a enum value
Or, we always use host memory to transfer data when the queues are different (for handling opencl cases)

I put the code here. we can use it check.
Compile: dpcpp <file>
Run: ./a.out <first_device_filter> <second_device_filter>

#include <CL/sycl.hpp>
#include <iostream>

int main(int argc, char *argv[])
{
    if (argc != 3) {
        std::cout << "Usage: " << argv[0] << " first_filter second_filter"
                  << std::endl;
        std::exit(1);
    }
    auto first_filter = sycl::ONEAPI::filter_selector(argv[1]);
    auto second_filter = sycl::ONEAPI::filter_selector(argv[2]);
    auto *first_queue =
        new sycl::queue{first_filter, sycl::property::queue::in_order{}};
    auto *second_queue =
        new sycl::queue{second_filter, sycl::property::queue::in_order{}};
    std::cout << "First queue device name: "
              << first_queue->get_device().get_info<sycl::info::device::name>()
              << std::endl;
    std::cout << "Second queue device name: "
              << second_queue->get_device().get_info<sycl::info::device::name>()
              << std::endl;
    const int n = 32;

    float data[n];
    for (int i = 0; i < n; i++) {
        data[i] = static_cast<float>(i) / n;
    }
    float *first_data = sycl::malloc_device<float>(n, *first_queue);
    float *first_data2 = sycl::malloc_device<float>(n, *first_queue);
    float *second_data = sycl::malloc_device<float>(n, *second_queue);
    float *second_data2 = sycl::malloc_device<float>(n, *second_queue);
    std::cout << "First queue test:" << std::endl;
    std::cout << "From Host to Device" << std::endl;
    first_queue->memcpy(first_data, data, n * sizeof(float)).wait();
    std::cout << "From Device to Device (self-queue)" << std::endl;
    first_queue->memcpy(first_data2, first_data, n * sizeof(float)).wait();
    std::cout << "From Device to Host" << std::endl;
    first_queue->memcpy(data, first_data2, n * sizeof(float)).wait();
    first_queue->wait_and_throw();
    std::cout << "Second queue test:" << std::endl;
    std::cout << "From Host to Device" << std::endl;
    second_queue->memcpy(second_data, data, n * sizeof(float)).wait();
    std::cout << "From Device to Device (self-queue)" << std::endl;
    second_queue->memcpy(second_data2, second_data, n * sizeof(float)).wait();
    std::cout << "From Device to Host" << std::endl;
    second_queue->memcpy(data, second_data2, n * sizeof(float)).wait();
    second_queue->wait_and_throw();
    bool same_device =
        (first_queue->get_device() == second_queue->get_device());
    bool same_context =
        (first_queue->get_context() == second_queue->get_context());
    std::cout << "device check: " << same_device << std::endl;
    std::cout << "context check: " << same_context << std::endl;
    if (same_device) {
        std::cout << "Access the data from different queue on the same device"
                  << std::endl;
        std::cout << "Access the second queue data from first queue"
                  << std::endl;
        first_queue->memcpy(second_data2, second_data, n * sizeof(float))
            .wait();
        std::cout << "Access the first queue data from second queue"
                  << std::endl;
        second_queue->memcpy(first_data2, first_data, n * sizeof(float)).wait();
    }
    std::cout << "Communication between first/second queue:" << std::endl;
    std::cout << "From second to first (handled by first queue)" << std::endl;
    first_queue->memcpy(first_data, second_data, n * sizeof(float)).wait();
    std::cout << "From fisrt to second (handled by second queue)" << std::endl;
    second_queue->memcpy(second_data, first_data, n * sizeof(float)).wait();
    first_queue->wait_and_throw();
    second_queue->wait_and_throw();
    sycl::free(first_data, first_queue->get_context());
    sycl::free(first_data2, first_queue->get_context());
    sycl::free(second_data, second_queue->get_context());
    sycl::free(second_data2, second_queue->get_context());
    return 0;
}

update the accessible check

[upsj]

A round-trip via host memory sounds sensible, I mean we do the same thing for ROCm <-> CUDA copies.

[yhmtsai]

I forgot to update level_zero information.
level_zero backend allows to transfer data on the same device, but opencl cannot.
get_device() are the same but get_context() are different

[tcojean]

LGTM. Some small comments.

[tcojean]

I think we should do a copy through the host here like in the ROCM <-> NVIDIA cases. We should also add a debug level message to state that this might not be optimized. I expect that later direct GPU copies will be supported?

[yhmtsai]

should we add the message?
for the GPU direct copies, it seems to use context to take care (identify) the memory.
If I pass the context to another one, they can communicate without segmentation fault.
but I can also pass the opencl context to level_zero queue and then they communicate without segfault, which is weird.
context should be under the platform but opencl and level zero should be the different platforms

[yhmtsai]

I use #if (GKO_VERBOSE_LEVEL >= 1) && !defined(NDEBUG) to enable output clog

[tcojean]

To make it work as you say, we would need a global context that is used by all executors? Or can you have one context per device and use the correct context so that the copy works?

[yhmtsai]

I think it should be a context per platform.
but I need to check the level_zero/opencl context behavior
two queues with the same context and then they can communicate

[yhmtsai]

rebase!

[sonarcloud]

Kudos, SonarCloud Quality Gate passed!

0 Bugs
0 Vulnerabilities
0 Security Hotspots
1 Code Smell

No Coverage information
0.0% Duplication