[OpenCL] cl_khr_integer_dot_product is much slower than dp4a emulation

[ProjectPhysX]

Hi all,

the new cl_khr_integer_dot_product addition in 2025.1 release is broken.

• On Windows (2025.19.3.0.17_230222), the __opencl_c_integer_dot_product_input_4x8bit and __opencl_c_integer_dot_product_input_4x8bit_packed feature macros are now present, but the dot(char4, char4)/dot_acc_sat(char4, char4, int) functions fail to compile with errors instructions in function CompilerException Failed to lookup symbol add_kernel JIT session error: Symbols not found: [ _Z3dotDv4_cS_ ] / [ _Z11dot_acc_satDv4_cS_i ].
• On Linux (2025.19.3.0.17_230222), both

  int dp4a(const char4 a, const char4 b, const int c) {
  	return c+dot(a, b); // 0.020 TIOPs/s
  }

  and

  int dp4a(const char4 a, const char4 b, const int c) {
  	return dot_acc_sat(a, b, c); // 0.015 TIOPs/s
  }

  perform much slower than the emulation variant

  int dp4a(const char4 a, const char4 b, const int c) {
  	return c+a.x*b.x+a.y*b.y+a.z*b.z+a.w*b.w; // 0.064 TIOPs/s
  }

  as measured on my i7-8700K CPU with my OpenCL-Benchmark. The full dp4a function implementation is here. The performance behavior is the same on AMD Ryzen 9 7950X.

Kind regards,
Moritz

[qichaogu]

For your question about integer dot product on Windows, please see the following example which demonstrates building kernels with dot functions works on Windows.

#include <CL/cl.h>
#include <CL/cl_ext.h>
#include <iostream>
#include <vector>

#define checkErr(err)                                                       \
  if (err != CL_SUCCESS)                                                    \
  {                                                                         \
    printf("Error code: %d, %s:%u\n", err, __FILE__, __LINE__);             \
    exit(-1);                                                               \
  }  

// OpenCL kernel source code
const char* kernelSource = R"(
__kernel void dot_product(__global char4* a, __global char4* b, __global int* result) {
    int gid = get_global_id(0);
    result[gid] = dot(a[gid], b[gid]);
}
)";

int main() {
    // Initialize data
    std::vector<cl_char4> a = {{1, 2, 3, 4}, {5, 6, 7, 8}};
    std::vector<cl_char4> b = {{1, 0, 1, 0}, {1, 0, 1, 0}};
    std::vector<int> result(a.size());

    // Get platform and device information
    cl_platform_id platform_id = nullptr;
    cl_device_id device_id = nullptr;
    cl_uint ret_num_devices;
    cl_uint ret_num_platforms;
    cl_int ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
    checkErr(ret)
    ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);
    checkErr(ret)

    // Create an OpenCL context
    cl_context context = clCreateContext(nullptr, 1, &device_id, nullptr, nullptr, &ret);
    checkErr(ret)

    // Create a command queue
    cl_command_queue command_queue = clCreateCommandQueueWithProperties(context, device_id, 0, &ret);
    checkErr(ret)

    // Create memory buffers on the device for each vector
    cl_mem a_mem = clCreateBuffer(context, CL_MEM_READ_ONLY, a.size() * sizeof(cl_char4), nullptr, &ret);
    checkErr(ret)
    cl_mem b_mem = clCreateBuffer(context, CL_MEM_READ_ONLY, b.size() * sizeof(cl_char4), nullptr, &ret);
    checkErr(ret)
    cl_mem result_mem = clCreateBuffer(context, CL_MEM_WRITE_ONLY, result.size() * sizeof(int), nullptr, &ret);
    checkErr(ret)

    // Copy the lists to their respective memory buffers
    ret = clEnqueueWriteBuffer(command_queue, a_mem, CL_TRUE, 0, a.size() * sizeof(cl_char4), a.data(), 0, nullptr, nullptr);
    checkErr(ret)
    ret = clEnqueueWriteBuffer(command_queue, b_mem, CL_TRUE, 0, b.size() * sizeof(cl_char4), b.data(), 0, nullptr, nullptr);
    checkErr(ret)

    // Create a program from the kernel source
    cl_program program = clCreateProgramWithSource(context, 1, &kernelSource, nullptr, &ret);
    checkErr(ret)

    // Build the program
    ret = clBuildProgram(program, 1, &device_id, "-cl-std=CL2.0", nullptr, nullptr);
    checkErr(ret)

    // Create the OpenCL kernel
    cl_kernel kernel = clCreateKernel(program, "dot_product", &ret);
    checkErr(ret)

    // Set the arguments of the kernel
    ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), &a_mem);
    checkErr(ret)
    ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), &b_mem);
    checkErr(ret)
    ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), &result_mem);
    checkErr(ret)

    // Execute the OpenCL kernel on the list
    size_t global_item_size = 2;
    size_t local_item_size = 1;
    ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, nullptr, &global_item_size, &local_item_size, 0, nullptr, nullptr);
    checkErr(ret)

    // Read the memory buffer result on the device to the local variable result
    ret = clEnqueueReadBuffer(command_queue, result_mem, CL_TRUE, 0, result.size() * sizeof(int), result.data(), 0, nullptr, nullptr);
    checkErr(ret)

    // Display the result
    for (size_t i = 0; i < result.size(); ++i) {
        std::cout << "Dot product of a[" << i << "] and b[" << i << "] = " << result[i] << std::endl;
    }

    // Clean up
    ret = clFlush(command_queue);
    ret = clFinish(command_queue);
    ret = clReleaseKernel(kernel);
    ret = clReleaseProgram(program);
    ret = clReleaseMemObject(a_mem);
    ret = clReleaseMemObject(b_mem);
    ret = clReleaseMemObject(result_mem);
    ret = clReleaseCommandQueue(command_queue);
    ret = clReleaseContext(context);
    checkErr(ret)

    return 0;
}

