add clang-format on source/

.clang-format

@@ -1,8 +1,63 @@
-Language: Cpp
-BasedOnStyle: LLVM
+---
+Language:        Cpp
+# BasedOnStyle:  Google
 AccessModifierOffset: -1
-BreakInheritanceList: BeforeComma
-BreakConstructorInitializers: BeforeComma
-Cpp11BracedListStyle: false
+ConstructorInitializerIndentWidth: 4
+AlignEscapedNewlinesLeft: true
+AlignTrailingComments: true
+AllowAllParametersOfDeclarationOnNextLine: true
+AllowShortBlocksOnASingleLine: false
+AllowShortCaseLabelsOnASingleLine: false
+AllowShortIfStatementsOnASingleLine: true
+AllowShortLoopsOnASingleLine: true
+AllowShortFunctionsOnASingleLine: All
+AlwaysBreakAfterDefinitionReturnType: false
+AlwaysBreakTemplateDeclarations: true
+AlwaysBreakBeforeMultilineStrings: true
+BreakBeforeBinaryOperators: None
+BreakBeforeTernaryOperators: true
+BreakConstructorInitializersBeforeComma: false
+BinPackParameters: true
+BinPackArguments: true
+ColumnLimit:     80
+ConstructorInitializerAllOnOneLineOrOnePerLine: true
+# Do not adapt pointer and reference alignment to the current style.
+# Just use PointerAlignment parameter unconditionally for consistency.
+DerivePointerAlignment: false
+ExperimentalAutoDetectBinPacking: false
+IndentCaseLabels: true
+IndentWrappedFunctionNames: false
+IndentFunctionDeclarationAfterType: false
+MaxEmptyLinesToKeep: 1
+KeepEmptyLinesAtTheStartOfBlocks: false
+NamespaceIndentation: None
+ObjCSpaceAfterProperty: false
+ObjCSpaceBeforeProtocolList: false
+PenaltyBreakBeforeFirstCallParameter: 1
+PenaltyBreakComment: 300
+PenaltyBreakString: 1000
+PenaltyBreakFirstLessLess: 120
+PenaltyExcessCharacter: 1000000
+PenaltyReturnTypeOnItsOwnLine: 200
 PointerAlignment: Left
-SpaceBeforeCpp11BracedList: true
+SpacesBeforeTrailingComments: 2
+Cpp11BracedListStyle: true
+Standard:        c++17
+IndentWidth:     2
+TabWidth:        8
+UseTab:          Never
+BreakBeforeBraces: Attach
+SpacesInParentheses: false
+SpacesInSquareBrackets: false
+SpacesInAngles:  false
+SpaceInEmptyParentheses: false
+SpacesInCStyleCastParentheses: false
+SpaceAfterCStyleCast: false
+SpacesInContainerLiterals: true
+SpaceBeforeAssignmentOperators: true
+ContinuationIndentWidth: 4
+CommentPragmas:  '^ IWYU pragma:'
+ForEachMacros:   [ foreach, Q_FOREACH, BOOST_FOREACH ]
+SpaceBeforeParens: ControlStatements
+DisableFormat:   false
+SortIncludes:   true

.github/workflows/CI.yml

@@ -52,3 +52,27 @@ jobs:
           name: spec-outputs
           path: |
             /tmp/out
+
+  check-clang-format:
+    if: github.event_name == 'pull_request'
+    runs-on: ubuntu-latest
+    steps:
+      - name: Checkout
+        uses: actions/checkout@v3
+        with:
+          fetch-depth: 0 # Need to check out base branch as well
+      - name: Install clang-format
+        run: sudo apt update; sudo apt install -y clang-format
+      # Inspired by SYCL-CTS
+      - name: Run clang-format on changed files
+        run: |
+          set -o errexit -o pipefail -o noclobber -o nounset
+          DIFF=$( git diff -U0 --no-color -C ./adoc/code/ ${{ github.event.pull_request.base.sha }} | ./adoc/scripts/clang-format-diff.py -p1 )
+          if [ ! -z "$DIFF" ]; then
+            echo 'The following changes are not formatted according to `clang-format`:' >> $GITHUB_STEP_SUMMARY
+            echo '```diff' >> $GITHUB_STEP_SUMMARY
+            echo "$DIFF" >> $GITHUB_STEP_SUMMARY
+            echo '```' >> $GITHUB_STEP_SUMMARY
+            echo "Not all files are formatted according to clang-format. See workflow summary for details."
+            exit 1 # Fail CI run
+          fi

