[SYCLomatic] Block Load headers core#1640
[SYCLomatic] Block Load headers core#1640zhimingwang36 merged 76 commits intooneapi-src:SYCLomaticfrom
Conversation
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
| for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++) { | ||
| new (&items[ITEM]) | ||
| InputT(block_itr[item.get_sub_group().get_local_range()[0] + | ||
| linear_tid + (ITEM * ITEMS_PER_THREAD)]); |
There was a problem hiding this comment.
A similar comment to above with regards to the stride. In this case, I believe instead of ITEMS_PER_THREAD it should be the subgroup size.
|
Thinking about the design in more depth, I do not think these APIs should be implemented as free functions. Rather, I think they should be implemented as classes which overload the function call operator or have a member function that invokes the load operation. My reasoning for this is that certain applications during migration may need to pass a "load" object to a separate function However, if they are implemented as free functions, then they would need to be passed as function pointers: which is not supported in the SYCL programming model. To make this work, What I think we should do is have two classes |
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
|
This PR looks to be in good shape for me. However, I think the outstanding comments in oneapi-src/SYCLomatic-test#619 need to be addressed, so we can have confidence in the implementation's correctness. |
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
| } | ||
|
|
||
| template <size_t ITEMS_PER_WORK_ITEM, load_algorithm ALGORITHM, typename InputT, | ||
| typename InputIteratorT, typename Item> |
There was a problem hiding this comment.
here are quite many item usage, ITEM/Item/iterms/ITERMS_PER_WORK_ITEM, maybe it should give a better name for "typename Item", it is a little confusing. or please add comments.
There was a problem hiding this comment.
thanks will add comments to clarify.
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h
Outdated
Show resolved
Hide resolved
| // InputT: typename parameter controlled at runtime from input sequence. | ||
| // InputIteratorT: typename parameter for iterator pointer controlled at | ||
| // runtime. | ||
| // Item : typename parameter resembling sycl::nd_item<3> . |
There was a problem hiding this comment.
can we rename to nd_item? (what is the diff b/w Item and sycl::nd_item)
There was a problem hiding this comment.
Yes it is same thing, but to maintain the same variable name across the file (similar for other apis), Item is used.
Co-authored-by: Wang, Zhiming <zhiming.wang@intel.com>
Co-authored-by: Wang, Zhiming <zhiming.wang@intel.com>
Co-authored-by: Wang, Zhiming <zhiming.wang@intel.com>
|
This PR is now tested with oneapi-src/SYCLomatic-test#619 and all tests are passing. |
danhoeflinger
left a comment
There was a problem hiding this comment.
LGTM,
There are some changes needed in the testing PR still, but with some local changes to those tests and passing runs, I'm confident enough in this PR now to approve.
In Progress PR for Load/Store header functions for Block API (related later to #1305 )
cc @yihanwg @danhoeflinger @mmichel11