This is an archive of the discontinued LLVM Phabricator instance.

[NVPTX] convert pointers in byval kernel arguments to global
ClosedPublic

Authored by jingyue on Jul 25 2015, 4:50 PM.

Details

Summary

For example, in

struct S {
  int *x;
  int *y;
};
__global__ void foo(S s) {
  int *b = s.y;
  // use b
}

"b" is guaranteed to point to global. NVPTX should emit ld.global/st.global for
accessing "b".

Diff Detail

Event Timeline

jingyue updated this revision to Diff 30639.Jul 25 2015, 4:50 PM
jingyue retitled this revision from to [NVPTX] convert pointers in byval kernel arguments to global.
jingyue updated this object.
jingyue added a reviewer: jholewinski.
jingyue added a subscriber: llvm-commits.
jingyue updated this object.Jul 25 2015, 9:44 PM
jholewinski accepted this revision.Jul 31 2015, 1:03 PM
jholewinski edited edge metadata.

Looks reasonable to me, with the minor comment.

Thanks!

lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp
210

Can you enable this for CUDA only?

This revision is now accepted and ready to land.Jul 31 2015, 1:03 PM
jingyue updated this revision to Diff 31170.Jul 31 2015, 2:37 PM
jingyue edited edge metadata.

As Justin pointed out, enable this optimization for CUDA only

jingyue closed this revision.Jul 31 2015, 2:44 PM