Mismatch of return value data type between function declaration and call-site

[ghost]

I'm on the latest HLC compiler from the branch hsail-stable-3.7.
Consider the following OpenCL program:

__attribute__((__noinline__)) bool IsEmpty()
{
    return true;
}

__kernel void Main()
{
    IsEmpty();
}

Compiling the program above with -O0 gives the following error:

The error is:

>         call    &IsEmpty (%IsEmpty) ();
>                          ^
input(27,17): Incompatible types of formal and actual arguments

ERROR:  The following command failed with return code 1.
        HSAILasm -o /tmp/hsa_finalizer-KenL4R/temp.hsail /tmp/cloc20050/temp.hsail

The corresponding intermediate assembly looks like:

module &__llvm_hsail_module:1:0:$full:$large:$near;

decl function &IsEmpty(arg_u32 %ret)();

function &IsEmpty(arg_u32 %IsEmpty)()
{

// BB#0:
    mov_b32    $s0, 1;
    st_arg_u32    $s0, [%IsEmpty];
    ret;
};

prog kernel &__OpenCL_Main_kernel(
    kernarg_u64 %__global_offset_0,
    kernarg_u64 %__global_offset_1,
    kernarg_u64 %__global_offset_2,
    kernarg_u64 %__printf_buffer,
    kernarg_u64 %__vqueue_pointer,
    kernarg_u64 %__aqlwrap_pointer)
{

    align(4) spill_u8 %__spillStack[4];
// BB#0:
    {
        arg_u8 %IsEmpty;
        call    &IsEmpty (%IsEmpty) ();
        ld_arg_u8    $s0, [%IsEmpty];
    }
    st_spill_align(4)_u32    $s0, [%__spillStack]; // 4-byte Folded Spill
    ret;
};

Apparently, the HSAIL code generator outputs an arg_u32 return value type for the IsEmpty() function, but the corresponding argument on the call-site is of type arg_u8.

[arsenm]

Can you post the LLVM IR you get for this?

[ghost]

Compiling the kernel with cloc.sh -opt 0 -ll test.cl gives the following LLVM IR (from test.ll):

; ModuleID = '/tmp/cloc9889/test.bc'
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024"
target triple = "spir64-unknown-unknown"

; Function Attrs: noinline nounwind
define spir_func zeroext i1 @IsEmpty() #0 {
  ret i1 true
}

; Function Attrs: nounwind
define spir_kernel void @__OpenCL_Main_kernel() #1 {
  %1 = call spir_func zeroext i1 @IsEmpty()
  ret void
}

attributes #0 = { noinline nounwind }
attributes #1 = { nounwind }

!opencl.kernels = !{!0}
!opencl.enable.FP_CONTRACT = !{}
!opencl.ocl.version = !{!6}
!opencl.spir.version = !{!6}

!0 = !{void ()* @__OpenCL_Main_kernel, !1, !2, !3, !4, !5}
!1 = !{!"kernel_arg_addr_space"}
!2 = !{!"kernel_arg_access_qual"}
!3 = !{!"kernel_arg_type"}
!4 = !{!"kernel_arg_type_qual"}
!5 = !{!"kernel_arg_base_type"}
!6 = !{i32 2, i32 0}

[geidav]

Any news one this issue?