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

[Justin Lebar via cfe-commits]

Thank you, Tom.  I will have a look.

On Wed, Jun 1, 2016 at 11:22 AM, Tom Stellard <tom at 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 at lists.llvm.org
>> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits