Skip to content

Commit 29993d2

Browse files
ranapratap55ranapratap55
committed
Minor modification to format changes
1 parent 901f395 commit 29993d2

File tree

1 file changed

+27
-25
lines changed

1 file changed

+27
-25
lines changed

‎clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp‎

Lines changed: 27 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -133,8 +133,8 @@ Value *EmitAMDGPUWorkGroupSize(CodeGenFunction &CGF, unsigned Index) {
133133
}
134134

135135
llvm::MDBuilder MDHelper(CGF.getLLVMContext());
136-
llvm::MDNode *RNode = MDHelper.createRange(
137-
APInt(16, 1), APInt(16, CGF.getTarget().getMaxOpenCLWorkGroupSize() + 1));
136+
llvm::MDNode *RNode = MDHelper.createRange(APInt(16, 1),
137+
APInt(16, CGF.getTarget().getMaxOpenCLWorkGroupSize() + 1));
138138
LD->setMetadata(llvm::LLVMContext::MD_range, RNode);
139139
LD->setMetadata(llvm::LLVMContext::MD_noundef,
140140
llvm::MDNode::get(CGF.getLLVMContext(), {}));
@@ -253,7 +253,8 @@ emitAMDGCNImageOverloadedReturnType(clang::CodeGen::CodeGenFunction &CGF,
253253
}
254254

255255
// Emit an intrinsic that has 1 float or double operand, and 1 integer.
256-
static Value *emitFPIntBuiltin(CodeGenFunction &CGF, const CallExpr *E,
256+
static Value *emitFPIntBuiltin(CodeGenFunction &CGF,
257+
const CallExpr *E,
257258
unsigned IntrinsicID) {
258259
llvm::Value *Src0 = CGF.EmitScalarExpr(E->getArg(0));
259260
llvm::Value *Src1 = CGF.EmitScalarExpr(E->getArg(1));
@@ -455,8 +456,8 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
455456
llvm::Value *Y = EmitScalarExpr(E->getArg(1));
456457
llvm::Value *Z = EmitScalarExpr(E->getArg(2));
457458

458-
llvm::Function *Callee =
459-
CGM.getIntrinsic(Intrinsic::amdgcn_div_scale, X->getType());
459+
llvm::Function *Callee = CGM.getIntrinsic(Intrinsic::amdgcn_div_scale,
460+
X->getType());
460461

461462
llvm::Value *Tmp = Builder.CreateCall(Callee, {X, Y, Z});
462463

@@ -476,8 +477,8 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
476477
llvm::Value *Src2 = EmitScalarExpr(E->getArg(2));
477478
llvm::Value *Src3 = EmitScalarExpr(E->getArg(3));
478479

479-
llvm::Function *F =
480-
CGM.getIntrinsic(Intrinsic::amdgcn_div_fmas, Src0->getType());
480+
llvm::Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_div_fmas,
481+
Src0->getType());
481482
llvm::Value *Src3ToBool = Builder.CreateIsNotNull(Src3);
482483
return Builder.CreateCall(F, {Src0, Src1, Src2, Src3ToBool});
483484
}
@@ -618,13 +619,13 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
618619
case AMDGPU::BI__builtin_amdgcn_frexp_expf: {
619620
Value *Src0 = EmitScalarExpr(E->getArg(0));
620621
Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_frexp_exp,
621-
{Builder.getInt32Ty(), Src0->getType()});
622+
{ Builder.getInt32Ty(), Src0->getType() });
622623
return Builder.CreateCall(F, Src0);
623624
}
624625
case AMDGPU::BI__builtin_amdgcn_frexp_exph: {
625626
Value *Src0 = EmitScalarExpr(E->getArg(0));
626627
Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_frexp_exp,
627-
{Builder.getInt16Ty(), Src0->getType()});
628+
{ Builder.getInt16Ty(), Src0->getType() });
628629
return Builder.CreateCall(F, Src0);
629630
}
630631
case AMDGPU::BI__builtin_amdgcn_fract:
@@ -670,8 +671,8 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
670671

