Thank you, Tom. I will have a look.
On Wed, Jun 1, 2016 at 11:22 AM, Tom Stellard <t...@stellard.net> wrote: > On Tue, May 31, 2016 at 09:27:13PM -0000, Justin Lebar via cfe-commits wrote: >> 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. >> > > Hi, > > The MachineInstr INLINEASM also needs to be marked as convergent, > otherwise you will run into the same problem with the > MachineInstr passes. > > -Tom > >> 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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
