<div dir="ltr">Forgot to reply to all. <div class="gmail_extra"><br><br><div class="gmail_quote">On Mon, Mar 24, 2014 at 12:30 PM, Justin Holewinski <span dir="ltr"><<a href="mailto:jholewinski@nvidia.com" target="_blank">jholewinski@nvidia.com</a>></span> wrote:<br>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<div bgcolor="#FFFFFF" text="#000000">
Okay, I see. That makes sense. Thanks!<div><div class="h5"><br>
<br>
<div>On 03/24/2014 03:22 PM, Jingyue Wu
wrote:<br>
</div>
<blockquote type="cite">
<div dir="ltr">
<div>We discussed this offline with Peter before. __shared__ is
a declaration qualifier instead of a type qualifier.
Therefore, clang will only generate (float *) in the AST,
because it wants to faithfully model the type system. it's
true that codegen creates an addrspace(3) allocation for 'a',
but "a" is still declared and used as float * in the AST. This
is why we fix codegen to addrspacecast every __shared__
variable. If __shared__ is part of the type system, I would
agree to add the __shared__ type in the AST and emit the code
you expected. </div>
<div><br>
</div>
Regarding "ExpectedAddrSpace != AddrSpace", as aforementioned,
AddrSpace is the allocation address space (3 for __shared__),
and ExpectedAddrSpace is the declared address space (0 in this
case). Therefore, they are different, and need an
addrspacecast.
<div>
<br>
</div>
<div>Jingyue</div>
</div>
<div class="gmail_extra"><br>
<br>
<div class="gmail_quote">On Mon, Mar 24, 2014 at 11:36 AM,
Justin Holewinski <span dir="ltr"><<a href="mailto:jholewinski@nvidia.com" target="_blank">jholewinski@nvidia.com</a>></span>
wrote:<br>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<div bgcolor="#FFFFFF" text="#000000"> I'm not sure I
understand the problem with <a href="http://address-space-conversion.cu" target="_blank">address-space-conversion.cu</a>. I
would expect the frontend to look at the pointer type when
determining the type of store to emit.<br>
<br>
For example, if there is code like the following:<br>
<br>
__shared__ float a;<br>
float *ptr = &a;<br>
*ptr = 2.0f;<br>
<br>
<br>
I would expect the frontend to create an addrspace(3)
allocation for 'a', store the value of 'addrspacecast
float addrspace(3)* to float*' into 'ptr', and then emit a
'store ..., float*' for the final assignment.<br>
<br>
But if the indirection through 'ptr' was not there, the
frontend would see an addrspace(3) pointer and emit the
corresponding store instead of first casting to
addrspace(0).<br>
<br>
Maybe I'm just misunderstanding what you're trying to do
here. Why is ExpectedAddrSpace != AddrSpace for the
examples in <a href="http://address-spaces.cu" target="_blank">address-spaces.cu</a>?
<div>
<div><br>
<br>
<br>
<div>On 03/24/2014 02:22 PM, Jingyue Wu wrote:<br>
</div>
<blockquote type="cite">
<div dir="ltr">Justin,
<div><br>
</div>
<div>I overlooked that you were referring
specifically to the last test case in <a href="http://address-spaces.cu" target="_blank">address-spaces.cu</a>.
<div><br>
</div>
<div>In that particular example, although the
previous code looks to emit more optimized
code, it only worked by chance because the
program only loads from and stores to "shared
int lk". If the test case had been more
complicated, e.g., "int *lkp = &lk" after
"shared int lk", the codegen would have
emitted a StoreInst with mismatched types, and
crashed just as in many cases in <a href="http://address-space-conversion.cu" target="_blank">address-space-conversion.cu</a>. </div>
<div><br>
</div>
<div>Jingyue</div>
</div>
</div>
<div class="gmail_extra"><br>
<br>
<div class="gmail_quote">On Mon, Mar 24, 2014 at
10:05 AM, Jingyue Wu <span dir="ltr"><<a href="mailto:jingyue@google.com" target="_blank">jingyue@google.com</a>></span>
wrote:<br>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<div dir="ltr">Right. We are aware of this
issue, and think it should be addressed in
the IR optimizer (similar to
InstCombineLoadCast and
InstCombineStoreToCast) instead of clang. Do
you think this is an appropriate approach?
Is this optimization general enough to stay
in the IR optimizer or target-dependent? <span><font color="#888888">
<div> <br>
</div>
<div>Jingyue</div>
</font></span></div>
<div>
<div>
<div class="gmail_extra"><br>
<br>
<div class="gmail_quote">On Mon, Mar 24,
2014 at 4:54 AM, Justin Holewinski <span dir="ltr"><<a href="mailto:justin.holewinski@gmail.com" target="_blank">justin.holewinski@gmail.com</a>></span>
wrote:<br>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<div dir="ltr">Hi Jingyue,
<div><br>
</div>
<div>I committed the addrspacecast
isel patterns to NVPTX. Also, I
wanted to point out that your
changes in the last test case in
this patch (<a href="http://address-spaces.cu" target="_blank">address-spaces.cu</a>)
represent changes that may lead
to performance degradation.
Specific address spaces should
be used whenever possible for
loads/stores. Casting
everything to a generic address
is still correct, but may lead
to additional indirections for
the hardware.</div>
</div>
<div class="gmail_extra">
<div>
<div><br>
<br>
<div class="gmail_quote">On
Fri, Mar 21, 2014 at 2:25
PM, Justin Holewinski <span dir="ltr"><<a href="mailto:jholewinski@nvidia.com" target="_blank">jholewinski@nvidia.com</a>></span>
wrote:<br>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<div bgcolor="#FFFFFF" text="#000000">
<div>addrspacecast
support in NVPTX is on
my todo list. I'll
try to put something
together in the next
few days.
<div>
<div><br>
<br>
On 3/21/14, 2:20
PM, Jingyue Wu
wrote:<br>
</div>
</div>
</div>
<div>
<div>
<blockquote type="cite">
<div dir="ltr">Hi,
<div><br>
</div>
<div>Static
local
variables in
CUDA can be
declared with
address space
qualifiers,
such as
__shared__.
Therefore, the
codegen needs
to potentially
addrspacecast
a static local
variable to
the type
expected by
its
declaration.
Peter did
something
similar for
global
variables in
r157167. </div>
<div><br>
</div>
<div>All clang
tests passed. </div>
<div><br>
</div>
<div>Justin: The
NVPTX backend
support for
addrspacecast
seems not
complete. We
can send you
follow-up
patches once
this one gets
in. </div>
<div><br>
</div>
<div> Jingyue</div>
</div>
</blockquote>
<br>
<br>
</div>
</div>
<span><font color="#888888">
<pre cols="72">--
Thanks,
Justin Holewinski</pre>
<div>
<hr> </div>
<div>This email
message is for the
sole use of the
intended
recipient(s) and
may contain
confidential
information. Any
unauthorized
review, use,
disclosure or
distribution is
prohibited. If
you are not the
intended
recipient, please
contact the sender
by reply email and
destroy all copies
of the original
message. </div>
<div>
<hr> </div>
</font></span></div>
</blockquote>
</div>
<br>
<br clear="all">
<div><br>
</div>
</div>
</div>
<span><font color="#888888">-- <br>
<br>
<div>Thanks,</div>
<div><br>
</div>
<div>Justin Holewinski</div>
</font></span></div>
</blockquote>
</div>
<br>
</div>
</div>
</div>
</blockquote>
</div>
<br>
</div>
</blockquote>
<br>
</div>
</div>
</div>
</blockquote>
</div>
<br>
</div>
</blockquote>
<br>
</div></div></div>
</blockquote></div><br></div></div>