adoc/code/algorithms.cpp

@@ -1,17 +1,17 @@
 // Copyright (c) 2011-2024 The Khronos Group, Inc.
 // SPDX-License-Identifier: Apache-2.0
 
-buffer<int> inputBuf { 1024 };
-buffer<int> outputBuf { 2 };
+buffer<int> inputBuf{1024};
+buffer<int> outputBuf{2};
 {
   // Initialize buffer on the host with 0, 1, 2, 3, ..., 1023
-  host_accessor a { inputBuf };
+  host_accessor a{inputBuf};
   std::iota(a.begin(), a.end(), 0);
 }
 
 myQueue.submit([&](handler& cgh) {
-  accessor inputValues { inputBuf, cgh, read_only };
-  accessor outputValues { outputBuf, cgh, write_only, no_init };
+  accessor inputValues{inputBuf, cgh, read_only};
+  accessor outputValues{outputBuf, cgh, write_only, no_init};
 
   cgh.parallel_for(nd_range<1>(range<1>(16), range<1>(16)), [=](nd_item<1> it) {
     // Apply a group algorithm to any number of values, described by an iterator
@@ -31,5 +31,5 @@ myQueue.submit([&](handler& cgh) {
   });
 });
 
-host_accessor a { outputBuf };
+host_accessor a{outputBuf};
 assert(a[0] == 523776 && a[1] == 120);

adoc/code/anatomy.cpp

@@ -3,10 +3,10 @@
 
 #include <iostream>
 #include <sycl/sycl.hpp>
-using namespace sycl; // (optional) avoids need for "sycl::" before SYCL names
+using namespace sycl;  // (optional) avoids need for "sycl::" before SYCL names
 
 int main() {
-  int data[1024]; // Allocate data to be worked on
+  int data[1024];  // Allocate data to be worked on
 
   // Create a default queue to enqueue work to the default device
   queue myQueue;
@@ -16,20 +16,20 @@ int main() {
   // because the destructor of resultBuf will wait
   {
     // Wrap our data variable in a buffer
-    buffer<int, 1> resultBuf { data, range<1> { 1024 } };
+        buffer<int, 1> resultBuf{data, range<1>{1024}};
 
     // Create a command group to issue commands to the queue
     myQueue.submit([&](handler& cgh) {
       // Request write access to the buffer without initialization
-      accessor writeResult { resultBuf, cgh, write_only, no_init };
+      accessor writeResult{resultBuf, cgh, write_only, no_init};
 
       // Enqueue a parallel_for task with 1024 work-items
       cgh.parallel_for(1024, [=](id<1> idx) {
         // Initialize each buffer element with its own rank number starting at 0
         writeResult[idx] = idx;
-      }); // End of the kernel function
-    });   // End of our commands for this queue
-  }       // End of scope, so we wait for work producing resultBuf to complete
+      });  // End of the kernel function
+    });    // End of our commands for this queue
+  }  // End of scope, so we wait for work producing resultBuf to complete
 
   // Print result
   for (int i = 0; i < 1024; i++)

adoc/code/aspectTraitExample.cpp

@@ -2,11 +2,12 @@
 // SPDX-License-Identifier: MIT
 
 #include <sycl/sycl.hpp>
-using namespace sycl; // (optional) avoids need for "sycl::" before SYCL names
+using namespace sycl;  // (optional) avoids need for "sycl::" before SYCL names
 
 constexpr int N = 512;
 
