[Libclc-dev] [PATCH] atomics: redefine atom_inc/atom_dec using atom_add/atom_sub
Jan Vesely via Libclc-dev
libclc-dev at lists.llvm.org
Mon Jun 26 04:18:13 PDT 2017
On Fri, 2017-06-23 at 20:33 -0500, Aaron Watry via Libclc-dev wrote:
> On Thu, Jun 22, 2017 at 10:25 AM, Jan Vesely <jan.vesely at rutgers.edu> wrote:
> > On Wed, 2017-06-21 at 11:06 -0500, Aaron Watry via Libclc-dev wrote:
> > > On Wed, Jun 21, 2017 at 9:16 AM, Jan Vesely <jan.vesely at rutgers.edu> wrote:
> > > > On Tue, 2017-06-20 at 21:11 -0500, Aaron Watry via Libclc-dev wrote:
> > > > > This is exactly what the atomic_[inc|dec] functions do, and it fixes
> > > > > kernel compilation failures using the OpenCL CTS.
> > > >
> > > > what are the compilation failures? I thought piglit had these ops
> > > > covered.
> > >
> > > Piglit has atomic_* covered, but not the CL 1.0 atom_* functions.
> >
> > ah right. guess we should add those as well (another time...).
>
> Yeah. In theory, we can add those to piglit. Since the CTS was
> open-sourced, I'm mostly using its failures as a worklist for now just
> to prevent duplication of work.
>
> >
> > >
> > > The issue I'm running into with the CTS is that atomic_* works fine,
> > > but I'm getting an ambiguous function reference for
> > > atom_inc/atom_dec/atom_or/atom_and.
> > >
> > > Example:
> > >
> > > #pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
> > > __kernel void test_atomic_fn(volatile __global uint *destMemory,
> > > __global uint *oldValues)
> > > {
> > > int tid = get_global_id(0);
> > > oldValues[tid] = atom_dec( &destMemory[0] );
> > > }
> > >
> > > Log:
> > > Build not successful for device "AMD PITCAIRN (DRM 2.49.0 /
> > > 4.11.0-041100rc8-generic, LLVM 5.0.0)", status: CL_BUILD_ERROR
> > > Build log for device "AMD PITCAIRN (DRM 2.49.0 /
> > > 4.11.0-041100rc8-generic, LLVM 5.0.0)" is:
> > > ------------
> > > input.cl:6:22: error: call to 'atom_dec' is ambiguous
> > > /usr/local/include/clc/cl_khr_global_int32_base_atomics/atom_dec.h:1:29:
> > > note: candidate function
> > > /usr/local/include/clc/cl_khr_global_int32_base_atomics/atom_dec.h:2:38:
> > > note: candidate function
> > > ------------
> >
> > this looks like clang treats both:
> > "volatile __global uint* -> __global unsigned int*"
> > and
> > "volatile __global uint* -> __globale int *"
> > as equal cost.
> >
> > Can you confirm that removing 'volatile' from destMemory hides the
> > issue?
>
> Yes, I can confirm that if I rebuild libclc using the latest upstream
> revision (and without any of my changes)
> and then remove the volatile keyword from the atom_* test kernels
> being compiled that the
> tests compile and run to completion successfully.
>
> >
> > >
> > > >
> > > > >
> > > > > Signed-off-by: Aaron Watry <awatry at gmail.com>
> > > > > ---
> > > > > generic/include/clc/cl_khr_global_int32_base_atomics/atom_dec.h | 3 +--
> > > > > generic/include/clc/cl_khr_global_int32_base_atomics/atom_inc.h | 3 +--
> > > > > generic/include/clc/cl_khr_local_int32_base_atomics/atom_dec.h | 3 +--
> > > > > generic/include/clc/cl_khr_local_int32_base_atomics/atom_inc.h | 3 +--
> > > > > generic/lib/SOURCES | 4 ----
> > > > > generic/lib/cl_khr_global_int32_base_atomics/atom_dec.cl | 9 ---------
> > > > > generic/lib/cl_khr_global_int32_base_atomics/atom_inc.cl | 9 ---------
> > > > > generic/lib/cl_khr_local_int32_base_atomics/atom_dec.cl | 9 ---------
> > > > > generic/lib/cl_khr_local_int32_base_atomics/atom_inc.cl | 9 ---------
> > > > > 9 files changed, 4 insertions(+), 48 deletions(-)
> > > > > delete mode 100644 generic/lib/cl_khr_global_int32_base_atomics/atom_dec.cl
> > > > > delete mode 100644 generic/lib/cl_khr_global_int32_base_atomics/atom_inc.cl
> > > > > delete mode 100644 generic/lib/cl_khr_local_int32_base_atomics/atom_dec.cl
> > > > > delete mode 100644 generic/lib/cl_khr_local_int32_base_atomics/atom_inc.cl
> > > > >
> > > > > diff --git a/generic/include/clc/cl_khr_global_int32_base_atomics/atom_dec.h b/generic/include/clc/cl_khr_global_int32_base_atomics/atom_dec.h
> > > > > index bbc872c..a520fe4 100644
> > > > > --- a/generic/include/clc/cl_khr_global_int32_base_atomics/atom_dec.h
> > > > > +++ b/generic/include/clc/cl_khr_global_int32_base_atomics/atom_dec.h
> > > > > @@ -1,2 +1 @@
> > > > > -_CLC_OVERLOAD _CLC_DECL int atom_dec(global int *p);
> > > > > -_CLC_OVERLOAD _CLC_DECL unsigned int atom_dec(global unsigned int *p);
> > > > > +#define atom_dec(p) atom_sub(p, 1)
> > > >
> > > > Does this not cause redefinition warning (since local version uses the
> > > > same define?
> > >
> > > Yes, it looks like it does. When building libclc and running the CTS
> > > tests, everything looks ok, but if I compile the atom_dec kernel
> > > manually via clang, I get:
> > > --------------------------------
> > > atom_dec_uint.cl:5:32: warning: passing 'volatile __global uint *'
> > > (aka 'volatile __global unsigned int *') to parameter of type
> > > '__global int *' discards qualifiers
> > > [-Wincompatible-pointer-types-discards-qualifiers]
> > > oldValues[tid] = atom_dec( &destMemory[0] );
> > > ^~~~~~~~~~~~~~
> > > /usr/local/include/clc/cl_khr_local_int32_base_atomics/atom_dec.h:9:30:
> > > note: expanded from macro 'atom_dec'
> > > #define atom_dec(p) atom_sub(p, 1)
> > > ^
> > > /usr/local/include/clc/cl_khr_global_int32_base_atomics/atom_sub.h:10:50:
> > > note: passing argument to parameter 'p' here
> > > _CLC_OVERLOAD _CLC_DECL int atom_sub(global int *p, int val);
> > > ^
> > > 1 warning generated.
> > > --------------------------------
> > >
> > > So yeah, I guess this is the wrong approach.
> > >
> > > I did just discover while poking around at the kernel that if I remove
> > > the "volatile" keyword from destMemory, then the kernel builds fine.
> > > The difference in the atomic_inc and the CL 1.0 atom_inc function
> > > declarations is basically that volatile keyword. If I re-add the
> > > volatile keyword to the kernel, and then also add it to the function
> > > declaration in atom_dec.h, the kernel also builds. Looking at POCL
> > > and Beignet they both just do a single blanket define in a top-level
> > > header file along the lines of:
> > > #define atom_dec atomic_dec
> > >
> > > That implies that both of those runtimes declare the global/local
> > > pointers as volatile even for the CL1.0 variants, which doesn't
> > > necessarily match the spec (although given how the CL 1.1/1.2 spec are
> > > written, they basically just say "the CL 1.0 atom_* functions are
> > > still supported), but it is what the CTS is testing. I don't
> > > currently have access to the definitions that nv/amd use in their
> > > closed-source binaries (and I haven't managed to find what ROCm uses,
> > > but I'm thinking that the implication is that when Khronos
> > > renamed/aliased atom_* to atomic_* in CL 1.1, they were implicitly
> > > retroactively adding the 'volatile' keyword to the function argument.
> > >
> > > I guess we could just remove the existing defines under
> > > cl_khr_[global|local]_int32_base_atomics and put a single set of
> > > defines in either clc.h, or in a new header under something like
> > > generic/include/clc/atomic/atom_functions.h that just defines the
> > > renames of the existing functions. Alternatively, we could just go
> > > back and add the volatile keyword to the pointer arguments in the
> > > atom_* headers/implementations.
> > >
> > > Thoughts?
> >
> > I checked that we also use similar define for atomic_inc/dec.
> > 9040bf38 addressed similar issue when calling atom_add.
> > I think adding volatile to atom_* functions would just hide the issue
> > (since test suites use volatile pointers).
> > see [0]:
> > "A conversion from a pointer of type ``T*`` to a pointer of type ``U*``, where
> > ``T`` and ``U`` are incompatible, is allowed, but is ranked below all other
> > types of conversions. Please note: ``U`` lacking qualifiers that are present
> > on ``T`` is sufficient for ``T`` and ``U`` to be incompatible."
> >
> > I think a proper fix needs to be done on clang side.
>
> Yes, the commit message for that llvm revision you linked sounds
> exactly like what is going on here.
>
> Ugh, looks like I need to read up on llvm development guidelines again
> (a codebase that I try to limit my exposure to for my own sanity),
> unless you want to look into that part. ;-)
Clang is not really my thing, and this issue is a rather low priority
for me. completing atomics support or adding register spilling for EG
are higher on my list.
You can try contacting Yaxun, he appeared to run into the same issue
and contributes to clang regularly.
regards,
Jan
>
> --Aaron
>
> >
> > Jan
> >
> > [0] https://reviews.llvm.org/D24113
> >
> > >
> > > --Aaron
> > >
> > > >
> > > > Jan
> > > >
> > > > > diff --git a/generic/include/clc/cl_khr_global_int32_base_atomics/atom_inc.h b/generic/include/clc/cl_khr_global_int32_base_atomics/atom_inc.h
> > > > > index 050747c..2b4436d 100644
> > > > > --- a/generic/include/clc/cl_khr_global_int32_base_atomics/atom_inc.h
> > > > > +++ b/generic/include/clc/cl_khr_global_int32_base_atomics/atom_inc.h
> > > > > @@ -1,2 +1 @@
> > > > > -_CLC_OVERLOAD _CLC_DECL int atom_inc(global int *p);
> > > > > -_CLC_OVERLOAD _CLC_DECL unsigned int atom_inc(global unsigned int *p);
> > > > > +#define atom_inc(p) atom_add(p, 1)
> > > > > diff --git a/generic/include/clc/cl_khr_local_int32_base_atomics/atom_dec.h b/generic/include/clc/cl_khr_local_int32_base_atomics/atom_dec.h
> > > > > index e74d8fc..a520fe4 100644
> > > > > --- a/generic/include/clc/cl_khr_local_int32_base_atomics/atom_dec.h
> > > > > +++ b/generic/include/clc/cl_khr_local_int32_base_atomics/atom_dec.h
> > > > > @@ -1,2 +1 @@
> > > > > -_CLC_OVERLOAD _CLC_DECL int atom_dec(local int *p);
> > > > > -_CLC_OVERLOAD _CLC_DECL unsigned int atom_dec(local unsigned int *p);
> > > > > +#define atom_dec(p) atom_sub(p, 1)
> > > > > diff --git a/generic/include/clc/cl_khr_local_int32_base_atomics/atom_inc.h b/generic/include/clc/cl_khr_local_int32_base_atomics/atom_inc.h
> > > > > index 718f1f2..2b4436d 100644
> > > > > --- a/generic/include/clc/cl_khr_local_int32_base_atomics/atom_inc.h
> > > > > +++ b/generic/include/clc/cl_khr_local_int32_base_atomics/atom_inc.h
> > > > > @@ -1,2 +1 @@
> > > > > -_CLC_OVERLOAD _CLC_DECL int atom_inc(local int *p);
> > > > > -_CLC_OVERLOAD _CLC_DECL unsigned int atom_inc(local unsigned int *p);
> > > > > +#define atom_inc(p) atom_add(p, 1)
> > > > > diff --git a/generic/lib/SOURCES b/generic/lib/SOURCES
> > > > > index 9e0157b..3386e8d 100644
> > > > > --- a/generic/lib/SOURCES
> > > > > +++ b/generic/lib/SOURCES
> > > > > @@ -8,8 +8,6 @@ atomic/atomic_xchg.cl
> > > > > atomic/atomic_impl.ll
> > > > > cl_khr_global_int32_base_atomics/atom_add.cl
> > > > > cl_khr_global_int32_base_atomics/atom_cmpxchg.cl
> > > > > -cl_khr_global_int32_base_atomics/atom_dec.cl
> > > > > -cl_khr_global_int32_base_atomics/atom_inc.cl
> > > > > cl_khr_global_int32_base_atomics/atom_sub.cl
> > > > > cl_khr_global_int32_base_atomics/atom_xchg.cl
> > > > > cl_khr_global_int32_extended_atomics/atom_and.cl
> > > > > @@ -19,8 +17,6 @@ cl_khr_global_int32_extended_atomics/atom_or.cl
> > > > > cl_khr_global_int32_extended_atomics/atom_xor.cl
> > > > > cl_khr_local_int32_base_atomics/atom_add.cl
> > > > > cl_khr_local_int32_base_atomics/atom_cmpxchg.cl
> > > > > -cl_khr_local_int32_base_atomics/atom_dec.cl
> > > > > -cl_khr_local_int32_base_atomics/atom_inc.cl
> > > > > cl_khr_local_int32_base_atomics/atom_sub.cl
> > > > > cl_khr_local_int32_base_atomics/atom_xchg.cl
> > > > > cl_khr_local_int32_extended_atomics/atom_and.cl
> > > > > diff --git a/generic/lib/cl_khr_global_int32_base_atomics/atom_dec.cl b/generic/lib/cl_khr_global_int32_base_atomics/atom_dec.cl
> > > > > deleted file mode 100644
> > > > > index cc24d2f..0000000
> > > > > --- a/generic/lib/cl_khr_global_int32_base_atomics/atom_dec.cl
> > > > > +++ /dev/null
> > > > > @@ -1,9 +0,0 @@
> > > > > -#include <clc/clc.h>
> > > > > -
> > > > > -#define IMPL(TYPE) \
> > > > > -_CLC_OVERLOAD _CLC_DEF TYPE atom_dec(global TYPE *p) { \
> > > > > - return atom_sub(p, (TYPE)1); \
> > > > > -}
> > > > > -
> > > > > -IMPL(int)
> > > > > -IMPL(unsigned int)
> > > > > diff --git a/generic/lib/cl_khr_global_int32_base_atomics/atom_inc.cl b/generic/lib/cl_khr_global_int32_base_atomics/atom_inc.cl
> > > > > deleted file mode 100644
> > > > > index 9193ae3..0000000
> > > > > --- a/generic/lib/cl_khr_global_int32_base_atomics/atom_inc.cl
> > > > > +++ /dev/null
> > > > > @@ -1,9 +0,0 @@
> > > > > -#include <clc/clc.h>
> > > > > -
> > > > > -#define IMPL(TYPE) \
> > > > > -_CLC_OVERLOAD _CLC_DEF TYPE atom_inc(global TYPE *p) { \
> > > > > - return atom_add(p, (TYPE)1); \
> > > > > -}
> > > > > -
> > > > > -IMPL(int)
> > > > > -IMPL(unsigned int)
> > > > > diff --git a/generic/lib/cl_khr_local_int32_base_atomics/atom_dec.cl b/generic/lib/cl_khr_local_int32_base_atomics/atom_dec.cl
> > > > > deleted file mode 100644
> > > > > index cfb3d80..0000000
> > > > > --- a/generic/lib/cl_khr_local_int32_base_atomics/atom_dec.cl
> > > > > +++ /dev/null
> > > > > @@ -1,9 +0,0 @@
> > > > > -#include <clc/clc.h>
> > > > > -
> > > > > -#define IMPL(TYPE) \
> > > > > -_CLC_OVERLOAD _CLC_DEF TYPE atom_dec(local TYPE *p) { \
> > > > > - return atom_sub(p, (TYPE)1); \
> > > > > -}
> > > > > -
> > > > > -IMPL(int)
> > > > > -IMPL(unsigned int)
> > > > > diff --git a/generic/lib/cl_khr_local_int32_base_atomics/atom_inc.cl b/generic/lib/cl_khr_local_int32_base_atomics/atom_inc.cl
> > > > > deleted file mode 100644
> > > > > index 8ea4738..0000000
> > > > > --- a/generic/lib/cl_khr_local_int32_base_atomics/atom_inc.cl
> > > > > +++ /dev/null
> > > > > @@ -1,9 +0,0 @@
> > > > > -#include <clc/clc.h>
> > > > > -
> > > > > -#define IMPL(TYPE) \
> > > > > -_CLC_OVERLOAD _CLC_DEF TYPE atom_inc(local TYPE *p) { \
> > > > > - return atom_add(p, (TYPE)1); \
> > > > > -}
> > > > > -
> > > > > -IMPL(int)
> > > > > -IMPL(unsigned int)
> > >
> > > _______________________________________________
> > > Libclc-dev mailing list
> > > Libclc-dev at lists.llvm.org
> > > http://lists.llvm.org/cgi-bin/mailman/listinfo/libclc-dev
>
> _______________________________________________
> Libclc-dev mailing list
> Libclc-dev at lists.llvm.org
> http://lists.llvm.org/cgi-bin/mailman/listinfo/libclc-dev
--
Jan Vesely <jan.vesely at rutgers.edu>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 833 bytes
Desc: This is a digitally signed message part
URL: <http://lists.llvm.org/pipermail/libclc-dev/attachments/20170626/c83e24bf/attachment.sig>
More information about the Libclc-dev
mailing list