671672
// FIXME-GFX10: How should 32 bit mask be handled?
672673
Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_icmp,
673-
{Builder.getInt64Ty(), Src0->getType()});
674-
return Builder.CreateCall(F, {Src0, Src1, Src2});
674+
{ Builder.getInt64Ty(), Src0->getType() });
675+
return Builder.CreateCall(F, { Src0, Src1, Src2 });
675676
}
676677
case AMDGPU::BI__builtin_amdgcn_fcmp:
677678
case AMDGPU::BI__builtin_amdgcn_fcmpf: {
@@ -681,8 +682,8 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
681682

682683
// FIXME-GFX10: How should 32 bit mask be handled?
683684
Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_fcmp,
684-
{Builder.getInt64Ty(), Src0->getType()});
685-
return Builder.CreateCall(F, {Src0, Src1, Src2});
685+
{ Builder.getInt64Ty(), Src0->getType() });
686+
return Builder.CreateCall(F, { Src0, Src1, Src2 });
686687
}
687688
case AMDGPU::BI__builtin_amdgcn_class:
688689
case AMDGPU::BI__builtin_amdgcn_classf:
@@ -694,12 +695,11 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
694695
Intrinsic::amdgcn_fmed3);
695696
case AMDGPU::BI__builtin_amdgcn_ds_append:
696697
case AMDGPU::BI__builtin_amdgcn_ds_consume: {
697-
Intrinsic::ID Intrin = BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_append
698-
? Intrinsic::amdgcn_ds_append
699-
: Intrinsic::amdgcn_ds_consume;
698+
Intrinsic::ID Intrin = BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_append ?
699+
Intrinsic::amdgcn_ds_append : Intrinsic::amdgcn_ds_consume;
700700
Value *Src0 = EmitScalarExpr(E->getArg(0));
701-
Function *F = CGM.getIntrinsic(Intrin, {Src0->getType()});
702-
return Builder.CreateCall(F, {Src0, Builder.getFalse()});
701+
Function *F = CGM.getIntrinsic(Intrin, { Src0->getType() });
702+
return Builder.CreateCall(F, { Src0, Builder.getFalse() });
703703
}
704704
case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_i32:
705705
case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_v2i32:
@@ -919,10 +919,12 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
919919

920920
// The builtins take these arguments as vec4 where the last element is
921921
// ignored. The intrinsic takes them as vec3.
922-
RayOrigin = Builder.CreateShuffleVector(RayOrigin, RayOrigin, {0, 1, 2});
923-
RayDir = Builder.CreateShuffleVector(RayDir, RayDir, {0, 1, 2});
924-
RayInverseDir =
925-
Builder.CreateShuffleVector(RayInverseDir, RayInverseDir, {0, 1, 2});
922+
RayOrigin = Builder.CreateShuffleVector(RayOrigin, RayOrigin,
923+
{0, 1, 2});
924+
RayDir =
925+
Builder.CreateShuffleVector(RayDir, RayDir, {0, 1, 2});
926+
RayInverseDir = Builder.CreateShuffleVector(RayInverseDir, RayInverseDir,
927+
{0, 1, 2});
926928

927929
Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_image_bvh_intersect_ray,
928930
{NodePtr->getType(), RayDir->getType()});
@@ -996,8 +998,8 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
996998
Value *Rtn = Builder.CreateExtractValue(Call, 0);
997999
Value *A = Builder.CreateExtractValue(Call, 1);
9981000
llvm::Type *RetTy = ConvertType(E->getType());
999-
Value *I0 =
1000-
Builder.CreateInsertElement(PoisonValue::get(RetTy), Rtn, (uint64_t)0);
1001+
Value *I0 = Builder.CreateInsertElement(PoisonValue::get(RetTy), Rtn,
1002+
(uint64_t)0);
10011003
// ds_bvh_stack_push8_pop2_rtn returns {i64, i32} but the builtin returns
10021004
// <2 x i64>, zext the second value.
10031005
if (A->getType()->getPrimitiveSizeInBits() <
@@ -1702,7 +1704,7 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
17021704
llvm::Value *Src1 = EmitScalarExpr(E->getArg(1));
17031705
llvm::Value *Src2 = EmitScalarExpr(E->getArg(2));
17041706
Function *F = CGM.getIntrinsic(Intrinsic::fshr, Src0->getType());
1705-
return Builder.CreateCall(F, {Src0, Src1, Src2});
1707+
return Builder.CreateCall(F, { Src0, Src1, Src2 });
17061708
}
17071709
case AMDGPU::BI__builtin_amdgcn_fence: {
17081710
ProcessOrderScopeAMDGCN(EmitScalarExpr(E->getArg(0)),

0 commit comments

Comments
 (0)