<html>
<head>
<meta http-equiv="Content-Type" content="text/html; charset=utf-8">
</head>
<body text="#000000" bgcolor="#FFFFFF">
<br>
<div class="moz-cite-prefix">On 06/21/2018 02:32 PM, Justin Lebar
wrote:<br>
</div>
<blockquote type="cite"
cite="mid:CAMuNMfpgm_OvRFOkG8hjcNCggdJHki=Ly-1spG1A2AkEGsExrQ@mail.gmail.com">
<meta http-equiv="Content-Type" content="text/html; charset=utf-8">
We already have a pass that vectorizes loads and stores in nvptx
and amdgpu. Not at my laptop, I forget the exact filename, but
it's called load-store vectorizer.</blockquote>
<br>
It's here: lib/Transforms/Vectorize/LoadStoreVectorizer.cpp<br>
<br>
I agree that, if LLVM can do this explicitly, it seems better.<br>
<br>
-Hal<br>
<br>
<blockquote type="cite"
cite="mid:CAMuNMfpgm_OvRFOkG8hjcNCggdJHki=Ly-1spG1A2AkEGsExrQ@mail.gmail.com">
<div><br>
</div>
<div>I think the question is, why is LSV not vectorizing this
code?</div>
<div><br>
</div>
<div>I think the answer is, llvm can't tell that the loads are
aligned. Ptxas can, but only because it's (apparently) doing
vectorization *after* it reesolves the shmem variables to
physical addresses. That is a cool trick, and llvm can't do it,
because llvm never sees the physical shmem addresses.</div>
<div><br>
</div>
<div>If you told llvm that the shmem variables were aligned to 16
bytes, LSV might do what you want here. llvm and ptxas should
be able to cooperate to give you the alignment you ask for in
the IR.</div>
<div><br>
</div>
<div>If that doesn't work I'd recommend trying to debug the LSV
code to see why it's not vectorizing. You can run `opt
-load-store-vectorizer -debug` -- or, my favorite way of doing
it is to run that command under <a href="http://rr-project.org"
moz-do-not-send="true">rr-project.org</a>.</div>
<div><br>
</div>
<div>It's possible that clang should opportunistically mark all
shmem variables over a certain size as align(16) so that this
happens automagically. That would kind of be a weird heuristic,
but maybe it makes sense. I don't think that would make sense
for LLVM to do that, though, so it wouldn't help you.</div>
<div><br>
</div>
<div>I think relying on LSV to do its job is better than messing
with the instruction order because the former is more powerful
-- it can vectorize in cases where ptxas would have a much
harder time.</div>
<div><br>
</div>
<div>Justin<br>
<div dir="auto"><br>
<div class="gmail_quote">
<div dir="ltr">On Thu, Jun 21, 2018, 7:48 AM Hal Finkel via
llvm-dev <<a href="mailto:llvm-dev@lists.llvm.org"
moz-do-not-send="true">llvm-dev@lists.llvm.org</a>>
wrote:<br>
</div>
<blockquote class="gmail_quote" style="margin:0 0 0
.8ex;border-left:1px #ccc solid;padding-left:1ex"><br>
On 06/21/2018 12:18 PM, Tim Besard via llvm-dev wrote:<br>
> Hi all,<br>
><br>
> I'm looking into the performance difference of a
benchmark compiled with<br>
> NVCC vs NVPTX (coming from Julia, not CUDA C) and I'm
seeing a<br>
> significant difference due to PTX instruction
ordering. The relevant<br>
> source code consists of two nested loops that get
fully unrolled, doing<br>
> some basic arithmetic with values loaded from shared
memory:<br>
><br>
>> #define BLOCK_SIZE 16<br>
>><br>
>> __shared__ float dia[BLOCK_SIZE][BLOCK_SIZE];<br>
>> __shared__ float
peri_col[BLOCK_SIZE][BLOCK_SIZE];<br>
>><br>
>> int idx = threadIdx.x - BLOCK_SIZE;<br>
>> for (int i = 0; i < BLOCK_SIZE; i++) {<br>
>> for (int j = 0; j < i; j++)<br>
>> peri_col[idx][i] -= peri_col[idx][j] *
dia[j][i];<br>
>> peri_col[idx][i] /= dia[i][i];<br>
>> }<br>
> NVCC emits PTX instructions where all loads from
shared memory are<br>
> packed together:<br>
><br>
>> ...<br>
>> ld.shared.f32 %f546, [kernel_dia+440];<br>
>> ld.shared.f32 %f545, [%r4+-996];<br>
>> ld.shared.f32 %f544, [kernel_dia+56];<br>
>> ld.shared.f32 %f543, [kernel_dia+88];<br>
>> ld.shared.f32 %f542, [kernel_dia+500];<br>
>> ld.shared.f32 %f541, [kernel_dia+84];<br>
>> ld.shared.f32 %f540, [%r4+-972];<br>
>> ld.shared.f32 %f539, [%r4+-1008];<br>
>> ld.shared.f32 %f538, [kernel_dia+496];<br>
>> ld.shared.f32 %f537, [kernel_dia+136];<br>
>> ld.shared.f32 %f536, [%r4+-976];<br>
>> ld.shared.f32 %f535, [kernel_dia+428];<br>
>> ... # hundreds of these<br>
> Even though this heavily bloats register usage (and
NVCC seems to do<br>
> this unconditionally, even with launch configurati<a
href="https://maps.google.com/?q=ons+whe&entry=gmail&source=g"
moz-do-not-send="true">ons whe</a>re this could<br>
> <a
href="https://maps.google.com/?q=hurt+p&entry=gmail&source=g"
moz-do-not-send="true">hurt p</a>erformance), i<a
href="https://maps.google.com/?q=t+allow&entry=gmail&source=g"
moz-do-not-send="true">t allow</a>s the CUDA PTX JIT to
emit 128-bit loads:<br>
><br>
>> LDS.128 R76, [0x2f0];<br>
>> LDS.128 R60, [0xa0];<br>
>> LDS.128 R72, [0x130];<br>
>> LDS.128 R96, [0x1b0];<br>
>> LDS.128 R92, [0x30];<br>
>> LDS.128 R116, [0x50];<br>
>> LDS.128 R108, [0x1f0];<br>
> LLVM preserves the operations more or less as they
are emitted by the<br>
> front-end, interleaving memory operations with
arithmetic. As a result,<br>
> the SASS code contains many more 32-bit loads, which
lowers performance<br>
> by ~10% on this specific benchmark.<br>
><br>
> What would be the best approach to improve generated
code? I can imagine<br>
> a late IR pass shuffling instructions around, but I
figured I'd ask to<br>
> see if this is a good approach and whether there's
existing work doing<br>
> similar transformations.<br>
<br>
You could make a custom pass, late IR or MI. You might
also be able to<br>
use the existing instruction-scheduling infrastructure.
You can<br>
implement ScheduleDAGMutation that does the clustering
that you'd like,<br>
or if the existing ones do what you want, use those. We
have preexisting<br>
createLoadClusterDAGMutation and
createStoreClusterDAGMutation<br>
functions. If you look at AMDGPU/AMDGPUTargetMachine.cpp,
you'll see<br>
calls like this:<br>
<br>
DAG->addMutation(createLoadClusterDAGMutation(DAG->TII,
DAG->TRI));<br>
<br>
and I think that you probably want to do the same.<br>
<br>
Also, you might want to override the subtarget's useAA()
method to<br>
return true (as this gives more freedom to the scheduler
to move memory<br>
accesses around to do this kind of clustering).<br>
<br>
-Hal<br>
<br>
><br>
> Thanks,<br>
<br>
-- <br>
Hal Finkel<br>
Lead, Compiler Technology and Programming Languages<br>
Leadership Computing Facility<br>
Argonne National Laboratory<br>
<br>
_______________________________________________<br>
LLVM Developers mailing list<br>
<a href="mailto:llvm-dev@lists.llvm.org" target="_blank"
moz-do-not-send="true">llvm-dev@lists.llvm.org</a><br>
<a
href="http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev"
rel="noreferrer" target="_blank" moz-do-not-send="true">http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev</a><br>
</blockquote>
</div>
</div>
</div>
</blockquote>
<br>
<pre class="moz-signature" cols="72">--
Hal Finkel
Lead, Compiler Technology and Programming Languages
Leadership Computing Facility
Argonne National Laboratory</pre>
</body>
</html>