[LLVMdev] [PATCH] R600 - Fix zero extend of i1

[Tom Stellard]

On Mon, Dec 30, 2013 at 11:50:23PM -0500, Jon Pry wrote:
> Hi,
> 
>     When trying to compile a trivial opencl kernel such as:
> 
> __kernel void if_eq(__global int * out, int arg0, int arg1){
>     out[0] = arg0==arg1?0:1;
> }
> 
>    Clang generates IR like:
> 
>   %1 = icmp eq i32 %arg0, %arg1
>   %. = zext i1 %1 to i32
> 
>    This eventually crashes ISel on R600. Attached patch adds a selector so
> it will compile.
> 
> Regards,
> 
> Jon Pry
> jonpry at gmail.com

> From a9800b30a7498241640cb31e8ed7508ffd22b569 Mon Sep 17 00:00:00 2001
> From: Jon Pry <jonpry at gmail.com>
> Date: Mon, 30 Dec 2013 23:43:28 -0500
> Subject: [PATCH] R600 - Fix zero extend of i1
> 

This patch looks good, but you need to add a test case.  You can add it
to the file test/CodeGen/R600/zero_extend.ll

-Tom

> ---
>  lib/Target/R600/SIInstructions.td |    5 +++++
>  1 file changed, 5 insertions(+)
> 
> diff --git a/lib/Target/R600/SIInstructions.td b/lib/Target/R600/SIInstructions.td
> index 3baa4cd..5ea08a1 100644
> --- a/lib/Target/R600/SIInstructions.td
> +++ b/lib/Target/R600/SIInstructions.td
> @@ -1807,6 +1807,11 @@ def : Pat <
>    (V_CNDMASK_B32_e64 (i32 0), (i32 -1), $src0)
>  >;
>  
> +def : Pat <
> +  (i32 (zext i1:$src0)),
> +  (V_CNDMASK_B32_e64 (i32 0), (i32 1), $src0)
> +>;
> +
>  // 1. Offset as 8bit DWORD immediate
>  def : Pat <
>    (SIload_constant i128:$sbase, IMM8bitDWORD:$offset),
> -- 
> 1.7.10.4
> 

> _______________________________________________
> LLVM Developers mailing list
> LLVMdev at cs.uiuc.edu         http://llvm.cs.uiuc.edu
> http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev