================
@@ -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
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
