kpet
Attempt at use ffmpeg opencl filters result in compilation failure (termux)
I have Android device with vendor-provided Vulkan driver (somewhere deep in system, I am not root so I can't change much)
in termux I have
clvk/stable,now 0.0.20230124.132813g9635d8e0 aarch64 [installed] Experimental implementation of OpenCL on Vulkan
it gives clinfo like this:
~ $ clinfo
Number of platforms 1
Platform Name clvk
Platform Vendor clvk
Platform Version OpenCL 3.0 clvk
Platform Profile FULL_PROFILE
Platform Extensions cl_khr_icd cl_khr_extended_versioning
Platform Extensions with Version cl_khr_icd 0x400000 (1.0.0)
cl_khr_extended_versioning 0x400000 (1.0.0)
Platform Numeric Version 0xc00000 (3.0.0)
Platform Extensions function suffix clvk
Platform Host timer resolution 0ns
Platform Name clvk
Number of devices 1
Device Name PowerVR Rogue GE8322
Device Vendor FIXME
Device Vendor ID 0x1010
Device Version OpenCL 3.0 CLVK on Vulkan v1.1.131 driver 5776728
Device UUID 32322038-3620-3130-3420-323138000000
Driver UUID 35373736-3732-3800-0000-000000000000
Valid Device LUID No
Device LUID 0000-000000000000
Device Node Mask 0
Device Numeric Version 0xc00000 (3.0.0)
Driver Version 3.0 CLVK on Vulkan v1.1.131 driver 5776728
Device OpenCL C Version OpenCL C 1.2 CLVK on Vulkan v1.1.131 driver 5776728
Device OpenCL C Numeric Version 0x402000 (1.2.0)
Device OpenCL C all versions OpenCL C 0x400000 (1.0.0)
OpenCL C 0x401000 (1.1.0)
OpenCL C 0x402000 (1.2.0)
OpenCL C 0xc00000 (3.0.0)
Device OpenCL C features __opencl_c_images 0xc00000 (3.0.0)
__opencl_c_read_write_images 0xc00000 (3.0.0)
__opencl_c_3d_image_writes 0xc00000 (3.0.0)
__opencl_c_atomic_order_acq_rel 0xc00000 (3.0.0)
__opencl_c_atomic_scope_device 0xc00000 (3.0.0)
__opencl_c_int64 0xc00000 (3.0.0)
Latest conformance test passed FIXME
Device Type GPU, Default
Device Profile FULL_PROFILE
Device Available Yes
Compiler Available Yes
Linker Available Yes
Max compute units 1
Max clock frequency 0MHz
Device Partition (core)
Max number of sub-devices 0
Supported partition types None
Supported affinity domains (n/a)
Max work item dimensions 3
Max work item sizes 512x512x64
Max work group size 512
Preferred work group size multiple (device) 16
Preferred work group size multiple (kernel) 16
Max sub-groups per work group 0
Preferred / native vector sizes
char 1 / 1
short 1 / 1
int 1 / 1
long 1 / 1
half 1 / 1 (cl_khr_fp16)
float 1 / 1
double 1 / 1 (n/a)
Half-precision Floating-point support (cl_khr_fp16)
Denormals No
Infinity and NANs Yes
Round to nearest Yes
Round to zero No
Round to infinity No
IEEE754-2008 fused multiply-add Yes
Support is emulated in software No
Single-precision Floating-point support (core)
Denormals No
Infinity and NANs Yes
Round to nearest Yes
Round to zero No
Round to infinity No
IEEE754-2008 fused multiply-add Yes
Support is emulated in software No
Correctly-rounded divide and sqrt operations No
Double-precision Floating-point support (n/a)
Address bits 32, Little-Endian
Global memory size 536870912 (512MiB)
Error Correction support No
Max memory allocation 134217728 (128MiB)
Unified memory for Host and Device Yes
Shared Virtual Memory (SVM) capabilities (core)
Coarse-grained buffer sharing No
Fine-grained buffer sharing No
Fine-grained system sharing No
Atomics No
Minimum alignment for any data type 128 bytes
Alignment of base address 1024 bits (128 bytes)
Preferred alignment for atomics
SVM 0 bytes
Global 0 bytes
Local 0 bytes
Atomic memory capabilities relaxed, acquire/release, work-group scope, device scope
Atomic fence capabilities relaxed, acquire/release, work-item scope, work-group scope, device scope
Max size for global variable 0
Preferred total size of global vars 0
Global Memory cache type None
Image support Yes
Max number of samplers per kernel 20
Max size for 1D images from buffer 4096 pixels
Max 1D or 2D image array size 2048 images
Base address alignment for 2D image buffers 0 bytes
Pitch alignment for 2D image buffers 0 pixels
Max 2D image size 4096x4096 pixels
Max 3D image size 2048x2048x2048 pixels
Max number of read image args 128
Max number of write image args 128
Max number of read/write image args 128
Pipe support No
Max number of pipe args 0
Max active pipe reservations 0
Max pipe packet size 0
Local memory type Local
Local memory size 16384 (16KiB)
Max number of constant args 8
Max constant buffer size 65536 (64KiB)
Generic address space support No
Max size of kernel argument 1024
Queue properties (on host)
Out-of-order execution No
Profiling Yes
Device enqueue capabilities (n/a)
Queue properties (on device)
Out-of-order execution No
Profiling No
Preferred size 0
Max size 0
Max queues on device 0
Max events on device 0
Prefer user sync for interop Yes
Profiling timer resolution 1ns
Execution capabilities
Run OpenCL kernels Yes
Run native kernels No
Non-uniform work-groups Yes
Work-group collective functions No
Sub-group independent forward progress No
IL version SPIR-V_1.0
ILs with version SPIR-V 0x400000 (1.0.0)
printf() buffer size 1048576 (1024KiB)
Built-in kernels (n/a)
Built-in kernels with version (n/a)
Device Extensions cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_byte_addressable_store cl_khr_extended_versioning cl_khr_create_command_queue cl_khr_il_program cl_khr_spirv_no_integer_wrap_decoration cl_arm_non_uniform_work_group_size cl_khr_suggested_local_work_size cl_khr_3d_image_writes cl_khr_device_uuid cl_khr_fp16
Device Extensions with Version cl_khr_global_int32_base_atomics 0x400000 (1.0.0)
cl_khr_global_int32_extended_atomics 0x400000 (1.0.0)
cl_khr_local_int32_base_atomics 0x400000 (1.0.0)
cl_khr_local_int32_extended_atomics 0x400000 (1.0.0)
cl_khr_byte_addressable_store 0x400000 (1.0.0)
cl_khr_extended_versioning 0x400000 (1.0.0)
cl_khr_create_command_queue 0x400000 (1.0.0)
cl_khr_il_program 0x400000 (1.0.0)
cl_khr_spirv_no_integer_wrap_decoration 0x400000 (1.0.0)
cl_arm_non_uniform_work_group_size 0x400000 (1.0.0)
cl_khr_suggested_local_work_size 0x400000 (1.0.0)
cl_khr_3d_image_writes 0x400000 (1.0.0)
cl_khr_device_uuid 0x400000 (1.0.0)
cl_khr_fp16 0x400000 (1.0.0)
NULL platform behavior
clGetPlatformInfo(NULL, CL_PLATFORM_NAME, ...) clvk
clGetDeviceIDs(NULL, CL_DEVICE_TYPE_ALL, ...) Success [clvk]
clCreateContext(NULL, ...) [default] Success [clvk]
clCreateContextFromType(NULL, CL_DEVICE_TYPE_DEFAULT) Success (1)
Platform Name clvk
Device Name PowerVR Rogue GE8322
clCreateContextFromType(NULL, CL_DEVICE_TYPE_CPU) No devices found in platform
clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU) Success (1)
Platform Name clvk
Device Name PowerVR Rogue GE8322
clCreateContextFromType(NULL, CL_DEVICE_TYPE_ACCELERATOR) No devices found in platform
clCreateContextFromType(NULL, CL_DEVICE_TYPE_CUSTOM) No devices found in platform
clCreateContextFromType(NULL, CL_DEVICE_TYPE_ALL) Success (1)
Platform Name clvk
Device Name PowerVR Rogue GE8322
ICD loader properties
ICD loader Name OpenCL ICD Loader
ICD loader Vendor OCL Icd free software
ICD loader Version 2.3.1
ICD loader Profile OpenCL 3.0
~ $
I try to use distribution ffmpeg like this:
CLVK_LOG=4 ffmpeg -init_hw_device opencl=gpu -filter_hw_device gpu -i ~/20210419_055507A.mp4 -vf "hwupload, boxblur_opencl, hwdownload,format=yuv420p" -f null /dev/null
it fails with compiler error:
[CLVK] retain: obj = 0xb4000075479025c0, refcount now 15 [CLVK] clCreateProgramWithSource: #line 1 "/home/builder/.termux-build/ffmpeg/src/libavfilter/opencl/avgblur.cl" /* * Copyright (c) 2018 Dylan Fernando * * This file is part of FFmpeg. * * FFmpeg is free software; you can redistribute it and/or * modify it under the terms of the GNU Lesser General Public * License as published by the Free Software Foundation; either * version 2.1 of the License, or (at your option) any later version. * * FFmpeg is distributed in the hope that it will be useful, * but WITHOUT ANY WARRANTY; without even the implied warranty of * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU * Lesser General Public License for more details. * * You should have received a copy of the GNU Lesser General Public * License along with FFmpeg; if not, write to the Free Software * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA */ __kernel void avgblur_horiz(__write_only image2d_t dst, __read_only image2d_t src, int rad) { const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_NEAREST); int2 loc = (int2)(get_global_id(0), get_global_id(1)); int2 size = (int2)(get_global_size(0), get_global_size(1)); int count = 0; float4 acc = (float4)(0,0,0,0); for (int xx = max(0, loc.x - rad); xx < min(loc.x + rad + 1, size.x); xx++) { count++; acc += read_imagef(src, sampler, (int2)(xx, loc.y)); } write_imagef(dst, loc, acc / count); } __kernel void avgblur_vert(__write_only image2d_t dst, __read_only image2d_t src, int radv) { const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_NEAREST); int2 loc = (int2)(get_global_id(0), get_global_id(1)); int2 size = (int2)(get_global_size(0), get_global_size(1)); int count = 0; float4 acc = (float4)(0,0,0,0); for (int yy = max(0, loc.y - radv); yy < min(loc.y + radv + 1, size.y); yy++) { count++; acc += read_imagef(src, sampler, (int2)(loc.x, yy)); } write_imagef(dst, loc, acc / count); } [CLVK] clBuildProgram: program = 0xb4000075c7902868, num_device = 1, device_list = 0xb4000074f78f6eb0, options = (null), pfn_notify = 0x0, user_data = 0x0
[CLVK] retain: obj = 0xb4000075c7902850, refcount now 2 release: obj = 0xb4000075b7966e20, refcount now 3 [CLVK] release: obj = 0xb4000075a7909a50, refcount now 1 [CLVK] release: obj = 0xb400007597902030, refcount now 9 [CLVK] release: obj = 0xb400007587900030, refcount now 0 [CLVK] release: obj = 0xb4000075479025c0, refcount now 14 [CLVK] release: obj = 0xb400007597902030, refcount now 8 [CLVK] release: obj = 0xb4000075b7966e20, refcount now 2 [CLVK] release: obj = 0xb400007587903cf0, refcount now 0 [CLVK] release: obj = 0xb4000075479025c0, refcount now 13 [CLVK] release: obj = 0xb400007597902030, refcount now 7 [CLVK] release: obj = 0xb4000075a7909a50, refcount now 0 [CLVK] do_build_inner_online: build_from_il 0 - build_to_ir 0 [CLVK] About to compile " -cl-single-precision-constant -cl-kernel-arg-info -fp64=0 -rewrite-packed-structs -spv-version=1.3 -max-pushconstant-size=128 -max-ubo-size=134217728 -global-offset -long-vector -module-constants-in-storage-buffer -cl-arm-non-uniform-work-group-size "
[CLVK] release: obj = 0xb4000075479025c0, refcount now 12 [CLVK] release: obj = 0xb400007587903bd0, refcount now 0 [CLVK] release: obj = 0xb4000075479025c0, refcount now 11 [CLVK] release: obj = 0xb400007597902030, refcount now 6 [CLVK] release: obj = 0xb4000075b7966e20, refcount now 1 [CLVK] release: obj = 0xb4000075879005d0, refcount now 0 [CLVK] release: obj = 0xb4000075479025c0, refcount now 10 [CLVK] release: obj = 0xb400007597902030, refcount now 5 [CLVK] release: obj = 0xb400007587902550, refcount now 0 [CLVK] release: obj = 0xb4000075479025c0, refcount now 9 [CLVK] release: obj = 0xb400007597902030, refcount now 4 [CLVK] release: obj = 0xb400007597902030, refcount now 3 [CLVK] SPIR-V binary is valid. [CLVK] check_capabilities: Program requires SPIR-V capability 1. [CLVK] check_capabilities: Program requires SPIR-V capability 56. [CLVK] check_capabilities: Device does not support SPIR-V capability 56.
[CLVK] Missing support for required SPIR-V capabilities. [CLVK] release: obj = 0xb4000075c7902850, refcount now 1 [Parsed_boxblur_opencl_1 @ 0xb4000075679027d0] Failed to build program: -11.
guess on this device it not supposed to work ...But I do not know what is "SPIR-V capability 56". May be more user-friendly error can be given?
Capability 56 is StorageImageWriteWithoutFormat (https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#Capability)
This is needed because of the write_imagef call in one of the kernels.
@Randrianasulu In case that was not clear from @rjodinchr's message, the error you're seeing is not a problem in clvk but clvk reporting that the Vulkan implementation on the device is missing a feature required for the kernels you are trying to compile.
Regarding error messages, I agree we could and should do better. I've created https://github.com/kpet/clvk/issues/498 to at least report names users could search for instead of numbers. It seems we also do not support all cases of checking which image formats support StorageImageWriteWithoutFormat, I've created https://github.com/kpet/clvk/issues/497.
