Skip to content

Commit d37ea02

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

File tree

8 files changed

+65
-61
lines changed

8 files changed

+65
-61
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def‎

Lines changed: 30 additions & 30 deletions
Original file line numberDiff line numberDiff line change
@@ -889,75 +889,75 @@ TARGET_BUILTIN(__builtin_amdgcn_cooperative_atomic_store_8x16B, "vV4i*V4iIicC*",
889889
// Image builtins
890890
//===----------------------------------------------------------------------===//
891891
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f32_i32, "V4fiiQtii", "nc", "image-insts")
892-
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f16_i32, "V4xiiQtii", "nc", "image-insts")
892+
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f16_i32, "V4hiiQtii", "nc", "image-insts")
893893
TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f32_i32, "V4fiiiQtii", "nc", "image-insts")
894-
TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f16_i32, "V4xiiiQtii", "nc", "image-insts")
894+
TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f16_i32, "V4hiiiQtii", "nc", "image-insts")
895895
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_f32_i32, "fiiiQtii", "nc", "image-insts")
896896
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f32_i32, "V4fiiiQtii", "nc", "image-insts")
897-
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f16_i32, "V4xiiiQtii", "nc", "image-insts")
897+
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f16_i32, "V4hiiiQtii", "nc", "image-insts")
898898
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_f32_i32, "fiiiiQtii", "nc", "image-insts")
899899
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
900-
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts")
900+
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
901901
TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
902-
TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts")
902+
TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
903903
TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
904-
TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts")
904+
TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
905905
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f32_i32, "V4fiiiQtii", "nc", "image-insts")
906-
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f16_i32, "V4xiiiQtii", "nc", "image-insts")
906+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f16_i32, "V4hiiiQtii", "nc", "image-insts")
907907
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
908-
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts")
908+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
909909
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_f32_i32, "fiiiiQtii", "nc", "image-insts")
910910
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
911-
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts")
911+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
912912
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_f32_i32, "fiiiiiQtii", "nc", "image-insts")
913913
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f32_i32, "V4fiiiiiQtii", "nc", "image-insts")
914-
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f16_i32, "V4xiiiiiQtii", "nc", "image-insts")
914+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f16_i32, "V4hiiiiiQtii", "nc", "image-insts")
915915
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f32_i32, "V4fiiiiiQtii", "nc", "image-insts")
916-
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f16_i32, "V4xiiiiiQtii", "nc", "image-insts")
916+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f16_i32, "V4hiiiiiQtii", "nc", "image-insts")
917917
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f32_i32, "V4fiiiiiQtii", "nc", "image-insts")
918-
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f16_i32, "V4xiiiiiQtii", "nc", "image-insts")
918+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f16_i32, "V4hiiiiiQtii", "nc", "image-insts")
919919
TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f32_i32, "vV4fiiQtii", "nc", "image-insts")
920-
TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f16_i32, "vV4xiiQtii", "nc", "image-insts")
920+
TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f16_i32, "vV4hiiQtii", "nc", "image-insts")
921921
TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f32_i32, "vV4fiiiQtii", "nc", "image-insts")
922-
TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f16_i32, "vV4xiiiQtii", "nc", "image-insts")
922+
TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f16_i32, "vV4hiiiQtii", "nc", "image-insts")
923923
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_f32_i32, "vfiiiQtii", "nc", "image-insts")
924924
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f32_i32, "vV4fiiiQtii", "nc", "image-insts")
925-
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f16_i32, "vV4xiiiQtii", "nc", "image-insts")
925+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f16_i32, "vV4hiiiQtii", "nc", "image-insts")
926926
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_f32_i32, "vfiiiiQtii", "nc", "image-insts")
927927
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
928-
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts")
928+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
929929
TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
930-
TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts")
930+
TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
931931
TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
932-
TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts")
932+
TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
933933
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f32_i32, "vV4fiiiQtii", "nc", "image-insts")
934-
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f16_i32, "vV4xiiiQtii", "nc", "image-insts")
934+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f16_i32, "vV4hiiiQtii", "nc", "image-insts")
935935
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
936-
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts")
936+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
937937
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_f32_i32, "vfiiiiQtii", "nc", "image-insts")
938938
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
939-
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts")
939+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
940940
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_f32_i32, "vfiiiiiQtii", "nc", "image-insts")
941941
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f32_i32, "vV4fiiiiiQtii", "nc", "image-insts")
942-
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f16_i32, "vV4xiiiiiQtii", "nc", "image-insts")
942+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f16_i32, "vV4hiiiiiQtii", "nc", "image-insts")
943943
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f32_i32, "vV4fiiiiiQtii", "nc", "image-insts")
944-
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f16_i32, "vV4xiiiiiQtii", "nc", "image-insts")
944+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f16_i32, "vV4hiiiiiQtii", "nc", "image-insts")
945945
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f32_i32, "vV4fiiiiiQtii", "nc", "image-insts")
946-
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f16_i32, "vV4xiiiiiQtii", "nc", "image-insts")
946+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f16_i32, "vV4hiiiiiQtii", "nc", "image-insts")
947947
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f32_f32, "V4fifQtV4ibii", "nc", "image-insts")
948-
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f16_f32, "V4xifQtV4ibii", "nc", "image-insts")
948+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f16_f32, "V4hifQtV4ibii", "nc", "image-insts")
949949
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f32_f32, "V4fiffQtV4ibii", "nc", "image-insts")
950-
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f16_f32, "V4xiffQtV4ibii", "nc", "image-insts")
950+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f16_f32, "V4hiffQtV4ibii", "nc", "image-insts")
951951
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_f32_f32, "fiffQtV4ibii", "nc", "image-insts")
952952
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f32_f32, "V4fiffQtV4ibii", "nc", "image-insts")
953-
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f16_f32, "V4xiffQtV4ibii", "nc", "image-insts")
953+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f16_f32, "V4hiffQtV4ibii", "nc", "image-insts")
954954
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_f32_f32, "fifffQtV4ibii", "nc", "image-insts")
955955
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f32_f32, "V4fifffQtV4ibii", "nc", "image-insts")
956-
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f16_f32, "V4xifffQtV4ibii", "nc", "image-insts")
956+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f16_f32, "V4hifffQtV4ibii", "nc", "image-insts")
957957
TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f32_f32, "V4fifffQtV4ibii", "nc", "image-insts")
958-
TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f16_f32, "V4xifffQtV4ibii", "nc", "image-insts")
958+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f16_f32, "V4hifffQtV4ibii", "nc", "image-insts")
959959
TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f32_f32, "V4fifffQtV4ibii", "nc", "image-insts")
960-
TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f16_f32, "V4xifffQtV4ibii", "nc", "image-insts")
960+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f16_f32, "V4hifffQtV4ibii", "nc", "image-insts")
961961
TARGET_BUILTIN(__builtin_amdgcn_image_sample_lz_1d_v4f32_f32, "V4fifQtV4ibii", "nc", "extended-image-insts")
962962
TARGET_BUILTIN(__builtin_amdgcn_image_sample_lz_1d_v4f16_f32, "V4xifQtV4ibii", "nc", "extended-image-insts")
963963
TARGET_BUILTIN(__builtin_amdgcn_image_sample_lz_1darray_v4f32_f32, "V4fiffQtV4ibii", "nc", "extended-image-insts")

