Merge pull request #102 from frasercrmck/regenerate-spirv

[UnitCL] Fix regenerate-spirv (with LLVM 15)
uxlfoundation · Aug 24, 2023 · e90e458 · e90e458
2 parents 1ca64b9 + c5598cf
commit e90e458
Show file tree

Hide file tree

Showing 21 changed files with 46 additions and 12 deletions.
diff --git a/source/cl/test/UnitCL/cmake/CompileKernelToSPIRVAsm.cmake b/source/cl/test/UnitCL/cmake/CompileKernelToSPIRVAsm.cmake
@@ -171,8 +171,9 @@ execute_process(
   COMMAND ${CLANG_PATH} -c -emit-llvm -target spir-unknown-unknown
     -cl-std=CL${SPIRV_CL_STD}
     -Xclang -finclude-default-header
+    -Werror
     ${DEFS_LIST} ${SPIRV_OPTIONS_LIST}
-    -O0 -Werror
+    -O0
     -o ${temp_bc32}
     ${INPUT_FILE}
   RESULT_VARIABLE clang_result
@@ -186,8 +187,9 @@ if(NOT clang_result EQUAL 0)
     ${CLANG_PATH} -c -emit-llvm -target spir-unknown-unknown
       -cl-std=CL${SPIRV_CL_STD}
       -Xclang -finclude-default-header
+      -Werror
       ${PRINTABLE_OPTIONS_LIST}
-      -O0 -Werror
+      -O0
       -o ${temp_bc32}
       ${INPUT_FILE}
     ${clang_error}")
@@ -229,8 +231,9 @@ execute_process(
   COMMAND ${CLANG_PATH} -c -emit-llvm -target spir64-unknown-unknown
     -cl-std=CL${SPIRV_CL_STD}
     -Xclang -finclude-default-header
+    -Werror
     ${DEFS_LIST} ${SPIRV_OPTIONS_LIST}
-    -O0 -Werror
+    -O0
     -o ${temp_bc64}
     ${INPUT_FILE}
   RESULT_VARIABLE clang_result
@@ -244,8 +247,9 @@ if(NOT clang_result EQUAL 0)
     ${CLANG_PATH} -c -emit-llvm -target spir64-unknown-unknown
       -cl-std=CL${SPIRV_CL_STD}
       -Xclang -finclude-default-header
+      -Werror
       ${PRINTABLE_OPTIONS_LIST}
-      -O0 -Werror
+      -O0
       -o ${temp_bc64}
       ${INPUT_FILE}
     ${clang_error}")

diff --git a/source/cl/test/UnitCL/kernels/barrier.11_barrier_with_align.cl b/source/cl/test/UnitCL/kernels/barrier.11_barrier_with_align.cl
@@ -14,6 +14,8 @@
 //
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+// SPIRV OPTIONS: -Wno-pointer-to-int-cast
+
 struct struct_needs_32bit_align {
   char x;
   int y;

diff --git a/source/cl/test/UnitCL/kernels/barrier.17_barrier_store_mask.cl b/source/cl/test/UnitCL/kernels/barrier.17_barrier_store_mask.cl
@@ -14,6 +14,8 @@
 //
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+// SPIRV OPTIONS: -Wno-pointer-to-int-cast
+
 __kernel void barrier_store_mask(int mask, __global int* output) {
   int global_id = get_global_id(0);
   barrier(CLK_LOCAL_MEM_FENCE);

diff --git a/source/cl/test/UnitCL/kernels/barrier.18_barrier_store_mask.cl b/source/cl/test/UnitCL/kernels/barrier.18_barrier_store_mask.cl
@@ -14,6 +14,8 @@
 //
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+// SPIRV OPTIONS: -Wno-pointer-to-int-cast
+
 __kernel void barrier_store_mask(int mask, __global int* output) {
   int global_id = get_global_id(0);
 

diff --git a/source/cl/test/UnitCL/kernels/barrier.19_barrier_store_mask.cl b/source/cl/test/UnitCL/kernels/barrier.19_barrier_store_mask.cl
@@ -14,6 +14,8 @@
 //
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+// SPIRV OPTIONS: -Wno-pointer-to-int-cast
+
 struct struct_needs_32bit_align {
   char x;
   int y;

diff --git a/source/cl/test/UnitCL/kernels/ext_async.01_simple_2d.cl b/source/cl/test/UnitCL/kernels/ext_async.01_simple_2d.cl
@@ -14,6 +14,10 @@
 //
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+// The LLVM compilers we use to build the SPIR-V tests don't have support for,
+// or aren't able to be told about, the extension that provides this builtin.
+// REQUIRES: nospirv
+
 __kernel void simple_2d(__local int *tmpA, __local int *tmpB,
                          __local int *tmpC, __global int *A, __global int *B,
                          __global int *C) {

diff --git a/source/cl/test/UnitCL/kernels/ext_async.02_simple_3d.cl b/source/cl/test/UnitCL/kernels/ext_async.02_simple_3d.cl
@@ -14,6 +14,10 @@
 //
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+// The LLVM compilers we use to build the SPIR-V tests don't have support for,
+// or aren't able to be told about, the extension that provides this builtin.
+// REQUIRES: nospirv
+
 __kernel void simple_3d(__local int *tmpA, __local int *tmpB,
                          __local int *tmpC, __global int *A, __global int *B,
                          __global int *C) {

diff --git a/source/cl/test/UnitCL/kernels/offline_type.cl.in b/source/cl/test/UnitCL/kernels/offline_type.cl.in
@@ -14,6 +14,11 @@
 //
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+// We only build these kernels for OpenCL. Note that if we did want to build
+// them for SPIR-V we'd have to ensure the right extensions are enabled for
+// upstream clang (half, double, etc)
+// REQUIRES: nospirv
+
 __kernel void type_${fulltype}(${fulltype} in_p,
                                global ${fulltype} *in_g,
                                constant ${fulltype} *in_c,

diff --git a/source/cl/test/UnitCL/kernels/printf.10_print_nan.cl b/source/cl/test/UnitCL/kernels/printf.10_print_nan.cl
@@ -24,7 +24,7 @@
 // To work around this issue, we disable the fp64 extension to ensure that the
 // generated SPIR-V does not contain this implicit conversion.
 
-// SPIRV OPTIONS: -Xclang;-cl-ext=-cl_khr_fp64
+// SPIRV OPTIONS: -Xclang;-cl-ext=-cl_khr_fp64;-Wno-format
 
 __kernel void print_nan(void) {
   // single specifiers

diff --git a/source/cl/test/UnitCL/kernels/printf.11_print_inf.cl b/source/cl/test/UnitCL/kernels/printf.11_print_inf.cl
@@ -24,7 +24,7 @@
 // To work around this issue, we disable the fp64 extension to ensure that the
 // generated SPIR-V does not contain this implicit conversion.
 
-// SPIRV OPTIONS: -Xclang;-cl-ext=-cl_khr_fp64
+// SPIRV OPTIONS: -Xclang;-cl-ext=-cl_khr_fp64;-Wno-format
 
 __kernel void print_inf(void) {
   // single specifiers

diff --git a/source/cl/test/UnitCL/kernels/printf.12_multiple_workgroups.cl b/source/cl/test/UnitCL/kernels/printf.12_multiple_workgroups.cl
@@ -14,6 +14,8 @@
 //
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+// SPIRV OPTIONS: -Wno-format
+
 __kernel void multiple_workgroups(void) {
   size_t i = get_local_id(0);
   if (i == 0) {

diff --git a/source/cl/test/UnitCL/kernels/printf.13_concurrent_printf.cl b/source/cl/test/UnitCL/kernels/printf.13_concurrent_printf.cl
@@ -14,6 +14,8 @@
 //
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+// SPIRV OPTIONS: -Wno-format
+
 __kernel void concurrent_printf(void) {
   size_t i = get_global_id(0);
   if (0 == (i % 2)) {

diff --git a/source/cl/test/UnitCL/kernels/printf.14_print_vector.cl b/source/cl/test/UnitCL/kernels/printf.14_print_vector.cl
@@ -14,6 +14,8 @@
 //
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+// SPIRV OPTIONS: -Wno-format
+
 kernel void print_vector(global int4* in, int flag) {
   size_t gid = get_global_id(0);
   if (gid < flag) {

diff --git a/source/cl/test/UnitCL/kernels/printf.15_floats.cl b/source/cl/test/UnitCL/kernels/printf.15_floats.cl
@@ -25,7 +25,7 @@
 // generated SPIR-V does not contain this implicit conversion.
 
 // DEFINITIONS: -DNUM_INPUTS=32
-// SPIRV OPTIONS: -Xclang;-cl-ext=-cl_khr_fp64
+// SPIRV OPTIONS: -Xclang;-cl-ext=-cl_khr_fp64;-Wno-format
 
 #ifndef NUM_INPUTS
 #define NUM_INPUTS 1

diff --git a/source/cl/test/UnitCL/kernels/printf.16_floats_vectors.cl b/source/cl/test/UnitCL/kernels/printf.16_floats_vectors.cl
@@ -14,6 +14,7 @@
 //
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+// SPIRV OPTIONS: -Wno-format
 // DEFINITIONS: -D NUM_INPUTS=16
 
 #ifndef NUM_INPUTS

diff --git a/source/cl/test/UnitCL/kernels/printf.17_float_formatting_aa.cl b/source/cl/test/UnitCL/kernels/printf.17_float_formatting_aa.cl
@@ -24,7 +24,7 @@
 // To work around this issue, we disable the fp64 extension to ensure that the
 // generated SPIR-V does not contain this implicit conversion.
 
-// SPIRV OPTIONS: -Xclang;-cl-ext=-cl_khr_fp64
+// SPIRV OPTIONS: -Xclang;-cl-ext=-cl_khr_fp64;-Wno-format
 
 kernel void float_formatting_aa(int size, global float* in) {
   // We are only printing from one workitem to make sure the order remains

diff --git a/source/cl/test/UnitCL/kernels/printf.17_float_formatting_ee.cl b/source/cl/test/UnitCL/kernels/printf.17_float_formatting_ee.cl
@@ -24,7 +24,7 @@
 // To work around this issue, we disable the fp64 extension to ensure that the
 // generated SPIR-V does not contain this implicit conversion.
 
-// SPIRV OPTIONS: -Xclang;-cl-ext=-cl_khr_fp64
+// SPIRV OPTIONS: -Xclang;-cl-ext=-cl_khr_fp64;-Wno-format
 
 kernel void float_formatting_ee(int size, global float* in) {
   // We are only printing from one workitem to make sure the order remains

diff --git a/source/cl/test/UnitCL/kernels/printf.17_float_formatting_ff.cl b/source/cl/test/UnitCL/kernels/printf.17_float_formatting_ff.cl
@@ -24,7 +24,7 @@
 // To work around this issue, we disable the fp64 extension to ensure that the
 // generated SPIR-V does not contain this implicit conversion.
 
-// SPIRV OPTIONS: -Xclang;-cl-ext=-cl_khr_fp64
+// SPIRV OPTIONS: -Xclang;-cl-ext=-cl_khr_fp64;-Wno-format
 
 kernel void float_formatting_ff(int size, global float* in) {
   // We are only printing from one workitem to make sure the order remains

diff --git a/source/cl/test/UnitCL/kernels/printf.17_float_formatting_gg.cl b/source/cl/test/UnitCL/kernels/printf.17_float_formatting_gg.cl
@@ -24,7 +24,7 @@
 // To work around this issue, we disable the fp64 extension to ensure that the
 // generated SPIR-V does not contain this implicit conversion.
 
-// SPIRV OPTIONS: -Xclang;-cl-ext=-cl_khr_fp64
+// SPIRV OPTIONS: -Xclang;-cl-ext=-cl_khr_fp64;-Wno-format
 
 kernel void float_formatting_gg(int size, global float* in) {
   // We are only printing from one workitem to make sure the order remains

diff --git a/source/cl/test/UnitCL/kernels/printf.19_print_halfs.cl b/source/cl/test/UnitCL/kernels/printf.19_print_halfs.cl
@@ -15,7 +15,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
 // REQUIRES: half
-// SPIRV OPTIONS: -Xclang;-cl-ext=-cl_khr_fp64,+cl_khr_fp16
+// SPIRV OPTIONS: -Xclang;-cl-ext=-cl_khr_fp64,+cl_khr_fp16;-Wno-format
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
 __kernel void print_halfs(__global half* a, __global half* b) {

diff --git a/source/cl/test/UnitCL/kernels/regression.57_attribute_aligned.cl b/source/cl/test/UnitCL/kernels/regression.57_attribute_aligned.cl
@@ -14,6 +14,8 @@
 //
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+// SPIRV OPTIONS: -Wno-pointer-to-int-cast
+
 struct struct_needs_1024byte_align {
   char c;
   short vec[3];