-template <bool HasFp16> class MyKernel {
+template <bool HasFp16>
+class MyKernel {
  public:
   void operator()(id<1> i) {
     if constexpr (HasFp16) {
@@ -22,11 +23,9 @@ int main() {
   myQueue.submit([&](handler& cgh) {
     device dev = myQueue.get_device();
     if (dev.has(aspect::fp16)) {
-      cgh.parallel_for(range { N },
-                       MyKernel<any_device_has_v<aspect::fp16>> {});
+      cgh.parallel_for(range{N}, MyKernel<any_device_has_v<aspect::fp16>>{});
     } else {
-      cgh.parallel_for(range { N },
-                       MyKernel<all_devices_have_v<aspect::fp16>> {});
+      cgh.parallel_for(range{N}, MyKernel<all_devices_have_v<aspect::fp16>>{});
     }
   });
 

adoc/code/basicParallelForItem.cpp

@@ -2,7 +2,7 @@
 // SPDX-License-Identifier: Apache-2.0
 
 myQueue.submit([&](handler& cgh) {
-  accessor acc { myBuffer, cgh, write_only };
+  accessor acc{myBuffer, cgh, write_only};
 
   cgh.parallel_for(range<1>(numWorkItems), [=](item<1> item) {
     // kernel argument type is item

adoc/code/basicparallelfor.cpp

@@ -2,7 +2,7 @@
 // SPDX-License-Identifier: Apache-2.0
 
 myQueue.submit([&](handler& cgh) {
-  accessor acc { myBuffer, cgh, write_only };
+  accessor acc{myBuffer, cgh, write_only};
 
   cgh.parallel_for(range<1>(numWorkItems),
                    [=](id<1> index) { acc[index] = 42.0f; });

adoc/code/bundle-builtin-kernel.cpp

@@ -2,7 +2,7 @@
 // SPDX-License-Identifier: Apache-2.0
 
 #include <sycl/sycl.hpp>
-using namespace sycl; // (optional) avoids need for "sycl::" before SYCL names
+using namespace sycl;  // (optional) avoids need for "sycl::" before SYCL names
 
 int main() {
   queue myQueue;
@@ -15,7 +15,7 @@ int main() {
   // Get an executable kernel_bundle containing all the built-in kernels
   // supported by the device.
   kernel_bundle<bundle_state::executable> myBundle =
-      get_kernel_bundle(myContext, { myDevice }, builtinKernelIds);
+      get_kernel_bundle(myContext, {myDevice}, builtinKernelIds);
 
   // Retrieve a kernel object that can be used to query for more information
   // about the built-in kernel or to submit it to a command group.  We assume
@@ -26,7 +26,7 @@ int main() {
   myQueue.submit([&](handler& cgh) {
     // Setting the arguments depends on the backend and the exact kernel used.
     cgh.set_args(...);
-    cgh.parallel_for(range { 1024 }, builtinKernel);
+    cgh.parallel_for(range{1024}, builtinKernel);
   });
 
   myQueue.wait();

adoc/code/bundle-kernel-introspection.cpp

@@ -2,9 +2,9 @@
 // SPDX-License-Identifier: Apache-2.0
 
 #include <sycl/sycl.hpp>
-using namespace sycl; // (optional) avoids need for "sycl::" before SYCL names
+using namespace sycl;  // (optional) avoids need for "sycl::" before SYCL names
 
-class MyKernel; // Forward declare the name of our kernel.
+class MyKernel;  // Forward declare the name of our kernel.
 
 int main() {
   size_t N = 1024;
@@ -15,7 +15,7 @@ int main() {
   // Get an executable kernel bundle containing our kernel.
   kernel_id kernelId = get_kernel_id<MyKernel>();
   auto myBundle =
-      get_kernel_bundle<bundle_state::executable>(myContext, { kernelId });
+      get_kernel_bundle<bundle_state::executable>(myContext, {kernelId});
 
   // Get the kernel's maximum work-group size when running on our device.
   kernel myKernel = myBundle.get_kernel(kernelId);
@@ -24,11 +24,11 @@ int main() {
 
   // Compute a good ND-range to use for iteration in the kernel
   // based on the maximum work-group size.
-  std::array<size_t, 11> divisors = { 1024, 512, 256, 128, 64, 32,
-                                      16,   8,   4,   2,   1 };
+  std::array<size_t, 11> divisors = {1024, 512, 256, 128, 64, 32,
+                                     16,   8,   4,   2,   1};
   size_t wgSize = *std::find_if(divisors.begin(), divisors.end(),
                                 [=](auto d) { return (d <= maxWgSize); });
-  nd_range myRange { range { N }, range { wgSize } };
+  nd_range myRange{range{N}, range{wgSize}};
 
   myQueue.submit([&](handler& cgh) {
     // Use the kernel bundle we queried, so we are sure the queried work-group

adoc/code/bundle-pre-compile.cpp

@@ -2,7 +2,7 @@
 // SPDX-License-Identifier: Apache-2.0
 
 #include <sycl/sycl.hpp>
-using namespace sycl; // (optional) avoids need for "sycl::" before SYCL names
+using namespace sycl;  // (optional) avoids need for "sycl::" before SYCL names
 
 int main() {
   queue myQueue;
@@ -18,7 +18,7 @@ int main() {
     // pre-compiled kernel from "myBundle".
     cgh.use_kernel_bundle(myBundle);
 
-    cgh.parallel_for(range { 1024 }, ([=](item index) {
+    cgh.parallel_for(range{1024}, ([=](item index) {
                        // kernel code
                      }));
   });

adoc/code/bundle-spec-constants.cpp

@@ -2,7 +2,7 @@
 // SPDX-License-Identifier: Apache-2.0
 
 #include <sycl/sycl.hpp>
-using namespace sycl; // (optional) avoids need for "sycl::" before SYCL names
+using namespace sycl;  // (optional) avoids need for "sycl::" before SYCL names
 
 // Forward declare names for our two kernels.
 class MyKernel1;
@@ -21,7 +21,7 @@ int main() {
 
   // Get the identifiers for our kernels, then get an input kernel bundle that
   // contains our two kernels.
-  auto kernelIds = { get_kernel_id<MyKernel1>(), get_kernel_id<MyKernel2>() };
+  auto kernelIds = {get_kernel_id<MyKernel1>(), get_kernel_id<MyKernel2>()};
   auto inputBundle =
       get_kernel_bundle<bundle_state::input>(myContext, kernelIds);
 
@@ -37,7 +37,7 @@ int main() {
     // Use the kernel bundle we built in this command group.
     cgh.use_kernel_bundle(exeBundle);
     cgh.parallel_for<MyKernel1>(
-        range { 1024 }, ([=](item index, kernel_handler kh) {
+        range{1024}, ([=](item index, kernel_handler kh) {
           // Read the value of the specialization constant.
           int w = kh.get_specialization_constant<width>();
           // ...
@@ -48,7 +48,7 @@ int main() {
     // This command group uses the same kernel bundle.
     cgh.use_kernel_bundle(exeBundle);
     cgh.parallel_for<MyKernel2>(
-        range { 1024 }, ([=](item index, kernel_handler kh) {
+        range{1024}, ([=](item index, kernel_handler kh) {
           int h = kh.get_specialization_constant<height>();
           // ...
         }));

adoc/code/deviceHas.cpp

@@ -10,13 +10,13 @@ class KernelFunctor {
 
  private:
   void foo() const {
-    half fp = 1.0; // No compiler diagnostic here
+    half fp = 1.0;  // No compiler diagnostic here
   }
 
   void bar() const {
     sycl::atomic_ref longAtomic(longValue);
-    longAtomic.fetchAdd(1); // ERROR: Compiler issues diagnostic because
-                            // "aspect::atomic64" missing from "device_has()"
+    longAtomic.fetchAdd(1);  // ERROR: Compiler issues diagnostic because
+                             // "aspect::atomic64" missing from "device_has()"
   }
 };
 
@@ -25,5 +25,5 @@ class KernelFunctor {
 // still check the device's aspects before submitting the kernel.
 if (myQueue.get_device().has(aspect::fp16)) {
   myQueue.submit(
-      [&](handler& h) { h.parallel_for(range { 16 }, KernelFunctor {}); });
+      [&](handler& h) { h.parallel_for(range{16}, KernelFunctor{}); });
 }

adoc/code/explicitcopy.cpp

@@ -4,19 +4,19 @@
 const size_t nElems = 10u;
 
 // Create a vector and fill it with values 0, 1, 2, 3, 4, 5, 6, 7, 8, 9
-std::vector<int> v { nElems };
+std::vector<int> v{nElems};
 std::iota(std::begin(v), std::end(v), 0);
 
 // Create a buffer with no associated user storage
-sycl::buffer<int, 1> b { range<1>(nElems) };
+sycl::buffer<int, 1> b{range<1>(nElems)};
 
 // Create a queue
 queue myQueue;
 
 myQueue.submit([&](handler& cgh) {
   // Retrieve a ranged write accessor to a global buffer with access to the
   // first half of the buffer
-  accessor acc { b, cgh, range<1>(nElems / 2), id<1>(0), write_only };
+  accessor acc{b, cgh, range<1>(nElems / 2), id<1>(0), write_only};
   // Copy the first five elements of the vector into the buffer associated with
   // the accessor
   cgh.copy(v.data(), acc);