Skip to content

Add sycl_khr_free_function_commands extension #644

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

Draft
Pennycook wants to merge 59 commits into
base: main
Choose a base branch
from
Draft

Add sycl_khr_free_function_commands extension #644

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
59 commits
Select commit Hold shift + click to select a range
47863fa
Add sycl_khr_free_function_commands extension
Pennycook Oct 17, 2024
ebe4dc0
Merge branch 'main' into khr_free_function_commands
Pennycook Dec 2, 2024
ce08652
Reword khr_free_function_commands comment
Pennycook Dec 5, 2024
372bb3b
Add periods to khr_free_function_commands comments
Pennycook Dec 5, 2024
9747f7a
Add + marks to code blocks containing ...
Pennycook Dec 6, 2024
2527e90
Require at least 1 reduction in *_reduce functions
Pennycook Dec 6, 2024
f63adeb
Replace "must be" with "is" in constraints
Pennycook Dec 6, 2024
47c08f8
Use bulleted list for multiple constraints
Pennycook Dec 6, 2024
15fd80a
Rewrite preconditions for USM copy functions
Pennycook Dec 6, 2024
9c62792
Fix typo in non-normative note
Pennycook Dec 6, 2024
4377549
Define kernel object overloads via equivalence
Pennycook Dec 6, 2024
c24fb13
Clarify dependencies for command_/event_barrier
Pennycook Dec 6, 2024
664a912
Clarify that event_barrier can be a no-op
Pennycook Dec 6, 2024
16111b2
Add missing invocation constructor
Pennycook Dec 6, 2024
88b540a
Restart numbering at 1 in each synopsis block
Pennycook Dec 6, 2024
2575901
Replace backticks with [code] environment
Pennycook Dec 6, 2024
8fa1ce2
Add + marks to code blocks containing ... again
Pennycook Dec 6, 2024
8d79af4
Fix grammar: "is" to "are"
Pennycook Dec 6, 2024
2ba2394
Fix formatting of bulleted lists
Pennycook Dec 6, 2024
e382dbc
Fix more instances of "is" that should be "are"
Pennycook Dec 6, 2024
a9bdc10
Remove unnecessary device-copyable constraint
Pennycook Dec 6, 2024
269706c
Remove khr::invocation from free_function_commands
Pennycook Dec 6, 2024
fa8a8f6
Remove empty issues section
Pennycook Dec 9, 2024
db380b4
Add missing constraints to fill overloads
Pennycook Dec 9, 2024
d26831a
Remove *_reduce functions for kernel objects
Pennycook Dec 12, 2024
75867f7
Fix copy-paste error in launch_task definition
Pennycook Jan 17, 2025
151f632
Remove unnecessary "is"
Pennycook Jan 20, 2025
32d11f5
Explain potential performance overhead of events
Pennycook Jan 20, 2025
165d07e
Add no-op note to command_barrier
Pennycook Jan 20, 2025
be56bb7
Weaken note about no-op from "is" to "may be"
Pennycook Jan 20, 2025
9897e6d
Fix copy-paste error in khr::copy
Pennycook Jan 22, 2025
b756d9e
Change KHR names to lower case
Pennycook Jan 23, 2025
e68a7c0
Remove sycl:: in free-function-command synopses
Pennycook Jan 24, 2025
96c101e
Add missing require() calls from queue overloads
Pennycook Jan 24, 2025
b279f37
Fix alignment of overload numbers
Pennycook Jan 24, 2025
d7ce65f
Fix parameter pack syntax
Pennycook Jan 24, 2025
19440ec
Add missing ptr parameter to memset
Pennycook Jan 28, 2025
c63e287
Use const queue& for khr_free_function_commands
Pennycook Apr 25, 2025
b314b77
Fix alignment of (1), (2) in khr_free_functions
Pennycook Apr 25, 2025
f1c9607
Add new line between extensions
Pennycook Apr 29, 2025
caf3b0a
Forbid calling anything after a free function
Pennycook May 19, 2025
ae40379
Add comments with references to old APIs
Pennycook May 27, 2025
a9289fe
Revert "Forbid calling anything after a free function"
Pennycook Jun 3, 2025
1e3737b
Remove all handler overloads
Pennycook Jun 3, 2025
f9ca360
Add khr::requirements class and overloads
Pennycook Jun 4, 2025
3057215
Allow kernel_bundle as a requirement
Pennycook Jun 13, 2025
3f66bb4
Use "status" instead of "state" for events
Pennycook Jun 17, 2025
bbd090e
Forbid requirements with multiple tracking objects
Pennycook Jun 17, 2025
5edf35d
Add blank line before bulleted list
Pennycook Jun 17, 2025
44ed92c
Expand Requirements... parameter pack correctly
Pennycook Jun 17, 2025
3073621
Support only non-deprecated accessors
Pennycook Jun 17, 2025
2afc9f0
Add exposition-only paragraph to function synopses
Pennycook Jun 17, 2025
f31095a
Add constraints to limit kernel bundles to kernels
Pennycook Jun 19, 2025
fcd2ca0
Add constraints to limit accessors to kernels
Pennycook Jun 19, 2025
1ee9698
Add constraints to limit accessor targets
Pennycook Jun 19, 2025
547c100
Use requirements in free function commands example
Pennycook Jul 18, 2025
896d71f
Merge branch 'main' into khr_free_function_commands
TApplencourt Sep 12, 2025
288047b
Testing
gmlueck Sep 23, 2025
a844ce8
Testing again
gmlueck Oct 8, 2025
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
1 change: 1 addition & 0 deletions adoc/chapters/architecture.adoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1901,6 +1901,7 @@ always matches the byte order of the devices.
This allows data to be copied between the host and the devices without any byte
swapping.

[[subsec:example.sycl.application]]
== Example SYCL application

Below is a more complex example application, combining some of the features
Expand Down
10 changes: 10 additions & 0 deletions adoc/config/rouge/lib/rouge/lexers/sycl.rb
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -436,6 +436,14 @@ class Sycl < Cpp
replace_me # Replace with list of actual keywords
)

# Exposition-only identifiers
sycl_exposition_only = %w(
register-events
register-accessors
register-kernel-bundle
has-tracking
)

# Here are some interesting tokens
# https://pygments.org/docs/tokens/ unused in C++ we can reuse
# in SYCL mode:
Expand Down Expand Up @@ -466,6 +474,8 @@ class Sycl < Cpp
# Insert some specific rules at the beginning of the statement
# rule of the C++ lexer
prepend :statements do
rule %r/(?:#{sycl_exposition_only.join('|')})\b/,
Generic::Emph
rule %r/(?:#{sycl_data_types.join('|')})\b/,
Keyword::Pseudo
rule %r/(?:#{sycl_functions.join('|')})\b/,
Expand Down
3 changes: 2 additions & 1 deletion adoc/config/rouge/lib/rouge/themes/sycl_spec.rb
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,7 +33,8 @@ class SYCLspec < Github
style Comment::Single, :fg => '#9acd32'
# Use a clearer white background
style Text, :bg => '#ffffff'

# Render exposition-only functions in italics to match ISO C++
style Generic::Emph, :fg => '#000000', :italic => true
Copy link
Contributor

@gmlueck gmlueck Jun 17, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The exposition-only function names don't seem to be rendered in italics in the PDF render.

Note the comment above saying that italics don't work in asciiidoctor-pdf. However, I'm not sure why this would be. Italics do seem to be working in general in the PDF render, so maybe it's some weird interaction with the rouge highlighter. I wonder if the real problem, though, is that we're just not installing the correct font.

Copy link
Contributor Author

@Pennycook Pennycook Jun 17, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, I think you're right. asciidoctor/asciidoctor-pdf#2350 (comment) suggests its because the default monospace font selected by asciidoctor for PDFs does not have an italic version.

I don't have any opinion at all about which fonts we should use. Is there a reason that we use different fonts for the HTML and PDF versions?

Copy link
Contributor

@gmlueck gmlueck Jun 17, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I haven't looked into it much. It seems like we use the default fonts mostly for both HTML and PDF. There's some tooling for KaTeX fonts in the HTML render, but I think this is related to the math package.

The asciidoc-pdf comment you linked above says to install a custom font if you want italic monospace. That might be worth exploring.

Copy link
Contributor Author

@Pennycook Pennycook Jun 17, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Okay, I spent some time exploring this.

Using a custom font is straightforward, and you just have to point AsciiDoctor to where the font lives via your theme file (which for us is config/themes/pdf-theme.yml). See https://docs.asciidoctor.org/pdf-converter/latest/theme/custom-font/ for full details. Unfortunately, getting things to work with the SYCL specification seems much harder than this.

I couldn't figure out how to pass extra arguments to our docker-based build system, and I also couldn't figure out whether we'd need/want to put the new fonts inside the docker image or if we'd have to import them somehow.

I tried to make progress just by replacing all instances of "M+ 1mn" with "Courier" (since that font is available and should support italics), but that leads to errors and warnings in the build:

asciidoctor: ERROR: the table cell on page 139 has been truncated; Asciidoctor PDF does not support table cell content that exceeds the height of a single page
asciidoctor: ERROR: the table cell on page 302 has been truncated; Asciidoctor PDF does not support table cell content that exceeds the height of a single page
asciidoctor: ERROR: the table cell on page 408 has been truncated; Asciidoctor PDF does not support table cell content that exceeds the height of a single page
asciidoctor: ERROR: the table cell on page 412 has been truncated; Asciidoctor PDF does not support table cell content that exceeds the height of a single page
asciidoctor: WARNING: The following text could not be fully converted to the Windows-1252 character set:
| acos(x) / π
asciidoctor: WARNING: The following text could not be fully converted to the Windows-1252 character set:
| acos(x[i]) / π
asciidoctor: WARNING: The following text could not be fully converted to the Windows-1252 character set:
| asin(x) / π
asciidoctor: WARNING: The following text could not be fully converted to the Windows-1252 character set:
| asin(x[i]) / π
asciidoctor: WARNING: The following text could not be fully converted to the Windows-1252 character set:
| atan(x) / π
asciidoctor: WARNING: The following text could not be fully converted to the Windows-1252 character set:
| atan(x[i]) / π
asciidoctor: WARNING: The following text could not be fully converted to the Windows-1252 character set:
| atan2(y, x) / π
asciidoctor: WARNING: The following text could not be fully converted to the Windows-1252 character set:
| atan2(y[i], x[i]) / π
asciidoctor: WARNING: The following text could not be fully converted to the Windows-1252 character set:
| cos(π * x)
asciidoctor: WARNING: The following text could not be fully converted to the Windows-1252 character set:
| cos(π * x[i])
asciidoctor: WARNING: The following text could not be fully converted to the Windows-1252 character set:
| sin(π * x)
asciidoctor: WARNING: The following text could not be fully converted to the Windows-1252 character set:
| sin(π * x[i])
asciidoctor: WARNING: The following text could not be fully converted to the Windows-1252 character set:
| tan(π * x)
asciidoctor: WARNING: The following text could not be fully converted to the Windows-1252 character set:
| tan(π * x[i])
asciidoctor: WARNING: The following text could not be fully converted to the Windows-1252 character set:
| (180 / π) * radi­ans
asciidoctor: WARNING: The following text could not be fully converted to the Windows-1252 character set:
| (180 / π) * radi­ans[i]
asciidoctor: WARNING: The following text could not be fully converted to the Windows-1252 character set:
| (π / 180) * degrees
asciidoctor: WARNING: The following text could not be fully converted to the Windows-1252 character set:
| (π / 180) * degrees[i]

I think we could fix the warnings by swapping π for something that's valid code, using an obvious placeholder constant like PI, or perhaps rewriting the block so that {pi} is not inside of [code] tags.

For the errors... I don't know what to do. As far as I can tell, changing the font is slightly adjusting the height of some table cells, which makes them big enough to fail to compile. This is probably a sign that our table cells are too large, and I hope it will fix itself when we have finished moving everything to the new format.

So... what should we do here? Should I just open a new issue to track supporting monospace italics in the PDF build? The exposition-only names are still obviously declared as such, both via a comment and the use of an invalid C++ name.

Copy link
Contributor

@gmlueck gmlueck Jun 18, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I couldn't figure out how to pass extra arguments to our docker-based build system, and I also couldn't figure out whether we'd need/want to put the new fonts inside the docker image or if we'd have to import them somehow.

I think we can pass additional command line options to Asciidoctor by changing the ADOCPDFOPTS variable in the Makefile. I also think we could ask for the docker image to be updated to contain new font files if we want. Would this help?

From your experiment, it seems like we either need to find a font that defines a "pi" character, or somehow change the font used for the code listing blocks without changing the font used for the inline [code] markup. (I think the "pi" characters are all in inline [code] markup.) I'm not sure if either of these are easy to do.

The table cell overflow on page 139 looks like the same one that @steffenlarsen found, which he fixed via #824. I'm not sure about the others because I don't see any table cells on those pages that are close to overflowing. Maybe they are just an artifact of the overflow on page 139? Or maybe your new font does something really weird when there is an unrenderable character, which causes the table cell to overflow?

If we could use the new font only for the code listing blocks (and not the inline [code] markup), that might also fix the table cell overflow problem. The overflow on page 139 is due to the [code] markup, not the code listing block.

If there is no easy way to fix the problems, I think we could live with the lack of italics in the code listings. I think most people use the HTML rendering. I was going to suggest that we could stop creating a PDF render entirely, but it seems like all the other Khronos specs have both HTML and PDF, so I guess this is a bad idea.

Copy link
Contributor Author

@Pennycook Pennycook Jun 19, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Just changing the font for source blocks causes a different set of table cells to overflow. I've tried changing the font size too, but a different set of table cells overflow.

I don't think I know enough about AsciiDoctor and how it handles fonts to fix this. The fact that I can't even see what went wrong (just an error message) is pretty frustrating.

I'm sure this is fixable, but I have no idea what I'm doing. It might be a better idea to get somebody like Jon Leech to take a look at it.

end
end
end
3 changes: 2 additions & 1 deletion adoc/extensions/index.adoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,4 +14,5 @@ include::sycl_khr_default_context.adoc[leveloffset=2]
include::sycl_khr_queue_empty_query.adoc[leveloffset=2]
include::sycl_khr_group_interface.adoc[leveloffset=2]
include::sycl_khr_max_work_group_queries.adoc[leveloffset=2]
include::sycl_khr_queue_flush.adoc[leveloffset=2]
include::sycl_khr_queue_flush.adoc[leveloffset=2]
include::sycl_khr_free_function_commands.adoc[leveloffset=2]
Loading