[Kernel][Hardware][AMD][ROCm] Fix rocm/attention.cu compilation on RO…

…Cm 6.0.3

In ROCm 6.0.3, __hip_bfloat16 doesn't support the + operator, while
hip_bfloat16 is OK.

Refer to ROCm/ROCm#2534

ERROR LOG:

vllm/build/temp.linux-x86_64-cpython-311/csrc/rocm/attention.hip:168:20:
error: invalid operands to binary expression
('__hip_bfloat16' and '__hip_bfloat16')
      res.b = t1.b + t2.b;
              ~~~~ ^ ~~~~
vllm/build/temp.linux-x86_64-cpython-311/csrc/rocm/attention.hip:683:15:
note: in instantiation of function template specialization
'addx4<__hip_bfloat16>' requested here
              addx4<scalar_t>(vout[qh][vh], vout_shared[qh][vh][laneid][w]);
              ^

Signed-off-by: Hollow Man <[email protected]>

csrc/rocm/attention.cu

@@ -163,7 +163,14 @@ __device__ __forceinline__ _B16x4 addx4(const _B16x4& inp1,
     for (int i = 0; i < 4; i++) {
       t1.u = inp1[i];
       t2.u = inp2[i];
-      res.b = t1.b + t2.b;
+      // See https://github.com/ROCm/ROCm/issues/2534
+      hip_bfloat16 t1b, t2b;
+      __hip_bfloat16 resb;
+      t1b.data = t1.b.data;
+      t2b.data = t2.b.data;
+      resb.data = (t1b + t2b).data;
+      res.b = resb;
+      // res.b = t1.b + t2.b;
       ret[i] = res.u;
     }
     return ret;