[CI][CLANG-FORMAT] Switch to clang-20-format

This keeps it in line with upstream dpc++

Author: Colin Davidson

.github/workflows/run_ock_internal_tests.yml

@@ -172,11 +172,11 @@ jobs:
 
       - name: setup-ubuntu-clang-format
         run:
-          pip install clang-format==19.1.0
+          pip install clang-format==20.1.0
 
       - name: run clang-format
         run: |
-          # we've installed clang-format-19 in the docker via pip, which just installs it as clang-format,
+          # we've installed clang-format-20 in the docker via pip, which just installs it as clang-format,
           # so just use clang-format-diff and -b clang-format directly
           git fetch origin ${{ github.base_ref }}
           git diff --no-color origin/${{ github.base_ref }} | \

CMakeLists.txt

@@ -197,7 +197,12 @@ find_package(PythonInterp 3.6 REQUIRED)
 # which perform static analysis and style checking on source files.
 # When updating the version here, also update that used in the merge request
 # config
-find_package(ClangTools 19 COMPONENTS clang-format clang-tidy)
+find_package(ClangTools 19 COMPONENTS clang-tidy)
+if(TARGET ClangTools::clang-tidy)
+  ca_option(CA_CLANG_TIDY_FLAGS STRING
+    "Semi-color separated list of clang-tidy flags" "")
+endif()
+find_package(ClangTools 20 COMPONENTS clang-format)
 if(TARGET ClangTools::clang-tidy)
   ca_option(CA_CLANG_TIDY_FLAGS STRING
     "Semi-color separated list of clang-tidy flags" "")

README.md