Build:
icpx integer_dot_product.cpp -o a.exe -lOpenCL
Run:

$ ./a.exe
Dot product of a[0] and b[0] = 4
Dot product of a[1] and b[1] = 12

For your question about performance on Linux, we'll investigate.

[ProjectPhysX]

Hi @qichaogu,

your sample program by default does not test the CPU Runtime - it default selects the dGPU (in my system Arc B580) which works.

If I select the CPU device (i7-8700K, Windows, 2025.19.3.0.17_230222), clBuildProgram fails with code -11 due to Symbols not found: [ _Z3dotDv4_cS_ ], just like my OpenCL-Benchmark.

On the iGPU (UHD 630, Windows, 31.0.101.2134), in your sample clBuildProgram also fails with code -11, due to this driver bug - here the error is error: undefined reference to '__builtin_spirv_OpDot_v4i8_v4i8' in function: '__builtin_spirv_OpDot_v4i8_v4i8' called by kernel: 'dp4a_kernel' error: backend compiler failed build.

You need a better test that checks on all OpenCL devices on the system, not just the default device of OpenCL Platform 0.

Regards,
Moritz

[qichaogu]

We are the OpenCL CPU runtime team, focusing on CPU.
To run the test on CPU, you can replace CL_DEVICE_TYPE_DEFAULT wtih CL_DEVICE_TYPE_CPU:
ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_CPU, 1, &device_id, &ret_num_devices);
My environment has only one CPU device, so either CL_DEVICE_TYPE_DEFAULT or CL_DEVICE_TYPE_CPU works.

$ sycl-ls
[opencl:cpu][opencl:0] Intel(R) OpenCL, Intel(R) Xeon(R) Platinum 8180M CPU @ 2.50GHz OpenCL 3.0 (Build 0) [2025.19.3.0.02_160000]

I also tested with the target CPU architecture of your CPU i7-8700K which supports at least AVX2 instruction set.

$ export CL_CONFIG_CPU_TARGET_ARCH=core-avx2
$ icpx integer_dot_product.cpp -o a.exe -lOpenCL
$ ./a.exe
Dot product of a[0] and b[0] = 4
Dot product of a[1] and b[1] = 12

[ProjectPhysX]

Hi @qichaogu,

please show me that it's actually running on the CPU on your side.

