<html>
<head>
<meta content="text/html; charset=UTF-8" http-equiv="Content-Type">
</head>
<body bgcolor="#FFFFFF" text="#000000">
I'm not sure I understand the problem with
address-space-conversion.cu. 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
address-spaces.cu?<br>
<br>
<br>
<div class="moz-cite-prefix">On 03/24/2014 02:22 PM, Jingyue Wu
wrote:<br>
</div>
<blockquote
cite="mid:CAMROOrFaLRniDmBFJc3YZ-AZ_JQgzAJNGC87v4p_VVR0SrFtzQ@mail.gmail.com"
type="cite">
<div dir="ltr">Justin,
<div><br>
</div>
<div>I overlooked that you were referring specifically to the
last test case in <a moz-do-not-send="true"
href="http://address-spaces.cu">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 moz-do-not-send="true"
href="http://address-space-conversion.cu">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 moz-do-not-send="true"
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 class="HOEnZb"><font
color="#888888">
<div>
<br>
</div>
<div>Jingyue</div>
</font></span></div>
<div class="HOEnZb">
<div class="h5">
<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
moz-do-not-send="true"
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 moz-do-not-send="true"
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 moz-do-not-send="true"
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>
</body>
</html>