Updated getTextureDescIndex() to lambda and minor changes

Author: ranapratap55

clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp

@@ -184,29 +184,7 @@ 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,
+static llvm::Value *loadTextureDescPtorAsVec8I32(CodeGenFunction &CGF,
                                                  llvm::Value *RsrcPtr) {
   auto &B = CGF.Builder;
   auto *VecTy = llvm::FixedVectorType::get(B.getInt32Ty(), 8);
@@ -215,15 +193,13 @@ static llvm::Value *LoadTextureDescPtorAsVec8I32(CodeGenFunction &CGF,
     return RsrcPtr;
 
   if (RsrcPtr->getType()->isIntegerTy(32)) {
-    unsigned AS = 8;
-    llvm::PointerType *VecPtrTy = llvm::PointerType::get(VecTy, AS);
+    llvm::PointerType *VecPtrTy = llvm::PointerType::get(CGF.getLLVMContext(), 8);
     llvm::Value *Ptr = B.CreateIntToPtr(RsrcPtr, VecPtrTy, "tex.rsrc.from.int");
     return B.CreateAlignedLoad(VecTy, Ptr, llvm::Align(32), "tex.rsrc.val");
   }
 
   if (RsrcPtr->getType()->isPointerTy()) {
-    unsigned AS = RsrcPtr->getType()->getPointerAddressSpace();
-    auto *VecPtrTy = llvm::PointerType::get(VecTy, AS);
+    auto *VecPtrTy = llvm::PointerType::get(CGF.getLLVMContext(), RsrcPtr->getType()->getPointerAddressSpace());
     llvm::Value *Typed = B.CreateBitCast(RsrcPtr, VecPtrTy, "tex.rsrc.typed");
     return B.CreateAlignedLoad(VecTy, Typed, llvm::Align(32), "tex.rsrc.val");
   }
@@ -232,40 +208,44 @@ static llvm::Value *LoadTextureDescPtorAsVec8I32(CodeGenFunction &CGF,
   if (DL.getTypeSizeInBits(RsrcPtr->getType()) == 256)
     return B.CreateBitCast(RsrcPtr, VecTy, "tex.rsrc.val");
 
-  RsrcPtr->getType()->print(llvm::errs());
-  llvm::report_fatal_error(": Unexpected texture resource argument form");
-}
-
-static unsigned GetTextureDescIndex(unsigned BuiltinID, const CallExpr *E) {
-  unsigned N = E->getNumArgs();
-  if (IsImageSampleBuiltIn(BuiltinID)) {
-    if (N < 5)
-      return (unsigned)-1;
-    return N - 5;
-  }
-
-  if (N < 3)
-    return (unsigned)-1;
-  return N - 3;
+  llvm::report_fatal_error("Unexpected texture resource argument form");
 }
 
 llvm::CallInst *
-EmitAMDGCNImageOverloadedReturnType(clang::CodeGen::CodeGenFunction &CGF,
+emitAMDGCNImageOverloadedReturnType(clang::CodeGen::CodeGenFunction &CGF,
                                     const clang::CallExpr *E,
                                     unsigned IntrinsicID, bool IsImageStore) {
+  auto findTextureDescIndex = [&CGF](const CallExpr *E) -> unsigned {
+    QualType TexQT = CGF.getContext().AMDGPUTextureTy;
+    for (unsigned I = 0, N = E->getNumArgs(); I < N; ++I) {
+      QualType ArgTy = E->getArg(I)->getType();
+      if (ArgTy == TexQT) {
+        return I;
+      }
+
+      if (ArgTy.getCanonicalType() == TexQT.getCanonicalType()) {
+        return I;
+      }
+    }
+    
+    return ~0U;
+  };
+
   clang::SmallVector<llvm::Value *, 10> Args;
-  unsigned RsrcIndex = GetTextureDescIndex(E->getBuiltinCallee(), E);
+  unsigned RsrcIndex = findTextureDescIndex(E);
+
+  if (RsrcIndex == ~0U) {
+    llvm::report_fatal_error("Invalid argument count for image builtin");
+  }
 
   for (unsigned I = 0; I < E->getNumArgs(); ++I) {
     llvm::Value *V = CGF.EmitScalarExpr(E->getArg(I));
     if (I == RsrcIndex)
-      V = LoadTextureDescPtorAsVec8I32(CGF, V);
+      V = loadTextureDescPtorAsVec8I32(CGF, V);
     Args.push_back(V);
   }
 
-  llvm::Type *RetTy = CGF.ConvertType(E->getType());
-  if (IsImageStore)
-    RetTy = CGF.VoidTy;
+  llvm::Type *RetTy = IsImageStore ? CGF.VoidTy: CGF.ConvertType(E->getType());
   llvm::CallInst *Call = CGF.Builder.CreateIntrinsic(RetTy, IntrinsicID, Args);
   return Call;
 }
@@ -1028,133 +1008,133 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
   }
   case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_load_1d, false);
   case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_load_1darray, false);
   case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_load_2d, false);
   case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_load_2darray, false);
   case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_load_3d, false);
   case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_load_cube, false);
   case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_load_mip_1d, false);
   case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_load_mip_1darray, false);
   case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_load_mip_2d, false);
   case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_load_mip_2darray, false);
   case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_load_mip_3d, false);
   case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_load_mip_cube, false);
   case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_store_1d, true);
   case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_store_1darray, true);
   case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_store_2d, true);
   case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_store_2darray, true);
   case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_store_3d, true);
   case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_store_cube, true);
   case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_store_mip_1d, true);
   case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_store_mip_1darray, true);
   case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_store_mip_2d, true);
   case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_store_mip_2darray, true);
   case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_store_mip_3d, true);
   case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_store_mip_cube, true);
   case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
   case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_sample_1d, false);
   case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
   case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_sample_1darray, false);
   case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
   case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
   case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_sample_2d, false);
   case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
   case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
   case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_sample_2darray, false);
   case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
   case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_sample_3d, false);
   case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
   case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32:
-    return EmitAMDGCNImageOverloadedReturnType(
+    return emitAMDGCNImageOverloadedReturnType(
         *this, E, Intrinsic::amdgcn_image_sample_cube, false);
   case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
   case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {

clang/lib/Sema/SemaAMDGPU.cpp

@@ -166,9 +166,9 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
     unsigned ArgCount = TheCall->getNumArgs() - 1;
     llvm::APSInt Result;
 
-    return ((SemaRef.BuiltinConstantArg(TheCall, 0, Result)) ||
+    return (SemaRef.BuiltinConstantArg(TheCall, 0, Result)) ||
             (SemaRef.BuiltinConstantArg(TheCall, ArgCount, Result)) ||
-            (SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result)));
+            (SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result));
   }
   case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
@@ -210,9 +210,9 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
     unsigned ArgCount = TheCall->getNumArgs() - 1;
     llvm::APSInt Result;
 
-    return ((SemaRef.BuiltinConstantArg(TheCall, 1, Result)) ||
+    return (SemaRef.BuiltinConstantArg(TheCall, 1, Result)) ||
             (SemaRef.BuiltinConstantArg(TheCall, ArgCount, Result)) ||
-            (SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result)));
+            (SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result));
   }
   default:
     return false;