[Libclc-dev] vload problems (was: Re: [PATCH 1/2] amdgcn: Implement {read_, write_, }mem_fence builtin)

[Jan Vesely via Libclc-dev]

On Thu, 2017-08-17 at 21:53 -0500, Aaron Watry wrote:
> On Wed, Aug 16, 2017 at 6:21 PM, Jan Vesely <jan.vesely at rutgers.edu> wrote:
> > On Sat, 2017-08-12 at 13:18 -0500, Aaron Watry wrote:
> > > On Fri, Aug 11, 2017 at 8:59 PM, Jan Vesely via Libclc-dev
> > > <libclc-dev at lists.llvm.org> wrote:
> > > > Signed-off-by: Jan Vesely <jan.vesely at rutgers.edu>
> > > > ---
> > > > 
> > > > __builtin_amdgcn_s_waitcnt path is compile tested only. I currently
> > > > don't have machine with GCN hw and LLVM > 4
> > > 
> > > Tested on GCN 1.0 (PITCAIRN) with LLVM 6.0 svn (from sometime last week).
> > > 
> > > I ran the local-memory/global-memory piglit tests, and the conformance
> > > basic/test_basic vload*, vstore*, and barrier tests. No change in
> > > pass-rate.
> > > 
> > > The only ones that fail are vload_private with char/uchar/short/ushort
> > > data types, but those failed before your series.
> > 
> > do you remember if those failures are new in LLVM 6? I just posted new
> > vload/vstore piglit tests, and those ran OK* on carrizo/iceland system
> > using LLVM 5.
> > 
> > *mostly. vload_half is missing from libclc to those test failed.
> > Unlike Turks, which fails ~40% of them (LLVM 6).
> 
> I don't recall running this test specifically before the 5.0rc
> branch-point, so I don't know if this is a new failure.  All of the
> vload tests for local/global/constant seem ok, and vload private is
> working for anything that has a 32-bit gentype or larger
> (int/long/float/double). It's just the vloads of
> char/uchar/short/ushort that are failing.  I've been looking at the
> vload code, and at least the CL C code we've got looks correct.
> 
> The CTS test allows you to tweak the number of data points that are
> tested in individual threads, and curiously, when the global size goes
> from 15 to 16, things start failing (at least for the char2 test).

If it's intermittent I'd suspect buffer manipulation rather than GPU
execution, but I might be wrong.
I mostly wanted to find out whether the new piglit tests hit the issue,
or the situation is more complicated.
Modified vstore tests managed to uncover fail in vstore-{u,}char-
private, but that's assertion failure, which I assume is different from
what you're seeing in CTS.

http://paul.rutgers.edu/~jv356/piglit/gcn-latest-3/problems.html

> 
> I've diffed the clover_dump.cl and .ll files for both, and they're
> identical.  I copied all of the test data into a piglit test, and the
> same CL code and input/output data passes in piglit... The piglit CL
> and .ll files are identical to what is executed by the CTS other than
> the giant piglit comment section at the top related to buffer/test
> setup.  Possibly a difference in API usage, or buffer alignments, or
> something.
> 
> I'm not entirely sure that the vload code is to blame here, but it's
> possible. It's working just fine for global/local/constant vload
> (which doesn't rule out any weirdness with private memory in GCN
> having possible alignment/register-size restrictions that I'm not
> aware of).

VI+ GCN parts support i16 instructions, and there was a similar bug
before: https://reviews.llvm.org/D30281

> I think I'm going to spend a little more time working on
> making sure that the create/write/read buffer pieces in clover (and
> all of our synchronization pieces) are solid before I spend too much
> more time here.

thanks. clover on GCN is not really high on my list. my gcn machine is
setup with rocm.
Moreover, upgrading to LLVM 5 regressed both pyrite (can't select
device) and bfgminer (flood of acpi errors in dmesg).

> 
> Ooh, fun fact...  I just tried to run the test on my Barts (6850/NI),
> and it hung the machine on the first test.  At least my GCN just fails
> the test...

this looks related to the ongoing clpeak reported bug, the kernel
driver is hit or miss when it comes to recovering hung GPU.

> 
> Anyway, I'll be out of communication for a few days.  New kid arriving
> tomorrow, so I'll be buried in parents/in-laws and lacking in sleep.

congratz, and no worries, there's no rush when it comes to clover :)

> 
> If you want to take a look at the .cl/ll files I used from piglit, I'm
> attaching them.

thanks,
Jan

> 
> --Aaron
> 
> > 
> > regards,
> > Jan
> > 
> > [SNIP]

-- 
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/20170818/6c90af04/attachment.sig>