Bug 99856 - OpenCL Hello world returns "unsupported call to function get_local_size"
Summary: OpenCL Hello world returns "unsupported call to function get_local_size"
Status: RESOLVED FIXED
Alias: None
Product: Mesa
Classification: Unclassified
Component: Other (show other bugs)
Version: 17.0
Hardware: Other All
: medium normal
Assignee: Jan Vesely
QA Contact: mesa-dev
URL:
Whiteboard:
Keywords:
Depends on:
Blocks: 99553
  Show dependency treegraph
 
Reported: 2017-02-19 03:24 UTC by Henrique Dante de Almeida
Modified: 2017-09-29 19:11 UTC (History)
3 users (show)

See Also:
i915 platform:
i915 features:


Attachments
OpenCL hello strace (207.17 KB, text/plain)
2017-02-25 21:34 UTC, Henrique Dante de Almeida
Details
Arch Linux libclc tahiti-amdgcn--.bc (5.03 MB, application/octet-stream)
2017-02-25 21:55 UTC, Henrique Dante de Almeida
Details
size_t is 32bit in mesa3d (1.39 KB, patch)
2017-02-26 12:42 UTC, Jan Vesely
Details | Splinter Review
size_t is 32bit in mesa3d (3.55 KB, patch)
2017-03-09 05:01 UTC, Jan Vesely
Details | Splinter Review
Restore support for llvm-3.9 (8.00 KB, patch)
2017-09-21 02:31 UTC, Jan Vesely
Details | Splinter Review

Note You need to log in before you can comment on or make changes to this bug.
Description Henrique Dante de Almeida 2017-02-19 03:24:23 UTC
https://bugs.archlinux.org/task/53018

When trying to execute opencl hello world from apple site, the following error appears:

./hello
Error: Failed to build program executable!
<unknown>:0:0: in function square void (float addrspace(1)*, float addrspace(1)*, i32): unsupported call to function get_local_size

Additional info:
* package version(s)
opencl-mesa 17.0.0-1

https://developer.apple.com/library/content/samplecode/OpenCL_Hello_World_Example/Listings/hello_c.html#//apple_ref/doc/uid/DTS40008187-hello_c-DontLinkElementID_4

* config and/or log files etc.

Number of platforms 1
Platform Name Clover
Platform Vendor Mesa
Platform Version OpenCL 1.1 Mesa 17.0.0
Platform Profile FULL_PROFILE
Platform Extensions cl_khr_icd
Platform Extensions function suffix MESA

Platform Name Clover
Number of devices 1
Device Name AMD CAPE VERDE (DRM 2.48.0 / 4.9.8-1-ARCH, LLVM 3.9.1)
Device Vendor AMD
Device Vendor ID 0x1002
Device Version OpenCL 1.1 Mesa 17.0.0
(...)
Comment 1 Vedran Miletić 2017-02-20 15:41:08 UTC
This function should be provided by libclc. Do you have it installed?
Comment 2 Henrique Dante de Almeida 2017-02-21 14:13:45 UTC
Yes:

Nome                 : libclc
Versão               : 0.2.0+334+520743b-1
Descrição            : Library requirements of the OpenCL C programming language
Arquitetura          : any
URL                  : http://libclc.llvm.org/
Licenças             : MIT
Grupos               : Nenhum
Provê                : Nenhum
Depende de           : Nenhum
Depend. opcionais    : Nenhum
Necessário para      : Nenhum
Opcional para        : Nenhum
Conflita com         : Nenhum
Substitui            : Nenhum
Tamanho instalado    : 35,96 MiB
Empacotador          : Laurent Carlier <lordheavym@gmail.com>
Data da compilação   : qui 27 out 2016 08:58:06 BRST
Data de instalação   : qua 02 nov 2016 16:03:36 BRST
Motivo da instalação : Instalado como dependência de outro pacote
Script de instalação : Não
Validado por         : Assinatura
Comment 3 Vedran Miletić 2017-02-21 14:38:29 UTC
On Fedora 25 with custom-compiled LLVM, libclc, and Mesa:

$ gcc -lOpenCL -lm hello.c
hello.c: In function ‘main’:
hello.c:269:5: warning: ‘clCreateCommandQueue’ is deprecated [-Wdeprecated-declarations]
     commands = clCreateCommandQueue(context, device_id, 0, &err);
     ^~~~~~~~