‎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)),

‎clang/test/CodeGen/builtins-image-load.c‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,8 @@
33

44
typedef int int4 __attribute__((ext_vector_type(4)));
55
typedef float float4 __attribute__((ext_vector_type(4)));
6-
typedef _Float16 half4 __attribute__((ext_vector_type(4)));
6+
typedef _Float16 half;
7+
typedef half half4 __attribute__((ext_vector_type(4)));
78

89
// CHECK-LABEL: define dso_local float @test_builtin_image_load_2d(
910
// CHECK-SAME: float noundef [[F32:%.*]], i32 noundef [[I32:%.*]], ptr [[TEX:%.*]]) #[[ATTR0:[0-9]+]] {

‎clang/test/CodeGen/builtins-image-store.c‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,8 @@
22
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1010 %s -emit-llvm -o - | FileCheck %s
33

44
typedef float float4 __attribute__((ext_vector_type(4)));
5-
typedef _Float16 half4 __attribute__((ext_vector_type(4)));
5+
typedef _Float16 half;
6+
typedef half half4 __attribute__((ext_vector_type(4)));
67

78
// CHECK-LABEL: define dso_local void @test_builtin_image_store_2d(
89
// CHECK-SAME: float noundef [[F32:%.*]], i32 noundef [[I32:%.*]], ptr [[TEX:%.*]]) #[[ATTR0:[0-9]+]] {

‎clang/test/SemaOpenCL/builtins-image-load-param-gfx1100-err.cl‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33

44
typedef int int4 __attribute__((ext_vector_type(4)));
55
typedef float float4 __attribute__((ext_vector_type(4)));
6-
typedef _Float16 half4 __attribute__((ext_vector_type(4)));
6+
typedef half half4 __attribute__((ext_vector_type(4)));
77

88
float test_builtin_image_load_2d(float f32, int i32, __amdgpu_texture_t tex) {
99

‎clang/test/SemaOpenCL/builtins-image-load-param-gfx942-err.cl‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33

44
typedef int int4 __attribute__((ext_vector_type(4)));
55
typedef float float4 __attribute__((ext_vector_type(4)));
6-
typedef _Float16 half4 __attribute__((ext_vector_type(4)));
6+
typedef half half4 __attribute__((ext_vector_type(4)));
77

88
float test_builtin_image_load_2d(float f32, int i32, __amdgpu_texture_t tex) {
99

‎clang/test/SemaOpenCL/builtins-image-store-param-gfx1100-err.cl‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22
// REQUIRES: amdgpu-registered-target
33

44
typedef float float4 __attribute__((ext_vector_type(4)));
5-
typedef _Float16 half4 __attribute__((ext_vector_type(4)));
5+
typedef half half4 __attribute__((ext_vector_type(4)));
66

77
void test_builtin_image_store_2d(float f32, int i32, __amdgpu_texture_t tex) {
88

‎clang/test/SemaOpenCL/builtins-image-store-param-gfx942-err.cl‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22
// REQUIRES: amdgpu-registered-target
33

44
typedef float float4 __attribute__((ext_vector_type(4)));
5-
typedef _Float16 half4 __attribute__((ext_vector_type(4)));
5+
typedef half half4 __attribute__((ext_vector_type(4)));
66

77
void test_builtin_image_store_2d(float f32, int i32, __amdgpu_texture_t tex) {
88

0 commit comments

Comments
 (0)