[llvm-bugs] [Bug 46199] New: [OpenCL] Invalid optimization of subgroup intrinsics

[via llvm-bugs]

https://bugs.llvm.org/show_bug.cgi?id=46199

            Bug ID: 46199
           Summary: [OpenCL] Invalid optimization of subgroup intrinsics
           Product: clang
           Version: trunk
          Hardware: PC
                OS: Windows NT
            Status: NEW
          Severity: enhancement
          Priority: P
         Component: OpenCL
          Assignee: unassignedclangbugs at nondot.org
          Reporter: piotr.fusik at intel.com
                CC: anastasia.stulova at arm.com, llvm-bugs at lists.llvm.org

The following OpenCL code:

kernel void test(global int *data)
{
    uint id = (uint) get_global_id(0);
    if (id < 4)
        data[id] = sub_group_elect();
    else
        data[id] = sub_group_elect();
}

when compiled with the current master
(f2c97656644e783622a6e60fe452b41ffe0f1d18) as follows: clang -cl-std=CL2.0
-include opencl-c.h -S -emit-llvm sub_group_elect_opt.cl

has an invalid optimization of combining the branches:

; Function Attrs: convergent norecurse nounwind uwtable
define dso_local spir_kernel void @test(i32* nocapture %data)
local_unnamed_addr #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5
!kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 {
entry:
  %call = tail call i64 @"?get_global_id@@$$J0YAKI at Z"(i32 0) #3
  %call2 = tail call i32 @"?sub_group_elect@@$$J0YAHXZ"() #4
  %idxprom = and i64 %call, 4294967295
  %arrayidx = getelementptr inbounds i32, i32* %data, i64 %idxprom
  store i32 %call2, i32* %arrayidx, align 4, !tbaa !8
  ret void
}

sub_group_elect is one of the many functions added in
https://reviews.llvm.org/D79781 that perform implicit communication within a
subgroup, i.e. threads implemented as SIMD on GPU. All the added functions are
affected.

https://reviews.llvm.org/D68994 is a proposal of addressing this problem with
the `convergent` attribute.

-- 
You are receiving this mail because:
You are on the CC list for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-bugs/attachments/20200604/923266ad/attachment.html>