[clang] [compiler-rt] [AMDGPU] Add UBSan checks for AMDGPU lane access builtins (PR #231485)

[Joseph Huber via cfe-commits]

================
@@ -683,12 +699,39 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
   case AMDGPU::BI__builtin_amdgcn_permlane64:
     return emitBuiltinWithOneOverloadedType<1>(*this, E,
                                                Intrinsic::amdgcn_permlane64);
-  case AMDGPU::BI__builtin_amdgcn_readlane:
-    return emitBuiltinWithOneOverloadedType<2>(*this, E,
-                                               Intrinsic::amdgcn_readlane);
-  case AMDGPU::BI__builtin_amdgcn_wave_shuffle:
-    return emitBuiltinWithOneOverloadedType<2>(*this, E,
-                                               Intrinsic::amdgcn_wave_shuffle);
+  case AMDGPU::BI__builtin_amdgcn_readlane: {
+    Value *Src = EmitScalarExpr(E->getArg(0));
+    Value *Lane = EmitScalarExpr(E->getArg(1));
+    // The lane index must be uniform across the wavefront.
+    emitBuiltinCheck(*this, E, BCK_NonUniformLane, [&] {
+      Value *First = Builder.CreateIntrinsic(Intrinsic::amdgcn_readfirstlane,
+                                             {Lane->getType()}, {Lane});
+      Value *Diff = Builder.CreateICmpNE(Lane, First);
----------------
jhuber6 wrote:

I don't think so? I don't think these are enabled in clang for AMDGPU yet.

https://github.com/llvm/llvm-project/pull/231485