In file included from hello.c:121:0:
/usr/include/CL/cl.h:1427:1: note: declared here
 clCreateCommandQueue(cl_context                     /* context */,
 ^~~~~~~~~~~~~~~~~~~~

$ ./hello
Computed '1024/1024' correct values!

I would say this is a distribution issue.
Comment 4 Aidan Thornton 2017-02-25 19:04:27 UTC
I've been having the same issue on Gentoo for a while with Mesa and libdrm from git master, and it appears someone's been having the same problem with the Debian-packaged version of Mesa too: https://bugs.debian.org/cgi-bin/bugreport.cgi?bug=848258

So if it is a distro issue it's one that affects multiple distros.  There's a comment on the Debian bug from a Debian developer/AMD employee complaining that it's an upstream bug too.
Comment 5 Henrique Dante de Almeida 2017-02-25 20:37:22 UTC
I'll try with an upstream environment and generate strace output soon
Comment 6 Jan Vesely 2017-02-25 20:40:12 UTC
(In reply to Henrique Dante de Almeida from comment #2)
> Yes:
> 
> Nome                 : libclc
> Versão               : 0.2.0+334+520743b-1
> Descrição            : Library requirements of the OpenCL C programming
> language
> Arquitetura          : any
> URL                  : http://libclc.llvm.org/
> Licenças             : MIT
> Grupos               : Nenhum
> Provê                : Nenhum
> Depende de           : Nenhum
> Depend. opcionais    : Nenhum
> Necessário para      : Nenhum
> Opcional para        : Nenhum
> Conflita com         : Nenhum
> Substitui            : Nenhum
> Tamanho instalado    : 35,96 MiB
> Empacotador          : Laurent Carlier <lordheavym@gmail.com>
> Data da compilação   : qui 27 out 2016 08:58:06 BRST
> Data de instalação   : qua 02 nov 2016 16:03:36 BRST
> Motivo da instalação : Instalado como dependência de outro pacote
> Script de instalação : Não
> Validado por         : Assinatura

can you post the list of files installed by this package?
Comment 7 Henrique Dante de Almeida 2017-02-25 20:43:57 UTC
libclc /usr/
libclc /usr/include/
libclc /usr/include/clc/
libclc /usr/include/clc/as_type.h
libclc /usr/include/clc/async/
libclc /usr/include/clc/async/async_work_group_copy.h
libclc /usr/include/clc/async/async_work_group_copy.inc
libclc /usr/include/clc/async/async_work_group_strided_copy.h
libclc /usr/include/clc/async/async_work_group_strided_copy.inc
libclc /usr/include/clc/async/gentype.inc
libclc /usr/include/clc/async/prefetch.h
libclc /usr/include/clc/async/prefetch.inc
libclc /usr/include/clc/async/wait_group_events.h
libclc /usr/include/clc/atomic/
libclc /usr/include/clc/atomic/atomic_add.h
libclc /usr/include/clc/atomic/atomic_and.h
libclc /usr/include/clc/atomic/atomic_cmpxchg.h
libclc /usr/include/clc/atomic/atomic_dec.h
libclc /usr/include/clc/atomic/atomic_decl.inc
libclc /usr/include/clc/atomic/atomic_inc.h
libclc /usr/include/clc/atomic/atomic_max.h
libclc /usr/include/clc/atomic/atomic_min.h
libclc /usr/include/clc/atomic/atomic_or.h
libclc /usr/include/clc/atomic/atomic_sub.h
libclc /usr/include/clc/atomic/atomic_xchg.h
libclc /usr/include/clc/atomic/atomic_xor.h
libclc /usr/include/clc/cl_khr_global_int32_base_atomics/
libclc /usr/include/clc/cl_khr_global_int32_base_atomics/atom_add.h
libclc /usr/include/clc/cl_khr_global_int32_base_atomics/atom_cmpxchg.h
libclc /usr/include/clc/cl_khr_global_int32_base_atomics/atom_dec.h
libclc /usr/include/clc/cl_khr_global_int32_base_atomics/atom_inc.h
libclc /usr/include/clc/cl_khr_global_int32_base_atomics/atom_sub.h
libclc /usr/include/clc/cl_khr_global_int32_base_atomics/atom_xchg.h
libclc /usr/include/clc/cl_khr_global_int32_extended_atomics/
libclc /usr/include/clc/cl_khr_global_int32_extended_atomics/atom_and.h
libclc /usr/include/clc/cl_khr_global_int32_extended_atomics/atom_max.h
libclc /usr/include/clc/cl_khr_global_int32_extended_atomics/atom_min.h
libclc /usr/include/clc/cl_khr_global_int32_extended_atomics/atom_or.h
libclc /usr/include/clc/cl_khr_global_int32_extended_atomics/atom_xor.h
libclc /usr/include/clc/cl_khr_local_int32_base_atomics/
libclc /usr/include/clc/cl_khr_local_int32_base_atomics/atom_add.h
libclc /usr/include/clc/cl_khr_local_int32_base_atomics/atom_cmpxchg.h
libclc /usr/include/clc/cl_khr_local_int32_base_atomics/atom_dec.h
libclc /usr/include/clc/cl_khr_local_int32_base_atomics/atom_inc.h
libclc /usr/include/clc/cl_khr_local_int32_base_atomics/atom_sub.h
libclc /usr/include/clc/cl_khr_local_int32_base_atomics/atom_xchg.h
libclc /usr/include/clc/cl_khr_local_int32_extended_atomics/
libclc /usr/include/clc/cl_khr_local_int32_extended_atomics/atom_and.h
libclc /usr/include/clc/cl_khr_local_int32_extended_atomics/atom_max.h
libclc /usr/include/clc/cl_khr_local_int32_extended_atomics/atom_min.h
libclc /usr/include/clc/cl_khr_local_int32_extended_atomics/atom_or.h
libclc /usr/include/clc/cl_khr_local_int32_extended_atomics/atom_xor.h
libclc /usr/include/clc/clc.h
libclc /usr/include/clc/clcfunc.h
libclc /usr/include/clc/clctypes.h
libclc /usr/include/clc/clcversion.h
libclc /usr/include/clc/common/
libclc /usr/include/clc/common/degrees.h
libclc /usr/include/clc/common/degrees.inc
libclc /usr/include/clc/common/mix.h
libclc /usr/include/clc/common/mix.inc
libclc /usr/include/clc/common/radians.h
libclc /usr/include/clc/common/radians.inc
libclc /usr/include/clc/common/sign.h
libclc /usr/include/clc/common/smoothstep.h
libclc /usr/include/clc/common/smoothstep.inc
libclc /usr/include/clc/common/step.h
libclc /usr/include/clc/common/step.inc
libclc /usr/include/clc/convert.h
libclc /usr/include/clc/float/
libclc /usr/include/clc/float/definitions.h
libclc /usr/include/clc/geometric/
libclc /usr/include/clc/geometric/cross.h
libclc /usr/include/clc/geometric/distance.h
libclc /usr/include/clc/geometric/distance.inc
libclc /usr/include/clc/geometric/dot.h
libclc /usr/include/clc/geometric/dot.inc
libclc /usr/include/clc/geometric/fast_distance.h
libclc /usr/include/clc/geometric/fast_distance.inc
libclc /usr/include/clc/geometric/fast_length.h
libclc /usr/include/clc/geometric/fast_length.inc
libclc /usr/include/clc/geometric/fast_normalize.h
libclc /usr/include/clc/geometric/fast_normalize.inc
libclc /usr/include/clc/geometric/floatn.inc
libclc /usr/include/clc/geometric/length.h
libclc /usr/include/clc/geometric/length.inc
libclc /usr/include/clc/geometric/normalize.h
libclc /usr/include/clc/geometric/normalize.inc
libclc /usr/include/clc/image/
libclc /usr/include/clc/image/image.h
libclc /usr/include/clc/image/image_defines.h
libclc /usr/include/clc/integer/
libclc /usr/include/clc/integer/abs.h
libclc /usr/include/clc/integer/abs.inc
libclc /usr/include/clc/integer/abs_diff.h
libclc /usr/include/clc/integer/abs_diff.inc
libclc /usr/include/clc/integer/add_sat.h
libclc /usr/include/clc/integer/add_sat.inc
libclc /usr/include/clc/integer/clz.h
libclc /usr/include/clc/integer/clz.inc
libclc /usr/include/clc/integer/definitions.h
libclc /usr/include/clc/integer/gentype.inc
libclc /usr/include/clc/integer/hadd.h
libclc /usr/include/clc/integer/hadd.inc
libclc /usr/include/clc/integer/integer-gentype.inc
libclc /usr/include/clc/integer/mad24.h
libclc /usr/include/clc/integer/mad24.inc
libclc /usr/include/clc/integer/mad_hi.h
libclc /usr/include/clc/integer/mad_sat.h
libclc /usr/include/clc/integer/mad_sat.inc
libclc /usr/include/clc/integer/mul24.h
libclc /usr/include/clc/integer/mul24.inc
libclc /usr/include/clc/integer/mul_hi.h
libclc /usr/include/clc/integer/mul_hi.inc
libclc /usr/include/clc/integer/rhadd.h
libclc /usr/include/clc/integer/rhadd.inc
libclc /usr/include/clc/integer/rotate.h
libclc /usr/include/clc/integer/rotate.inc
libclc /usr/include/clc/integer/sub_sat.h
libclc /usr/include/clc/integer/sub_sat.inc
libclc /usr/include/clc/integer/upsample.h
libclc /usr/include/clc/math/
libclc /usr/include/clc/math/acos.h
libclc /usr/include/clc/math/acos.inc
libclc /usr/include/clc/math/acosh.h
libclc /usr/include/clc/math/acosh.inc
libclc /usr/include/clc/math/acospi.h
libclc /usr/include/clc/math/acospi.inc
libclc /usr/include/clc/math/asin.h
libclc /usr/include/clc/math/asin.inc
libclc /usr/include/clc/math/asinh.h
libclc /usr/include/clc/math/asinh.inc
libclc /usr/include/clc/math/asinpi.h
libclc /usr/include/clc/math/asinpi.inc
libclc /usr/include/clc/math/atan.h
libclc /usr/include/clc/math/atan.inc
libclc /usr/include/clc/math/atan2.h
libclc /usr/include/clc/math/atan2.inc
libclc /usr/include/clc/math/atan2pi.h
libclc /usr/include/clc/math/atan2pi.inc
libclc /usr/include/clc/math/atanh.h
libclc /usr/include/clc/math/atanh.inc
libclc /usr/include/clc/math/atanpi.h
libclc /usr/include/clc/math/atanpi.inc
libclc /usr/include/clc/math/binary_decl.inc
libclc /usr/include/clc/math/binary_intrin.inc
libclc /usr/include/clc/math/cbrt.h
libclc /usr/include/clc/math/cbrt.inc
libclc /usr/include/clc/math/ceil.h
libclc /usr/include/clc/math/clc_nextafter.h
libclc /usr/include/clc/math/copysign.h
libclc /usr/include/clc/math/copysign.inc
libclc /usr/include/clc/math/cos.h
libclc /usr/include/clc/math/cos.inc
libclc /usr/include/clc/math/cosh.h
libclc /usr/include/clc/math/cosh.inc
libclc /usr/include/clc/math/cospi.h
libclc /usr/include/clc/math/cospi.inc
libclc /usr/include/clc/math/erf.h
libclc /usr/include/clc/math/erfc.h
libclc /usr/include/clc/math/exp.h
libclc /usr/include/clc/math/exp10.h
libclc /usr/include/clc/math/exp2.h
libclc /usr/include/clc/math/exp2.inc
libclc /usr/include/clc/math/fabs.h
libclc /usr/include/clc/math/fdim.h
libclc /usr/include/clc/math/fdim.inc
libclc /usr/include/clc/math/floor.h
libclc /usr/include/clc/math/fma.h
libclc /usr/include/clc/math/fmax.h
libclc /usr/include/clc/math/fmin.h
libclc /usr/include/clc/math/fmod.h
libclc /usr/include/clc/math/fmod.inc
libclc /usr/include/clc/math/fract.h
libclc /usr/include/clc/math/fract.inc
libclc /usr/include/clc/math/frexp.h
libclc /usr/include/clc/math/frexp.inc
libclc /usr/include/clc/math/gentype.inc
libclc /usr/include/clc/math/half_rsqrt.h
libclc /usr/include/clc/math/half_sqrt.h
libclc /usr/include/clc/math/hypot.h
libclc /usr/include/clc/math/hypot.inc
libclc /usr/include/clc/math/ilogb.h
libclc /usr/include/clc/math/ilogb.inc
libclc /usr/include/clc/math/ldexp.h
libclc /usr/include/clc/math/ldexp.inc
libclc /usr/include/clc/math/lgamma.h
libclc /usr/include/clc/math/lgamma.inc
libclc /usr/include/clc/math/lgamma_r.h
libclc /usr/include/clc/math/lgamma_r.inc
libclc /usr/include/clc/math/log.h
libclc /usr/include/clc/math/log.inc
libclc /usr/include/clc/math/log10.h
libclc /usr/include/clc/math/log1p.h
libclc /usr/include/clc/math/log1p.inc
libclc /usr/include/clc/math/log2.h
libclc /usr/include/clc/math/log2.inc
libclc /usr/include/clc/math/mad.h
libclc /usr/include/clc/math/mad.inc
libclc /usr/include/clc/math/modf.h
libclc /usr/include/clc/math/modf.inc
libclc /usr/include/clc/math/native_cos.h
libclc /usr/include/clc/math/native_divide.h
libclc /usr/include/clc/math/native_exp.h
libclc /usr/include/clc/math/native_exp10.h
libclc /usr/include/clc/math/native_exp2.h
libclc /usr/include/clc/math/native_log.h
libclc /usr/include/clc/math/native_log.inc
libclc /usr/include/clc/math/native_log2.h
libclc /usr/include/clc/math/native_log2.inc
libclc /usr/include/clc/math/native_powr.h
libclc /usr/include/clc/math/native_sin.h
libclc /usr/include/clc/math/native_sqrt.h
libclc /usr/include/clc/math/nextafter.h
libclc /usr/include/clc/math/pow.h
libclc /usr/include/clc/math/pown.h
libclc /usr/include/clc/math/rint.h
libclc /usr/include/clc/math/round.h
libclc /usr/include/clc/math/rsqrt.h
libclc /usr/include/clc/math/sin.h
libclc /usr/include/clc/math/sin.inc
libclc /usr/include/clc/math/sincos.h
libclc /usr/include/clc/math/sincos.inc
libclc /usr/include/clc/math/sinpi.h
libclc /usr/include/clc/math/sinpi.inc
libclc /usr/include/clc/math/sqrt.h
libclc /usr/include/clc/math/sqrt.inc
libclc /usr/include/clc/math/tan.h
libclc /usr/include/clc/math/tan.inc
libclc /usr/include/clc/math/tanh.h
libclc /usr/include/clc/math/tanh.inc
libclc /usr/include/clc/math/ternary_intrin.inc
libclc /usr/include/clc/math/tgamma.h
libclc /usr/include/clc/math/tgamma.inc
libclc /usr/include/clc/math/trunc.h
libclc /usr/include/clc/math/unary_decl.inc
libclc /usr/include/clc/math/unary_intrin.inc
libclc /usr/include/clc/relational/
libclc /usr/include/clc/relational/all.h
libclc /usr/include/clc/relational/any.h
libclc /usr/include/clc/relational/binary_decl.inc
libclc /usr/include/clc/relational/bitselect.h
libclc /usr/include/clc/relational/bitselect.inc
libclc /usr/include/clc/relational/floatn.inc
libclc /usr/include/clc/relational/isequal.h
libclc /usr/include/clc/relational/isfinite.h
libclc /usr/include/clc/relational/isgreater.h
libclc /usr/include/clc/relational/isgreaterequal.h
libclc /usr/include/clc/relational/isinf.h
libclc /usr/include/clc/relational/isless.h
libclc /usr/include/clc/relational/islessequal.h
libclc /usr/include/clc/relational/islessgreater.h
libclc /usr/include/clc/relational/isnan.h
libclc /usr/include/clc/relational/isnormal.h
libclc /usr/include/clc/relational/isnotequal.h
libclc /usr/include/clc/relational/isordered.h
libclc /usr/include/clc/relational/isunordered.h
libclc /usr/include/clc/relational/select.h
libclc /usr/include/clc/relational/signbit.h
libclc /usr/include/clc/relational/unary_decl.inc
libclc /usr/include/clc/shared/
libclc /usr/include/clc/shared/clamp.h
libclc /usr/include/clc/shared/clamp.inc
libclc /usr/include/clc/shared/max.h
libclc /usr/include/clc/shared/max.inc
libclc /usr/include/clc/shared/min.h
libclc /usr/include/clc/shared/min.inc
libclc /usr/include/clc/shared/vload.h
libclc /usr/include/clc/shared/vstore.h
libclc /usr/include/clc/synchronization/
libclc /usr/include/clc/synchronization/barrier.h
libclc /usr/include/clc/synchronization/cl_mem_fence_flags.h
libclc /usr/include/clc/workitem/
libclc /usr/include/clc/workitem/get_global_id.h
libclc /usr/include/clc/workitem/get_global_offset.h
libclc /usr/include/clc/workitem/get_global_size.h
libclc /usr/include/clc/workitem/get_group_id.h
libclc /usr/include/clc/workitem/get_local_id.h
libclc /usr/include/clc/workitem/get_local_size.h
libclc /usr/include/clc/workitem/get_num_groups.h
libclc /usr/include/clc/workitem/get_work_dim.h
libclc /usr/lib/
libclc /usr/lib/clc/
libclc /usr/lib/clc/amdgcn--amdhsa.bc
libclc /usr/lib/clc/aruba-r600--.bc
libclc /usr/lib/clc/barts-r600--.bc
libclc /usr/lib/clc/bonaire-amdgcn--.bc
libclc /usr/lib/clc/bonaire-amdgcn--amdhsa.bc
libclc /usr/lib/clc/bonaire-amdgcn-mesa-mesa3d.bc
libclc /usr/lib/clc/caicos-r600--.bc
libclc /usr/lib/clc/carrizo-amdgcn--.bc
libclc /usr/lib/clc/carrizo-amdgcn--amdhsa.bc
libclc /usr/lib/clc/carrizo-amdgcn-mesa-mesa3d.bc
libclc /usr/lib/clc/cayman-r600--.bc
libclc /usr/lib/clc/cedar-r600--.bc
libclc /usr/lib/clc/cypress-r600--.bc
libclc /usr/lib/clc/fiji-amdgcn--.bc
libclc /usr/lib/clc/fiji-amdgcn--amdhsa.bc
libclc /usr/lib/clc/fiji-amdgcn-mesa-mesa3d.bc
libclc /usr/lib/clc/hainan-amdgcn--.bc
libclc /usr/lib/clc/hainan-amdgcn-mesa-mesa3d.bc
libclc /usr/lib/clc/hawaii-amdgcn--.bc
libclc /usr/lib/clc/hawaii-amdgcn--amdhsa.bc
libclc /usr/lib/clc/hawaii-amdgcn-mesa-mesa3d.bc
libclc /usr/lib/clc/hemlock-r600--.bc
libclc /usr/lib/clc/iceland-amdgcn--.bc
libclc /usr/lib/clc/iceland-amdgcn--amdhsa.bc
libclc /usr/lib/clc/iceland-amdgcn-mesa-mesa3d.bc
libclc /usr/lib/clc/juniper-r600--.bc
libclc /usr/lib/clc/kabini-amdgcn--.bc
libclc /usr/lib/clc/kabini-amdgcn--amdhsa.bc
libclc /usr/lib/clc/kabini-amdgcn-mesa-mesa3d.bc
libclc /usr/lib/clc/kaveri-amdgcn--.bc
libclc /usr/lib/clc/kaveri-amdgcn--amdhsa.bc
libclc /usr/lib/clc/kaveri-amdgcn-mesa-mesa3d.bc
libclc /usr/lib/clc/mullins-amdgcn--.bc
libclc /usr/lib/clc/mullins-amdgcn--amdhsa.bc
libclc /usr/lib/clc/mullins-amdgcn-mesa-mesa3d.bc
libclc /usr/lib/clc/nvptx--nvidiacl.bc
libclc /usr/lib/clc/nvptx64--nvidiacl.bc
libclc /usr/lib/clc/oland-amdgcn--.bc
libclc /usr/lib/clc/oland-amdgcn-mesa-mesa3d.bc
libclc /usr/lib/clc/palm-r600--.bc
libclc /usr/lib/clc/pitcairn-amdgcn--.bc
libclc /usr/lib/clc/pitcairn-amdgcn-mesa-mesa3d.bc
libclc /usr/lib/clc/polaris10-amdgcn--.bc
libclc /usr/lib/clc/polaris10-amdgcn--amdhsa.bc
libclc /usr/lib/clc/polaris10-amdgcn-mesa-mesa3d.bc
libclc /usr/lib/clc/polaris11-amdgcn--.bc
libclc /usr/lib/clc/polaris11-amdgcn--amdhsa.bc
libclc /usr/lib/clc/polaris11-amdgcn-mesa-mesa3d.bc
libclc /usr/lib/clc/redwood-r600--.bc
libclc /usr/lib/clc/stoney-amdgcn--.bc
libclc /usr/lib/clc/stoney-amdgcn--amdhsa.bc
libclc /usr/lib/clc/stoney-amdgcn-mesa-mesa3d.bc
libclc /usr/lib/clc/subnormal_disable.bc
libclc /usr/lib/clc/subnormal_use_default.bc
libclc /usr/lib/clc/sumo-r600--.bc
libclc /usr/lib/clc/sumo2-r600--.bc
libclc /usr/lib/clc/tahiti-amdgcn--.bc
libclc /usr/lib/clc/tahiti-amdgcn-mesa-mesa3d.bc
libclc /usr/lib/clc/tonga-amdgcn--.bc
libclc /usr/lib/clc/tonga-amdgcn--amdhsa.bc
libclc /usr/lib/clc/tonga-amdgcn-mesa-mesa3d.bc
libclc /usr/lib/clc/turks-r600--.bc
libclc /usr/lib/clc/verde-amdgcn--.bc
libclc /usr/lib/clc/verde-amdgcn-mesa-mesa3d.bc
libclc /usr/share/
libclc /usr/share/licenses/
libclc /usr/share/licenses/libclc/
libclc /usr/share/licenses/libclc/LICENSE
libclc /usr/share/pkgconfig/
libclc /usr/share/pkgconfig/libclc.pc
Comment 8 Henrique Dante de Almeida 2017-02-25 20:45:41 UTC
Should I have verde-r600-* ? I'm not using amdgpu.
Comment 9 Henrique Dante de Almeida 2017-02-25 21:05:44 UTC
I'm currently unable to use amdgpu with kernel 4.9. I'm getting this when exiting X and unable to shutdown cleanly:

fev 25 18:00:44 dragonmount kernel: amdgpu 0000:01:00.0: couldn't schedule ib on ring <sdma0>
fev 25 18:00:44 dragonmount kernel: [drm:amdgpu_job_run [amdgpu]] *ERROR* Error scheduling IBs (
fev 25 18:00:44 dragonmount kernel: [drm:amd_sched_main [amdgpu]] *ERROR* Failed to run job!
fev 25 18:00:44 dragonmount kernel: amdgpu 0000:01:00.0: couldn't schedule ib on ring <sdma0>
fev 25 18:00:44 dragonmount kernel: [drm:amdgpu_job_run [amdgpu]] *ERROR* Error scheduling IBs (
fev 25 18:00:44 dragonmount kernel: [drm:amd_sched_main [amdgpu]] *ERROR* Failed to run job!
f


https://bbs.archlinux.org/viewtopic.php?id=222476


Can we fix this for the radeon driver for now ?
Comment 10 Jan Vesely 2017-02-25 21:18:08 UTC
(In reply to Henrique Dante de Almeida from comment #8)
> Should I have verde-r600-* ? I'm not using amdgpu.

no. verde-amdgcn- is correct.

the Linux kernel driver does not matter. the failure is in gpu-kernel build phase.

can you run strace with the failing program (append as attachment, pls).
mesa is probably looking for libclc in wrong place.

another possibility is that the kernel is built using -cl-opt-disable option, but that should not be the case unless the program explicitly requests it.
Comment 11 Henrique Dante de Almeida 2017-02-25 21:34:18 UTC
Created attachment 129915 [details]
OpenCL hello strace
Comment 12 Jan Vesely 2017-02-25 21:49:54 UTC
(In reply to Henrique Dante de Almeida from comment #11)
> Created attachment 129915 [details]
> OpenCL hello strace

thank you.
libclc is located and loaded correctly.
can you attach /usr/lib/clc/tahiti-amdgcn--.bc ?
Comment 13 Henrique Dante de Almeida 2017-02-25 21:55:13 UTC
Created attachment 129916 [details]
Arch Linux libclc tahiti-amdgcn--.bc
Comment 14 Jan Vesely 2017-02-25 23:20:08 UTC
I think the problem is that the libclc implementation contains uninlined calls (function calls are not supported an AMD GPUs):

; Function Attrs: alwaysinline nounwind
define linkonce_odr i32 @get_global_id(i32) local_unnamed_addr #12 {
  switch i32 %0, label %get_group_id.exit [
    i32 0, label %get_group_id.exit.thread
    i32 1, label %get_group_id.exit.thread1
    i32 2, label %get_group_id.exit.thread2
  ]

get_group_id.exit.thread:                         ; preds = %1
  %2 = tail call i32 @llvm.amdgcn.workgroup.id.x() #14
  %3 = tail call i32 bitcast (i64 (i32)* @get_local_size to i32 (i32)*)(i32 0) #18

attributes #12 = { alwaysinline nounwind ...
attributes #18 = { nobuiltin nounwind }

libclc commit 520743b generates this code when compiled using llvm-3.9.
This does not happen with LLVM-4/5.
Comment 15 Jan Vesely 2017-02-26 12:42:37 UTC
Created attachment 129929 [details] [review]
size_t is 32bit in mesa3d

building libclc (commit 520743) with this patch applied should produce a working library
Comment 16 Henrique Dante de Almeida 2017-02-26 23:12:45 UTC
Weird ! So was this a bug in IR source code ?
Comment 17 Jan Vesely 2017-02-27 14:37:08 UTC
(In reply to Henrique Dante de Almeida from comment #16)
> Weird ! So was this a bug in IR source code ?

more like a combination of llvm/libclc. the function would ideally be rewritten in clc.
can you confirm that the patch works (i.e.the resulting libclc works in your setup) ? I could only test basic kernel build.
thank you
Comment 18 Aidan Thornton 2017-02-27 15:44:24 UTC
While that patch fixes the problem with get_local_size at least for me, it appears that get_global_size and get_num_groups have the exact same problem and require the same fix. The simple hello world example doesn't require them but this breaks more complex OpenCL code which does with a almost identical error. Making the same change to both of these appears to give me a working libclc on Cape Verde. Hopefully it will for Henrique too.
Comment 19 Mig 2017-02-27 18:40:16 UTC
Hello,
I built the libclc package with the patch applied with the latest llvm/clang (5.0.0svn) on arch, and getting the following error

1. Device: AMD TONGA (DRM 3.8.0 / 4.9.11-1-ARCH, LLVM 3.9.1)
Error: Failed to build program executable!
<unknown>:0:0: in function square void (float addrspace(1)*, float addrspace(1)*, i32): unsupported call to function get_local_size

How can I fix this?
Comment 20 Mig 2017-02-27 18:56:19 UTC
That was the old error, sorry. After applying the patch the error is:
. Device: AMD TONGA (DRM 3.8.0 / 4.9.11-1-ARCH, LLVM 3.9.1)
'+fp64-fp16-denormals' is not a recognized feature for this target (ignoring feature)
'+fp64-fp16-denormals' is not a recognized feature for this target (ignoring feature)
'+fp64-fp16-denormals' is not a recognized feature for this target (ignoring feature)
'+fp64-fp16-denormals' is not a recognized feature for this target (ignoring feature)
Error: Failed to build program executable!
<unknown>:0:0: in function square void (float addrspace(1)*, float addrspace(1)*, i32): unsupported call to function get_global_id
Comment 21 Jan Vesely 2017-02-27 20:25:51 UTC
(In reply to Mig from comment #19)
> Hello,
> I built the libclc package with the patch applied with the latest llvm/clang
> (5.0.0svn) on arch, and getting the following error
> 
> 1. Device: AMD TONGA (DRM 3.8.0 / 4.9.11-1-ARCH, LLVM 3.9.1)

you should not mix llvm versions. libclc should be compiled with the same version that you use with mesa/clover (3.9.1 in you case)
Comment 22 Henrique Dante de Almeida 2017-02-28 22:10:16 UTC
Jan, the patch fixes the problem and hello world works here. Should I report the bug in llvm bug tracker ?
Comment 23 Mig 2017-02-28 22:45:30 UTC
I recompiled with the default llvm compiler 3.9.1 and it works for me too!

Number of platforms: 2 
1. Device: AMD TONGA (DRM 3.8.0 / 4.9.11-1-ARCH, LLVM 3.9.1)
success: got back 1 binaries, total size 8434
binary 0: size 8434 dumped to square0.gallium_bin
Computed '1024/1024' correct values!
Comment 24 Jan Vesely 2017-03-01 17:58:49 UTC
(In reply to Henrique Dante de Almeida from comment #22)
> Jan, the patch fixes the problem and hello world works here. Should I report
> the bug in llvm bug tracker ?

I'm not sure if that'd help. There is no release_39 stable branch for libclc. Even if there was, afaik there is no plan for another 3.9 release.
distros will probably have to apply that patch individually (Aidan's version might be better as it fixes more instances of the problem).

another option would be to rewrite all 3 files to clc, but that'd need mesa changes

I'm not sure what the preferred solution would be.

@Tom, any ideas?
Comment 25 Jan Vesely 2017-03-09 05:01:21 UTC
Created attachment 130136 [details] [review]
size_t is 32bit in mesa3d

more complete patch
Comment 26 Jan Vesely 2017-09-15 22:04:15 UTC
There are no further fixes planned for clover on llvm-3.9.
sorry
Comment 27 Jan Vesely 2017-09-21 02:31:34 UTC
Created attachment 134395 [details] [review]
Restore support for llvm-3.9

Hi,

the attached patch restores support for llvm-3.9 to libclc HEAD and fixes the size_t problem.
I don't have the right combination of HW + system running llvm-3.9 so any testing is appreciated.
Comment 28 Jan Vesely 2017-09-29 19:11:05 UTC
Fixed in libclc >=r314543


Use of freedesktop.org services, including Bugzilla, is subject to our Code of Conduct. How we collect and use information is described in our Privacy Policy.