intel / intel/llvm

[ESIMD] Discussion topics, possible enhancements

Open
#5,367 0 comments 5 reactions 0 assignees View on GitHub
confirmed enhancement
Dominant language
LLVM
Stars
1.5k
Forks
854
Avg merge
3d 17h
Merged PRs (30d)
137

Description

This is to record code review suggestions from various people.

1. Inter-relation with `std::simd`.

- @rolandschulz suggests __not__ to imeplement `std::simd` for Intel graphics based on
`sycl::intel::gpu`, but directly on top of clang vectors.
> Implementing it on top of `gpu::simd` is the strategy which requires a reimplementation
> of `std::simd` features missing from `gpu::simd`. On the other hand there is an existing
> implementation directly on top of gcc/clang vector: https://github.com/VcDevel/std-simd.
> We should explore how we can reuse most of that implementation also for the GPU.

- @mattkretz suggests to consider making `std::simd` good enough for GPU programming
instead of creating a highly specific extension.
> A SIMD type that doesn't allow you to reach full performance is a failure
> (or hopefully just buggy and needs to be fixed). Performance is the reason why we use it.

...

> I guess what I'd like to discuss is: If you deviate from the Parallelism TS 2, why?
> Is it to simplify your implementation, is it because of missing functionality, is it
> because of performance, or? This is important feedback to the C++ committee when merging
> the TS into the IS is discussed.

2. Invocation API
`Cgh.single_task([=] () EXPLICIT_SIMD`
- @Ruyk suggests to use different handler API to invoke ESIMD kernels rather
than using attributes to mark lambda/functor type:
> Well we could have a single "parallel_for" for everything but sometimes is
> better to distinguish entry points, specially if the programming model is very
> different. That macro in the lambda is easy to miss, specially if the kernel
> lambda is defined elsewhere.

...

> One of my concerns is that you could have the kernel separated from the
> single_task point of entry, and you may not even know what kind of kernel is
> that you are dispatching, just a SYCL kernel. Without a way for a C++
> developer to identify the usage of this different API inside, they will be
> reduced to random runtime errors if the wrong kernel is used in the wrong device.

- @keryell suggests to use kernel properties and avoid macros:
https://www.youtube.com/watch?v=Fp8DuVWesT4 and https://gitlab.khronos.org/sycl/Specification/issues/296

- @kbobrovs mentioned that there was another suggestion from @rolandschulz to
use some extra "executor" template parameter, which would be set to "ESIMD"
for ESIMD kernels.

3. Enabling this extension for other architectures, such as x86, with extracting
and clearly marking generic and target-dependent API portions

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.