]> granicus.if.org Git - clang/commit
[CUDA] Mark all CUDA device-side function defs, decls, and calls as convergent.
authorJustin Lebar <jlebar@google.com>
Wed, 24 Feb 2016 21:55:11 +0000 (21:55 +0000)
committerJustin Lebar <jlebar@google.com>
Wed, 24 Feb 2016 21:55:11 +0000 (21:55 +0000)
commit5931844b43db04042de37f652a9a3077b53a646b
tree22ce02d87132c02aac57876165d1b2701a5e4fd8
parenta4841b214ebe36b744a229a82d701a1f0cfe6dcd
[CUDA] Mark all CUDA device-side function defs, decls, and calls as convergent.

Summary:
This is important for e.g. the following case:

  void sync() { __syncthreads(); }
  void foo() {
    do_something();
    sync();
    do_something_else():
  }

Without this change, if the optimizer does not inline sync() (which it
won't because __syncthreads is also marked as noduplicate, for now
anyway), it is free to perform optimizations on sync() that it would not
be able to perform on __syncthreads(), because sync() is not marked as
convergent.

Similarly, we need a notion of convergent calls, since in the case when
we can't statically determine a call's target(s), we need to know
whether it's safe to perform optimizations around the call.

This change is conservative; the optimizer will remove these attrs where
it can, see r260318, r260319.

Reviewers: majnemer

Subscribers: cfe-commits, jhen, echristo, tra

Differential Revision: http://reviews.llvm.org/D17056

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@261779 91177308-0d34-0410-b5e6-96231b3b80d8
lib/CodeGen/CGCall.cpp
test/CodeGenCUDA/convergent.cu [new file with mode: 0644]
test/CodeGenCUDA/device-var-init.cu