@@ -217,8 +217,7 @@ static llvm::Value *LoadTextureDescPtorAsVec8I32(CodeGenFunction &CGF,
217
217
if (RsrcPtr->getType ()->isIntegerTy (32 )) {
218
218
unsigned AS = 8 ;
219
219
llvm::PointerType *VecPtrTy = llvm::PointerType::get (VecTy, AS);
220
- llvm::Value *Ptr =
221
- B.CreateIntToPtr (RsrcPtr, VecPtrTy, " tex.rsrc.from.int" );
220
+ llvm::Value *Ptr = B.CreateIntToPtr (RsrcPtr, VecPtrTy, " tex.rsrc.from.int" );
222
221
return B.CreateAlignedLoad (VecTy, Ptr, llvm::Align (32 ), " tex.rsrc.val" );
223
222
}
224
223
@@ -240,22 +239,24 @@ static llvm::Value *LoadTextureDescPtorAsVec8I32(CodeGenFunction &CGF,
240
239
static unsigned GetTextureDescIndex (unsigned BuiltinID, const CallExpr *E) {
241
240
unsigned N = E->getNumArgs ();
242
241
if (IsImageSampleBuiltIn (BuiltinID)) {
243
- if (N < 5 ) return (unsigned )-1 ;
242
+ if (N < 5 )
243
+ return (unsigned )-1 ;
244
244
return N - 5 ;
245
245
}
246
-
247
- if (N < 3 ) return (unsigned )-1 ;
246
+
247
+ if (N < 3 )
248
+ return (unsigned )-1 ;
248
249
return N - 3 ;
249
250
}
250
251
251
- llvm::CallInst *EmitAMDGCNImageOverloadedReturnType (clang::CodeGen::CodeGenFunction &CGF,
252
- const clang::CallExpr *E ,
253
- unsigned IntrinsicID ,
254
- bool IsImageStore) {
252
+ llvm::CallInst *
253
+ EmitAMDGCNImageOverloadedReturnType ( clang::CodeGen::CodeGenFunction &CGF ,
254
+ const clang::CallExpr *E ,
255
+ unsigned IntrinsicID, bool IsImageStore) {
255
256
clang::SmallVector<llvm::Value *, 10 > Args;
256
257
unsigned RsrcIndex = GetTextureDescIndex (E->getBuiltinCallee (), E);
257
258
258
- for (unsigned I = 0 ; I < E->getNumArgs (); ++I){
259
+ for (unsigned I = 0 ; I < E->getNumArgs (); ++I) {
259
260
llvm::Value *V = CGF.EmitScalarExpr (E->getArg (I));
260
261
if (I == RsrcIndex)
261
262
V = LoadTextureDescPtorAsVec8I32 (CGF, V);
@@ -1028,133 +1029,133 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1028
1029
case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
1029
1030
case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
1030
1031
return EmitAMDGCNImageOverloadedReturnType (
1031
- *this , E, Intrinsic::amdgcn_image_load_1d, false );
1032
+ *this , E, Intrinsic::amdgcn_image_load_1d, false );
1032
1033
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
1033
1034
case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
1034
1035
return EmitAMDGCNImageOverloadedReturnType (
1035
- *this , E, Intrinsic::amdgcn_image_load_1darray, false );
1036
+ *this , E, Intrinsic::amdgcn_image_load_1darray, false );
1036
1037
case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
1037
1038
case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
1038
1039
case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
1039
1040
return EmitAMDGCNImageOverloadedReturnType (
1040
- *this , E, Intrinsic::amdgcn_image_load_2d, false );
1041
+ *this , E, Intrinsic::amdgcn_image_load_2d, false );
1041
1042
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
1042
1043
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
1043
1044
case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
1044
1045
return EmitAMDGCNImageOverloadedReturnType (
1045
- *this , E, Intrinsic::amdgcn_image_load_2darray, false );
1046
+ *this , E, Intrinsic::amdgcn_image_load_2darray, false );
1046
1047
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
1047
1048
case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
1048
1049
return EmitAMDGCNImageOverloadedReturnType (
1049
- *this , E, Intrinsic::amdgcn_image_load_3d, false );
1050
+ *this , E, Intrinsic::amdgcn_image_load_3d, false );
1050
1051
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
1051
1052
case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
1052
1053
return EmitAMDGCNImageOverloadedReturnType (
1053
- *this , E, Intrinsic::amdgcn_image_load_cube, false );
1054
+ *this , E, Intrinsic::amdgcn_image_load_cube, false );
1054
1055
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
1055
1056
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
1056
1057
return EmitAMDGCNImageOverloadedReturnType (
1057
- *this , E, Intrinsic::amdgcn_image_load_mip_1d, false );
1058
+ *this , E, Intrinsic::amdgcn_image_load_mip_1d, false );
1058
1059
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
1059
1060
case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
1060
1061
return EmitAMDGCNImageOverloadedReturnType (
1061
- *this , E, Intrinsic::amdgcn_image_load_mip_1darray, false );
1062
+ *this , E, Intrinsic::amdgcn_image_load_mip_1darray, false );
1062
1063
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
1063
1064
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
1064
1065
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
1065
1066
return EmitAMDGCNImageOverloadedReturnType (
1066
- *this , E, Intrinsic::amdgcn_image_load_mip_2d, false );
1067
+ *this , E, Intrinsic::amdgcn_image_load_mip_2d, false );
1067
1068
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
1068
1069
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
1069
1070
case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
1070
1071
return EmitAMDGCNImageOverloadedReturnType (
1071
- *this , E, Intrinsic::amdgcn_image_load_mip_2darray, false );
1072
+ *this , E, Intrinsic::amdgcn_image_load_mip_2darray, false );
1072
1073
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
1073
1074
case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
1074
1075
return EmitAMDGCNImageOverloadedReturnType (
1075
- *this , E, Intrinsic::amdgcn_image_load_mip_3d, false );
1076
+ *this , E, Intrinsic::amdgcn_image_load_mip_3d, false );
1076
1077
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
1077
1078
case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
1078
1079
return EmitAMDGCNImageOverloadedReturnType (
1079
- *this , E, Intrinsic::amdgcn_image_load_mip_cube, false );
1080
+ *this , E, Intrinsic::amdgcn_image_load_mip_cube, false );
1080
1081
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
1081
1082
case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
1082
1083
return EmitAMDGCNImageOverloadedReturnType (
1083
- *this , E, Intrinsic::amdgcn_image_store_1d, true );
1084
+ *this , E, Intrinsic::amdgcn_image_store_1d, true );
1084
1085
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
1085
1086
case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
1086
1087
return EmitAMDGCNImageOverloadedReturnType (
1087
- *this , E, Intrinsic::amdgcn_image_store_1darray, true );
1088
+ *this , E, Intrinsic::amdgcn_image_store_1darray, true );
1088
1089
case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
1089
1090
case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
1090
1091
case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
1091
1092
return EmitAMDGCNImageOverloadedReturnType (
1092
- *this , E, Intrinsic::amdgcn_image_store_2d, true );
1093
+ *this , E, Intrinsic::amdgcn_image_store_2d, true );
1093
1094
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
1094
1095
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
1095
1096
case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
1096
1097
return EmitAMDGCNImageOverloadedReturnType (
1097
- *this , E, Intrinsic::amdgcn_image_store_2darray, true );
1098
+ *this , E, Intrinsic::amdgcn_image_store_2darray, true );
1098
1099
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
1099
1100
case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
1100
1101
return EmitAMDGCNImageOverloadedReturnType (
1101
- *this , E, Intrinsic::amdgcn_image_store_3d, true );
1102
+ *this , E, Intrinsic::amdgcn_image_store_3d, true );
1102
1103
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
1103
1104
case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
1104
1105
return EmitAMDGCNImageOverloadedReturnType (
1105
- *this , E, Intrinsic::amdgcn_image_store_cube, true );
1106
+ *this , E, Intrinsic::amdgcn_image_store_cube, true );
1106
1107
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
1107
1108
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
1108
1109
return EmitAMDGCNImageOverloadedReturnType (
1109
- *this , E, Intrinsic::amdgcn_image_store_mip_1d, true );
1110
+ *this , E, Intrinsic::amdgcn_image_store_mip_1d, true );
1110
1111
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
1111
1112
case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
1112
1113
return EmitAMDGCNImageOverloadedReturnType (
1113
- *this , E, Intrinsic::amdgcn_image_store_mip_1darray, true );
1114
+ *this , E, Intrinsic::amdgcn_image_store_mip_1darray, true );
1114
1115
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
1115
1116
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
1116
1117
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
1117
1118
return EmitAMDGCNImageOverloadedReturnType (
1118
- *this , E, Intrinsic::amdgcn_image_store_mip_2d, true );
1119
+ *this , E, Intrinsic::amdgcn_image_store_mip_2d, true );
1119
1120
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
1120
1121
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
1121
1122
case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
1122
1123
return EmitAMDGCNImageOverloadedReturnType (
1123
- *this , E, Intrinsic::amdgcn_image_store_mip_2darray, true );
1124
+ *this , E, Intrinsic::amdgcn_image_store_mip_2darray, true );
1124
1125
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
1125
1126
case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
1126
1127
return EmitAMDGCNImageOverloadedReturnType (
1127
- *this , E, Intrinsic::amdgcn_image_store_mip_3d, true );
1128
+ *this , E, Intrinsic::amdgcn_image_store_mip_3d, true );
1128
1129
case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
1129
- case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32:
1130
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32:
1130
1131
return EmitAMDGCNImageOverloadedReturnType (
1131
- *this , E, Intrinsic::amdgcn_image_store_mip_cube, true );
1132
+ *this , E, Intrinsic::amdgcn_image_store_mip_cube, true );
1132
1133
case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
1133
1134
case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
1134
1135
return EmitAMDGCNImageOverloadedReturnType (
1135
- *this , E, Intrinsic::amdgcn_image_sample_1d, false );
1136
+ *this , E, Intrinsic::amdgcn_image_sample_1d, false );
1136
1137
case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
1137
1138
case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
1138
1139
return EmitAMDGCNImageOverloadedReturnType (
1139
- *this , E, Intrinsic::amdgcn_image_sample_1darray, false );
1140
+ *this , E, Intrinsic::amdgcn_image_sample_1darray, false );
1140
1141
case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
1141
1142
case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
1142
1143
case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
1143
1144
return EmitAMDGCNImageOverloadedReturnType (
1144
- *this , E, Intrinsic::amdgcn_image_sample_2d, false );
1145
+ *this , E, Intrinsic::amdgcn_image_sample_2d, false );
1145
1146
case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
1146
1147
case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
1147
1148
case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
1148
1149
return EmitAMDGCNImageOverloadedReturnType (
1149
- *this , E, Intrinsic::amdgcn_image_sample_2darray, false );
1150
+ *this , E, Intrinsic::amdgcn_image_sample_2darray, false );
1150
1151
case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
1151
1152
case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
1152
1153
return EmitAMDGCNImageOverloadedReturnType (
1153
- *this , E, Intrinsic::amdgcn_image_sample_3d, false );
1154
+ *this , E, Intrinsic::amdgcn_image_sample_3d, false );
1154
1155
case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
1155
1156
case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32:
1156
1157
return EmitAMDGCNImageOverloadedReturnType (
1157
- *this , E, Intrinsic::amdgcn_image_sample_cube, false );
1158
+ *this , E, Intrinsic::amdgcn_image_sample_cube, false );
1158
1159
case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
1159
1160
case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {
1160
1161
llvm::FixedVectorType *VT = FixedVectorType::get (Builder.getInt32Ty (), 8 );
0 commit comments