diff --git a/clang/test/CodeGen/X86/amx_api.c b/clang/test/CodeGen/X86/amx_api.c
--- a/clang/test/CodeGen/X86/amx_api.c
+++ b/clang/test/CodeGen/X86/amx_api.c
@@ -1,6 +1,9 @@
 // RUN: %clang_cc1 %s -flax-vector-conversions=none -ffreestanding -triple=x86_64-unknown-unknown  -target-feature +avx512f  -target-feature +amx-int8  \
 // RUN: -target-feature +amx-bf16 -emit-llvm -o - -Werror -pedantic | FileCheck %s --check-prefixes=CHECK
 
+// RUN: %clang_cc1 %s -flax-vector-conversions=none -ffreestanding -triple=x86_64-unknown-unknown  -target-feature +avx512f  -target-feature +amx-int8  \
+// RUN: -target-feature +amx-bf16 -O2 -emit-llvm -o - -Werror -pedantic | FileCheck %s --check-prefixes=CHECK2
+
 #include <immintrin.h>
 
 char buf[1024];
@@ -31,6 +34,15 @@
   __tile_stored(buf, STRIDE, c);
 }
 
+// Not fold the bitcast const vector into amx intrisic.
+void test_tile_init(short row, short col) {
+  __tile1024i c = {row, col, {1, 2, 3}};
+  __tile_stored(buf, STRIDE, c);
+  //CHECK2-LABEL: @test_tile_init
+  //CHECK2: {{%.*}} = bitcast <256 x i32> <i32 1, i32 2, i32 3, i32 0,
+  //CHECK2-NEXT: call void @llvm.x86.tilestored64.internal({{.*}}, x86_amx {{%.*}})
+}
+
 void test_tile_loadd(short row, short col) {
   //CHECK-LABEL: @test_tile_loadd
   //CHECK: call x86_amx @llvm.x86.tileloadd64.internal
diff --git a/llvm/lib/Analysis/ConstantFolding.cpp b/llvm/lib/Analysis/ConstantFolding.cpp
--- a/llvm/lib/Analysis/ConstantFolding.cpp
+++ b/llvm/lib/Analysis/ConstantFolding.cpp
@@ -104,10 +104,16 @@
   assert(CastInst::castIsValid(Instruction::BitCast, C, DestTy) &&
          "Invalid constantexpr bitcast!");
 
+  // We won't fold bitcast for tile type, becasue there is no way to
+  // assigne a tmm reg from a constant. We manually generate tilestore
+  // and tileload at pass "Lower AMX type".
+  if (DestTy->isX86_AMXTy())
+    return nullptr;
+
   // Catch the obvious splat cases.
-  if (C->isNullValue() && !DestTy->isX86_MMXTy() && !DestTy->isX86_AMXTy())
+  if (C->isNullValue() && !DestTy->isX86_MMXTy())
     return Constant::getNullValue(DestTy);
-  if (C->isAllOnesValue() && !DestTy->isX86_MMXTy() && !DestTy->isX86_AMXTy() &&
+  if (C->isAllOnesValue() && !DestTy->isX86_MMXTy() &&
       !DestTy->isPtrOrPtrVectorTy()) // Don't get ones for ptr types!
     return Constant::getAllOnesValue(DestTy);
 
diff --git a/llvm/lib/Transforms/Scalar/SCCP.cpp b/llvm/lib/Transforms/Scalar/SCCP.cpp
--- a/llvm/lib/Transforms/Scalar/SCCP.cpp
+++ b/llvm/lib/Transforms/Scalar/SCCP.cpp
@@ -826,7 +826,7 @@
   if (Constant *OpC = getConstant(OpSt)) {
     // Fold the constant as we build.
     Constant *C = ConstantFoldCastOperand(I.getOpcode(), OpC, I.getType(), DL);
-    if (isa<UndefValue>(C))
+    if (!C || isa<UndefValue>(C))
       return;
     // Propagate constant value
     markConstant(&I, C);