[DCU]fix mip and compile issues

Signed-off-by: Tangao <2205747538@qq.com>

tests/pytorch/distributed/test_fusible_ops_with_userbuffers.py

@@ -432,7 +432,7 @@ def test_fuser_ops_with_userbuffers(
     command = []
     if tex.ubuf_built_with_mpi():
         python_exe = pathlib.Path(sys.executable).resolve()
-        command.extend(("mpirun", "-np", str(world_size), "--oversubscribe", "--quiet", python_exe))
+        command.extend(("mpirun", "-np", str(world_size), "--oversubscribe", "--allow-run-as-root", python_exe))
     else:
         command.extend(("torchrun", f"--nproc_per_node={world_size}"))
 

transformer_engine/common/dropout/dropout.cu

@@ -44,7 +44,17 @@ __device__ __forceinline__ uint32_t bytewise_less_than(uint32_t a, uint32_t b) {
   // Bitwise logical op to get answer in MSBs
   // Equivalent logic: result = (a == b) ? !result : b
 #ifdef __HIP_PLATFORM_AMD__
-  result = (a == b) ? !result : b;
+  // Use HIP vector types for byte-wise parallel comparison
+  union { uint32_t u32; uint8_t bytes[4]; } a_union, b_union;
+  a_union.u32 = a;
+  b_union.u32 = b;
+
+  uint32_t mask = 0;
+  mask |= (a_union.bytes[0] < b_union.bytes[0]) ? 0x80000000U : 0;
+  mask |= (a_union.bytes[1] < b_union.bytes[1]) ? 0x00800000U : 0;
+  mask |= (a_union.bytes[2] < b_union.bytes[2]) ? 0x00008000U : 0;
+  mask |= (a_union.bytes[3] < b_union.bytes[3]) ? 0x00000080U : 0;
+  result = mask;
 #else
   asm("lop3.b32 %0, %1, %2, %3, 0x4D;\n\t" : "=r"(result) : "r"(a), "r"(b), "r"(result));
 #endif