llvm dev - Dec 2013 - [LLVMdev] [PATCH] R600 - Fix zero extend of i1

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
-------------- next part --------------
An HTML attachment was scrubbed...
URL:
<http://lists.llvm.org/pipermail/llvm-dev/attachments/20131230/057c13c5/attachment.html>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: 0001-R600-Fix-zero-extend-of-i1.patch
Type: text/x-patch
Size: 770 bytes
Desc: not available
URL:
<http://lists.llvm.org/pipermail/llvm-dev/attachments/20131230/057c13c5/attachment.bin>

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

> 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
Version 2 of patch attached which includes test case.


-Jon
-------------- next part --------------