From cd98ef873f55e5950249241df5d828ad96e07aa7 Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Tue, 21 Jul 2026 15:15:26 -0500 Subject: [PATCH 01/34] detect when data-noescape should be applied to a code block --- Scripts/LessonChecker.py | 90 ++++++++++++++++++++++++++++++++-------- 1 file changed, 72 insertions(+), 18 deletions(-) diff --git a/Scripts/LessonChecker.py b/Scripts/LessonChecker.py index 7810d062..3a95ae06 100644 --- a/Scripts/LessonChecker.py +++ b/Scripts/LessonChecker.py @@ -45,15 +45,24 @@ class Mode(IntFlag): SILENT = auto() +class State(IntFlag): + NONE = 0 + CODE = auto() + MARK = auto() + ERROR = auto() + + +# ------------------------------------------------------------------------------------------ + + @dataclass class SYCLAParser(HTMLParser): lesson_name: str mode: Mode code_blocks: list[tuple[int, str]] = field(default_factory=list) output: list[str] = field(default_factory=list) - is_code: bool = False - is_mark: bool = False - parser_error: bool = False + + state: State = State.NONE SELF_CLOSING_TAGS = { "area", @@ -87,6 +96,24 @@ def msg_action(self): def output_joined(self): return "".join(self.output) + @property + def parser_error(self): + return State.ERROR in self.state + + @property + def last_code_idx(self): + if State.CODE in self.state: + # reverse search output for last code tag + return next( + ( + i + for i in range(len(self.output) - 1, -1, -1) + # < is important, because the word "code" can apepar in a comment, just like it does in this one + if " in without data-noescape prop in lesson {self.lesson_name} line num: {self.getpos()[0]}" + ) - case "pre" if self.is_code: + # Add data-noescape prop to code + code_tag = self.output[last_code_idx] + print(code_tag) + self.output[last_code_idx] = ( + code_tag[:-1] + " data-noescape" + code_tag[-1] + ) - self.parser_error = True + case "pre" if State.CODE in self.state: + + self.state |= State.ERROR self._warn( f"misalligned and
 tags in lesson {self.lesson_name} line num: {self.getpos()[0]}"
                 )
@@ -130,8 +178,8 @@ def handle_starttag(self, tag, attrs):
 
             case _:
                 # Any other unescaped
-                if self.is_code:
-                    self.parser_error = True
+                if State.CODE in self.state:
+                    self.state |= State.ERROR
                     self._warn(
                         f" annotations are the only tags allowed in  blocks! Violation in lesson {self.lesson_name} line num: {self.getpos()} tag: {tag}"
                     )
@@ -147,12 +195,12 @@ def handle_starttag(self, tag, attrs):
         self.output.append(self._encode_attrs(tag, attrs))
 
     def handle_data(self, data):
-        if self.is_code and bool(self.mode & Mode.EXTRACT):
+        if State.CODE in self.state and bool(self.mode & Mode.EXTRACT):
             pos, e_data = self.code_blocks[-1]
             self.code_blocks[-1] = (pos, e_data + data)
 
         if data == "&":
-            self.parser_error = True
+            self.state |= State.ERROR
             self._warn(
                 f"Unescaped &! Violation in lesson {self.lesson_name} line num: {self.getpos()}"
             )
@@ -166,7 +214,7 @@ def handle_entityref(self, name):
             case "lt" | "gt" | "amp":
                 raw_entity = f"&{name};"
             case _:
-                self.parser_error = True
+                self.state |= State.ERROR
                 self._warn(
                     f"Impropperly escaped entity ref! Violation in lesson {self.lesson_name} line num: {self.getpos()} ref: {name}"
                 )
@@ -175,7 +223,7 @@ def handle_entityref(self, name):
 
         self.output.append(raw_entity)
 
-        if self.is_code and bool(self.mode & Mode.EXTRACT):
+        if State.CODE in self.state and bool(self.mode & Mode.EXTRACT):
             converted_char = html.unescape(raw_entity)
             pos, e_data = self.code_blocks[-1]
             self.code_blocks[-1] = (pos, e_data + converted_char)
@@ -185,15 +233,19 @@ def handle_startendtag(self, tag, attrs):
 
     def handle_endtag(self, tag):
         if tag in self.SELF_CLOSING_TAGS:
-            self.parser_error = True
+            self.state |= State.ERROR
             self._warn(
                 f"Self closing tag does not require an endtag Violation in lesson {self.lesson_name} line num: {self.getpos()} tag: {tag}"
             )
             return
 
-        if tag == "code":
-            self.is_code = False
-            self.is_mark = False
+        # if tag == "code":
+
+        match tag:
+            case "code":
+                self.state &= ~State.CODE
+            case "mark":
+                self.state &= ~State.MARK
 
         self.output.append(f"")
 
@@ -275,7 +327,9 @@ def verify_html(lesson, file_str, args) -> tuple[str, list, bool]:
             with open(f"{out_base}/CMakeLists.txt", "w") as cmake:
                 cmake.write(cmake_executables)
 
-            if args.autofix:
+            # only write back to file if there is sonething to fix
+            if args.autofix and error:
+                print("saving back because fixed a thing")
                 write_back_path = l_path
 
                 with open(write_back_path, "w") as out:

From cfbf8f261a52040424be68c11fdf64b4955ff693 Mon Sep 17 00:00:00 2001
From: Michael Lance 
Date: Tue, 21 Jul 2026 16:07:58 -0500
Subject: [PATCH 02/34] update lesson checker and tests

---
 Scripts/LessonChecker.py      | 14 ++++++++++++--
 Scripts/test_LessonChecker.py | 29 ++++++++++++++++++++++++-----
 2 files changed, 36 insertions(+), 7 deletions(-)

diff --git a/Scripts/LessonChecker.py b/Scripts/LessonChecker.py
index 3a95ae06..cafb8c95 100644
--- a/Scripts/LessonChecker.py
+++ b/Scripts/LessonChecker.py
@@ -138,8 +138,19 @@ def handle_starttag(self, tag, attrs):
                 self.state |= State.CODE
 
                 if ("class", "language-cpp") not in attrs:
+                    self.state |= State.ERROR
+                    self._warn(
+                        f"highlight langauge class missing in lesson {self.lesson_name} line num: {self.getpos()[0]}"
+                    )
                     attrs.append(("class", "language-cpp"))
 
+                if ("class", "code-100pc") in attrs:
+                    self.state |= State.ERROR
+                    self._warn(
+                        f"Old stling in lesson {self.lesson_name} line num: {self.getpos()[0]}"
+                    )
+                    attrs.remove(("class", "code-100pc"))
+
                 self.code_blocks.append((self.getpos()[0], ""))
 
             case "mark":
@@ -160,7 +171,6 @@ def handle_starttag(self, tag, attrs):
 
                     # Add data-noescape prop to code
                     code_tag = self.output[last_code_idx]
-                    print(code_tag)
                     self.output[last_code_idx] = (
                         code_tag[:-1] + " data-noescape" + code_tag[-1]
                     )
@@ -329,7 +339,7 @@ def verify_html(lesson, file_str, args) -> tuple[str, list, bool]:
 
             # only write back to file if there is sonething to fix
             if args.autofix and error:
-                print("saving back because fixed a thing")
+                print("Saving fixes back to lesson")
                 write_back_path = l_path
 
                 with open(write_back_path, "w") as out:
diff --git a/Scripts/test_LessonChecker.py b/Scripts/test_LessonChecker.py
index a3bdd0f2..de84a291 100644
--- a/Scripts/test_LessonChecker.py
+++ b/Scripts/test_LessonChecker.py
@@ -18,18 +18,27 @@
         "extract": False,
         "verify": True,
         "fix": False,
-        "silent": True,
+        "silent": False,
         "expected_return": True,
     },
     {
         "name": "detect_mark_in_code",
-        "body": "
message
",
+        "body": '
message
',
         "extract": False,
         "verify": True,
         "fix": False,
         "silent": True,
         "expected_return": False,
     },
+    {
+        "name": "detect_mark_in_code_missing_escape",
+        "body": '
message
',
+        "extract": False,
+        "verify": True,
+        "fix": False,
+        "silent": True,
+        "expected_return": True,
+    },
     {
         "name": "parse_start_end_tag",
         "body": "
", @@ -77,7 +86,7 @@ }, { "name": "extract_code_block", - "body": "
int main() { return 0; }
", + "body": '
int main() { return 0; }
', "extract": True, "verify": False, "fix": False, @@ -97,7 +106,7 @@ }, { "name": "extract_proper_lt_and_gt", - "body": "
#include<sycl/sycl.hpp>
", + "body": '
#include<sycl/sycl.hpp>
', "extract": True, "verify": False, "fix": False, @@ -140,11 +149,21 @@ "body": '
&
', "extract": False, "verify": False, - "fix": False, + "fix": True, "silent": False, "expected_return": True, "expected_output": '
&
', }, + { + "name": "fix_old_styling_in_code", + "body": '
', + "extract": False, + "verify": False, + "fix": True, + "silent": False, + "expected_return": True, + "expected_output": '
', + }, ] From d56b4833790160475aa32e84e61e61902b407d39 Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Tue, 21 Jul 2026 16:09:28 -0500 Subject: [PATCH 03/34] autofix Data_and_Dependence --- .../Data_and_Dependencies/index.html | 56 +++++++++---------- 1 file changed, 25 insertions(+), 31 deletions(-) diff --git a/Lesson_Materials/Data_and_Dependencies/index.html b/Lesson_Materials/Data_and_Dependencies/index.html index 312af96a..d0c4910e 100644 --- a/Lesson_Materials/Data_and_Dependencies/index.html +++ b/Lesson_Materials/Data_and_Dependencies/index.html @@ -2,17 +2,11 @@ - - - - - + + + + + @@ -117,10 +111,10 @@
-
+							

 sycl::buffer buf {data, sycl::range{1024}};
 
-gpuQueue.submit([&](sycl::handler &cgh) {
+gpuQueue.submit([&](sycl::handler &cgh) {
   sycl::accessor acc {buf, cgh};
 
   cgh.parallel_for<kernel_a>(sycl::range{1024},
@@ -129,7 +123,7 @@
     });
 });
 
-gpuQueue.submit([&](sycl::handler &cgh) {
+gpuQueue.submit([&](sycl::handler &cgh) {
   sycl::accessor acc{buf, cgh};
 
   cgh.parallel_for<kernel_b>(sycl::range{1024},
@@ -155,10 +149,10 @@
 					
-
+							

 sycl::buffer buf {data, sycl::range{1024}};
 
-gpuQueue.submit([&](sycl::handler &cgh) {
+gpuQueue.submit([&](sycl::handler &cgh) {
   sycl::accessor acc {buf, cgh};
 
   cgh.parallel_for<kernel_a>(sycl::range{1024},
@@ -167,7 +161,7 @@
     });
 });
 
-gpuQueue.submit([&](sycl::handler &cgh) {
+gpuQueue.submit([&](sycl::handler &cgh) {
   sycl::accessor acc {buf, cgh};
 
   cgh.parallel_for<kernel_b>(sycl::range{1024},
@@ -194,10 +188,10 @@
 					
-
+							

 sycl::buffer buf {data, sycl::range{1024}};
 
-gpuQueue.submit([&](sycl::handler &cgh) {
+gpuQueue.submit([&](sycl::handler &cgh) {
   sycl::accessor acc{buf, cgh};
 
   cgh.parallel_for<my_kernel>(sycl::range{1024},
@@ -206,7 +200,7 @@
     });
 });
 
-gpuQueue.submit([&](sycl::handler &cgh) {
+gpuQueue.submit([&](sycl::handler &cgh) {
   sycl::accessor acc{buf, cgh};
 
   cgh.parallel_for<my_kernel>(sycl::range{1024},
@@ -232,10 +226,10 @@
 					
-
+							

 buf = sycl::buffer(data, sycl::range{1024});
 
-gpuQueue.submit([&](sycl::handler &cgh) {
+gpuQueue.submit([&](sycl::handler &cgh) {
   sycl::accessor acc {buf, cgh};
 
   cgh.parallel_for<my_kernel>(sycl::range{1024},
@@ -244,7 +238,7 @@
     });
 });
 
-gpuQueue.submit([&](sycl::handler &cgh) {
+gpuQueue.submit([&](sycl::handler &cgh) {
   sycl::accessor acc {buf, cgh};
 
   cgh.parallel_for<my_kernel>(sycl::range{1024},
@@ -269,7 +263,7 @@
 					
-
+							

 auto devicePtr =
   sycl::malloc_device<int>(1024, gpuQueue);
 
@@ -308,7 +302,7 @@
 					
-
+							

 auto devicePtr =
   sycl::malloc_device<int>(1024, gpuQueue);
 
@@ -343,7 +337,7 @@
 					
-
+							

 auto devicePtr =
   sycl::malloc_device<int>(1024, gpuQueue);
 
@@ -396,11 +390,11 @@
 					
-
+							

 sycl::buffer bufA {dataA, sycl::range{1024}};
 sycl::buffer bufB {dataB, sycl::range{1024}};
 
-gpuQueue.submit([&](sycl::handler &cgh) {
+gpuQueue.submit([&](sycl::handler &cgh) {
   sycl::accessor accA {bufA, cgh};
 
   cgh.parallel_for<kernel_a>(sycl::range{1024},
@@ -409,7 +403,7 @@
     });
 });
 
-gpuQueue.submit([&](sycl::handler &cgh) {
+gpuQueue.submit([&](sycl::handler &cgh) {
   sycl::accessor accB {bufB, cgh};
 
   cgh.parallel_for<kernel_b>(sycl::range{1024},
@@ -435,7 +429,7 @@
 						#### Concurrent data flow with USM
 					
-
+						

 auto devicePtrA = sycl::malloc_device<int>(1024, gpuQueue);
 auto devicePtrB = sycl::malloc_device<int>(1024, gpuQueue);
 
@@ -469,7 +463,7 @@
 						#### Concurrent data flow with USM
 					
-
+						

 auto devicePtrA = sycl::malloc_device<int>(1024, gpuQueue);
 auto devicePtrB = sycl::malloc_device<int>(1024, gpuQueue);
 

From ea86c0417c22cf8246c0335ce9a23b59cdb11821 Mon Sep 17 00:00:00 2001
From: Michael Lance 
Date: Tue, 21 Jul 2026 16:29:57 -0500
Subject: [PATCH 04/34] autofix Data_Parallelism

---
 Lesson_Materials/Data_Parallelism/index.html | 24 ++++++++------------
 1 file changed, 9 insertions(+), 15 deletions(-)

diff --git a/Lesson_Materials/Data_Parallelism/index.html b/Lesson_Materials/Data_Parallelism/index.html
index 1aae25ad..368823fe 100644
--- a/Lesson_Materials/Data_Parallelism/index.html
+++ b/Lesson_Materials/Data_Parallelism/index.html
@@ -2,17 +2,11 @@
 
 
   
-    
-    
-    
-    
-    
+    
+    
+    
+    
+    
     
     
     
@@ -257,8 +251,8 @@
 							   * Runtime may not always have enough information to choose the best-performing size
 						
- ND-Range
- SYCL-Range + ND-Range
+ SYCL-Range
@@ -432,8 +426,8 @@ event cpy_evt_1 = gpuQueue.copy(hA.data(), dA, hA.size()); event cpy_evt_2 = gpuQueue.copy(hB.data(), dB, hB.size()); -gpuQueue.parallel_for(range<1>(dA.size()) {cpy_evt_1, cpy_evt_2}, - [=](id<1> i){ +gpuQueue.parallel_for(range&<1>(dA.size()) {cpy_evt_1, cpy_evt_2}, + [=](id&lt1> i){ out[i] = inA[i] + inB[i]; }); From a0defa6f5ad79c95ddcd4537495cabe5cd22c165 Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Tue, 21 Jul 2026 16:40:33 -0500 Subject: [PATCH 05/34] only create cmakelists on extract --- Scripts/LessonChecker.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/Scripts/LessonChecker.py b/Scripts/LessonChecker.py index cafb8c95..b382ca0b 100644 --- a/Scripts/LessonChecker.py +++ b/Scripts/LessonChecker.py @@ -334,8 +334,8 @@ def verify_html(lesson, file_str, args) -> tuple[str, list, bool]: with open(wpath, "w") as out: out.write(code) - with open(f"{out_base}/CMakeLists.txt", "w") as cmake: - cmake.write(cmake_executables) + with open(f"{out_base}/CMakeLists.txt", "w") as cmake: + cmake.write(cmake_executables) # only write back to file if there is sonething to fix if args.autofix and error: From 41d24dbadb8b37ea0536b65ddf4ac0580d768dce Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Tue, 21 Jul 2026 16:43:16 -0500 Subject: [PATCH 06/34] autofix Device_Discovery --- Lesson_Materials/Device_Discovery/index.html | 44 +++++++++----------- 1 file changed, 19 insertions(+), 25 deletions(-) diff --git a/Lesson_Materials/Device_Discovery/index.html b/Lesson_Materials/Device_Discovery/index.html index c0398174..613e9e09 100644 --- a/Lesson_Materials/Device_Discovery/index.html +++ b/Lesson_Materials/Device_Discovery/index.html @@ -2,17 +2,11 @@ - - - - - + + + + + @@ -94,7 +88,7 @@
-
+							

 auto platforms = platform::get_platforms();
 
 							
@@ -115,7 +109,7 @@
-
+							

 auto intelDevices = intelPlatform.get_devices();
 
 							
@@ -137,7 +131,7 @@
-
+							

 auto devices = device::get_devices();
 
 							
@@ -176,7 +170,7 @@
-
+							

 auto gpuDevice = device(gpu_selector_v); 
 							
@@ -185,7 +179,7 @@
- * A device selector takes a parameter of type `const device &` and gives it a "score". + * A device selector takes a parameter of type `const device &` and gives it a "score". * Used to query all devices and return the one with the highest "score". * A device with a negative score will never be chosen.
@@ -197,7 +191,7 @@
-
+							

 auto chosenDevice = device();
 auto chosenDevice = device(default_selector_v);
 							
@@ -218,8 +212,8 @@ #### Creating a custom device selector
-
-int my_gpu_selector(const device& dev) {
+						

+int my_gpu_selector(const device& dev) {
 
 }
 						
@@ -235,8 +229,8 @@ #### Creating a custom device selector
-
-int my_gpu_selector(const device& dev) {
+						

+int my_gpu_selector(const device& dev) {
   if (dev.is_gpu()){
     return 1;
   }
@@ -257,8 +251,8 @@
 						#### Creating a custom device selector
 					
-
-int my_gpu_selector(const device& dev) {
+						

+int my_gpu_selector(const device& dev) {
   if (dev.is_gpu()){
     return 1;
   }
@@ -284,7 +278,7 @@
 					
-
+							

 auto plt = dev.get_platform();
 auto platformName
   = dev.get_info<info::device::name>();
@@ -307,7 +301,7 @@
 					
-
+							

 bool supportsFp16 = dev.has(aspect::fp16);
 							
From b84beb5796bd4b778cf60d4df84ae00e60959f34 Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Tue, 21 Jul 2026 16:44:39 -0500 Subject: [PATCH 07/34] autofix Enqueueing_a_Kernel --- .../Enqueueing_a_Kernel/index.html | 74 +++++++++---------- 1 file changed, 34 insertions(+), 40 deletions(-) diff --git a/Lesson_Materials/Enqueueing_a_Kernel/index.html b/Lesson_Materials/Enqueueing_a_Kernel/index.html index 0113f7fa..5d552f88 100644 --- a/Lesson_Materials/Enqueueing_a_Kernel/index.html +++ b/Lesson_Materials/Enqueueing_a_Kernel/index.html @@ -2,17 +2,11 @@ - - - - - + + + + + @@ -109,8 +103,8 @@
-
-gpuQueue.submit([&](handler &cgh){
+							

+gpuQueue.submit([&](handler &cgh){
   
   /* Command group function */
   
@@ -131,8 +125,8 @@
 					
-
-gpuQueue.submit([&](handler &cgh){
+							

+gpuQueue.submit([&](handler &cgh){
   
   /* Command group function */
   
@@ -153,8 +147,8 @@
 					
-
-gpuQueue.submit([&](handler &cgh){
+							

+gpuQueue.submit([&](handler &cgh){
 
   /* Command group function */
   
@@ -205,10 +199,10 @@
 					
-
+							

 class my_kernel;
 
-gpuQueue.submit([&](handler &cgh){
+gpuQueue.submit([&](handler &cgh){
 
   cgh.single_task<my_kernel>([=]() {
     /* kernel code */
@@ -228,10 +222,10 @@
 				
-
+							

 class my_kernel;
 
-gpuQueue.submit([&](handler &cgh){
+gpuQueue.submit([&](handler &cgh){
 								
   cgh.single_task<my_kernel>([=]() {
     /* kernel code */
@@ -250,10 +244,10 @@
 				
-
+							

 class my_kernel;
 
-gpuQueue.submit([&](handler &cgh){
+gpuQueue.submit([&](handler &cgh){
 								
   cgh.single_task<my_kernel>([=]() {
     /* kernel code */
@@ -272,10 +266,10 @@
 				
-
+							

 class my_kernel;
 
-gpuQueue.submit([&](handler &cgh){
+gpuQueue.submit([&](handler &cgh){
 								
   cgh.single_task<my_kernel>([=]() {
     /* kernel code */
@@ -295,7 +289,7 @@
 					
#### SYCL kernel function rules
-
+
* Must be defined using a C++ lambda or function object, they cannot be a function pointer or std::function. * Must always capture or store members by-value. * SYCL kernel functions declared with a lambda ~~must be named using a forward declarable C++ type, declared in global scope~~ can be anonymous since SYCL 2020! @@ -307,7 +301,7 @@
#### SYCL kernel function restrictions
-
+
* No dynamic allocation * No dynamic polymorphism * No function pointers @@ -321,11 +315,11 @@
-
+							

 class my_kernel;
 
 queue gpuQueue;
-gpuQueue.submit([&](handler &cgh){
+gpuQueue.submit([&](handler &cgh){
 
   cgh.single_task<my_kernel>([=]() {
     /* kernel code */
@@ -345,7 +339,7 @@
 					
-
+							

 struct my_kernel { 
   void operator()() const {
     /* kernel function */
@@ -367,16 +361,16 @@
 					
-
+							

 struct my_kernel { 
   void operator()() const {
     /* kernel function */
   }
 };
 							
-
+							

 queue gpuQueue;
-gpuQueue.submit([&](handler &cgh){
+gpuQueue.submit([&](handler &cgh){
 								
   cgh.single_task(my_kernel{}); 
 }).wait();
@@ -393,7 +387,7 @@
 					
#### Streams
-
+
* A `stream` can be used in a kernel function to print text to the console from the device, similarly to how you would with `std::cout`. * The `stream` is a buffered output stream so the output may not appear until the kernel function is complete. * The `stream` is useful for debugging, but should not be relied on in performance critical code. @@ -405,8 +399,8 @@ #### Streams
-
-sycl::stream(size_t bufferSize, size_t workItemBufferSize, handler &cgh);
+						

+sycl::stream(size_t bufferSize, size_t workItemBufferSize, handler &cgh);
 						
@@ -423,11 +417,11 @@
-
+							

 class my_kernel;
 
 queue gpuQueue;
-gpuQueue.submit([&](handler &cgh){
+gpuQueue.submit([&](handler &cgh){
 
   auto os = sycl::stream(1024, 1024, cgh);
 
@@ -450,11 +444,11 @@
 					
-
+							

 class my_kernel;
 
 queue gpuQueue;
-gpuQueue.submit([&](handler &cgh){
+gpuQueue.submit([&](handler &cgh){
 
   auto os = sycl::stream(1024, 1024, cgh);
 

From c22668a60e117894fc9377ce0e03e73170451c8b Mon Sep 17 00:00:00 2001
From: Michael Lance 
Date: Wed, 22 Jul 2026 11:07:32 -0500
Subject: [PATCH 08/34] better handling for partial escape sequences

---
 Scripts/LessonChecker.py | 4 +++-
 1 file changed, 3 insertions(+), 1 deletion(-)

diff --git a/Scripts/LessonChecker.py b/Scripts/LessonChecker.py
index b382ca0b..6569d932 100644
--- a/Scripts/LessonChecker.py
+++ b/Scripts/LessonChecker.py
@@ -11,6 +11,7 @@
 from pathlib import Path
 import argparse
 import sys
+import re
 
 # ------------------------------------------------------------------------------------------
 # Config
@@ -219,10 +220,11 @@ def handle_data(self, data):
         self.output.append(data)
 
     def handle_entityref(self, name):
-
         match name:
             case "lt" | "gt" | "amp":
                 raw_entity = f"&{name};"
+            case str() if m := re.search(r"lt(.+)", name):
+                raw_entity = f"<{m.group(1)};"
             case _:
                 self.state |= State.ERROR
                 self._warn(

From f8986c564fb72adf7fbe6524962144ec9c2df784 Mon Sep 17 00:00:00 2001
From: Michael Lance 
Date: Wed, 22 Jul 2026 17:13:31 -0500
Subject: [PATCH 09/34] improved handling for unescaped characters in bodies

---
 Scripts/LessonChecker.py      | 17 +++++++++++------
 Scripts/test_LessonChecker.py | 11 +++++++++++
 2 files changed, 22 insertions(+), 6 deletions(-)

diff --git a/Scripts/LessonChecker.py b/Scripts/LessonChecker.py
index 6569d932..de52f39b 100644
--- a/Scripts/LessonChecker.py
+++ b/Scripts/LessonChecker.py
@@ -190,6 +190,7 @@ def handle_starttag(self, tag, attrs):
             case _:
                 # Any other unescaped
                 if State.CODE in self.state:
+                    print(tag)
                     self.state |= State.ERROR
                     self._warn(
                         f" annotations are the only tags allowed in  blocks! Violation in lesson {self.lesson_name} line num: {self.getpos()} tag: {tag}"
@@ -210,12 +211,13 @@ def handle_data(self, data):
             pos, e_data = self.code_blocks[-1]
             self.code_blocks[-1] = (pos, e_data + data)
 
-        if data == "&":
+        if any((viol_char := char) in {"&", "<", ">"} for char in data):
             self.state |= State.ERROR
             self._warn(
-                f"Unescaped &! Violation in lesson {self.lesson_name} line num: {self.getpos()}"
+                f"Unescaped character {viol_char} in lesson {self.lesson_name} line num: {self.getpos()[0]}"
             )
-            data = "&"
+
+        data = data.replace("&", "&").replace("<", "<").replace(">", ">")
 
         self.output.append(data)
 
@@ -225,6 +227,11 @@ def handle_entityref(self, name):
                 raw_entity = f"&{name};"
             case str() if m := re.search(r"lt(.+)", name):
                 raw_entity = f"<{m.group(1)};"
+
+                self.state |= State.ERROR
+                self._warn(
+                    f"Partial escape sequence found in lesson {self.lesson_name} line num: {self.getpos()[0]}"
+                )
             case _:
                 self.state |= State.ERROR
                 self._warn(
@@ -251,8 +258,6 @@ def handle_endtag(self, tag):
             )
             return
 
-        # if tag == "code":
-
         match tag:
             case "code":
                 self.state &= ~State.CODE
@@ -268,7 +273,7 @@ def handle_decl(self, decl):
         self.output.append(f"")
 
 
-# ------------------------------------------------------------------------------------------
+# -----------------------------------------------------------------------------------------
 
 
 def verify_html(lesson, file_str, args) -> tuple[str, list, bool]:
diff --git a/Scripts/test_LessonChecker.py b/Scripts/test_LessonChecker.py
index de84a291..cc4b2e0c 100644
--- a/Scripts/test_LessonChecker.py
+++ b/Scripts/test_LessonChecker.py
@@ -164,6 +164,17 @@
         "expected_return": True,
         "expected_output": '
', }, + { + "name": "fix_partial_escape_sequences", + "body": '
&
', + "extract": False, + "verify": False, + "fix": True, + "silent": False, + "expected_return": True, + "expected_output": '
&
', + }, + ] From c95819353d6aceadc4b3c50b9427efabd97f3473 Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Wed, 22 Jul 2026 17:13:58 -0500 Subject: [PATCH 10/34] autofix Fast_Track --- Lesson_Materials/Fast_Track/index.html | 98 ++++++++++++-------------- 1 file changed, 46 insertions(+), 52 deletions(-) diff --git a/Lesson_Materials/Fast_Track/index.html b/Lesson_Materials/Fast_Track/index.html index 30ed5bbd..5fe199a4 100644 --- a/Lesson_Materials/Fast_Track/index.html +++ b/Lesson_Materials/Fast_Track/index.html @@ -2,17 +2,11 @@ - - - - - + + + + + @@ -31,7 +25,7 @@
-
+
## Learning Objectives
@@ -158,7 +152,7 @@
-
+							

 auto gpuDevice = device(gpu_selector_v); 
 							
@@ -167,7 +161,7 @@
- * A device selector takes a parameter of type `const device &` and gives it a "score". + * A device selector takes a parameter of type `const device &` and gives it a "score". * Used to query all devices and return the one with the highest "score". * A device with a negative score will never be chosen.
@@ -228,12 +222,12 @@ * Some computation on the GPU. * Migration of the result back to the CPU.
- +
- #### SYCL Buffers & Accessors + #### SYCL Buffers & Accessors
* The buffer/accessor model separates the storage and access of data @@ -246,7 +240,7 @@
- #### SYCL Buffers & Accessors + #### SYCL Buffers & Accessors
@@ -261,7 +255,7 @@
- #### SYCL Buffers & Accessors + #### SYCL Buffers & Accessors
@@ -280,17 +274,17 @@
-
-buffer<float, 1> bufA(dA.data(), range<1>(dA.size())); 
-buffer<float, 1> bufB(dB.data(), range<1>(dB.size())); 
-buffer<float, 1> bufO(dO.data(), range<1>(dO.size()));
+							

+buffer<float;, 1> bufA(dA.data(), range<1;>(dA.size())); 
+buffer<float;, 1> bufB(dB.data(), range<1;>(dB.size())); 
+buffer<float;, 1> bufO(dO.data(), range<1;>(dO.size()));
 
-gpuQueue.submit([&](handler &cgh){
+gpuQueue.submit([&](handler &cgh){
   sycl::accessor inA{bufA, cgh, sycl::read_only};
   sycl::accessor inB{bufB, cgh, sycl::read_only};
   sycl::accessor out{bufO, cgh, sycl::write_only};
-  cgh.parallel_for<add>(range<1>(dA.size()), 
-    [=](id<1> i){ 
+  cgh.parallel_for<add;>(range<1;>(dA.size()), 
+    [=](id<1;> i){ 
     out[i] = inA[i] + inB[i];
   });
 });
@@ -309,11 +303,11 @@
 						#### USM: Malloc_device
 					
-
-void* malloc_device(size_t numBytes, const queue& syclQueue, const property_list &propList = {});
+						

+void* malloc_device(size_t numBytes, const queue& syclQueue, const property_list &propList = {});
 
 template <typename T>
-T* malloc_device(size_t count, const queue& syclQueue, const property_list &propList = {});
+T* malloc_device(size_t count, const queue& syclQueue, const property_list &propList = {});
 						
@@ -330,8 +324,8 @@ #### USM: Free
-
-void free(void* ptr, queue& syclQueue);
+						

+void free(void* ptr, queue& syclQueue);
 						
@@ -347,8 +341,8 @@ #### USM: Memcpy
-
-event queue::memcpy(void* dest, const void* src, size_t numBytes, const std::vector &depEvents);
+						

+event queue::memcpy(void* dest, const void* src, size_t numBytes, const std::vector<event> &depEvents);
 						
@@ -434,24 +428,24 @@
-
+							

 							
-cgh.parallel_for<kernel>(range<1>(1024), 
-  [=](id<1> idx){
+cgh.parallel_for<kernel;>(range<1;>(1024), 
+  [=](id<1;> idx){
     /* kernel function code */
 });
 							
-
+							

 							
-cgh.parallel_for<kernel>(range<1>(1024), 
-  [=](item<1> item){
+cgh.parallel_for<kernel;>(range<1;>(1024), 
+  [=](item<1;> item){
     /* kernel function code */
 });
 							
-
+							

 							
-cgh.parallel_for<kernel>(nd_range<1>(range<1>(1024), 
-  range<1>(32)),[=](nd_item<1> ndItem){
+cgh.parallel_for<kernel;>(nd_range<1;>(range<1;>(1024), 
+  range<1;>(32)),[=](nd_item<1;> ndItem){
     /* kernel function code */
 });
 							
@@ -489,21 +483,21 @@
#### Function object
-
- ```c++ +
+

 class MyKernel {
-  sycl::accessor input_;
+  sycl::accessor<float> input_;
   float* output_;
 
-  MyKernel(sycl::buffer buf, float* output, sycl::handler& h)
+  MyKernel(sycl::buffer<float> buf, float* output, sycl::handler& h)
     : input{buf.get_access(h)}, output_{output} {}
 
   // const is required
-  void operator()(sycl::item<1> i) const {
+  void operator()(sycl::item<1> i) const {
     ; // computation here
   }
 };
-            ```
+						
The members are accessible on the device inside the function call @@ -515,20 +509,20 @@
#### Lambda function
-
- ```c++ -sycl::buffer buf = /* normal init */; +
+

+sycl::buffer<float, 1> buf = /* normal init */;
 float * output = sycl::malloc_device(/* params */);
 
 ... queue submit as normal ...
 
 auto acc = buf.get_access(h);
-auto func = [=](sycl::item<1> i) {
+auto func = [=](sycl::item<1> i) {
   acc[i] = someVal;
   output[i.get_global_linear_id()] = someOtherVal;
 });
 handler.parallel_for(range, func);
-            ```
+						
The variables used implicitly are captured by value and are usable in From 2e79aabd8e7113145ac3ae99d07a0acf9beaa028 Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Wed, 22 Jul 2026 17:15:32 -0500 Subject: [PATCH 11/34] autofix Enqueueing_a_Kernel II (better auto fix) --- Lesson_Materials/Enqueueing_a_Kernel/index.html | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/Lesson_Materials/Enqueueing_a_Kernel/index.html b/Lesson_Materials/Enqueueing_a_Kernel/index.html index 5d552f88..ff09e468 100644 --- a/Lesson_Materials/Enqueueing_a_Kernel/index.html +++ b/Lesson_Materials/Enqueueing_a_Kernel/index.html @@ -453,15 +453,15 @@ auto os = sycl::stream(1024, 1024, cgh); cgh.single_task<my_kernel>([=]() { - os << "Hello world!\n"; + os << "Hello world!\n"; }); }).wait();
* Next we capture the `stream` in the kernel function's lambda expression. - * Then we can print `"Hello World!"` to the console using the `<<` operator. - * This is where the work-item size comes in, this is the cache available to store text on the right-hand-size of the `<<` operator. + * Then we can print `"Hello World!"` to the console using the `<<` operator. + * This is where the work-item size comes in, this is the cache available to store text on the right-hand-size of the `<<` operator.
From ba3a9cccb7f125bad30f52d49f17ecc8a4770a6f Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Wed, 22 Jul 2026 17:17:33 -0500 Subject: [PATCH 12/34] autofix Hanlding_Errors --- Lesson_Materials/Handling_Errors/index.html | 21 +++++++-------------- 1 file changed, 7 insertions(+), 14 deletions(-) diff --git a/Lesson_Materials/Handling_Errors/index.html b/Lesson_Materials/Handling_Errors/index.html index a6cc4700..5f78abfe 100644 --- a/Lesson_Materials/Handling_Errors/index.html +++ b/Lesson_Materials/Handling_Errors/index.html @@ -2,17 +2,11 @@ - - - - - + + + + + @@ -83,8 +77,7 @@
-

+            

 #include <sycl/sycl.hpp>
 using namespace sycl;
 
@@ -185,7 +178,7 @@
     q.wait_and_throw();
   } catch (const exception& e) {
   std::cout << "Exception caught: " << e.what() 
-     << std::endl;
+     << std::endl;
   }
 }
 							
From 39f112c0eb6465d159493eae99b00f948275b182 Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Wed, 22 Jul 2026 17:18:05 -0500 Subject: [PATCH 13/34] autofix Image_Convolution --- Lesson_Materials/Image_Convolution/index.html | 20 +++++++------------ 1 file changed, 7 insertions(+), 13 deletions(-) diff --git a/Lesson_Materials/Image_Convolution/index.html b/Lesson_Materials/Image_Convolution/index.html index 1475dc1d..21f02d81 100644 --- a/Lesson_Materials/Image_Convolution/index.html +++ b/Lesson_Materials/Image_Convolution/index.html @@ -2,17 +2,11 @@ - - - - - + + + + + @@ -152,7 +146,7 @@ #### Input/output image locations
- +
 auto inputImageFile = "../Code_Exercises/Images/dogs.png";
 auto outputImageFile = ../Code_Exercises/Images/blurred_dogs.png";
@@ -172,7 +166,7 @@
 						#### Convolution filters
 					
- +
 auto filter = util::generate_filter(util::filter_type filterType, int width);
 						
From fb95e210e5a75be2fc833ca55beb827cf104e669 Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Wed, 22 Jul 2026 17:19:17 -0500 Subject: [PATCH 14/34] autofix Introduction_To_USM --- .../Introduction_to_USM/index.html | 18 ++++++------------ 1 file changed, 6 insertions(+), 12 deletions(-) diff --git a/Lesson_Materials/Introduction_to_USM/index.html b/Lesson_Materials/Introduction_to_USM/index.html index e8ec3b84..6afaff6c 100644 --- a/Lesson_Materials/Introduction_to_USM/index.html +++ b/Lesson_Materials/Introduction_to_USM/index.html @@ -2,17 +2,11 @@ - - - - - + + + + + @@ -158,7 +152,7 @@ Each SYCL platform and its device(s) will support different variants of USM and different kinds of memory allocation
-
if (dev.has(sycl::aspect::usm_device_allocations))
+
if (dev.has(sycl::aspect::usm_device_allocations))
From 0a83c92664e116c171e7bc1c70495b0a0cd9d99c Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Wed, 22 Jul 2026 17:20:08 -0500 Subject: [PATCH 15/34] autofix Local_Memory_Tiling --- .../Local_Memory_Tiling/index.html | 22 +++++++------------ 1 file changed, 8 insertions(+), 14 deletions(-) diff --git a/Lesson_Materials/Local_Memory_Tiling/index.html b/Lesson_Materials/Local_Memory_Tiling/index.html index 167d639a..c729287a 100644 --- a/Lesson_Materials/Local_Memory_Tiling/index.html +++ b/Lesson_Materials/Local_Memory_Tiling/index.html @@ -2,17 +2,11 @@ - - - - - + + + + + @@ -103,7 +97,7 @@ #### Local accessors
-
+						

 auto scratchpad = sycl::accessor<int, 1, sycl::access::target::local>(sycl::range{workGroupSize}, cgh);
 						
@@ -244,7 +238,7 @@ #### Group_barrier
-
+						

 sycl::group_barrier(item.get_group());
 						
@@ -259,7 +253,7 @@
#### Local memory image convolution performance
-
+
![SYCL](../../Static/images/image_convolution_performance_local_mem.png "SYCL")
From 0b38a2af9b2f19a78f6d0f17e3e26b78ecbe1541 Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Wed, 22 Jul 2026 17:21:20 -0500 Subject: [PATCH 16/34] autofix Managing_Data --- Lesson_Materials/Managing_Data/index.html | 62 ++++++++++------------- 1 file changed, 28 insertions(+), 34 deletions(-) diff --git a/Lesson_Materials/Managing_Data/index.html b/Lesson_Materials/Managing_Data/index.html index 0ad04956..88fb8b80 100644 --- a/Lesson_Materials/Managing_Data/index.html +++ b/Lesson_Materials/Managing_Data/index.html @@ -2,17 +2,11 @@ - - - - - + + + + + @@ -60,7 +54,7 @@ * Some computation on the GPU. * Migration of the result back to the CPU.
- +
@@ -71,7 +65,7 @@ * Memory transfers between CPU and GPU are a bottleneck. * We want to minimize these transfers, when possible.
- +
@@ -95,7 +89,7 @@ #### Using USM - Malloc Device
-
+						

 // Allocate memory on device
 T *device_ptr = sycl::malloc_device<T>(n, myQueue);
 
@@ -123,7 +117,7 @@
 						#### Using USM - Malloc Shared
 					
-
+						

 // Allocate shared memory 
 T *shared_ptr = sycl::malloc_shared<T>(n, myQueue);
 
@@ -146,7 +140,7 @@
 				
 				
- #### SYCL Buffers & Accessors + #### SYCL Buffers & Accessors
* SYCL provides an API which takes care of allocations and `memcpy`s, as well as some other things. @@ -154,7 +148,7 @@
- #### SYCL Buffers & Accessors + #### SYCL Buffers & Accessors
* The buffer/accessor model separates the storage and access of data @@ -167,7 +161,7 @@
- #### SYCL Buffers & Accessors + #### SYCL Buffers & Accessors
@@ -184,7 +178,7 @@
- #### SYCL Buffers & Accessors + #### SYCL Buffers & Accessors
@@ -199,7 +193,7 @@
- #### SYCL Buffers & Accessors + #### SYCL Buffers & Accessors
@@ -214,7 +208,7 @@
- #### SYCL Buffers & Accessors + #### SYCL Buffers & Accessors
@@ -230,7 +224,7 @@
- #### SYCL Buffers & Accessors + #### SYCL Buffers & Accessors
@@ -245,15 +239,15 @@
- #### SYCL Buffers & Accessors + #### SYCL Buffers & Accessors
-
+						

 T var = 42;
 
 {
   // Create buffer pointing to var.
-  auto buf = sycl::buffer{&var, sycl::range<1>{1}};
+  auto buf = sycl::buffer{&var, sycl::range<1>{1}};
 
   // ...
   // Do some computation on device. Use accessors to access buffer
@@ -295,7 +289,7 @@
 						#### Constructing an accessor
 					
-
+						

 auto acc = sycl::accessor{bufA, cgh};
 						
@@ -320,7 +314,7 @@ #### Specifying the access mode
-
+						

 auto readAcc = sycl::accessor{bufA, cgh, sycl::read_only};
 auto writeAcc = sycl::accessor{bufB, cgh, sycl::write_only};
 						
@@ -339,7 +333,7 @@ #### Specifying no initialization
-
+						

 auto acc = sycl::accessor{buf, cgh, sycl::no_init};
 						
@@ -355,15 +349,15 @@ #### Using Accessors
-
+						

 T var = 42;
 
 {
   // Create buffer pointing to var.
-  auto bufA = sycl::buffer{&var, sycl::range<1>{1}};
-  auto bufB = sycl::buffer{&var, sycl::range<1>{1}};
+  auto bufA = sycl::buffer{&var, sycl::range<1>{1}};
+  auto bufB = sycl::buffer{&var, sycl::range<1>{1}};
 
-  q.submit([&](sycl::handler &cgh) {
+  q.submit([&](sycl::handler &cgh) {
 	auto accA = sycl::accessor{bufA, cgh, sycl::read_only};
 	auto accB = sycl::accessor{bufA, cgh, sycl::no_init};
 
@@ -387,8 +381,8 @@
 						#### operator[]
 					
-
-gpuQueue.submit([&](handler &cgh){
+						

+gpuQueue.submit([&](handler &cgh){
   auto inA = sycl::accessor{bufA, cgh, sycl::read_only};
   auto inB = sycl::accessor{bufB, cgh, sycl::read_only};
   auto out = sycl::accessor{bufO, cgh, sycl::write_only};

From 007f295907e3987cb6891a1d39b2808fb919c6c9 Mon Sep 17 00:00:00 2001
From: Michael Lance 
Date: Wed, 22 Jul 2026 17:21:50 -0500
Subject: [PATCH 17/34] autofix Matrix_Transpose

---
 Lesson_Materials/Matrix_Transpose/index.html | 20 +++++++-------------
 1 file changed, 7 insertions(+), 13 deletions(-)

diff --git a/Lesson_Materials/Matrix_Transpose/index.html b/Lesson_Materials/Matrix_Transpose/index.html
index abbce19c..bf78a092 100644
--- a/Lesson_Materials/Matrix_Transpose/index.html
+++ b/Lesson_Materials/Matrix_Transpose/index.html
@@ -2,17 +2,11 @@
 
 
   
-    
-    
-    
-    
-    
+    
+    
+    
+    
+    
     
     
     
@@ -183,7 +177,7 @@
 						#### Local accessors
 					
-
+						

 auto scratchpad = sycl::local_accessor<int, dims>(sycl::range{workGroupSize}, cgh);
 						
@@ -324,7 +318,7 @@ #### Group_barrier
-
+						

 sycl::group_barrier(item.get_group());
 						
From 76ac842f5ca218d6a24b7b7e81ace3b1fa702513 Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Wed, 22 Jul 2026 17:24:09 -0500 Subject: [PATCH 18/34] autofix More_SYCL_Features --- .../More_SYCL_Features/index.html | 22 +++++++------------ 1 file changed, 8 insertions(+), 14 deletions(-) diff --git a/Lesson_Materials/More_SYCL_Features/index.html b/Lesson_Materials/More_SYCL_Features/index.html index ed70306b..d9b5451d 100644 --- a/Lesson_Materials/More_SYCL_Features/index.html +++ b/Lesson_Materials/More_SYCL_Features/index.html @@ -2,17 +2,11 @@ - - - - - + + + + + @@ -227,7 +221,7 @@
-

+							

 q.parallel_for([=](sycl::item<1> it) {
   sycl::atomic_ref<T,
           sycl::memory_order_relaxed,
@@ -288,7 +282,7 @@
 					
-

+							

 q.submit([&](sycl::handler &cgh) {
   // Output of reduction will be in ptr
   auto sumReduction = sycl::reduction(ptr,
@@ -315,7 +309,7 @@
 					
-

+							

 q.submit([&](sycl::handler &cgh) {
   // Output of reduction will be in ptr
   auto maxReduction = sycl::reduction(ptr,

From 6fcd1b186ae9773bf6a5a246c1686762ee91415d Mon Sep 17 00:00:00 2001
From: Michael Lance 
Date: Wed, 22 Jul 2026 17:25:43 -0500
Subject: [PATCH 19/34] autofix Multiple_Devices

---
 Lesson_Materials/Multiple_Devices/index.html | 22 +++++++-------------
 1 file changed, 8 insertions(+), 14 deletions(-)

diff --git a/Lesson_Materials/Multiple_Devices/index.html b/Lesson_Materials/Multiple_Devices/index.html
index 88c76ba7..1ab4e0ae 100644
--- a/Lesson_Materials/Multiple_Devices/index.html
+++ b/Lesson_Materials/Multiple_Devices/index.html
@@ -2,17 +2,11 @@
 
 
   
-    
-    
-    
-    
-    
+    
+    
+    
+    
+    
     
     
     
@@ -102,7 +96,7 @@
 						#### Creating an implicit context
 					
-
+						

 auto defaultQueue = queue{};
 						
@@ -118,7 +112,7 @@ #### Creating a context from devices
-
+						

 auto sharedContext = context{{cpuDevice, gpuDevice}};
 						
@@ -133,7 +127,7 @@ #### Creating a context from a platform
-
+						

 auto sharedContext = context{intelPlatform};
 						
From db88afc198f1c5dbe0ff22d118afaec964c6f69d Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Wed, 22 Jul 2026 17:28:23 -0500 Subject: [PATCH 20/34] autofix oneMath_gemm --- Lesson_Materials/oneMath_gemm/index.html | 36 ++++++++++-------------- 1 file changed, 15 insertions(+), 21 deletions(-) diff --git a/Lesson_Materials/oneMath_gemm/index.html b/Lesson_Materials/oneMath_gemm/index.html index 23c973e7..0da36689 100644 --- a/Lesson_Materials/oneMath_gemm/index.html +++ b/Lesson_Materials/oneMath_gemm/index.html @@ -2,17 +2,11 @@ - - - - - + + + + + @@ -25,7 +19,7 @@
-

+

oneAPI Math Library (oneMath)

@@ -59,13 +53,13 @@

* **Intel CPU/GPU**: Intel Math Kernels Library (oneMKL) * **NVIDIA GPU**: cuBLAS, cuSOLVER, cuRAND, cuFFT * **AMD GPU**: rocBLAS, rocSOLVER, rocRAND, rocFFT - * Imagine being able to use all of them with *single source code* → **oneMath** + * Imagine being able to use all of them with *single source code* **oneMath**

-

+

oneAPI and oneMath

@@ -80,14 +74,14 @@

* wrapper implementation dispatching SYCL API calls to a multitude of implementations, both generic and vendor-specific

- +
- +
-

+

oneMath library backends

@@ -98,7 +92,7 @@

#### Run-time dispatching

-

+						

 #include <oneapi/math.hpp>
 
 sycl::queue q{myDeviceSelector};
@@ -125,8 +119,8 @@ 

#### Compile-time dispatching

-

-#include <oneapi/math.hpp>
+						

+#include <oneapi;/math.hpp>
 
 sycl::queue cpu_queue{sycl::cpu_selector_v};
 
@@ -134,7 +128,7 @@ 

sycl::buffer<T,1> b{b_host, k*n}; sycl::buffer<T,1> c{c_host, m*n}; -oneapi::math::backend_selector<oneapi::math::backend::mklcpu> cpu_selector(cpu_queue); +oneapi::math::backend_selector<oneapi;::math::backend::mklcpu> cpu_selector(cpu_queue); // Select the Intel oneMKL CPU backend specifically ^^^^^^ oneapi::math::blas::column_major::gemm(cpu_selector, ..., m, n, k, ..., a, ..., b, ..., c, ... ); From 2e2c4e26d7f2e1cfc55dd1cff9b60d4b2a2c4f0e Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Wed, 22 Jul 2026 17:29:33 -0500 Subject: [PATCH 21/34] autofix Using_USM --- Lesson_Materials/Using_USM/index.html | 92 +++++++++++++-------------- 1 file changed, 43 insertions(+), 49 deletions(-) diff --git a/Lesson_Materials/Using_USM/index.html b/Lesson_Materials/Using_USM/index.html index cbecaa8a..104a5268 100644 --- a/Lesson_Materials/Using_USM/index.html +++ b/Lesson_Materials/Using_USM/index.html @@ -2,17 +2,11 @@ - - - - - + + + + + @@ -53,7 +47,7 @@
#### USM Allocation Types
-
+
| Type | Description | Access host | Access device | Nominal location | |--------|---------------------------------|----------------|----------------|----------------------| | device | device global allocations | ✗ | ✓ | device | @@ -67,11 +61,11 @@ #### Malloc_device
-
-void* malloc_device(size_t numBytes, const queue& syclQueue, const property_list &propList = {});
+						

+void* malloc_device(size_t numBytes, const queue& syclQueue, const property_list &propList = {});
 
 template <typename T>
-T* malloc_device(size_t count, const queue& syclQueue, const property_list &propList = {});
+T* malloc_device(size_t count, const queue& syclQueue, const property_list &propList = {});
 						
@@ -88,8 +82,8 @@ #### Free
-
-void free(void* ptr, queue& syclQueue);
+						

+void free(void* ptr, queue& syclQueue);
 						
@@ -105,8 +99,8 @@ #### Memcpy
-
-event queue::memcpy(void* dest, const void* src, size_t numBytes, const std::vector &depEvents);
+						

+event queue::memcpy(void* dest, const void* src, size_t numBytes, const std::vector<event> &depEvents);
 						
@@ -120,13 +114,13 @@
- #### Memset & fill + #### Memset & fill
-
-event queue::memset(void* ptr, int value, size_t numBytes, const std::vector &depEvents);
+						

+event queue::memset(void* ptr, int value, size_t numBytes, const std::vector<event> &depEvents);
 
-event queue::fill(void* ptr, const T& pattern, size_t count, const std::vector &depEvents);
+event queue::fill(void* ptr, const T& pattern, size_t count, const std::vector<event> &depEvents);
 						
@@ -142,12 +136,12 @@ #### Putting it all together
-
+						

 int square_number(int x){
 	
   auto myQueue = queue{};
 
-  myQueue.submit([&](handler &cgh){
+  myQueue.submit([&](handler &cgh){
     cgh.single_task<square_number>([=](){
       /* square some number */
     });
@@ -167,12 +161,12 @@
 						#### Putting it all together
 					
-
+						

 int square_number(int x){
 	
   auto myQueue = queue{usm_selector{}};
 
-  myQueue.submit([&](handler &cgh){
+  myQueue.submit([&](handler &cgh){
     cgh.single_task<square_number>([=](){
       /* square some number */
     });
@@ -192,14 +186,14 @@
 						#### Putting it all together
 					
-
+						

 int square_number(int x){
 	
   auto myQueue = queue{usm_selector{}};
 
   auto devicePtr = malloc_device<int>(1, myQueue);
 
-  myQueue.submit([&](handler &cgh){
+  myQueue.submit([&](handler &cgh){
     cgh.single_task<square_number>([=](){
       /* square some number */
     });
@@ -220,16 +214,16 @@
 						#### Putting it all together
 					
-
+						

 int square_number(int x){
 
   auto myQueue = queue{usm_selector{}};
 
   auto devicePtr = malloc_device<int>(1, myQueue);
 
-  myQueue.memcpy(devicePtr, &x, sizeof(int)).wait();
+  myQueue.memcpy(devicePtr, &x, sizeof(int)).wait();
 
-  myQueue.submit([&](handler &cgh){
+  myQueue.submit([&](handler &cgh){
     cgh.single_task<square>([=](){
       /* square some number */
     });
@@ -250,16 +244,16 @@
 						#### Putting it all together
 					
-
+						

 int square_number(int x){
 
   auto myQueue = queue{usm_selector{}};
 
   auto devicePtr = malloc_device<int>(1, myQueue);
 
-  myQueue.memcpy(devicePtr, &x, sizeof(int)).wait();
+  myQueue.memcpy(devicePtr, &x, sizeof(int)).wait();
 
-  myQueue.submit([&](handler &cgh){
+  myQueue.submit([&](handler &cgh){
     cgh.single_task<square>([=](){
       *devicePtr = (*devicePtr) * (*devicePtr);
     });
@@ -279,22 +273,22 @@
 						#### Putting it all together
 					
-
+						

 int square_number(int x){
 
   auto myQueue = queue{usm_selector{}};
 
   auto devicePtr = malloc_device<int>(1, myQueue);
 
-  myQueue.memcpy(devicePtr, &x, sizeof(int)).wait();
+  myQueue.memcpy(devicePtr, &x, sizeof(int)).wait();
 
-  myQueue.submit([&](handler &cgh){
+  myQueue.submit([&](handler &cgh){
     cgh.single_task<square>([=](){
       *devicePtr = (*devicePtr) * (*devicePtr);
     });
   }).wait();
 
-  myQueue.memcpy(&x, devicePtr, sizeof(int)).wait();
+  myQueue.memcpy(&x, devicePtr, sizeof(int)).wait();
 
   return x;
 }
@@ -310,22 +304,22 @@
 						#### Putting it all together
 					
-
+						

 int square_number(int x){
 
   auto myQueue = queue{usm_selector{}};
 
   auto devicePtr = malloc_device<int>(1, myQueue);
 
-  myQueue.memcpy(devicePtr, &x, sizeof(int)).wait();
+  myQueue.memcpy(devicePtr, &x, sizeof(int)).wait();
 
-  myQueue.submit([&](handler &cgh){
+  myQueue.submit([&](handler &cgh){
     cgh.single_task<square>([=](){
       *devicePtr = (*devicePtr) * (*devicePtr);
     });
   }).wait();
 
-  myQueue.memcpy(&x, devicePtr, sizeof(int)).wait();
+  myQueue.memcpy(&x, devicePtr, sizeof(int)).wait();
 
   free(devicePtr, myQueue);
 
@@ -343,12 +337,12 @@
 						#### Queue shortcuts
 					
-
+						

 template <typename KernelName, typename KernelType>
-event queue::single_task(const KernelType &KernelFunc);
+event queue::single_task(const KernelType &KernelFunc);
 
 template <typename KernelName, typename KernelType, int Dims>
-event queue::parallel_for(range GlobalRange, const KernelType &KernelFunc);
+event queue::parallel_for(range<Dims> GlobalRange, const KernelType &KernelFunc);
 						
@@ -362,20 +356,20 @@ #### With the queue shortcut
-
+						

 int square_number(int x){
 
   auto myQueue = queue{usm_selector{}};
 
   auto devicePtr = malloc_device<int>(1, myQueue);
 
-  myQueue.memcpy(devicePtr, &x, sizeof(int)).wait();
+  myQueue.memcpy(devicePtr, &x, sizeof(int)).wait();
 
   myQueue.single_task<square>([=](){
     *devicePtr = (*devicePtr) * (*devicePtr);
   }).wait();
 
-  myQueue.memcpy(&x, devicePtr, sizeof(int)).wait();
+  myQueue.memcpy(&x, devicePtr, sizeof(int)).wait();
 
   free(devicePtr, myQueue);
 

From 3b5564dddf0789446346a5439e4425a6c1c07f3e Mon Sep 17 00:00:00 2001
From: Michael Lance 
Date: Wed, 22 Jul 2026 17:30:24 -0500
Subject: [PATCH 22/34] autofix Vectors

---
 Lesson_Materials/Vectors/index.html | 40 ++++++++++++-----------------
 1 file changed, 17 insertions(+), 23 deletions(-)

diff --git a/Lesson_Materials/Vectors/index.html b/Lesson_Materials/Vectors/index.html
index d93e17c7..55919767 100644
--- a/Lesson_Materials/Vectors/index.html
+++ b/Lesson_Materials/Vectors/index.html
@@ -2,17 +2,11 @@
 
 
   
-    
-    
-    
-    
-    
+    
+    
+    
+    
+    
     
     
     
@@ -112,7 +106,7 @@
 						#### Vec class
 					
-
+						

 template <typename dataT, int numElements>
 class vec;
 						
@@ -129,7 +123,7 @@ #### Aliases
-
+						

 using float4 = vec<float, 4>;
 ...
 						
@@ -144,18 +138,18 @@ #### Vec constructors
-
+						

 auto f4 = sycl::float4{1.0f, 2.0f, 3.0f, 4.0f}; // {1.0f, 2.0f, 3.0f, 4.0f}
 						
-
+						

 auto f2 = sycl::float4{2.0f, 3.0f}; // {2.0f, 3.0f}
 auto f4 = sycl::float4{1.0f, f2, 4.0f}; // {1.0f, 2.0f, 3.0f, 4.0f}
 						
-
+						

 auto f4 = sycl::float4{0.0f};  // {0.0f, 0.0f, 0.0f, 0.0f}
 						
@@ -170,7 +164,7 @@ #### Vec operators
-
+						

 auto f4a = sycl::float4{1.0f, 2.0f, 3.0f, 4.0f}; // {1.0f, 2.0f, 3.0f, 4.0f}
 
 auto f4b = sycl::float4{2.0f}; // {2.0f, 2.0f, 2.0f, 2.0f}
@@ -188,7 +182,7 @@
 						#### Vec types
 					
-
+						

 sycl::int2
 sycl::int3
 sycl::int4
@@ -209,13 +203,13 @@
 						#### Swizzles
 					
-
+						

 auto f4 = sycl::float4{1.0f, 2.0f, 3.0f, 4.0f}; // {1.0f, 2.0f, 3.0f, 4.0f}
 auto f2 = f4.swizzle<0, 3>(); // {1.0f, 4.0f}
 						
-
+						

 auto f4 = sycl::float4{1.0f, 2.0f, 3.0f, 4.0f}; // {1.0f, 2.0f, 3.0f, 4.0f}
 f4.swizzle<1, 2>() = sycl::float2{9.0f, 9.0f}; // f4 becomes {1.0f, 9.0f, 9.0f, 4.0f}
 						
@@ -230,13 +224,13 @@ #### Simple swizzles
-
+						

 auto f4 = sycl::float4{1.0f, 2.0f, 3.0f, 4.0f}; // {1.0f, 2.0f, 3.0f, 4.0f}
 auto f2 = f4.xw(); // {1.0f, 4.0f}
 						
-
+						

 auto f4 = sycl::float4{1.0f, 2.0f, 3.0f, 4.0f}; // {1.0f, 2.0f, 3.0f, 4.0f}
 f4.yz() = sycl::float2{9.0f, 9.0f}; // f4 becomes {1.0f, 9.0f, 9.0f, 4.0f}
 						
@@ -250,7 +244,7 @@
#### Vectorized image convolution performance
-
+
![SYCL](../../Static/images/image_convolution_performance_vectorized.png "SYCL")

From caaab98138b9d741ce163081e0f0b5eed5214105 Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Thu, 23 Jul 2026 09:47:52 -0500 Subject: [PATCH 23/34] fix lessonchecker --- Scripts/LessonChecker.py | 7 ++++--- Scripts/test_LessonChecker.py | 7 +++---- 2 files changed, 7 insertions(+), 7 deletions(-) diff --git a/Scripts/LessonChecker.py b/Scripts/LessonChecker.py index de52f39b..0d118cf6 100644 --- a/Scripts/LessonChecker.py +++ b/Scripts/LessonChecker.py @@ -190,7 +190,6 @@ def handle_starttag(self, tag, attrs): case _: # Any other unescaped if State.CODE in self.state: - print(tag) self.state |= State.ERROR self._warn( f" annotations are the only tags allowed in blocks! Violation in lesson {self.lesson_name} line num: {self.getpos()} tag: {tag}" @@ -222,11 +221,13 @@ def handle_data(self, data): self.output.append(data) def handle_entityref(self, name): + print(name) match name: case "lt" | "gt" | "amp": raw_entity = f"&{name};" - case str() if m := re.search(r"lt(.+)", name): - raw_entity = f"<{m.group(1)};" + + case str() if m := re.search(r"(lt|gt|amp)(.+)", name): + raw_entity = f"&{m.group(1)};{m.group(2)}" self.state |= State.ERROR self._warn( diff --git a/Scripts/test_LessonChecker.py b/Scripts/test_LessonChecker.py index cc4b2e0c..d1010bc3 100644 --- a/Scripts/test_LessonChecker.py +++ b/Scripts/test_LessonChecker.py @@ -110,7 +110,7 @@ "extract": True, "verify": False, "fix": False, - "silent": True, + "silent": False, "expected_return": False, "expected_code": "#include", }, @@ -166,15 +166,14 @@ }, { "name": "fix_partial_escape_sequences", - "body": '
&
', + "body": "
<1
", "extract": False, "verify": False, "fix": True, "silent": False, "expected_return": True, - "expected_output": '
&
', + "expected_output": "
<1
", }, - ] From 18c55e8003821247f6eae362ee9d8a4fe060d8c0 Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Thu, 23 Jul 2026 09:57:38 -0500 Subject: [PATCH 24/34] remove random
--- Lesson_Materials/Data_Parallelism/index.html | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/Lesson_Materials/Data_Parallelism/index.html b/Lesson_Materials/Data_Parallelism/index.html index 368823fe..79a1c97f 100644 --- a/Lesson_Materials/Data_Parallelism/index.html +++ b/Lesson_Materials/Data_Parallelism/index.html @@ -251,7 +251,7 @@ * Runtime may not always have enough information to choose the best-performing size
- ND-Range
+ ND-Range SYCL-Range
From cd9d97f04abcfb0be17a23668f15d7a3b4de7a95 Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Thu, 23 Jul 2026 10:07:04 -0500 Subject: [PATCH 25/34] remove debug statement from lesson checker --- Lesson_Materials/Fast_Track/index.html | 98 ++++++++++++++------------ Scripts/LessonChecker.py | 1 - 2 files changed, 52 insertions(+), 47 deletions(-) diff --git a/Lesson_Materials/Fast_Track/index.html b/Lesson_Materials/Fast_Track/index.html index 5fe199a4..30ed5bbd 100644 --- a/Lesson_Materials/Fast_Track/index.html +++ b/Lesson_Materials/Fast_Track/index.html @@ -2,11 +2,17 @@ - - - - - + + + + + @@ -25,7 +31,7 @@
-
+
## Learning Objectives
@@ -152,7 +158,7 @@
-

+							
 auto gpuDevice = device(gpu_selector_v); 
 							
@@ -161,7 +167,7 @@
- * A device selector takes a parameter of type `const device &` and gives it a "score". + * A device selector takes a parameter of type `const device &` and gives it a "score". * Used to query all devices and return the one with the highest "score". * A device with a negative score will never be chosen.
@@ -222,12 +228,12 @@ * Some computation on the GPU. * Migration of the result back to the CPU.
- +
- #### SYCL Buffers & Accessors + #### SYCL Buffers & Accessors
* The buffer/accessor model separates the storage and access of data @@ -240,7 +246,7 @@
- #### SYCL Buffers & Accessors + #### SYCL Buffers & Accessors
@@ -255,7 +261,7 @@
- #### SYCL Buffers & Accessors + #### SYCL Buffers & Accessors
@@ -274,17 +280,17 @@
-

-buffer<float;, 1> bufA(dA.data(), range<1;>(dA.size())); 
-buffer<float;, 1> bufB(dB.data(), range<1;>(dB.size())); 
-buffer<float;, 1> bufO(dO.data(), range<1;>(dO.size()));
+							
+buffer<float, 1> bufA(dA.data(), range<1>(dA.size())); 
+buffer<float, 1> bufB(dB.data(), range<1>(dB.size())); 
+buffer<float, 1> bufO(dO.data(), range<1>(dO.size()));
 
-gpuQueue.submit([&](handler &cgh){
+gpuQueue.submit([&](handler &cgh){
   sycl::accessor inA{bufA, cgh, sycl::read_only};
   sycl::accessor inB{bufB, cgh, sycl::read_only};
   sycl::accessor out{bufO, cgh, sycl::write_only};
-  cgh.parallel_for<add;>(range<1;>(dA.size()), 
-    [=](id<1;> i){ 
+  cgh.parallel_for<add>(range<1>(dA.size()), 
+    [=](id<1> i){ 
     out[i] = inA[i] + inB[i];
   });
 });
@@ -303,11 +309,11 @@
 						#### USM: Malloc_device
 					
-

-void* malloc_device(size_t numBytes, const queue& syclQueue, const property_list &propList = {});
+						
+void* malloc_device(size_t numBytes, const queue& syclQueue, const property_list &propList = {});
 
 template <typename T>
-T* malloc_device(size_t count, const queue& syclQueue, const property_list &propList = {});
+T* malloc_device(size_t count, const queue& syclQueue, const property_list &propList = {});
 						
@@ -324,8 +330,8 @@ #### USM: Free
-

-void free(void* ptr, queue& syclQueue);
+						
+void free(void* ptr, queue& syclQueue);
 						
@@ -341,8 +347,8 @@ #### USM: Memcpy
-

-event queue::memcpy(void* dest, const void* src, size_t numBytes, const std::vector<event> &depEvents);
+						
+event queue::memcpy(void* dest, const void* src, size_t numBytes, const std::vector &depEvents);
 						
@@ -428,24 +434,24 @@
-

+							
 							
-cgh.parallel_for<kernel;>(range<1;>(1024), 
-  [=](id<1;> idx){
+cgh.parallel_for<kernel>(range<1>(1024), 
+  [=](id<1> idx){
     /* kernel function code */
 });
 							
-

+							
 							
-cgh.parallel_for<kernel;>(range<1;>(1024), 
-  [=](item<1;> item){
+cgh.parallel_for<kernel>(range<1>(1024), 
+  [=](item<1> item){
     /* kernel function code */
 });
 							
-

+							
 							
-cgh.parallel_for<kernel;>(nd_range<1;>(range<1;>(1024), 
-  range<1;>(32)),[=](nd_item<1;> ndItem){
+cgh.parallel_for<kernel>(nd_range<1>(range<1>(1024), 
+  range<1>(32)),[=](nd_item<1> ndItem){
     /* kernel function code */
 });
 							
@@ -483,21 +489,21 @@
#### Function object
-
-

+					
+ ```c++ class MyKernel { - sycl::accessor<float> input_; + sycl::accessor input_; float* output_; - MyKernel(sycl::buffer<float> buf, float* output, sycl::handler& h) + MyKernel(sycl::buffer buf, float* output, sycl::handler& h) : input{buf.get_access(h)}, output_{output} {} // const is required - void operator()(sycl::item<1> i) const { + void operator()(sycl::item<1> i) const { ; // computation here } }; -
+ ```
The members are accessible on the device inside the function call @@ -509,20 +515,20 @@
#### Lambda function
-
-

-sycl::buffer<float, 1> buf = /* normal init */;
+					
+ ```c++ +sycl::buffer buf = /* normal init */; float * output = sycl::malloc_device(/* params */); ... queue submit as normal ... auto acc = buf.get_access(h); -auto func = [=](sycl::item<1> i) { +auto func = [=](sycl::item<1> i) { acc[i] = someVal; output[i.get_global_linear_id()] = someOtherVal; }); handler.parallel_for(range, func); -
+ ```
The variables used implicitly are captured by value and are usable in diff --git a/Scripts/LessonChecker.py b/Scripts/LessonChecker.py index 0d118cf6..d2009a35 100644 --- a/Scripts/LessonChecker.py +++ b/Scripts/LessonChecker.py @@ -221,7 +221,6 @@ def handle_data(self, data): self.output.append(data) def handle_entityref(self, name): - print(name) match name: case "lt" | "gt" | "amp": raw_entity = f"&{name};" From e33fb33bfa1d9c4037799aaf9b601095b1ce189b Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Thu, 23 Jul 2026 10:07:47 -0500 Subject: [PATCH 26/34] re run fast track autofix --- Lesson_Materials/Fast_Track/index.html | 82 ++++++++++++-------------- 1 file changed, 38 insertions(+), 44 deletions(-) diff --git a/Lesson_Materials/Fast_Track/index.html b/Lesson_Materials/Fast_Track/index.html index 30ed5bbd..229a7ff7 100644 --- a/Lesson_Materials/Fast_Track/index.html +++ b/Lesson_Materials/Fast_Track/index.html @@ -2,17 +2,11 @@ - - - - - + + + + + @@ -31,7 +25,7 @@
-
+
## Learning Objectives
@@ -158,7 +152,7 @@
-
+							

 auto gpuDevice = device(gpu_selector_v); 
 							
@@ -167,7 +161,7 @@
- * A device selector takes a parameter of type `const device &` and gives it a "score". + * A device selector takes a parameter of type `const device &` and gives it a "score". * Used to query all devices and return the one with the highest "score". * A device with a negative score will never be chosen.
@@ -228,12 +222,12 @@ * Some computation on the GPU. * Migration of the result back to the CPU.
- +
- #### SYCL Buffers & Accessors + #### SYCL Buffers & Accessors
* The buffer/accessor model separates the storage and access of data @@ -246,7 +240,7 @@
- #### SYCL Buffers & Accessors + #### SYCL Buffers & Accessors
@@ -261,7 +255,7 @@
- #### SYCL Buffers & Accessors + #### SYCL Buffers & Accessors
@@ -280,17 +274,17 @@
-
-buffer<float, 1> bufA(dA.data(), range<1>(dA.size())); 
-buffer<float, 1> bufB(dB.data(), range<1>(dB.size())); 
-buffer<float, 1> bufO(dO.data(), range<1>(dO.size()));
+							

+buffer<float, 1> bufA(dA.data(), range<1>(dA.size())); 
+buffer<float, 1> bufB(dB.data(), range<1>(dB.size())); 
+buffer<float, 1> bufO(dO.data(), range<1>(dO.size()));
 
-gpuQueue.submit([&](handler &cgh){
+gpuQueue.submit([&](handler &cgh){
   sycl::accessor inA{bufA, cgh, sycl::read_only};
   sycl::accessor inB{bufB, cgh, sycl::read_only};
   sycl::accessor out{bufO, cgh, sycl::write_only};
-  cgh.parallel_for<add>(range<1>(dA.size()), 
-    [=](id<1> i){ 
+  cgh.parallel_for<add>(range<1>(dA.size()), 
+    [=](id<1> i){ 
     out[i] = inA[i] + inB[i];
   });
 });
@@ -309,11 +303,11 @@
 						#### USM: Malloc_device
 					
-
-void* malloc_device(size_t numBytes, const queue& syclQueue, const property_list &propList = {});
+						

+void* malloc_device(size_t numBytes, const queue& syclQueue, const property_list &propList = {});
 
 template <typename T>
-T* malloc_device(size_t count, const queue& syclQueue, const property_list &propList = {});
+T* malloc_device(size_t count, const queue& syclQueue, const property_list &propList = {});
 						
@@ -330,8 +324,8 @@ #### USM: Free
-
-void free(void* ptr, queue& syclQueue);
+						

+void free(void* ptr, queue& syclQueue);
 						
@@ -347,8 +341,8 @@ #### USM: Memcpy
-
-event queue::memcpy(void* dest, const void* src, size_t numBytes, const std::vector &depEvents);
+						

+event queue::memcpy(void* dest, const void* src, size_t numBytes, const std::vector<event> &depEvents);
 						
@@ -434,24 +428,24 @@
-
+							

 							
-cgh.parallel_for<kernel>(range<1>(1024), 
-  [=](id<1> idx){
+cgh.parallel_for<kernel>(range<1>(1024), 
+  [=](id<1> idx){
     /* kernel function code */
 });
 							
-
+							

 							
-cgh.parallel_for<kernel>(range<1>(1024), 
-  [=](item<1> item){
+cgh.parallel_for<kernel>(range<1>(1024), 
+  [=](item<1> item){
     /* kernel function code */
 });
 							
-
+							

 							
-cgh.parallel_for<kernel>(nd_range<1>(range<1>(1024), 
-  range<1>(32)),[=](nd_item<1> ndItem){
+cgh.parallel_for<kernel>(nd_range<1>(range<1>(1024), 
+  range<1>(32)),[=](nd_item<1> ndItem){
     /* kernel function code */
 });
 							
@@ -495,11 +489,11 @@ sycl::accessor input_; float* output_; - MyKernel(sycl::buffer buf, float* output, sycl::handler& h) + MyKernel(sycl::buffer buf, float* output, sycl::handler& h) : input{buf.get_access(h)}, output_{output} {} // const is required - void operator()(sycl::item<1> i) const { + void operator()(sycl::item<1> i) const { ; // computation here } }; @@ -523,7 +517,7 @@ ... queue submit as normal ... auto acc = buf.get_access(h); -auto func = [=](sycl::item<1> i) { +auto func = [=](sycl::item<1> i) { acc[i] = someVal; output[i.get_global_linear_id()] = someOtherVal; }); From 10eb4f0ca3588e0eb6de2f643ebfb76a2bed1dfb Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Thu, 23 Jul 2026 13:26:49 -0500 Subject: [PATCH 27/34] check for missing data-line-numbers --- Scripts/LessonChecker.py | 5 +++++ Scripts/test_LessonChecker.py | 22 +++++++++++----------- 2 files changed, 16 insertions(+), 11 deletions(-) diff --git a/Scripts/LessonChecker.py b/Scripts/LessonChecker.py index d2009a35..155043d6 100644 --- a/Scripts/LessonChecker.py +++ b/Scripts/LessonChecker.py @@ -152,6 +152,11 @@ def handle_starttag(self, tag, attrs): ) attrs.remove(("class", "code-100pc")) + if not any(attr[0] == "data-line-numbers" for attr in attrs): + self.state |= State.ERROR + self._warn(f"data-line-numbers missing in lesson {self.lesson_name} line num: {self.getpos()[0]}") + attrs.append(("data-line-numbers", None)) + self.code_blocks.append((self.getpos()[0], "")) case "mark": diff --git a/Scripts/test_LessonChecker.py b/Scripts/test_LessonChecker.py index d1010bc3..7851a673 100644 --- a/Scripts/test_LessonChecker.py +++ b/Scripts/test_LessonChecker.py @@ -23,7 +23,7 @@ }, { "name": "detect_mark_in_code", - "body": '
message
',
+        "body": '
message
',
         "extract": False,
         "verify": True,
         "fix": False,
@@ -86,7 +86,7 @@
     },
     {
         "name": "extract_code_block",
-        "body": '
int main() { return 0; }
', + "body": '
int main() { return 0; }
', "extract": True, "verify": False, "fix": False, @@ -106,7 +106,7 @@ }, { "name": "extract_proper_lt_and_gt", - "body": '
#include<sycl/sycl.hpp>
', + "body": '
#include<sycl/sycl.hpp>
', "extract": True, "verify": False, "fix": False, @@ -122,37 +122,37 @@ "fix": True, "silent": True, "expected_return": True, - "expected_output": '
', + "expected_output": '
', }, { "name": "fix_improper_escaped_ref", - "body": '
&evt
', + "body": '
&evt
', "extract": False, "verify": False, "fix": True, "silent": True, "expected_return": True, - "expected_output": '
&evt
', + "expected_output": '
&evt
', }, { "name": "fix_improper_escaped_gators", - "body": '
#include
', + "body": '
#include
', "extract": False, "verify": False, "fix": True, "silent": False, "expected_return": True, - "expected_output": '
#include<sycl/sycl.hpp>
', + "expected_output": '
#include<sycl/sycl.hpp>
', }, { "name": "fix_improper_escaped_ampersand", - "body": '
&
', + "body": '
&
', "extract": False, "verify": False, "fix": True, "silent": False, "expected_return": True, - "expected_output": '
&
', + "expected_output": '
&
', }, { "name": "fix_old_styling_in_code", @@ -162,7 +162,7 @@ "fix": True, "silent": False, "expected_return": True, - "expected_output": '
', + "expected_output": '
', }, { "name": "fix_partial_escape_sequences", From 3831f0338b1c8cb2d0a807ffb10c291e77e8a878 Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Thu, 23 Jul 2026 13:28:57 -0500 Subject: [PATCH 28/34] add line numbers to all lessons --- .../Advanced_Data_Flow/index.html | 30 +- Lesson_Materials/Buffer_Accessor/index.html | 304 ++++++++++++++++++ .../Data_and_Dependencies/index.html | 20 +- Lesson_Materials/Device_Discovery/index.html | 20 +- .../Enqueueing_a_Kernel/index.html | 28 +- Lesson_Materials/Fast_Track/index.html | 16 +- Lesson_Materials/Image_Convolution/index.html | 4 +- .../Introduction_to_USM/index.html | 2 +- .../Local_Memory_Tiling/index.html | 4 +- Lesson_Materials/Managing_Data/index.html | 16 +- Lesson_Materials/Matrix_Transpose/index.html | 4 +- .../More_SYCL_Features/index.html | 6 +- Lesson_Materials/Multiple_Devices/index.html | 6 +- Lesson_Materials/Using_USM/index.html | 26 +- Lesson_Materials/Vectors/index.html | 22 +- Lesson_Materials/oneMath_gemm/index.html | 4 +- 16 files changed, 405 insertions(+), 107 deletions(-) create mode 100644 Lesson_Materials/Buffer_Accessor/index.html diff --git a/Lesson_Materials/Advanced_Data_Flow/index.html b/Lesson_Materials/Advanced_Data_Flow/index.html index 22f47c87..10d96720 100644 --- a/Lesson_Materials/Advanced_Data_Flow/index.html +++ b/Lesson_Materials/Advanced_Data_Flow/index.html @@ -2,17 +2,11 @@ - - - - - + + + + + @@ -61,7 +55,7 @@ #### Initial data (USM)
-

+            

 auto devicePtr = sycl::malloc_device(sizeInBytes, gpuQueue);
 
 gpuQueue.memcpy(devicePtr, initialData, sizeInBytes).wait();
@@ -80,7 +74,7 @@
             #### Initial data (USM)
           
-

+            

 auto sharedData = sycl::malloc_shared(sizeInBytes, gpuQueue);
 						
@@ -96,7 +90,7 @@ #### Initial data (buffer/accessor)
-

+            

 auto buf = sycl::buffer{initialData, sycl::range{size}};
 						
@@ -115,7 +109,7 @@ #### Use_host_pointer property
-

+            

 auto buf = sycl::buffer{initialData, sycl::range{size},
   {sycl::property::buffer::use_host_ptr{}}};
 						
@@ -146,7 +140,7 @@ #### Set_final_data
-

+            

 auto buf = sycl::buffer{initialData, sycl::range{size}};
 
 buf.set_final_data(finalData);
@@ -165,7 +159,7 @@
             #### Set_final_data
           
-

+            

 auto buf = sycl::buffer{initialData, sycl::range{size}};
 
 buf.set_final_data(nullptr);
@@ -186,7 +180,7 @@
             #### Uninitialized buffers (buffer/accessor)
           
-

+            

 auto buf = sycl::buffer<T>{sycl::range{size}};
 						
diff --git a/Lesson_Materials/Buffer_Accessor/index.html b/Lesson_Materials/Buffer_Accessor/index.html new file mode 100644 index 00000000..e4bb8d58 --- /dev/null +++ b/Lesson_Materials/Buffer_Accessor/index.html @@ -0,0 +1,304 @@ + + + + + + + + + + + + + + + + + +
+
+ +
+
+ ## The Buffer Accessor Model +
+
+ +
+
+ ## Learning Objectives + * Learn about the buffer/accessor model + * Learn how to manage data with the buffer/accessor model +
+
+ +
+
+ #### SYCL Buffers & Accessors +
+
+ * The buffer/accessor model separates the storage and access of data + * A SYCL buffer manages data across the host and any number of devices + * A SYCL accessor requests access to data on the host or on a device for a specific SYCL kernel function + * Accessors are also used to access data within a SYCL kernel function + * This means they are declared in the host code but captured by and then accessed within a SYCL kernel function +
+
+ +
+
+ #### SYCL Buffers & Accessors +
+
+
+ * A SYCL buffer can be constructed with a pointer to host memory + * For the lifetime of the buffer this memory is owned by the SYCL runtime + * When a buffer object is constructed it will not allocate or copy to device memory at first + * This will only happen once the SYCL runtime knows the data needs to be accessed and where it needs to be accessed +
+
+ ![Buffer Host Memory](../../Static/images/buffer-hostmemory.png "Buffer Host Memory") +
+
+
+ +
+
+ #### SYCL Buffers & Accessors +
+
+
+ * Constructing an accessor specifies a request to access the data managed by the buffer + * There are a range of different types of accessor which provide different ways to access data +
+
+ ![Buffer Host Memory Accessor](../../Static/images/buffer-hostmemory-accessor.png "Buffer Host Memory Accessor") +
+
+
+ +
+
+ #### SYCL Buffers & Accessors +
+
+
+ * When an accessor is constructed it is associated with a command group via the handler object + * This connects the buffer that is being accessed, the way in which it’s being accessed and the device that the command group is being submitted to +
+
+ ![Buffer Host Memory Accessor CG](../../Static/images/buffer-hostmemory-accessor-cg.png "Buffer Host Memory Accessor CG") +
+
+
+ +
+
+ #### SYCL Buffers & Accessors +
+
+
+ * Once the SYCL scheduler selects the command group to be executed it must first satisfy its data dependencies + * If necessary, this includes allocating and copying the data to the device accessing that data + * If the most recent copy of the data is already on the device then the runtime will not copy again +
+
+ ![Buffer Host Memory Accessor CG Device](../../Static/images/buffer-hostmemory-accessor-cg-device.png "Buffer Host Memory Accessor CG Device") +
+
+
+ +
+
+ #### SYCL Buffers & Accessors +
+
+
+ * Data will remain in device memory after kernels finish executing until another accessor requests access in a different device or on the host + * When the buffer object is destroyed it will wait for any outstanding work that is accessing the data to complete and then copy back to the original host memory +
+
+ ![Buffer Destroyed](../../Static/images/buffer-destroyed.png "Buffer Destroyed") +
+
+
+ +
+
+ #### SYCL Buffers & Accessors +
+
+

+#include<sycl/sycl.hpp> 
+
+int main() {
+  int var = 42; 
+
+  auto buf = sycl::buffer{&var, sycl::range<1>{1}};
+
+}
+T var = 42;
+
+{
+  // Create buffer pointing to var.
+  auto buf = sycl::buffer{&var, sycl::range<1>{1}};
+
+  // ...
+  // Do some computation on device. Use accessors to access buffer
+  // ...
+  
+} // var updated here
+
+assert(var != 42);
+						
+
+
+ * A buffer is associated with a type, range and + dimensionality. Dimensionality must be either 1, 2 or + 3. + * Usually type and dimensionality can be inferred using + CTAD. + * If a buffer is associated with some allocation in host + memory, the host memory will be updated only once the + buffer goes out of scope. +
+
+
+
+ #### Accessor class +
+
+ * There are many different ways to use the `accessor` + class. + * Accessing data on a device. + * Accessing data immediately in the host application. + * Allocating local memory. + * For now we are going to focus on accessing data on a + device. +
+
+ +
+
+ #### Constructing an accessor +
+
+

+auto acc = sycl::accessor{bufA, cgh};
+						
+
+
+ * There are many ways to construct an `accessor`. + * Accessors are complicated templates with many template arguments, + representing the data type, dimensionality, access mode etc. + * The `accessor` class supports CTAD so it's not + necessary to specify all of the template arguments. + * The most common way to construct an `accessor` is from + a `buffer` and a `handler` associated with the command + group function you are within. + * The element type and dimensionality are inferred from + the `buffer`. + * The `access_mode` is defaulted to + `access_mode::read_write`. +
+
+ +
+
+ #### Specifying the access mode +
+
+

+auto readAcc = sycl::accessor{bufA, cgh, sycl::read_only};
+auto writeAcc = sycl::accessor{bufB, cgh, sycl::write_only};
+						
+
+
+ * When constructing an `accessor` you will likely also + want to specify the `access_mode` + * You can do this by passing one of the CTAD tags: + * `read_only` will result in `access_mode::read`. + * `write_only` will result in `access_mode::write`. +
+
+ +
+
+ #### Specifying no initialization +
+
+

+auto acc = sycl::accessor{buf, cgh, sycl::no_init};
+						
+
+
+ * When constructing an `accessor` you may also want to + discard the original data of a `buffer`. + * You can do this by passing the `no_init` property. +
+
+ +
+
+ #### Using Accessors +
+
+

+T var = 42;
+
+{
+  // Create buffer pointing to var.
+  auto bufA = sycl::buffer{&var, sycl::range<1>{1}};
+  auto bufB = sycl::buffer{&var, sycl::range<1>{1}};
+
+  q.submit([&](sycl::handler &cgh) {
+	auto accA = sycl::accessor{bufA, cgh, sycl::read_only};
+	auto accB = sycl::accessor{bufA, cgh, sycl::no_init};
+
+  cgh.single_task<mykernel>(...); // Do some work
+  });
+  
+} // var updated here
+
+assert(var != 42);
+						
+
+
+ * Buffers and accessors take care of memory + migration, as well as dependency analysis. + * More to come later on dependencies. +
+
+ +
+
+ #### operator[] +
+
+

+gpuQueue.submit([&](handler &cgh){
+  auto inA = sycl::accessor{bufA, cgh, sycl::read_only};
+  auto inB = sycl::accessor{bufB, cgh, sycl::read_only};
+  auto out = sycl::accessor{bufO, cgh, sycl::write_only};
+  cgh.single_task<mykernel>([=]{
+    out[0] = inA[0] + inB[0];
+  }); 
+});
+						
+
+
+ * As well as specifying data dependencies an `accessor` + can also be used to access the data from within a kernel + function. + * You can do this by calling `operator[]` on the + `accessor`. + * `operator[]` for USM pointers must take a `size_t`, + whereas `operator[]` for accessors can take a + multi-dimensional `sycl::id` or a `size_t`. +
+
+
+
+ + + diff --git a/Lesson_Materials/Data_and_Dependencies/index.html b/Lesson_Materials/Data_and_Dependencies/index.html index d0c4910e..a43af4ea 100644 --- a/Lesson_Materials/Data_and_Dependencies/index.html +++ b/Lesson_Materials/Data_and_Dependencies/index.html @@ -111,7 +111,7 @@
-

+							

 sycl::buffer buf {data, sycl::range{1024}};
 
 gpuQueue.submit([&](sycl::handler &cgh) {
@@ -149,7 +149,7 @@
 					
-

+							

 sycl::buffer buf {data, sycl::range{1024}};
 
 gpuQueue.submit([&](sycl::handler &cgh) {
@@ -188,7 +188,7 @@
 					
-

+							

 sycl::buffer buf {data, sycl::range{1024}};
 
 gpuQueue.submit([&](sycl::handler &cgh) {
@@ -226,7 +226,7 @@
 					
-

+							

 buf = sycl::buffer(data, sycl::range{1024});
 
 gpuQueue.submit([&](sycl::handler &cgh) {
@@ -263,7 +263,7 @@
 					
-

+							

 auto devicePtr =
   sycl::malloc_device<int>(1024, gpuQueue);
 
@@ -302,7 +302,7 @@
 					
-

+							

 auto devicePtr =
   sycl::malloc_device<int>(1024, gpuQueue);
 
@@ -337,7 +337,7 @@
 					
-

+							

 auto devicePtr =
   sycl::malloc_device<int>(1024, gpuQueue);
 
@@ -390,7 +390,7 @@
 					
-

+							

 sycl::buffer bufA {dataA, sycl::range{1024}};
 sycl::buffer bufB {dataB, sycl::range{1024}};
 
@@ -429,7 +429,7 @@
 						#### Concurrent data flow with USM
 					
-

+						

 auto devicePtrA = sycl::malloc_device<int>(1024, gpuQueue);
 auto devicePtrB = sycl::malloc_device<int>(1024, gpuQueue);
 
@@ -463,7 +463,7 @@
 						#### Concurrent data flow with USM
 					
-

+						

 auto devicePtrA = sycl::malloc_device<int>(1024, gpuQueue);
 auto devicePtrB = sycl::malloc_device<int>(1024, gpuQueue);
 
diff --git a/Lesson_Materials/Device_Discovery/index.html b/Lesson_Materials/Device_Discovery/index.html
index 613e9e09..eb2c5aae 100644
--- a/Lesson_Materials/Device_Discovery/index.html
+++ b/Lesson_Materials/Device_Discovery/index.html
@@ -88,7 +88,7 @@
 					
-

+							

 auto platforms = platform::get_platforms();
 
 							
@@ -109,7 +109,7 @@
-

+							

 auto intelDevices = intelPlatform.get_devices();
 
 							
@@ -131,7 +131,7 @@
-

+							

 auto devices = device::get_devices();
 
 							
@@ -170,7 +170,7 @@
-

+							

 auto gpuDevice = device(gpu_selector_v); 
 							
@@ -191,7 +191,7 @@
-

+							

 auto chosenDevice = device();
 auto chosenDevice = device(default_selector_v);
 							
@@ -212,7 +212,7 @@ #### Creating a custom device selector
-

+						

 int my_gpu_selector(const device& dev) {
 
 }
@@ -229,7 +229,7 @@
 						#### Creating a custom device selector
 					
-

+						

 int my_gpu_selector(const device& dev) {
   if (dev.is_gpu()){
     return 1;
@@ -251,7 +251,7 @@
 						#### Creating a custom device selector
 					
-

+						

 int my_gpu_selector(const device& dev) {
   if (dev.is_gpu()){
     return 1;
@@ -278,7 +278,7 @@
 					
-

+							

 auto plt = dev.get_platform();
 auto platformName
   = dev.get_info<info::device::name>();
@@ -301,7 +301,7 @@
 					
-

+							

 bool supportsFp16 = dev.has(aspect::fp16);
 							
diff --git a/Lesson_Materials/Enqueueing_a_Kernel/index.html b/Lesson_Materials/Enqueueing_a_Kernel/index.html index ff09e468..d5e6ec14 100644 --- a/Lesson_Materials/Enqueueing_a_Kernel/index.html +++ b/Lesson_Materials/Enqueueing_a_Kernel/index.html @@ -103,7 +103,7 @@
-

+							

 gpuQueue.submit([&](handler &cgh){
   
   /* Command group function */
@@ -125,7 +125,7 @@
 					
-

+							

 gpuQueue.submit([&](handler &cgh){
   
   /* Command group function */
@@ -147,7 +147,7 @@
 					
-

+							

 gpuQueue.submit([&](handler &cgh){
 
   /* Command group function */
@@ -199,7 +199,7 @@
 					
-

+							

 class my_kernel;
 
 gpuQueue.submit([&](handler &cgh){
@@ -222,7 +222,7 @@
 				
-

+							

 class my_kernel;
 
 gpuQueue.submit([&](handler &cgh){
@@ -244,7 +244,7 @@
 				
-

+							

 class my_kernel;
 
 gpuQueue.submit([&](handler &cgh){
@@ -266,7 +266,7 @@
 				
-

+							

 class my_kernel;
 
 gpuQueue.submit([&](handler &cgh){
@@ -315,7 +315,7 @@
 					
-

+							

 class my_kernel;
 
 queue gpuQueue;
@@ -339,7 +339,7 @@
 					
-

+							

 struct my_kernel { 
   void operator()() const {
     /* kernel function */
@@ -361,14 +361,14 @@
 					
-

+							

 struct my_kernel { 
   void operator()() const {
     /* kernel function */
   }
 };
 							
-

+							

 queue gpuQueue;
 gpuQueue.submit([&](handler &cgh){
 								
@@ -399,7 +399,7 @@
 						#### Streams
 					
-

+						

 sycl::stream(size_t bufferSize, size_t workItemBufferSize, handler &cgh);
 						
@@ -417,7 +417,7 @@
-

+							

 class my_kernel;
 
 queue gpuQueue;
@@ -444,7 +444,7 @@
 					
-

+							

 class my_kernel;
 
 queue gpuQueue;
diff --git a/Lesson_Materials/Fast_Track/index.html b/Lesson_Materials/Fast_Track/index.html
index 229a7ff7..afe5e1a3 100644
--- a/Lesson_Materials/Fast_Track/index.html
+++ b/Lesson_Materials/Fast_Track/index.html
@@ -152,7 +152,7 @@
 					
-

+							

 auto gpuDevice = device(gpu_selector_v); 
 							
@@ -274,7 +274,7 @@
-

+							

 buffer<float, 1> bufA(dA.data(), range<1>(dA.size())); 
 buffer<float, 1> bufB(dB.data(), range<1>(dB.size())); 
 buffer<float, 1> bufO(dO.data(), range<1>(dO.size()));
@@ -303,7 +303,7 @@
 						#### USM: Malloc_device
 					
-

+						

 void* malloc_device(size_t numBytes, const queue& syclQueue, const property_list &propList = {});
 
 template <typename T>
@@ -324,7 +324,7 @@
 						#### USM: Free
 					
-

+						

 void free(void* ptr, queue& syclQueue);
 						
@@ -341,7 +341,7 @@ #### USM: Memcpy
-

+						

 event queue::memcpy(void* dest, const void* src, size_t numBytes, const std::vector<event> &depEvents);
 						
@@ -428,21 +428,21 @@
-

+							

 							
 cgh.parallel_for<kernel>(range<1>(1024), 
   [=](id<1> idx){
     /* kernel function code */
 });
 							
-

+							

 							
 cgh.parallel_for<kernel>(range<1>(1024), 
   [=](item<1> item){
     /* kernel function code */
 });
 							
-

+							

 							
 cgh.parallel_for<kernel>(nd_range<1>(range<1>(1024), 
   range<1>(32)),[=](nd_item<1> ndItem){
diff --git a/Lesson_Materials/Image_Convolution/index.html b/Lesson_Materials/Image_Convolution/index.html
index 21f02d81..8ca99947 100644
--- a/Lesson_Materials/Image_Convolution/index.html
+++ b/Lesson_Materials/Image_Convolution/index.html
@@ -146,7 +146,7 @@
 						#### Input/output image locations
 					
- +
 auto inputImageFile = "../Code_Exercises/Images/dogs.png";
 auto outputImageFile = ../Code_Exercises/Images/blurred_dogs.png";
@@ -166,7 +166,7 @@
 						#### Convolution filters
 					
- +
 auto filter = util::generate_filter(util::filter_type filterType, int width);
 						
diff --git a/Lesson_Materials/Introduction_to_USM/index.html b/Lesson_Materials/Introduction_to_USM/index.html index 6afaff6c..1522463b 100644 --- a/Lesson_Materials/Introduction_to_USM/index.html +++ b/Lesson_Materials/Introduction_to_USM/index.html @@ -152,7 +152,7 @@ Each SYCL platform and its device(s) will support different variants of USM and different kinds of memory allocation
-
if (dev.has(sycl::aspect::usm_device_allocations))
+
if (dev.has(sycl::aspect::usm_device_allocations))
diff --git a/Lesson_Materials/Local_Memory_Tiling/index.html b/Lesson_Materials/Local_Memory_Tiling/index.html index c729287a..2bc9912c 100644 --- a/Lesson_Materials/Local_Memory_Tiling/index.html +++ b/Lesson_Materials/Local_Memory_Tiling/index.html @@ -97,7 +97,7 @@ #### Local accessors
-

+						

 auto scratchpad = sycl::accessor<int, 1, sycl::access::target::local>(sycl::range{workGroupSize}, cgh);
 						
@@ -238,7 +238,7 @@ #### Group_barrier
-

+						

 sycl::group_barrier(item.get_group());
 						
diff --git a/Lesson_Materials/Managing_Data/index.html b/Lesson_Materials/Managing_Data/index.html index 88fb8b80..13044605 100644 --- a/Lesson_Materials/Managing_Data/index.html +++ b/Lesson_Materials/Managing_Data/index.html @@ -89,7 +89,7 @@ #### Using USM - Malloc Device
-

+						

 // Allocate memory on device
 T *device_ptr = sycl::malloc_device<T>(n, myQueue);
 
@@ -117,7 +117,7 @@
 						#### Using USM - Malloc Shared
 					
-

+						

 // Allocate shared memory 
 T *shared_ptr = sycl::malloc_shared<T>(n, myQueue);
 
@@ -242,7 +242,7 @@
 						#### SYCL Buffers & Accessors
 					
-

+						

 T var = 42;
 
 {
@@ -289,7 +289,7 @@
 						#### Constructing an accessor
 					
-

+						

 auto acc = sycl::accessor{bufA, cgh};
 						
@@ -314,7 +314,7 @@ #### Specifying the access mode
-

+						

 auto readAcc = sycl::accessor{bufA, cgh, sycl::read_only};
 auto writeAcc = sycl::accessor{bufB, cgh, sycl::write_only};
 						
@@ -333,7 +333,7 @@ #### Specifying no initialization
-

+						

 auto acc = sycl::accessor{buf, cgh, sycl::no_init};
 						
@@ -349,7 +349,7 @@ #### Using Accessors
-

+						

 T var = 42;
 
 {
@@ -381,7 +381,7 @@
 						#### operator[]
 					
-

+						

 gpuQueue.submit([&](handler &cgh){
   auto inA = sycl::accessor{bufA, cgh, sycl::read_only};
   auto inB = sycl::accessor{bufB, cgh, sycl::read_only};
diff --git a/Lesson_Materials/Matrix_Transpose/index.html b/Lesson_Materials/Matrix_Transpose/index.html
index bf78a092..7daaa44f 100644
--- a/Lesson_Materials/Matrix_Transpose/index.html
+++ b/Lesson_Materials/Matrix_Transpose/index.html
@@ -177,7 +177,7 @@
 						#### Local accessors
 					
-

+						

 auto scratchpad = sycl::local_accessor<int, dims>(sycl::range{workGroupSize}, cgh);
 						
@@ -318,7 +318,7 @@ #### Group_barrier
-

+						

 sycl::group_barrier(item.get_group());
 						
diff --git a/Lesson_Materials/More_SYCL_Features/index.html b/Lesson_Materials/More_SYCL_Features/index.html index d9b5451d..6bacd8c3 100644 --- a/Lesson_Materials/More_SYCL_Features/index.html +++ b/Lesson_Materials/More_SYCL_Features/index.html @@ -221,7 +221,7 @@
-

+							

 q.parallel_for([=](sycl::item<1> it) {
   sycl::atomic_ref<T,
           sycl::memory_order_relaxed,
@@ -282,7 +282,7 @@
 					
-

+							

 q.submit([&](sycl::handler &cgh) {
   // Output of reduction will be in ptr
   auto sumReduction = sycl::reduction(ptr,
@@ -309,7 +309,7 @@
 					
-

+							

 q.submit([&](sycl::handler &cgh) {
   // Output of reduction will be in ptr
   auto maxReduction = sycl::reduction(ptr,
diff --git a/Lesson_Materials/Multiple_Devices/index.html b/Lesson_Materials/Multiple_Devices/index.html
index 1ab4e0ae..00d546bf 100644
--- a/Lesson_Materials/Multiple_Devices/index.html
+++ b/Lesson_Materials/Multiple_Devices/index.html
@@ -96,7 +96,7 @@
 						#### Creating an implicit context
 					
-

+						

 auto defaultQueue = queue{};
 						
@@ -112,7 +112,7 @@ #### Creating a context from devices
-

+						

 auto sharedContext = context{{cpuDevice, gpuDevice}};
 						
@@ -127,7 +127,7 @@ #### Creating a context from a platform
-

+						

 auto sharedContext = context{intelPlatform};
 						
diff --git a/Lesson_Materials/Using_USM/index.html b/Lesson_Materials/Using_USM/index.html index 104a5268..2ac0dfa4 100644 --- a/Lesson_Materials/Using_USM/index.html +++ b/Lesson_Materials/Using_USM/index.html @@ -61,7 +61,7 @@ #### Malloc_device
-

+						

 void* malloc_device(size_t numBytes, const queue& syclQueue, const property_list &propList = {});
 
 template <typename T>
@@ -82,7 +82,7 @@
 						#### Free
 					
-

+						

 void free(void* ptr, queue& syclQueue);
 						
@@ -99,7 +99,7 @@ #### Memcpy
-

+						

 event queue::memcpy(void* dest, const void* src, size_t numBytes, const std::vector<event> &depEvents);
 						
@@ -117,7 +117,7 @@ #### Memset & fill
-

+						

 event queue::memset(void* ptr, int value, size_t numBytes, const std::vector<event> &depEvents);
 
 event queue::fill(void* ptr, const T& pattern, size_t count, const std::vector<event> &depEvents);
@@ -136,7 +136,7 @@
 						#### Putting it all together
 					
-

+						

 int square_number(int x){
 	
   auto myQueue = queue{};
@@ -161,7 +161,7 @@
 						#### Putting it all together
 					
-

+						

 int square_number(int x){
 	
   auto myQueue = queue{usm_selector{}};
@@ -186,7 +186,7 @@
 						#### Putting it all together
 					
-

+						

 int square_number(int x){
 	
   auto myQueue = queue{usm_selector{}};
@@ -214,7 +214,7 @@
 						#### Putting it all together
 					
-

+						

 int square_number(int x){
 
   auto myQueue = queue{usm_selector{}};
@@ -244,7 +244,7 @@
 						#### Putting it all together
 					
-

+						

 int square_number(int x){
 
   auto myQueue = queue{usm_selector{}};
@@ -273,7 +273,7 @@
 						#### Putting it all together
 					
-

+						

 int square_number(int x){
 
   auto myQueue = queue{usm_selector{}};
@@ -304,7 +304,7 @@
 						#### Putting it all together
 					
-

+						

 int square_number(int x){
 
   auto myQueue = queue{usm_selector{}};
@@ -337,7 +337,7 @@
 						#### Queue shortcuts
 					
-

+						

 template <typename KernelName, typename KernelType>
 event queue::single_task(const KernelType &KernelFunc);
 
@@ -356,7 +356,7 @@
 						#### With the queue shortcut
 					
-

+						

 int square_number(int x){
 
   auto myQueue = queue{usm_selector{}};
diff --git a/Lesson_Materials/Vectors/index.html b/Lesson_Materials/Vectors/index.html
index 55919767..b3bfe015 100644
--- a/Lesson_Materials/Vectors/index.html
+++ b/Lesson_Materials/Vectors/index.html
@@ -106,7 +106,7 @@
 						#### Vec class
 					
-

+						

 template <typename dataT, int numElements>
 class vec;
 						
@@ -123,7 +123,7 @@ #### Aliases
-

+						

 using float4 = vec<float, 4>;
 ...
 						
@@ -138,18 +138,18 @@ #### Vec constructors
-

+						

 auto f4 = sycl::float4{1.0f, 2.0f, 3.0f, 4.0f}; // {1.0f, 2.0f, 3.0f, 4.0f}
 						
-

+						

 auto f2 = sycl::float4{2.0f, 3.0f}; // {2.0f, 3.0f}
 auto f4 = sycl::float4{1.0f, f2, 4.0f}; // {1.0f, 2.0f, 3.0f, 4.0f}
 						
-

+						

 auto f4 = sycl::float4{0.0f};  // {0.0f, 0.0f, 0.0f, 0.0f}
 						
@@ -164,7 +164,7 @@ #### Vec operators
-

+						

 auto f4a = sycl::float4{1.0f, 2.0f, 3.0f, 4.0f}; // {1.0f, 2.0f, 3.0f, 4.0f}
 
 auto f4b = sycl::float4{2.0f}; // {2.0f, 2.0f, 2.0f, 2.0f}
@@ -182,7 +182,7 @@
 						#### Vec types
 					
-

+						

 sycl::int2
 sycl::int3
 sycl::int4
@@ -203,13 +203,13 @@
 						#### Swizzles
 					
-

+						

 auto f4 = sycl::float4{1.0f, 2.0f, 3.0f, 4.0f}; // {1.0f, 2.0f, 3.0f, 4.0f}
 auto f2 = f4.swizzle<0, 3>(); // {1.0f, 4.0f}
 						
-

+						

 auto f4 = sycl::float4{1.0f, 2.0f, 3.0f, 4.0f}; // {1.0f, 2.0f, 3.0f, 4.0f}
 f4.swizzle<1, 2>() = sycl::float2{9.0f, 9.0f}; // f4 becomes {1.0f, 9.0f, 9.0f, 4.0f}
 						
@@ -224,13 +224,13 @@ #### Simple swizzles
-

+						

 auto f4 = sycl::float4{1.0f, 2.0f, 3.0f, 4.0f}; // {1.0f, 2.0f, 3.0f, 4.0f}
 auto f2 = f4.xw(); // {1.0f, 4.0f}
 						
-

+						

 auto f4 = sycl::float4{1.0f, 2.0f, 3.0f, 4.0f}; // {1.0f, 2.0f, 3.0f, 4.0f}
 f4.yz() = sycl::float2{9.0f, 9.0f}; // f4 becomes {1.0f, 9.0f, 9.0f, 4.0f}
 						
diff --git a/Lesson_Materials/oneMath_gemm/index.html b/Lesson_Materials/oneMath_gemm/index.html index 0da36689..81954582 100644 --- a/Lesson_Materials/oneMath_gemm/index.html +++ b/Lesson_Materials/oneMath_gemm/index.html @@ -92,7 +92,7 @@

#### Run-time dispatching

-

+						

 #include <oneapi/math.hpp>
 
 sycl::queue q{myDeviceSelector};
@@ -119,7 +119,7 @@ 

#### Compile-time dispatching

-

+						

 #include <oneapi;/math.hpp>
 
 sycl::queue cpu_queue{sycl::cpu_selector_v};

From a8db5ca0e73b3b8078be0e284a66a8892788a45c Mon Sep 17 00:00:00 2001
From: Michael Lance 
Date: Thu, 23 Jul 2026 13:38:37 -0500
Subject: [PATCH 29/34] address feedback

---
 Scripts/LessonChecker.py | 20 +++++++++++++-------
 1 file changed, 13 insertions(+), 7 deletions(-)

diff --git a/Scripts/LessonChecker.py b/Scripts/LessonChecker.py
index 155043d6..a9021cf8 100644
--- a/Scripts/LessonChecker.py
+++ b/Scripts/LessonChecker.py
@@ -108,9 +108,9 @@ def last_code_idx(self):
             return next(
                 (
                     i
-                    for i in range(len(self.output) - 1, -1, -1)
-                    # < is important, because the word "code" can apepar in a comment, just like it does in this one
-                    if ""} for char in data):
+        viol_chars = [char for char in data if char in {"&", "<", ">"}]
+
+        if viol_chars:
+            unique_viols = ", ".join(sorted(set(viol_chars)))
+
             self.state |= State.ERROR
             self._warn(
-                f"Unescaped character {viol_char} in lesson {self.lesson_name} line num: {self.getpos()[0]}"
+                f"Unescaped character(s) [{unique_viols}] in lesson {self.lesson_name} line num: {self.getpos()[0]}"
             )
 
         data = data.replace("&", "&").replace("<", "<").replace(">", ">")

From 94abec49479300c1a1587ddbce84aad07b9f794f Mon Sep 17 00:00:00 2001
From: Thomas Applencourt 
Date: Thu, 23 Jul 2026 14:06:31 -0500
Subject: [PATCH 30/34] Apply suggestion from @TApplencourt

---
 Scripts/LessonChecker.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/Scripts/LessonChecker.py b/Scripts/LessonChecker.py
index a9021cf8..7b255fa6 100644
--- a/Scripts/LessonChecker.py
+++ b/Scripts/LessonChecker.py
@@ -177,7 +177,7 @@ def handle_starttag(self, tag, attrs):
                         f" in  without data-noescape prop in lesson {self.lesson_name} line num: {self.getpos()[0]}"
                     )
 
-                    # Add data-noescape prop to code
+                    # `` -> ``
                     code_tag = self.output[last_code_idx]
                     self.output[last_code_idx] = (
                         code_tag[:-1] + " data-noescape" + code_tag[-1]

From 6cb8975f10556bf22a672d213bdd2c0bace96c58 Mon Sep 17 00:00:00 2001
From: Michael Lance 
Date: Thu, 23 Jul 2026 14:14:46 -0500
Subject: [PATCH 31/34] remove extra lesson oops

---
 Lesson_Materials/Buffer_Accessor/index.html | 304 --------------------
 1 file changed, 304 deletions(-)
 delete mode 100644 Lesson_Materials/Buffer_Accessor/index.html

diff --git a/Lesson_Materials/Buffer_Accessor/index.html b/Lesson_Materials/Buffer_Accessor/index.html
deleted file mode 100644
index e4bb8d58..00000000
--- a/Lesson_Materials/Buffer_Accessor/index.html
+++ /dev/null
@@ -1,304 +0,0 @@
-
-
-
-  
-    
-    
-    
-    
-    
-    
-    
-    
-    
-    
-    
-  
-	
-		
-
- -
-
- ## The Buffer Accessor Model -
-
- -
-
- ## Learning Objectives - * Learn about the buffer/accessor model - * Learn how to manage data with the buffer/accessor model -
-
- -
-
- #### SYCL Buffers & Accessors -
-
- * The buffer/accessor model separates the storage and access of data - * A SYCL buffer manages data across the host and any number of devices - * A SYCL accessor requests access to data on the host or on a device for a specific SYCL kernel function - * Accessors are also used to access data within a SYCL kernel function - * This means they are declared in the host code but captured by and then accessed within a SYCL kernel function -
-
- -
-
- #### SYCL Buffers & Accessors -
-
-
- * A SYCL buffer can be constructed with a pointer to host memory - * For the lifetime of the buffer this memory is owned by the SYCL runtime - * When a buffer object is constructed it will not allocate or copy to device memory at first - * This will only happen once the SYCL runtime knows the data needs to be accessed and where it needs to be accessed -
-
- ![Buffer Host Memory](../../Static/images/buffer-hostmemory.png "Buffer Host Memory") -
-
-
- -
-
- #### SYCL Buffers & Accessors -
-
-
- * Constructing an accessor specifies a request to access the data managed by the buffer - * There are a range of different types of accessor which provide different ways to access data -
-
- ![Buffer Host Memory Accessor](../../Static/images/buffer-hostmemory-accessor.png "Buffer Host Memory Accessor") -
-
-
- -
-
- #### SYCL Buffers & Accessors -
-
-
- * When an accessor is constructed it is associated with a command group via the handler object - * This connects the buffer that is being accessed, the way in which it’s being accessed and the device that the command group is being submitted to -
-
- ![Buffer Host Memory Accessor CG](../../Static/images/buffer-hostmemory-accessor-cg.png "Buffer Host Memory Accessor CG") -
-
-
- -
-
- #### SYCL Buffers & Accessors -
-
-
- * Once the SYCL scheduler selects the command group to be executed it must first satisfy its data dependencies - * If necessary, this includes allocating and copying the data to the device accessing that data - * If the most recent copy of the data is already on the device then the runtime will not copy again -
-
- ![Buffer Host Memory Accessor CG Device](../../Static/images/buffer-hostmemory-accessor-cg-device.png "Buffer Host Memory Accessor CG Device") -
-
-
- -
-
- #### SYCL Buffers & Accessors -
-
-
- * Data will remain in device memory after kernels finish executing until another accessor requests access in a different device or on the host - * When the buffer object is destroyed it will wait for any outstanding work that is accessing the data to complete and then copy back to the original host memory -
-
- ![Buffer Destroyed](../../Static/images/buffer-destroyed.png "Buffer Destroyed") -
-
-
- -
-
- #### SYCL Buffers & Accessors -
-
-

-#include<sycl/sycl.hpp> 
-
-int main() {
-  int var = 42; 
-
-  auto buf = sycl::buffer{&var, sycl::range<1>{1}};
-
-}
-T var = 42;
-
-{
-  // Create buffer pointing to var.
-  auto buf = sycl::buffer{&var, sycl::range<1>{1}};
-
-  // ...
-  // Do some computation on device. Use accessors to access buffer
-  // ...
-  
-} // var updated here
-
-assert(var != 42);
-						
-
-
- * A buffer is associated with a type, range and - dimensionality. Dimensionality must be either 1, 2 or - 3. - * Usually type and dimensionality can be inferred using - CTAD. - * If a buffer is associated with some allocation in host - memory, the host memory will be updated only once the - buffer goes out of scope. -
-
-
-
- #### Accessor class -
-
- * There are many different ways to use the `accessor` - class. - * Accessing data on a device. - * Accessing data immediately in the host application. - * Allocating local memory. - * For now we are going to focus on accessing data on a - device. -
-
- -
-
- #### Constructing an accessor -
-
-

-auto acc = sycl::accessor{bufA, cgh};
-						
-
-
- * There are many ways to construct an `accessor`. - * Accessors are complicated templates with many template arguments, - representing the data type, dimensionality, access mode etc. - * The `accessor` class supports CTAD so it's not - necessary to specify all of the template arguments. - * The most common way to construct an `accessor` is from - a `buffer` and a `handler` associated with the command - group function you are within. - * The element type and dimensionality are inferred from - the `buffer`. - * The `access_mode` is defaulted to - `access_mode::read_write`. -
-
- -
-
- #### Specifying the access mode -
-
-

-auto readAcc = sycl::accessor{bufA, cgh, sycl::read_only};
-auto writeAcc = sycl::accessor{bufB, cgh, sycl::write_only};
-						
-
-
- * When constructing an `accessor` you will likely also - want to specify the `access_mode` - * You can do this by passing one of the CTAD tags: - * `read_only` will result in `access_mode::read`. - * `write_only` will result in `access_mode::write`. -
-
- -
-
- #### Specifying no initialization -
-
-

-auto acc = sycl::accessor{buf, cgh, sycl::no_init};
-						
-
-
- * When constructing an `accessor` you may also want to - discard the original data of a `buffer`. - * You can do this by passing the `no_init` property. -
-
- -
-
- #### Using Accessors -
-
-

-T var = 42;
-
-{
-  // Create buffer pointing to var.
-  auto bufA = sycl::buffer{&var, sycl::range<1>{1}};
-  auto bufB = sycl::buffer{&var, sycl::range<1>{1}};
-
-  q.submit([&](sycl::handler &cgh) {
-	auto accA = sycl::accessor{bufA, cgh, sycl::read_only};
-	auto accB = sycl::accessor{bufA, cgh, sycl::no_init};
-
-  cgh.single_task<mykernel>(...); // Do some work
-  });
-  
-} // var updated here
-
-assert(var != 42);
-						
-
-
- * Buffers and accessors take care of memory - migration, as well as dependency analysis. - * More to come later on dependencies. -
-
- -
-
- #### operator[] -
-
-

-gpuQueue.submit([&](handler &cgh){
-  auto inA = sycl::accessor{bufA, cgh, sycl::read_only};
-  auto inB = sycl::accessor{bufB, cgh, sycl::read_only};
-  auto out = sycl::accessor{bufO, cgh, sycl::write_only};
-  cgh.single_task<mykernel>([=]{
-    out[0] = inA[0] + inB[0];
-  }); 
-});
-						
-
-
- * As well as specifying data dependencies an `accessor` - can also be used to access the data from within a kernel - function. - * You can do this by calling `operator[]` on the - `accessor`. - * `operator[]` for USM pointers must take a `size_t`, - whereas `operator[]` for accessors can take a - multi-dimensional `sycl::id` or a `size_t`. -
-
-
-
- - - From 8a01d4bda18a3a7a2846f93f75daf71eedad519b Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Thu, 23 Jul 2026 14:25:42 -0500 Subject: [PATCH 32/34] refix onemath --- Lesson_Materials/oneMath_gemm/index.html | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/Lesson_Materials/oneMath_gemm/index.html b/Lesson_Materials/oneMath_gemm/index.html index 81954582..f42290d1 100644 --- a/Lesson_Materials/oneMath_gemm/index.html +++ b/Lesson_Materials/oneMath_gemm/index.html @@ -120,7 +120,7 @@


-#include <oneapi;/math.hpp>
+#include <oneapi/math.hpp>
 
 sycl::queue cpu_queue{sycl::cpu_selector_v};
 
@@ -128,7 +128,7 @@ 

sycl::buffer<T,1> b{b_host, k*n}; sycl::buffer<T,1> c{c_host, m*n}; -oneapi::math::backend_selector<oneapi;::math::backend::mklcpu> cpu_selector(cpu_queue); +oneapi::math::backend_selector<oneapi::math::backend::mklcpu> cpu_selector(cpu_queue); // Select the Intel oneMKL CPU backend specifically ^^^^^^ oneapi::math::blas::column_major::gemm(cpu_selector, ..., m, n, k, ..., a, ..., b, ..., c, ... ); From 31c239df6db17cacbc9c9ca618cbe9bdddd6a069 Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Thu, 23 Jul 2026 14:38:53 -0500 Subject: [PATCH 33/34] refix onemath x 2 --- Lesson_Materials/oneMath_gemm/index.html | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/Lesson_Materials/oneMath_gemm/index.html b/Lesson_Materials/oneMath_gemm/index.html index f42290d1..7a68236e 100644 --- a/Lesson_Materials/oneMath_gemm/index.html +++ b/Lesson_Materials/oneMath_gemm/index.html @@ -53,7 +53,7 @@

* **Intel CPU/GPU**: Intel Math Kernels Library (oneMKL) * **NVIDIA GPU**: cuBLAS, cuSOLVER, cuRAND, cuFFT * **AMD GPU**: rocBLAS, rocSOLVER, rocRAND, rocFFT - * Imagine being able to use all of them with *single source code* **oneMath** + * Imagine being able to use all of them with *single source code* → **oneMath**

From 1389b822ed723c5a6653157cf7cf1a15631caa13 Mon Sep 17 00:00:00 2001 From: Michael Lance Date: Thu, 23 Jul 2026 15:00:02 -0500 Subject: [PATCH 34/34] add test for charref --- Scripts/LessonChecker.py | 3 +++ Scripts/test_LessonChecker.py | 9 +++++++++ 2 files changed, 12 insertions(+) diff --git a/Scripts/LessonChecker.py b/Scripts/LessonChecker.py index 7b255fa6..39948100 100644 --- a/Scripts/LessonChecker.py +++ b/Scripts/LessonChecker.py @@ -258,6 +258,9 @@ def handle_entityref(self, name): pos, e_data = self.code_blocks[-1] self.code_blocks[-1] = (pos, e_data + converted_char) + def handle_charref(self, name): + self.output.append(f"&#{name};") + def handle_startendtag(self, tag, attrs): self.output.append(self._encode_attrs(tag, attrs)) diff --git a/Scripts/test_LessonChecker.py b/Scripts/test_LessonChecker.py index 7851a673..99b28264 100644 --- a/Scripts/test_LessonChecker.py +++ b/Scripts/test_LessonChecker.py @@ -84,6 +84,15 @@ "silent": True, "expected_return": False, }, + { + "name": "validate_charref", + "body": "
💩
", + "extract": False, + "verify": True, + "fix": False, + "silent": True, + "expected_return": False, + }, { "name": "extract_code_block", "body": '
int main() { return 0; }
',