================
@@ -181,6 +184,92 @@ static Value *EmitAMDGCNBallotForExec(CodeGenFunction
&CGF, const CallExpr *E,
return Call;
}
+static bool IsImageSampleBuiltIn(unsigned BuiltinID) {
+ switch (BuiltinID) {
+ case clang::AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
+ case clang::AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
+ case clang::AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
+ case clang::AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
+ case clang::AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
+ case clang::AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
+ case clang::AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
+ case clang::AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
+ case clang::AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
+ case clang::AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
+ case clang::AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
+ case clang::AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
+ case clang::AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
+ case clang::AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32:
+ return true;
+ default:
+ return false;
+ }
+}
+
+static llvm::Value *LoadTextureDescPtorAsVec8I32(CodeGenFunction &CGF,
+ llvm::Value *RsrcPtr) {
+ auto &B = CGF.Builder;
+ auto *VecTy = llvm::FixedVectorType::get(B.getInt32Ty(), 8);
+
+ if (RsrcPtr->getType() == VecTy)
+ return RsrcPtr;
+
+ if (RsrcPtr->getType()->isIntegerTy(32)) {
+ unsigned AS = 8;
----------------
ranapratap55 wrote:
Updated.
https://github.com/llvm/llvm-project/pull/140210
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
