From 649d574e524c8b40fbdfc17c76242e7c4331c45e Mon Sep 17 00:00:00 2001 From: Justin Lebar Date: Tue, 31 May 2016 21:27:13 +0000 Subject: [PATCH] [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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@271336 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/CodeGen/CGStmt.cpp | 8 ++++++++ test/CodeGenCUDA/convergent.cu | 6 ++++++ 2 files changed, 14 insertions(+) diff --git a/lib/CodeGen/CGStmt.cpp b/lib/CodeGen/CGStmt.cpp index 6945ec9ef6..ff70bbc866 100644 --- a/lib/CodeGen/CGStmt.cpp +++ b/lib/CodeGen/CGStmt.cpp @@ -2054,6 +2054,14 @@ void CodeGenFunction::EmitAsmStmt(const AsmStmt &S) { llvm::ConstantAsMetadata::get(Loc))); } + if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) { + // Conservatively, mark all inline asm blocks in CUDA as convergent + // (meaning, they may call an intrinsically convergent op, such as bar.sync, + // and so can't have certain optimizations applied around them). + Result->addAttribute(llvm::AttributeSet::FunctionIndex, + llvm::Attribute::Convergent); + } + // Extract all of the register value results from the asm. std::vector RegResults; if (ResultRegTypes.size() == 1) { diff --git a/test/CodeGenCUDA/convergent.cu b/test/CodeGenCUDA/convergent.cu index d2e75f7e2f..6827c57d29 100644 --- a/test/CodeGenCUDA/convergent.cu +++ b/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]] = { -- 2.50.1