[PATCH] R600 - Fix zero extend of i1

Project Infrastructure LLVM Dev List Archives
#1

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@gmail.com

0001-R600-Fix-zero-extend-of-i1.patch (770 Bytes)

#2

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@gmail.com

From a9800b30a7498241640cb31e8ed7508ffd22b569 Mon Sep 17 00:00:00 2001
From: Jon Pry <jonpry@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

#3

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

0001-R600-Fix-zero-extend-of-i1-V2.txt (2.1 KB)

#4

It’s a good idea to include the : after the function name in the -LABEL checks. It helps avoid problems when comments refer to the function name, or if you have several functions that start with the same thing (i.e. test1 and test10)

#5

It’s a good idea to include the : after the function name in the -LABEL checks. It helps avoid problems when comments refer to the function name, or if you have several functions that start with the same thing (i.e. test1 and test10)

Fair enough.

0001-R600-Fix-zero-extend-of-i1-V3.txt (2.11 KB)

#6

> It’s a good idea to include the : after the function name in the -LABEL checks. It helps avoid problems when comments refer to the function name, or if you have several functions that start with the same thing (i.e. test1 and test10)

Fair enough.

From 4729d848f7db69bf361c5cab4689af1fd96f98a8 Mon Sep 17 00:00:00 2001
From: Jon Pry <jonpry@gmail.com>
Date: Wed, 1 Jan 2014 21:37:17 -0500
Subject: [PATCH] R600 - Fix zero extend of i1

---
lib/Target/R600/SIInstructions.td | 5 +++++
test/CodeGen/R600/zero_extend.ll | 21 ++++++++++++++++++---
2 files changed, 23 insertions(+), 3 deletions(-)

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),
diff --git a/test/CodeGen/R600/zero_extend.ll b/test/CodeGen/R600/zero_extend.ll
index 481b3b3..1d1f47c 100644
--- a/test/CodeGen/R600/zero_extend.ll
+++ b/test/CodeGen/R600/zero_extend.ll
@@ -1,14 +1,14 @@
; RUN: llc < %s -march=r600 -mcpu=redwood | FileCheck %s --check-prefix=R600-CHECK
; RUN: llc < %s -march=r600 -mcpu=SI -verify-machineinstrs | FileCheck %s --check-prefix=SI-CHECK

-; R600-CHECK: @test
+; R600-CHECK: @test1:

The function name checks should be R600-CHECK-LABEL, otherwise LGTM.

-Tom

P.S.

In the future, patches should be sent to llvm-commits.