I did these modifications to your sample, to select the CPU:

    // Get platform and device information
    const int select_platform = 2;
    const int select_device = 0;
    cl_platform_id platform_ids[10];
    cl_device_id device_ids[10];
    cl_uint ret_num_devices;
    cl_uint ret_num_platforms;
    cl_int ret = clGetPlatformIDs(10, platform_ids, &ret_num_platforms);
    checkErr(ret)
    ret = clGetDeviceIDs(platform_ids[select_platform], CL_DEVICE_TYPE_DEFAULT, 10, device_ids, &ret_num_devices);
    checkErr(ret)
    cl_device_id device_id = device_ids[select_device];
    char name[1024];
    clGetDeviceInfo(device_id, CL_DEVICE_NAME, sizeof(name), &name, nullptr);
    std::cout << name << "\n";

And to check OpenCL compiler output:

    // Build the program
    ret = clBuildProgram(program, 1, &device_id, "-cl-std=CL2.0", nullptr, nullptr);
    char build_log[10000];
    clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(build_log), build_log, nullptr);
    std::cout << build_log << "\n";
    checkErr(ret)

I'm compiling in Visual Studio Community 2019 on Windows 10 with MSVC.
OpenCL CPU Runtime is 2025.19.3.0.17_230222.

This is the output I get:

C:\Users\Moritz\Desktop\Dokumente\Projekte\C++\Entwicklung\OpenCL-Wrapper\bin>OpenCL-Wrapper.exe
Intel(R) Core(TM) i7-8700K CPU @ 3.70GHz
Compilation started
Compilation done
Linking started
Linking done
Device build started
Options used by backend compiler: -cl-std=CL2.0
Failed to build device program
remark: <unknown>:0:0: Kernel was 8-way vectorized2 gathers1 scatters
remark: <unknown>:0:0: 50 instructions in function
CompilerException Failed to lookup symbol dot_product
JIT session error: Symbols not found: [ _Z3dotDv4_cS_ ]
Failed to materialize symbols: { (main, { dot_product }) }

Error code: -11, C:\Users\Moritz\Desktop\Dokumente\Projekte\C++\Entwicklung\OpenCL-Wrapper\src\main.cpp:74

dot(char4, char4) is broken in the compiler.

Regards,
Moritz

[qichaogu]

Are there multiple OpenCL CPU runtime versions in your environment?

I improved the example with listing all available OpenCL platforms and devices and allowing user to select the platform and device to run the application, and now I'm able to reproduce your errors with an older version of OpenCL CPU runtime that does not have the cl_khr_integer_dot_product extension enabled.

Improved example:

#include <CL/cl.h>
#include <CL/cl_ext.h>
#include <iostream>
#include <vector>

#define checkErr(err)                                                       \
  if (err != CL_SUCCESS)                                                    \
  {                                                                         \
    printf("Error code: %d, %s:%u\n", err, __FILE__, __LINE__);             \
    exit(-1);                                                               \
  }  

// OpenCL kernel source code
const char* kernelSource = R"(
__kernel void dot_product(__global char4* a, __global char4* b, __global int* result) {
    int gid = get_global_id(0);
    result[gid] = dot(a[gid], b[gid]);
}
)";