@@ -45,13 +45,13 @@ To install the dependencies on Ubuntu, open the terminal and run:
 
 ### Recommended packages
 * [Ninja](https://ninja-build.org/)
-* [clang-format](https://clang.llvm.org/docs/ClangFormat.html) 16
+* [clang-format](https://clang.llvm.org/docs/ClangFormat.html) 20
 * [lit](https://llvm.org/docs/CommandGuide/lit.html)
 
 To install the recommended packages, run:
 ```sh
    $ sudo apt install -y ninja-build doxygen python3-pip
-   $ sudo pip3 install lit virtualenv cmakelint clang-format==19.1.0
+   $ sudo pip3 install lit virtualenv cmakelint clang-format==20.1.0
 ```
 
 ### Compiling oneAPI Construction Kit

clik/examples/clik_async/barrier_print/device_barrier_print.c

@@ -1,50 +0,0 @@
-// Copyright (C) Codeplay Software Limited
-//
-// Licensed under the Apache License, Version 2.0 (the "License") with LLVM
-// Exceptions; you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//     https://github.com/uxlfoundation/oneapi-construction-kit/blob/main/LICENSE.txt
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
-// WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
-// License for the specific language governing permissions and limitations
-// under the License.
-//
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-#include "device_barrier_print.h"
-
-// In this example, the kernel has two computation steps represented by printf
-// calls. A barrier is used to ensure step A has been performed by all
-// work-items in the work-group before any work-item in the group starts step B.
-//
-// With a work-group size of four, the output of this example will look like:
-//
-// Kernel part A (tid = 0)
-// Kernel part A (tid = 1)
-// Kernel part A (tid = 2)
-// Kernel part A (tid = 3)
-// Kernel part B (tid = 0)
-// Kernel part B (tid = 1)
-// Kernel part B (tid = 2)
-// Kernel part B (tid = 3)
-//
-// Without a barrier between the two steps, the output would look like:
-//
-// Kernel part A (tid = 0)
-// Kernel part B (tid = 0)
-// Kernel part A (tid = 1)
-// Kernel part B (tid = 1)
-// Kernel part A (tid = 2)
-// Kernel part B (tid = 2)
-// Kernel part A (tid = 3)
-// Kernel part B (tid = 3)
-
-__kernel void barrier_print(exec_state_t *ctx) {
-  uint tid = get_global_id(0, ctx);
-  print(ctx, "Kernel part A (tid = %d)\n", tid);
-  barrier(ctx);
-  print(ctx, "Kernel part B (tid = %d)\n", tid);
-}

clik/examples/clik_async/barrier_sum/device_barrier_sum.c

@@ -1,43 +0,0 @@
-// Copyright (C) Codeplay Software Limited
-//
-// Licensed under the Apache License, Version 2.0 (the "License") with LLVM
-// Exceptions; you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//     https://github.com/uxlfoundation/oneapi-construction-kit/blob/main/LICENSE.txt
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
-// WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
-// License for the specific language governing permissions and limitations
-// under the License.
-//
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-#include "device_barrier_sum.h"
-
-// In this example, each work-item in the kernel computes the sum of a few
-// values. These values are the same for all items in the work-group. In order
-// to reduce the number of memory operations, a small local array is used to
-// share these values between the items in the group. Each item only needs to
-// load one item from the source buffer, however a barrier is needed to ensure
-// all values have been copied before computing the sum.
-
-__kernel void barrier_sum(__global uint *src, __global uint *dst,
-                          __local uint *src_tile, exec_state_t *ctx) {
-  uint tid = get_global_id(0, ctx);
-
-  // Copy values from the source buffer to a local tile.
-  uint lid = get_local_id(0, ctx);
-  src_tile[lid] = src[tid];
-
-  // Wait for all items in the group to have finished copying values.
-  barrier(ctx);
-
-  // Sum values from the tile and write the result to the output buffer.
-  uint sum = 0;
-  for (uint i = 0; i < get_local_size(0, ctx); i++) {
-    sum += src_tile[i];
-  }
-  dst[tid] = sum;
-}

clik/examples/clik_async/blur/device_blur.c

@@ -1,64 +0,0 @@
-// Copyright (C) Codeplay Software Limited
-//
-// Licensed under the Apache License, Version 2.0 (the "License") with LLVM
-// Exceptions; you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//     https://github.com/uxlfoundation/oneapi-construction-kit/blob/main/LICENSE.txt
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
-// WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
-// License for the specific language governing permissions and limitations
-// under the License.
-//
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-#include "device_blur.h"
-
-__kernel void copy_and_pad_hor(__global uint *src, __global uint *dst,
-                               exec_state_t *item) {
-  // Figure out array slice to copy
-  uint tid = get_global_id(0, item);
-  uint src_idx = WIDTH * tid;
-  uint src_idx_end = src_idx + WIDTH;
-  uint dst_idx = EXTENDED_WIDTH * (tid + 1);
-
-  // Duplicate the first element
-  dst[dst_idx++] = src[src_idx];
-
-  // Copy over the middle
-  while (src_idx < src_idx_end) {
-    dst[dst_idx++] = src[src_idx++];
-  }
-
-  // Duplicate the last element
-  dst[dst_idx] = src[src_idx - 1];
-}
-
-__kernel void pad_vert(__global uint *dst, exec_state_t *item) {
-  uint tid = get_global_id(0, item);
-  uint gap = HEIGHT * EXTENDED_WIDTH;
-
-  // Duplicate the topmost and the bottommost pixel.
-  dst[tid] = dst[tid + EXTENDED_WIDTH];
-  dst[tid + gap + EXTENDED_WIDTH] = dst[tid + gap];
-}
-
-__kernel void blur(__global uint *src, __global uint *dst, exec_state_t *item) {
-  uint x = get_global_id(0, item);
-  uint y = get_global_id(1, item);
-
-  // Calculate row starting indices.
-  uint base0 = EXTENDED_WIDTH * y + x;
-  uint base1 = EXTENDED_WIDTH + base0;
-  uint base2 = EXTENDED_WIDTH + base1;
-
-  // Add together 9 pixels around destination.
-  uint total = src[base0] + src[base0 + 1] + src[base0 + 2] + src[base1] +
-               src[base1 + 1] + src[base1 + 2] + src[base2] + src[base2 + 1] +
-               src[base2 + 2];
-
-  // Store the mean of the 9 pixeld into the destination buffer.
-  dst[WIDTH * y + x] = total / 9;
-}

clik/examples/clik_async/concatenate_dma/device_concatenate_dma.c

@@ -1,41 +0,0 @@
-// Copyright (C) Codeplay Software Limited
-//
-// Licensed under the Apache License, Version 2.0 (the "License") with LLVM
-// Exceptions; you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//     https://github.com/uxlfoundation/oneapi-construction-kit/blob/main/LICENSE.txt
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
-// WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
-// License for the specific language governing permissions and limitations
-// under the License.
-//
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-#include "device_concatenate_dma.h"
-
-// Carry out the computation for one work-item.
-__kernel void concatenate_dma(__global uint *src1, __global uint *src2,
-                              __global uint *dst, uint block_size,
-                              exec_state_t *ctx) {
-  uint tid = get_global_id(0, ctx);
-  uint offset = tid * block_size;
-  uint size = get_global_size(0, ctx) * block_size;
-
-  // Enqueue a DMA operation from the first input buffer to the first half of
-  // the output buffer.
-  uintptr_t xfer1 =
-      start_dma(&dst[offset], &src1[offset], block_size * sizeof(uint), ctx);
-
-  // Enqueue a DMA operation from the second input buffer to the second half of
-  // the output buffer, without waiting for the first DMA operation to finish.
-  uintptr_t xfer2 = start_dma(&dst[size + offset], &src2[offset],
-                              block_size * sizeof(uint), ctx);
-
-  // Wait for both DMA operations to finish. Waiting for a DMA operation only
-  // returns when that operation as well as all other operations enqueued before
-  // it are completed.
-  wait_dma(xfer2, ctx);
-}

clik/examples/clik_async/hello/device_hello.c

@@ -1,22 +0,0 @@
-// Copyright (C) Codeplay Software Limited
-//
-// Licensed under the Apache License, Version 2.0 (the "License") with LLVM
-// Exceptions; you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//     https://github.com/uxlfoundation/oneapi-construction-kit/blob/main/LICENSE.txt
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
-// WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
-// License for the specific language governing permissions and limitations
-// under the License.
-//
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-#include "device_hello.h"
-
-__kernel void hello_async(exec_state_t *ctx) {
-  print(ctx, "Hello from clik_async! tid=%d, lid=%d, gid=%d\n",
-        get_global_id(0, ctx), get_local_id(0, ctx), get_group_id(0, ctx));
-}

clik/examples/clik_async/matrix_multiply/device_matrix_multiply.c

@@ -1,28 +0,0 @@
-// Copyright (C) Codeplay Software Limited
-//
-// Licensed under the Apache License, Version 2.0 (the "License") with LLVM
-// Exceptions; you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//     https://github.com/uxlfoundation/oneapi-construction-kit/blob/main/LICENSE.txt
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
-// WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
-// License for the specific language governing permissions and limitations
-// under the License.
-//
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-#include "device_matrix_multiply.h"
-
-__kernel void matrix_multiply(__global float *a, __global float *b,
-                              __global float *c, uint m, exec_state_t *item) {
-  uint col = get_global_id(0, item);
-  uint row = get_global_id(1, item);
-  float sum = 0.0f;
-  for (int i = 0; i < m; i++) {
-    sum += a[(row * m) + i] * b[(i * m) + col];
-  }
-  c[(row * m) + col] = sum;
-}

clik/examples/clik_async/matrix_multiply_tiled/device_matrix_multiply_tiled.c

@@ -1,44 +0,0 @@
-// Copyright (C) Codeplay Software Limited
-//
-// Licensed under the Apache License, Version 2.0 (the "License") with LLVM
-// Exceptions; you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//     https://github.com/uxlfoundation/oneapi-construction-kit/blob/main/LICENSE.txt
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
-// WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
-// License for the specific language governing permissions and limitations
-// under the License.
-//
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-#include "device_matrix_multiply_tiled.h"
-
-__kernel void matrix_multiply(__global float *a, __global float *b,
-                              __global float *c, uint m, exec_state_t *ctx) {
-  uint col = get_global_id(0, ctx);
-  uint row = get_global_id(1, ctx);
-  uint lx = get_local_id(0, ctx);
-  uint ly = get_local_id(1, ctx);
-
-  __local_variable float tile_a[TS * TS];
-  __local_variable float tile_b[TS * TS];
-  float sum = 0.0f;
-  for (int i = 0; i < m; i += TS) {
-    // Load one tile from A and B into local memory.
-    tile_a[(ly * TS) + lx] = a[(row * m) + i + lx];
-    tile_b[(ly * TS) + lx] = b[((i + ly) * m) + col];
-    barrier(ctx);
-
-    // Compute the partial sum for the loaded tile.
-    for (int j = 0; j < TS; j++) {
-      sum += tile_a[(ly * TS) + j] * tile_b[(j * TS) + lx];
-    }
-
-    // Wait for all items to be finished before computing the next tile.
-    barrier(ctx);
-  }
-  c[(row * m) + col] = sum;
-}