diff options
| author | Justin Lebar <jlebar@google.com> | 2016-05-31 21:27:13 +0000 |
|---|---|---|
| committer | Justin Lebar <jlebar@google.com> | 2016-05-31 21:27:13 +0000 |
| commit | f179364341ef49440757a481b377627d17db3aaa (patch) | |
| tree | 91f6ce66a817e125e164c2fffc16b171150cb443 /clang/test/CodeGenCUDA | |
| parent | 0aeb313e792649aa063e91783e3d8675cfdfd5be (diff) | |
| download | bcm5719-llvm-f179364341ef49440757a481b377627d17db3aaa.tar.gz bcm5719-llvm-f179364341ef49440757a481b377627d17db3aaa.zip | |
[CUDA] Conservatively mark inline asm as convergent.
Summary:
This is particularly important because a some convergent CUDA intrinsics
(e.g. __shfl_down) are implemented in terms of inline asm.
Reviewers: tra
Subscribers: cfe-commits
Differential Revision: http://reviews.llvm.org/D20836
llvm-svn: 271336
Diffstat (limited to 'clang/test/CodeGenCUDA')
| -rw-r--r-- | clang/test/CodeGenCUDA/convergent.cu | 6 |
1 files changed, 6 insertions, 0 deletions
diff --git a/clang/test/CodeGenCUDA/convergent.cu b/clang/test/CodeGenCUDA/convergent.cu index d2e75f7e2fa..6827c57d29f 100644 --- a/clang/test/CodeGenCUDA/convergent.cu +++ b/clang/test/CodeGenCUDA/convergent.cu @@ -25,6 +25,11 @@ __host__ __device__ void baz(); __host__ __device__ void bar() { // DEVICE: call void @_Z3bazv() [[CALL_ATTR:#[0-9]+]] baz(); + // DEVICE: call i32 asm "trap;", "=l"() [[ASM_ATTR:#[0-9]+]] + int x; + asm ("trap;" : "=l"(x)); + // DEVICE: call void asm sideeffect "trap;", ""() [[ASM_ATTR:#[0-9]+]] + asm volatile ("trap;"); } // DEVICE: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]] @@ -32,6 +37,7 @@ __host__ __device__ void bar() { // DEVICE-SAME: convergent // DEVICE-SAME: } // DEVICE: attributes [[CALL_ATTR]] = { convergent } +// DEVICE: attributes [[ASM_ATTR]] = { convergent // HOST: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]] // HOST: attributes [[BAZ_ATTR]] = { |
