[Libclc-dev] [PATCH] atomics: redefine atom_inc/atom_dec using atom_add/atom_sub
Aaron Watry via Libclc-dev
libclc-dev at lists.llvm.org
Fri Jun 23 18:33:55 PDT 2017
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. ;-)
--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
More information about the Libclc-dev
mailing list