[SYCLomatic] Block Load headers core #1640

abhilash1910 · 2024-01-18T12:23:55Z

In Progress PR for Load/Store header functions for Block API (related later to #1305 )
cc @yihanwg @danhoeflinger @mmichel11

clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h

mmichel11 · 2024-02-01T19:19:16Z

clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h

+  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)]);


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.

mmichel11 · 2024-02-02T18:59:11Z

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
/ routine in a kernel. If they are implemented as classes, then an object can be passed to the function where it could then be used:

load_striped<...> obj;
my_func(obj);

However, if they are implemented as free functions, then they would need to be passed as function pointers:

auto func_obj = &load_striped<...>;
my_func(func_obj);

which is not supported in the SYCL programming model. To make this work, load_striped would have to be called directly from my_func which I think may be difficult from the migration perspective.

What I think we should do is have two classes work_group_load and sub_group_load and then have some enum type dictating the strategy (ex: striped, blocked, etc). We could then have have specializations of this class to implement each strategy and maybe have some helpers to reduce code duplication.

clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h

yihanwg

LGTM

mmichel11 · 2024-04-30T02:38:12Z

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

zhimingwang36 · 2024-05-06T03:17:29Z

clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h

+}
+
+template <size_t ITEMS_PER_WORK_ITEM, load_algorithm ALGORITHM, typename InputT,
+          typename InputIteratorT, typename Item>


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.

thanks will add comments to clarify.

clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h

zhimingwang36 · 2024-05-07T08:48:13Z

clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h

+// 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> .


can we rename to nd_item? (what is the diff b/w Item and sycl::nd_item)

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 <[email protected]>

abhilash1910 · 2024-05-08T17:09:08Z

This PR is now tested with oneapi-src/SYCLomatic-test#619 and all tests are passing.
Tested on application side as well, working as expected.

danhoeflinger

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.

mmichel11

LGTM

load_helpers

bde928e

abhilash1910 requested a review from a team as a code owner January 18, 2024 12:23

abhilash1910 requested review from tomflinda and AndyCHHuang January 18, 2024 12:23

abhilash1910 marked this pull request as draft January 18, 2024 12:24

add load apis

8717079

abhilash1910 changed the title ~~[SYCLomatic] Block Load/Store headers core~~ [SYCLomatic] Block Load headers core Jan 29, 2024

clang-format

1549642

abhilash1910 marked this pull request as ready for review January 29, 2024 15:35

yihanwg suggested changes Jan 30, 2024

View reviewed changes

clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h 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

abhilash1910 added 5 commits January 30, 2024 09:06

reviews 1

41d994e

fix lit

99df7d9

update warp striped logic

789bd18

rename warp to subgroup

7872841

clang-format

ec7a718

mmichel11 reviewed Jan 30, 2024

View reviewed changes

abhilash1910 added 4 commits January 31, 2024 13:47

review commits 2

fe4c38e

fix lit

a4e2316

Merge branch 'oneapi-src:SYCLomatic' into block_load_store_headers

1a0b447

Merge branch 'oneapi-src:SYCLomatic' into block_load_store_headers

c5e4fad

abhilash1910 mentioned this pull request Feb 1, 2024

[SYClomatic-test] Add help function for Block Load header oneapi-src/SYCLomatic-test#619

Merged

clang-format

95edd0e

mmichel11 reviewed Feb 1, 2024

View reviewed changes

danhoeflinger reviewed Feb 1, 2024

View reviewed changes

mmichel11 reviewed Feb 1, 2024

View reviewed changes

abhilash1910 added 3 commits February 2, 2024 11:32

update review 1

1419253

review commit 2

6f99026

fix dtype

eb5539a

Merge branch 'oneapi-src:SYCLomatic' into block_load_store_headers

8ce2b68

yihanwg reviewed Apr 7, 2024

View reviewed changes

clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h Outdated Show resolved Hide resolved

use size_t

fdc2f2f

yihanwg approved these changes Apr 29, 2024

View reviewed changes

fix issues related to tests 619

95db67f

danhoeflinger reviewed Apr 30, 2024

View reviewed changes

clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h Outdated Show resolved Hide resolved

abhilash1910 added 2 commits April 30, 2024 18:39

remove ALGORITHM parameter

141ace7

format fix

540db29

mmichel11 reviewed May 1, 2024

View reviewed changes

clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h Outdated Show resolved Hide resolved

use local_range

ebf6237

zhimingwang36 reviewed May 6, 2024

View reviewed changes

abhilash1910 added 3 commits May 6, 2024 09:02

add comments

7f9d4e6

format

cb87b67

format

71d4047

zhimingwang36 reviewed May 7, 2024

View reviewed changes

clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h Outdated Show resolved Hide resolved

zhimingwang36 reviewed May 7, 2024

View reviewed changes

clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h Outdated Show resolved Hide resolved

zhimingwang36 reviewed May 7, 2024

View reviewed changes

clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h Outdated Show resolved Hide resolved

zhimingwang36 reviewed May 7, 2024

View reviewed changes

abhilash1910 and others added 6 commits May 7, 2024 14:44

Update clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h

bd24713

Co-authored-by: Wang, Zhiming <[email protected]>

Update clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h

26a3ae2

Co-authored-by: Wang, Zhiming <[email protected]>

Update clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h

6222566

Co-authored-by: Wang, Zhiming <[email protected]>

Merge branch 'oneapi-src:SYCLomatic' into block_load_store_headers

236c7ba

format

89cf7d3

fix issue in referencing

c4f0ca0

Merge branch 'oneapi-src:SYCLomatic' into block_load_store_headers

ec4f8ae

danhoeflinger approved these changes May 8, 2024

View reviewed changes

mmichel11 approved these changes May 9, 2024

View reviewed changes

Update dpcpp_extensions.h

d382f60

zhimingwang36 merged commit 01f2906 into oneapi-src:SYCLomatic May 9, 2024

[SYCLomatic] Block Load headers core #1640

[SYCLomatic] Block Load headers core #1640

Uh oh!

Conversation

abhilash1910 commented Jan 18, 2024

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

mmichel11 Feb 1, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

mmichel11 commented Feb 2, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

yihanwg left a comment

Choose a reason for hiding this comment

Uh oh!

mmichel11 commented Apr 30, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

Uh oh!

zhimingwang36 May 6, 2024

Choose a reason for hiding this comment

Uh oh!

abhilash1910 May 6, 2024

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

zhimingwang36 May 7, 2024

Choose a reason for hiding this comment

Uh oh!

abhilash1910 May 7, 2024

Choose a reason for hiding this comment

Uh oh!

abhilash1910 commented May 8, 2024

Uh oh!

danhoeflinger left a comment

Choose a reason for hiding this comment

Uh oh!

mmichel11 left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

mmichel11 Feb 1, 2024 •

edited

Loading

mmichel11 commented Feb 2, 2024 •

edited

Loading

mmichel11 commented Apr 30, 2024 •

edited

Loading