int main() {
    // Initialize data
    std::vector<cl_char4> a = {{1, 2, 3, 4}, {5, 6, 7, 8}};
    std::vector<cl_char4> b = {{1, 0, 1, 0}, {1, 0, 1, 0}};
    std::vector<int> result(a.size());

    // Get platform and device information
    cl_uint num_platforms = 0;
    cl_int ret = clGetPlatformIDs(0, nullptr, &num_platforms);
    checkErr(ret)
    std::vector<cl_platform_id> platforms(num_platforms);
    ret = clGetPlatformIDs(num_platforms, platforms.data(), nullptr);
    checkErr(ret)
    std::cout << "Number of platforms: " << num_platforms << std::endl;
    std::vector<std::vector<cl_device_id>> platform_devices(num_platforms);
    for (cl_uint i = 0; i < num_platforms; ++i) {
        std::cout << "Platform " << i << ":" << std::endl;
        cl_uint num_devices;
        ret = clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, 0, nullptr, &num_devices);
        checkErr(ret)
        //std::vector<cl_device_id> devices(num_devices);
        platform_devices[i].resize(num_devices);
        ret = clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, num_devices, platform_devices[i].data(), nullptr);
        checkErr(ret)
        std::cout << "  Number of devices: " << num_devices << std::endl;
        for (cl_uint j = 0; j < num_devices; ++j) {
            std::cout << "  Device " << j << ":" << std::endl;
            char buffer[1024];
            ret = clGetDeviceInfo(platform_devices[i][j], CL_DEVICE_NAME, sizeof(buffer), &buffer, nullptr);
            checkErr(ret)
            std::cout << "    Name: " << buffer << "\n";
            ret = clGetDeviceInfo(platform_devices[i][j], CL_DRIVER_VERSION, sizeof(buffer), buffer, nullptr);
            checkErr(ret)
            std::cout << "    Driver Version: " << buffer << std::endl;
        }
    }

    // User to select platform and device
    int platform_index = 0;
    int device_index = 0;
    std::cout << "----------------------------------------\n";
    std::cout << "Select platform: ";
    std::cin >> platform_index;
    cl_platform_id platform_id = platforms[platform_index];
    std::cout << "Select device: ";
    std::cin >> device_index;
    cl_device_id device_id = platform_devices[platform_index][device_index];
    std::cout << "----------------------------------------\n";

    // Create an OpenCL context
    cl_context context = clCreateContext(nullptr, 1, &device_id, nullptr, nullptr, &ret);
    checkErr(ret)

    // Create a command queue
    cl_command_queue command_queue = clCreateCommandQueueWithProperties(context, device_id, 0, &ret);
    checkErr(ret)

    // Create memory buffers on the device for each vector
    cl_mem a_mem = clCreateBuffer(context, CL_MEM_READ_ONLY, a.size() * sizeof(cl_char4), nullptr, &ret);
    checkErr(ret)
    cl_mem b_mem = clCreateBuffer(context, CL_MEM_READ_ONLY, b.size() * sizeof(cl_char4), nullptr, &ret);
    checkErr(ret)
    cl_mem result_mem = clCreateBuffer(context, CL_MEM_WRITE_ONLY, result.size() * sizeof(int), nullptr, &ret);
    checkErr(ret)

    // Copy the lists to their respective memory buffers
    ret = clEnqueueWriteBuffer(command_queue, a_mem, CL_TRUE, 0, a.size() * sizeof(cl_char4), a.data(), 0, nullptr, nullptr);
    checkErr(ret)
    ret = clEnqueueWriteBuffer(command_queue, b_mem, CL_TRUE, 0, b.size() * sizeof(cl_char4), b.data(), 0, nullptr, nullptr);
    checkErr(ret)

    // Create a program from the kernel source
    cl_program program = clCreateProgramWithSource(context, 1, &kernelSource, nullptr, &ret);
    checkErr(ret)

    // Build the program
    ret = clBuildProgram(program, 1, &device_id, "-cl-std=CL2.0", nullptr, nullptr);
    char build_log[10000];
    clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(build_log), build_log, nullptr);
    std::cout << build_log << "\n";
    checkErr(ret)

    // Create the OpenCL kernel
    cl_kernel kernel = clCreateKernel(program, "dot_product", &ret);
    checkErr(ret)

    // Set the arguments of the kernel
    ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), &a_mem);
    checkErr(ret)
    ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), &b_mem);
    checkErr(ret)
    ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), &result_mem);
    checkErr(ret)

    // Execute the OpenCL kernel on the list
    size_t global_item_size = 2;
    size_t local_item_size = 1;
    ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, nullptr, &global_item_size, &local_item_size, 0, nullptr, nullptr);
    checkErr(ret)

    // Read the memory buffer result on the device to the local variable result
    ret = clEnqueueReadBuffer(command_queue, result_mem, CL_TRUE, 0, result.size() * sizeof(int), result.data(), 0, nullptr, nullptr);
    checkErr(ret)

    // Display the result
    for (size_t i = 0; i < result.size(); ++i) {
        std::cout << "Dot product of a[" << i << "] and b[" << i << "] = " << result[i] << std::endl;
    }

    // Clean up
    ret = clFlush(command_queue);
    ret = clFinish(command_queue);
    ret = clReleaseKernel(kernel);
    ret = clReleaseProgram(program);
    ret = clReleaseMemObject(a_mem);
    ret = clReleaseMemObject(b_mem);
    ret = clReleaseMemObject(result_mem);
    ret = clReleaseCommandQueue(command_queue);
    ret = clReleaseContext(context);
    checkErr(ret)

    return 0;
}

Select the version 2025.19.3.0.17_230222 (platform 0, device 0), passed:

>a.exe
Number of platforms: 2
Platform 0:
  Number of devices: 1
  Device 0:
    Name: Intel(R) Xeon(R) Platinum 8180M CPU @ 2.50GHz
    Driver Version: 2025.19.3.0.17_230222
Platform 1:
  Number of devices: 1
  Device 0:
    Name: Intel(R) Xeon(R) Platinum 8180M CPU @ 2.50GHz
    Driver Version: 2024.18.10.0.08_160000
----------------------------------------
Select platform: 0
Select device: 0
----------------------------------------
Compilation started
Compilation done
Linking started
Linking done
Device build started
Options used by backend compiler: -cl-std=CL2.0
Device build done
remark: <unknown>:0:0: Kernel was 8-way vectorized2 gathers1 scatters
remark: <unknown>:0:0: 33 instructions in function
Kernel "dot_product" was not vectorized
Done.
Dot product of a[0] and b[0] = 4
Dot product of a[1] and b[1] = 12

Select an older version 2024.18.10.0.08_160000 (platform 1, device 0), failed with same errors as yours:

>a.exe
Number of platforms: 2
Platform 0:
  Number of devices: 1
  Device 0:
    Name: Intel(R) Xeon(R) Platinum 8180M CPU @ 2.50GHz
    Driver Version: 2025.19.3.0.17_230222
Platform 1:
  Number of devices: 1
  Device 0:
    Name: Intel(R) Xeon(R) Platinum 8180M CPU @ 2.50GHz
    Driver Version: 2024.18.10.0.08_160000
----------------------------------------
Select platform: 1
Select device: 0
----------------------------------------
Compilation started
Compilation done
Linking started
Linking done
Device build started
Options used by backend compiler: -cl-std=CL2.0
Failed to build device program
remark: <unknown>:0:0: Kernel was 8-way vectorized2 gathers1 scatters
remark: <unknown>:0:0: 50 instructions in function
CompilerException Failed to lookup symbol dot_product
JIT session error: Symbols not found: [ _Z3dotDv4_cS_ ]
Failed to materialize symbols: { (main, { dot_product }) }


Error code: -11, integer_dot_product.cpp:101

Could you try this example to make sure you are using the expected OpenCL CPU runtime version?

[ProjectPhysX]

Hi @qichaogu,

I have installed w_opencl_runtime_p_2025.1.0.972.exe from Intel's official download page - that is version 2025.19.3.0.17_230222.

I have just tested it again on another system, with 13th Gen Intel(R) Core(TM) i7-13700K, and get the same error Symbols not found: [ _Z3dotDv4_cS_ ].

I have not installed older versions of the OpenCL CPU Runtime side-by-side.

In the old version before 2025.1 update, __opencl_c_integer_dot_product_input_4x8bit was not enabled, and the dot(char4, char4) path would not be chosen and instead the fallback emulation - everything was fine.

In 2025.1 version, __opencl_c_integer_dot_product_input_4x8bit is present so dot(char4, char4) gets selected, but dot(char4, char4) is broken and fails to compile.

Here is my dp4a implementation for reference:

int dp4a(const char4 a, const char4 b, const int c) { // 4-wide byte dot product and accumulate
)+"#if cl_nv_compute_capability>=61"+R( // use hardware-supported dp4a on Nvidia Pascal or newer GPUs with inline PTX assembly
	int d;)+"asm(\"dp4a.s32.s32\t%0,%1,%2,%3;\":\"=r\"(d):\"r\"(as_int(a)),\"r\"(as_int(b)),\"r\"(c));"+R(return d;
)+"#elif defined(__opencl_c_integer_dot_product_input_4x8bit)"+R( // use hardware-supported dp4a on some Intel GPUs
	return c+dot(a, b); // dot_acc_sat(a, b, c); is slow
)+"#elif __has_builtin(__builtin_amdgcn_sdot4)"+R( // use hardware-supported dp4a on some AMD GPUs
    return __builtin_amdgcn_sdot4(as_int(a), as_int(b), c, false);
)+"#elif defined(cl_arm_integer_dot_product_accumulate_int8)"+R( // use hardware-supported dp4a on some ARM GPUs
    return arm_dot_acc(a, b, c);
)+"#else"+R( // fallback emulation (compilers will turn this into hardware-supported dp4a instruction if available)
	return c+a.x*b.x+a.y*b.y+a.z*b.z+a.w*b.w;
)+"#endif"+R(
}

Your sample will of course also fail for older OpenCL CPU Runtime versions, because you do not check for __opencl_c_integer_dot_product_input_4x8bit feature flag in the OpenCL kernel.

I'm asking you to download and install w_opencl_runtime_p_2025.1.0.972.exe from Intel's official download page and test with that. How is this broken version put up as official release?

Regards,
Moritz

[qichaogu]

With the installation of the w_opencl_runtime_p_2025.1.0.972.exe, the example also ran successfully.
I ran it on a Windows laptop and select the CPU device. Here's the output:

>a.exe
Number of platforms: 2
Platform 0:
  Number of devices: 1
  Device 0:
    Name: Intel(R) Graphics
    Driver Version: 32.0.101.5763
Platform 1:
  Number of devices: 1
  Device 0:
    Name: Intel(R) Core(TM) Ultra 7 165U
    Driver Version: 2025.19.3.0.17_230222
----------------------------------------
Select platform: 1
Select device: 0
----------------------------------------
Compilation started
Compilation done
Linking started
Linking done
Device build started
Options used by backend compiler: -cl-std=CL2.0
Device build done
remark: <unknown>:0:0: Kernel was 8-way vectorized2 gathers1 scatters
remark: <unknown>:0:0: 33 instructions in function
Kernel "dot_product" was not vectorized
Done.
Dot product of a[0] and b[0] = 4
Dot product of a[1] and b[1] = 12

Can you run the example and share your output?

As described in the download page, please make sure you have administrative permission on the Windows machine on which you're installing the OpenCL CPU runtime.

After installing, please double check:

1. The directory of Intel OpenCL CPU runtime is included in PATH:
   
2. The path of Intel OpenCL CPU runtime is set to the registry HKEY_LOCAL_MACHINE\SOFTWARE\Khronos\OpenCL\Vendors:
   

[ProjectPhysX]

Hi @qichaogu,

thanks, the regkey and environment variables were present.

Issue was different one: The 2025.1 installer didn't fully remove the old installation from the system. Old installation didn't show up as OpenCL device anymore after installing 2025.1, but was still present in installed Windows Apps. I had to remove it manually, and now dp4a in 2025.1 runtime works as intended.
Previous releases of the installer had always removed the old version during installation. Why new installer doesn't do it anymore? This should be at least added to download page as a note.

With installation issue resolved, I can benchmark the new native dp4a on Windows too:

CPU	native dp4a c+dot(a, b)	native dp4a dot_acc_sat(a, b, c)	emulated dp4a c+a.x*b.x+a.y*b.y+a.z*b.z+a.w*b.w
i7-13700K	0.037 TIOPs/s	0.027 TIOPs/s	0.195 TIOPs/s
i7-8700K	0.020 TIOPs/s	0.015 TIOPs/s	0.064 TIOPs/s
Intel Xeon 6960P	0.407 TIOPs/s	?	3.334 TIOPs/s

So the performance issue is the same on Windows/Linux: newly added native dp4a is much slower than emulated dp4a.

Regards,
Moritz