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.
For example, if there is code like the following:
__shared__ float a;
float *ptr = &a;
*ptr = 2.0f;
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.
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).
Maybe I'm just misunderstanding what you're trying to do here. Why is
ExpectedAddrSpace != AddrSpace for the examples in address-spaces.cu?
On 03/24/2014 02:22 PM, Jingyue Wu wrote:
Justin,
I overlooked that you were referring specifically to the last test
case in address-spaces.cu <http://address-spaces.cu>.
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 address-space-conversion.cu
<http://address-space-conversion.cu>.
Jingyue
On Mon, Mar 24, 2014 at 10:05 AM, Jingyue Wu <[email protected]
<mailto:[email protected]>> wrote:
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?
Jingyue
On Mon, Mar 24, 2014 at 4:54 AM, Justin Holewinski
<[email protected] <mailto:[email protected]>>
wrote:
Hi Jingyue,
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 (address-spaces.cu <http://address-spaces.cu>)
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.
On Fri, Mar 21, 2014 at 2:25 PM, Justin Holewinski
<[email protected] <mailto:[email protected]>> wrote:
addrspacecast support in NVPTX is on my todo list. I'll
try to put something together in the next few days.
On 3/21/14, 2:20 PM, Jingyue Wu wrote:
Hi,
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.
All clang tests passed.
Justin: The NVPTX backend support for addrspacecast seems
not complete. We can send you follow-up patches once this
one gets in.
Jingyue
--
Thanks,
Justin Holewinski
------------------------------------------------------------------------
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.
------------------------------------------------------------------------
--
Thanks,
Justin Holewinski
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
