Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Single-pass scan kernel template #1320

Merged
merged 78 commits into from
Apr 24, 2024
Merged

Single-pass scan kernel template #1320

Show file tree
Hide file tree
Changes from 76 commits
Commits
Show all changes
78 commits
Select commit Hold shift + click to select a range
58c2639
Start of single-pass scan kernel template
adamfidel Aug 18, 2023
e0a676d
Fix hang in inclusive scan
adamfidel Aug 24, 2023
956f139
Debug statements for scan kernel template
adamfidel Aug 31, 2023
deb92cb
Update scan kernel template test
adamfidel Sep 6, 2023
fd0af78
Merge remote-tracking branch 'antcarcomp01/dev/adamfidel/scan_kt' int…
adamfidel Sep 13, 2023
590b1c0
Only have a single work-item per group query for previous tile status
adamfidel Sep 14, 2023
5f4069a
First attempt at parallel lookback
adamfidel Sep 18, 2023
6d4aa3d
Working cooperative lookback
adamfidel Sep 22, 2023
7e32a6f
Fix correctness issue with non-power-of-2 sizes
adamfidel Oct 25, 2023
1fcdba4
Code cleanup and bug fixes
adamfidel Nov 17, 2023
db77c7d
Remove accidental debug statement
adamfidel Nov 17, 2023
0228724
Test with floats
adamfidel Nov 27, 2023
b911bc0
Fix incorrect values for size_t scan
adamfidel Dec 13, 2023
8dacc2f
Add support for sizes > 2^30
adamfidel Dec 18, 2023
69b5afe
Uglify
adamfidel Dec 19, 2023
d88abb1
Remove debug code
adamfidel Dec 19, 2023
53e17b1
Merge branch 'main' into dev/adamfidel/scan_kt
adamfidel Jan 8, 2024
ce2bb8a
Move single-pass scan test to same place as ESIMD radix sort test
adamfidel Jan 8, 2024
d67d579
Move kernel template to other kernel template directory
adamfidel Jan 8, 2024
6d7929a
Restructure scan KT tests to be similar to sort KT tests
adamfidel Jan 8, 2024
4b01e64
Delete old test code
adamfidel Jan 8, 2024
6e97ed2
clang-format
adamfidel Jan 8, 2024
063ac92
Re-arrange assert for 64-bit atomics
adamfidel Jan 8, 2024
ec5a47c
Rename entry from single_pass_inclusive_scan -> inclusive_scan
adamfidel Jan 11, 2024
217a848
Rewrite kernel submission of init kernel to use provided kernel name
adamfidel Jan 11, 2024
28ab915
Rewrite main kernel to use user-provided kernel name
adamfidel Jan 11, 2024
d48e6d6
Move input values to local memory to process them there
adamfidel Jan 12, 2024
c2a6686
Merge remote-tracking branch 'github/dev/adamfidel/scan_kt' into dev/…
adamfidel Jan 12, 2024
9cf1ec1
fix issues caused by merge
adamfidel Jan 12, 2024
9d621a9
Address various pull request comments
adamfidel Jan 12, 2024
605cc42
KT scan tests generation through CMakeLists.txt (#1351)
SergeyKopienko Jan 17, 2024
c747f75
Remove unnecessary chunking logic
adamfidel Jan 17, 2024
b5136ce
Merge remote-tracking branch 'github/dev/adamfidel/scan_kt' into dev/…
adamfidel Jan 17, 2024
96c25ef
Separate out flags and value to support signed integers
adamfidel Jan 17, 2024
5a7626e
Use acq_rel instead of seq_cst
adamfidel Jan 17, 2024
de2061b
Fix duplicate kernel name errors
adamfidel Jan 23, 2024
d78db9c
Fix tests using oneapi::dpl::begin/end
adamfidel Jan 24, 2024
d4f019f
Fill padding with identity values to fix errors in test
adamfidel Jan 29, 2024
64edea3
Merge remote-tracking branch 'github/main' into dev/adamfidel/scan_kt
adamfidel Feb 1, 2024
e272b18
Assert that malloc_device returned non-null ptrs
adamfidel Feb 16, 2024
cdbec2d
Merge remote-tracking branch 'github/main' into dev/adamfidel/scan_kt
adamfidel Feb 16, 2024
3e57dbe
Pass ranges by forward ref
adamfidel Feb 16, 2024
e55c491
Use sycl::reqd_sub_group_size instead of intel::reqd_sub_group_size
adamfidel Feb 20, 2024
b944c2b
Replace raw unroll pragma with _ONEDPL_PRAGMA_UNROLL macro
adamfidel Feb 26, 2024
49dea3f
Call single-group scan if possible
adamfidel Feb 28, 2024
5262700
Change namespace from kt::igpu to kt::gpu
adamfidel Feb 29, 2024
e7c32cb
Use single malloc for values array
adamfidel Mar 13, 2024
5a6352f
Use only a single malloc if the value type is the same as the flag st…
adamfidel Mar 13, 2024
85c730c
Add missing free for status value array
adamfidel Mar 13, 2024
f135477
Fix include guard
adamfidel Mar 18, 2024
4e65b42
Rename _FlagType to _FlagStorageType
adamfidel Mar 18, 2024
f8f8508
Replace group.leader() with is first work-item in subgroup
adamfidel Mar 20, 2024
c0e22e3
Simplify device memory allocation into a single device_malloc
adamfidel Mar 22, 2024
404fcc2
Fix all_view test
adamfidel Apr 1, 2024
0e7a87c
Support passing sycl::buffer directly
adamfidel Apr 3, 2024
d6b60a6
Use __dpl_bit_ceil
adamfidel Apr 3, 2024
a899426
Promote types for sizes to 64 bits
adamfidel Apr 3, 2024
a8caee1
Merge commit '6a45be7078a636676b6a128a142e5d02213722ca' into dev/adam…
adamfidel Apr 9, 2024
41adacc
Add new device_backend_tag for call to single-group scan
adamfidel Apr 9, 2024
247bffb
Fix duplicate kernel name error with single-group scan
adamfidel Apr 9, 2024
80e6cd4
clang format
adamfidel Apr 9, 2024
0d0efeb
Fix alignment of status values
adamfidel Apr 17, 2024
b530eec
Rename [build/run]-igpu-tests to [build/run]-scan-kt-tests
adamfidel Apr 17, 2024
849db06
Remove template for queue type
adamfidel Apr 17, 2024
d076f11
Fix access mode for input
adamfidel Apr 17, 2024
eb6fb65
Add missing includes
adamfidel Apr 17, 2024
f1e0709
::std -> std:: and replacing assert with exception
adamfidel Apr 17, 2024
265ab8f
Fix number of elements vs number of bytes
adamfidel Apr 19, 2024
fe18dac
Improve test data generation, especially for multiplies
adamfidel Apr 19, 2024
e96c8e9
Merge remote-tracking branch 'github/dev/adamfidel/scan_kt' into dev/…
adamfidel Apr 19, 2024
0df3640
Fix CMake target without any constant params
adamfidel Apr 19, 2024
33c8982
Adding a few includes for completeness
adamfidel Apr 22, 2024
e431f01
Address PR comments
adamfidel Apr 23, 2024
6bcc32d
Merge remote-tracking branch 'github/main' into dev/adamfidel/scan_kt
adamfidel Apr 23, 2024
5f2fea6
Use new get_new_kernel_params function
adamfidel Apr 23, 2024
978e0fa
Correctly pass kernel name with optional_kernel_name
adamfidel Apr 23, 2024
804c8c7
clang-format + Moving around get_new_kernel_params
adamfidel Apr 23, 2024
7cd2d63
clang-format
adamfidel Apr 23, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 3 additions & 1 deletion include/oneapi/dpl/experimental/kernel_templates
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,7 +15,9 @@
#include "kt/kernel_param.h"

#if __has_include(<sycl/ext/intel/esimd.hpp>)
# include "kt/esimd_radix_sort.h"
# include "kt/esimd_radix_sort.h"
#endif

#include "kt/single_pass_scan.h"

#endif // _ONEDPL_KERNEL_TEMPLATES
Loading
Loading