r271336 - [CUDA] Conservatively mark inline asm as convergent.

[Tom Stellard via cfe-commits]

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 at lists.llvm.org
> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits