AMDGPU: Implement tensor load and store instructions for gfx1250

 A few update based on review comments:

1. fix clang format
2. use ClangBuiltin in intrinsic defs to avoid the boilerplate in
   TargetBuiltins/AMDGPU.cpp
3. Fold to return of bool expression

Author: changpeng

clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp

@@ -621,32 +621,6 @@ 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)));
-    llvm::Function *F = CGM.getIntrinsic(IID, {});
-    return Builder.CreateCall(F, {Args});
-  }
   case AMDGPU::BI__builtin_amdgcn_load_to_lds: {
     // Should this have asan instrumentation?
     return emitBuiltinWithOneOverloadedType<5>(*this, E,

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

@@ -3606,10 +3606,10 @@ class AMDGPUTensorLoadStoreD2:
     "", [SDNPMemOperand]
   >;
 
-def int_amdgcn_tensor_load_to_lds : AMDGPUTensorLoadStore;
-def int_amdgcn_tensor_store_from_lds : AMDGPUTensorLoadStore;
-def int_amdgcn_tensor_load_to_lds_d2 : AMDGPUTensorLoadStoreD2;
-def int_amdgcn_tensor_store_from_lds_d2 : AMDGPUTensorLoadStoreD2;
+def int_amdgcn_tensor_load_to_lds : ClangBuiltin<"__builtin_amdgcn_tensor_load_to_lds">, AMDGPUTensorLoadStore;
+def int_amdgcn_tensor_store_from_lds : ClangBuiltin<"__builtin_amdgcn_tensor_store_from_lds">, AMDGPUTensorLoadStore;
+def int_amdgcn_tensor_load_to_lds_d2 : ClangBuiltin<"__builtin_amdgcn_tensor_load_to_lds_d2">, AMDGPUTensorLoadStoreD2;
+def int_amdgcn_tensor_store_from_lds_d2 : ClangBuiltin<"__builtin_amdgcn_tensor_store_from_lds_d2">, AMDGPUTensorLoadStoreD2;
 
 /// Emit an addrspacecast without null pointer checking.
 /// Should only be inserted by a pass based on analysis of an addrspacecast's src.

llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

@@ -4289,10 +4289,8 @@ bool AMDGPUAsmParser::validateTensorR128(const MCInst &Inst) {
     return true;
 
   int R128Idx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::r128);
-  if (R128Idx >= 0 && Inst.getOperand(R128Idx).getImm())
-    return false;
 
-  return true;
+  return R128Idx < 0 || !Inst.getOperand(R128Idx).getImm();
 }
 
 static bool IsRevOpcode(const unsigned Opcode)
@@ -5219,7 +5217,7 @@ bool AMDGPUAsmParser::validateInstruction(const MCInst &Inst,
   }
   if (!validateTensorR128(Inst)) {
     Error(getImmLoc(AMDGPUOperand::ImmTyD16, Operands),
-      "instruction must set modifier r128=0");
+          "instruction must set modifier r128=0");
     return false;
   }
   if (!validateMIMGMSAA(Inst)) {