HIP kernels with bool parameters produce non-conformant SPIR-V

[linehill]

HIP kernels with bool parameters fail to run on rusticl. This is due to the produced SPIR-V being non-conformant respect to the OpenCL SPIR-V environment specification. The SPIR-V kernels have parameters of OpTypeBool type but they are not allowed in kernels.

$ cat bool-kernel-arg.hip
#include <hip/hip_runtime.h>

__global__ void k(bool cond) { printf("cond=%d\n", cond); }

int main() {
  k<<<1, 1>>>(true);
  if (hipGetLastError() != hipSuccess)
    return 1;
  if (hipDeviceSynchronize() != hipSuccess)
    return 2;
  return 0;
}
$ ../install/bin/hipcc bool-kernel-arg.hip -O2 -o bool-kernel-arg
$ CHIP_LOGLEVEL=info ./bool-kernel-arg
CHIP info [TID 31602] [1715603385.982596711] : CHIP_PLATFORM=0
CHIP info [TID 31602] [1715603385.982717400] : CHIP_DEVICE_TYPE=default
CHIP info [TID 31602] [1715603385.982754591] : CHIP_DEVICE=0
CHIP info [TID 31602] [1715603385.982776661] : CHIP_BE=opencl
CHIP info [TID 31602] [1715603385.982803903] : CHIP_DUMP_SPIRV=off
CHIP info [TID 31602] [1715603385.982844175] : CHIP_JIT_FLAGS_OVERRIDE=-cl-kernel-arg-info -cl-std=CL3.0
CHIP info [TID 31602] [1715603385.982882198] : CHIP_L0_COLLECT_EVENTS_TIMEOUT=0
CHIP info [TID 31602] [1715603385.982905412] : CHIP_L0_EVENT_TIMEOUT=0
CHIP info [TID 31602] [1715603385.982921474] : CHIP_SKIP_UNINIT=off
CHIP info [TID 31602] [1715603387.961502576] : OpenCL Devices of type default with SPIR-V_1 support:
AMD Radeon Pro VII (radeonsi, vega20, LLVM 17.0.6, DRM 3.57, 6.5.0-28-generic)  is supported.

CHIP warning [TID 31602] [1715603387.961755173] : The device might not support subgroup size 32, warp-size sensitive kernels might not work correctly.
CHIP error [TID 31602] [1715603387.965727841] : hipErrorTbd (CL_INVALID_ARG_SIZE clSetKernelArg failed) in /mnt/md0/linehill/ws-chip-spv-3/chipstar/src/backend/OpenCL/CHIPBackendOpenCL.cc:1752:operator()

CHIP error [TID 31602] [1715603387.965848329] : Caught Error: hipErrorTbd