Support for program-scope variables

Global variables in program-scope are now supported, along with static global variables in function-scope, for both OpenCL C source and SPIR-V compilation. The implementation passes the basic/test_basic test of the OpenCL-CTS, and has been tested with client applications through chipStar.

global float testGlobalVar[128];

__kernel void test1 (__global const float *a) {
  size_t i = get_global_id(0) % 128;
  testGlobalVar[i] += a[i];
}

__kernel void test2 (__global const float *a) {
  size_t i = get_global_id(0) % 128;
  testGlobalVar[i] *= a[i];
}

__kernel void test3 (__global float *out) {
  size_t i = get_global_id(0) % 128;
  out[i] = testGlobalVar[i];
}

Support for generic address space

Generic AS is now supported, for both OpenCL C source and SPIR-V compilation. PoCL now passes the generic_address_space/test_generic_address_space test of the OpenCL-CTS, and has been tested with CUDA/HIP applications through chipStar.

int isOdd(int *val) {
  return val[0] % 2;
}

__kernel void test3 (__global int *in1, __local int *in2, __global int *out) {
  size_t i = get_global_id(0);
  out[i] = isOdd(in1+i) + isOdd(in2+(i % 128)];
}

Initial support for cl_khr_subgroups

The default is a single subgroup that always executes the whole X-dimension's WIs. Independent forward progress is not yet supported, but it's not needed for CTS compliance, due to the corner case of only one SG in flight.

Additionally, there is partial implementation for cl_khr_subgroup_shuffle, cl_intel_subgroups and cl_khr_subgroup_ballot with caveats:

• cl_khr_subgroup_shuffle: Passes the CTS, but only because it doesn't test non-uniform(lock-step) behavior, see: https://github.com/KhronosGroup/OpenCL-CTS/issues/1236
• cl_khr_subgroup_ballot: sub_group_ballot() works for uniform calls, the rest are unimplemented.
• cl_intel_subgroups: The block reads/writes are unimplemented.

Initial support for cl_intel_required_subgroup_size

This extension allows the programmer to specify the required subgroup size for a kernel function. This can be important for algorithm correctness in some cases. It's used by chipStar to implement fixed width warps when needed. The programmer can specify the size with a new kernel attribute: __attribute__((intel_reqd_sub_group_size(<int>)))

PoCL additionally implements CL_DEVICE_SUB_GROUP_SIZES_INTEL parameter for clGetDeviceInfo API, however CL_​KERNEL_​SPILL_​MEM_​SIZE_​INTEL and CL_​KERNEL_​COMPILE_​SUB_​GROUP_​SIZE_​INTEL for clGetKernelWorkGroupInfo API are not yet implemented.

Initial support for cl_khr_fp16

PoCL now has partial support for cl_khr_fp16 when compiled with Clang/LLVM 16+. The implementation relies on Clang, and may result in emulation (promoting to fp32) if the CPU does not support the required instruction set. In Clang/LLVM 16+, the following targets have native fp16 support: 32-bit and 64-bit ARM (depending on vendor), x86-64 with AVX512-FP16. Currently only implemented for a part of builtin library functions, those that are implemented with either an expression, or a Clang builtin.

Support for cl_intel_unified_shared_memory

This extension, together with SPIR-V support and other new features, allows using PoCL as an OpenCL backend for SYCL runtimes. This works with the both CPU driver (tested on x86-64 & ARM64) and the Level Zero driver. Vincent A. Arcila has contributed a guide for building PoCL as SYCL runtime backend on ARM.

Additionally, there is a new testsuite integrated into PoCL for testing USM support, intel-compute-samples. These are tests from https://github.com/intel/compute-samples and PoCL currently passes 78% of the tests (12 tests failed out of 54).