Author: jlebar Date: Tue May 31 16:27:13 2016 New Revision: 271336 URL: http://llvm.org/viewvc/llvm-project?rev=271336&view=rev Log: [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 Modified: cfe/trunk/lib/CodeGen/CGStmt.cpp cfe/trunk/test/CodeGenCUDA/convergent.cu Modified: cfe/trunk/lib/CodeGen/CGStmt.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGStmt.cpp?rev=271336&r1=271335&r2=271336&view=diff ============================================================================== --- cfe/trunk/lib/CodeGen/CGStmt.cpp (original) +++ cfe/trunk/lib/CodeGen/CGStmt.cpp Tue May 31 16:27:13 2016 @@ -2054,6 +2054,14 @@ void CodeGenFunction::EmitAsmStmt(const 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<llvm::Value*> RegResults; if (ResultRegTypes.size() == 1) { Modified: cfe/trunk/test/CodeGenCUDA/convergent.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/convergent.cu?rev=271336&r1=271335&r2=271336&view=diff ============================================================================== --- cfe/trunk/test/CodeGenCUDA/convergent.cu (original) +++ cfe/trunk/test/CodeGenCUDA/convergent.cu Tue May 31 16:27:13 2016 @@ -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]] = { _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits