[clang] [llvm] AMDGPU: Implement tensor load and store instructions for gfx1250 (PR #146636)

[Changpeng Fang via cfe-commits]

================
@@ -621,6 +621,32 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
     llvm::Function *F = CGM.getIntrinsic(IID, {LoadTy});
     return Builder.CreateCall(F, {Addr});
   }
+  case AMDGPU::BI__builtin_amdgcn_tensor_load_to_lds:
+  case AMDGPU::BI__builtin_amdgcn_tensor_load_to_lds_d2:
+  case AMDGPU::BI__builtin_amdgcn_tensor_store_from_lds:
+  case AMDGPU::BI__builtin_amdgcn_tensor_store_from_lds_d2: {
+    Intrinsic::ID IID;
+    switch (BuiltinID) {
+    case AMDGPU::BI__builtin_amdgcn_tensor_load_to_lds:
+      IID = Intrinsic::amdgcn_tensor_load_to_lds;
+      break;
+    case AMDGPU::BI__builtin_amdgcn_tensor_load_to_lds_d2:
+      IID = Intrinsic::amdgcn_tensor_load_to_lds_d2;
+      break;
+    case AMDGPU::BI__builtin_amdgcn_tensor_store_from_lds:
+      IID = Intrinsic::amdgcn_tensor_store_from_lds;
+      break;
+    case AMDGPU::BI__builtin_amdgcn_tensor_store_from_lds_d2:
+      IID = Intrinsic::amdgcn_tensor_store_from_lds_d2;
+      break;
+    }
+
+    SmallVector<Value *, 5> Args;
+    for (int i = 0, e = E->getNumArgs(); i != e; ++i)
+      Args.push_back(EmitScalarExpr(E->getArg(i)));
----------------
changpeng wrote:

> Can just assign size at the start and assign each index instead of push_back

Do you mean we can use Args[i] = ...? 
I though we do not know the array size at compile time, so push_back can grow. Also, we always sue push_back in similar places. 

https://github.com/llvm/llvm-project/pull/146636