[SYCL] Change lowering of 'cl::sycl::select' into SPIR-V

[AlexeySachkov]

Previously, OpSelect from SPIR-V core spec was used, but comparing to
the SYCL definition of 'select' built-in, it behaves
differently and even takes different arguments.

select instruction from OpenCL Extended Instruction Set must be used
instead.

Description of OpSelect from SPIR-V core spec:

Condition must be a scalar or vector of Boolean type.
If Condition is a scalar and true, the result is Object 1. If
Condition is a scalar and false, the result is Object 2.

If Condition is a vector, Result Type must be a vector with the same
number of components as Condition and the result is a mix of Object 1
and Object 2: When a component of Condition is true, the corresponding
component in the result is taken from Object 1, otherwise it is taken
from Object 2.

Description of select ExtInst:

For each component of a vector type, the result is a if the most
significant bit of c is zero, otherwise it is b.

For a scalar type, the result is a if c is zero, otherwise it is b.

c must be integer or vector(2,3,4,8,16) of integer values.

The latter perfectly matches both SYCL and OpenCL specs.

Note: previous implementation emulate ExtInst select behavior over
OpSelect by evaluating MSB of each vector compontent in C++ code, so, it
is functionally correct. However, it uses ext-vectors of booleans, which are
unsupported by OpenCL and confuse underlying OpenCL compilers sometimes
(crashes and hangs experienced in complex applications because of this).

[alexbatashev]

LGTM, my comment can be ignored