fix ci build error and fix llama 65b error

Author: RichardWooSJTU

paddle/fluid/framework/tensor_util.cc

@@ -307,7 +307,7 @@ void TensorCopySync(const phi::DenseTensor& src,
     return;
   }
 
-  VLOG(0) << "TensorCopySync " << src.dims() << " from " << src.place()
+  VLOG(3) << "TensorCopySync " << src.dims() << " from " << src.place()
           << " to " << dst_place;
   src.check_memory_size();
   dst->Resize(src.dims());

paddle/fluid/pybind/inference_api.cc

@@ -345,10 +345,15 @@ void PaddleTensorShareExternalData(paddle_infer::Tensor &tensor,     // NOLINT
         static_cast<int64_t *>(paddle_tensor.data<int64_t>()),
         shape,
         ToPaddleInferPlace(paddle_tensor.place().GetType()));
+  } else if (paddle_tensor.dtype() == phi::DataType::UINT8) {
+    tensor.ShareExternalData(
+        static_cast<uint8_t *>(paddle_tensor.data()),
+        shape,
+        ToPaddleInferPlace(paddle_tensor.place().GetType()));
   } else {
     PADDLE_THROW(platform::errors::Unimplemented(
         "Unsupported data type. Now share_external_data only supports INT32, "
-        "INT64, FLOAT32, FLOAT16, BFLOAT16 and BOOL."));
+        "INT64, UINT8, FLOAT32, FLOAT16, BFLOAT16 and BOOL."));
   }
 }
 

paddle/phi/kernels/fusion/gpu/mmha_util.cu.h

@@ -1633,6 +1633,7 @@ inline __device__ uint8_t round_tmp(float16 val) {
   return static_cast<uint8_t>(quant_value + 128.0);
 }
 
+#ifdef ENABLE_BF16
 template <>
 inline __device__ uint8_t round_tmp(__nv_bfloat16 val) {
   float quant_value =
@@ -1641,6 +1642,7 @@ inline __device__ uint8_t round_tmp(__nv_bfloat16 val) {
   quant_value = quant_value < -127.0f ? -127.0f : quant_value;
   return static_cast<uint8_t>(quant_value + 128.0);
 }
+#endif
 
 template <>
 inline __device__ uint16_t round_tmp(float2 val) {
@@ -1726,6 +1728,7 @@ inline __device__ uint64_t round_tmp(uint4 val) {
   return ret;
 }
 
+#ifdef ENABLE_BF16
 template <>
 inline __device__ uint16_t round_tmp(__nv_bfloat162 val) {
   union {
@@ -1760,6 +1763,7 @@ inline __device__ uint64_t round_tmp(bf16_8_t val) {
   int16[3] = round_tmp<uint16_t, __nv_bfloat162>(val.w);
   return int64;
 }
+#endif
 
 inline __device__ float2 rotary_embedding_coefficient(const int zid,
                                                       const int rot_embed_dim,

paddle/phi/kernels/gpu/c_embedding_kernel.cu

@@ -41,16 +41,15 @@ __global__ void CEmbedding(T* out,
   CUDA_KERNEL_LOOP(i, limit) {
     size_t row = i / columns;
     size_t col = i % columns;
-    auto id = ids[row];
+    auto id = static_cast<int64_t>(ids[row]);
 
-    PADDLE_ENFORCE(
-        id >= 0 && (vocab_size < 0 || id < vocab_size),
-        "The index is out of bounds, "
-        "please check whether the dimensions of index and "
-        "input meet the requirements. It should "
-        "be less than [%d] and greater than or equal to 0, but received [%d]",
-        vocab_size,
-        id);
+    // PADDLE_ENFORCE(
+    //     id >= 0 && (vocab_size < 0 || id < vocab_size),
+    //     "The index is out of bounds, "
+    //     "please check whether the dimensions of index and "
+    //     "input meet the requirements. It should "
+    //     "be less than [%d] and greater than or equal to 0, but received
+    //     [%d]", vocab_size, id);
     if (id >= start_idx && id < end_idx) {
       auto real_idx = id - start_idx;
       out[i] = table[real_idx * columns + col];

test/legacy_test/test_block_multihead_attention.py

@@ -88,6 +88,7 @@ def naive_attention_impl(
     scale=1.0,
     cache_k_dequant_scales=None,
     cache_v_dequant_scales=None,
+    use_cachekv_int8="None",
 ):
     batch = query.shape[0]
     heads = query.shape[1]
@@ -98,13 +99,18 @@ def naive_attention_impl(
     key = key.reshape([batch, kv_head, 1, seq_len, head_dim])
     key = paddle.tile(key, [1, 1, heads // kv_head, 1, 1])
     key = key.reshape([batch, heads, seq_len, head_dim])
+
+    if use_cachekv_int8 == "dynamic":
+        unsqueeze_shape = [2, 3]
+    elif use_cachekv_int8 == "static":
+        unsqueeze_shape = [0, 2, 3]
     if pre_cache_k is not None:
         key = paddle.concat([pre_cache_k, key], axis=2)
     if cache_k is not None:
         if cache_k_dequant_scales is not None:
             dequant_cache_k = (
                 (cache_k.astype('float32') - 128.0)
-                * cache_k_dequant_scales.unsqueeze([0, 2, 3])
+                * cache_k_dequant_scales.unsqueeze(unsqueeze_shape)
             ).astype(key.dtype)
             key = paddle.concat([dequant_cache_k, key], axis=2)
         else:
@@ -119,7 +125,7 @@ def naive_attention_impl(
         if cache_v_dequant_scales is not None:
             dequant_cache_v = (
                 (cache_v.astype('float32') - 128.0)
-                * cache_v_dequant_scales.unsqueeze([0, 2, 3])
+                * cache_v_dequant_scales.unsqueeze(unsqueeze_shape)
             ).astype(value.dtype)
             value = paddle.concat([dequant_cache_v, value], axis=2)
         else:
@@ -1306,6 +1312,13 @@ def test_all(self):
         )
 
 
+@unittest.skipIf(
+    not core.is_compiled_with_cuda()
+    or get_cuda_version() < 11040
+    or not is_sm_supported,
+    "core is not compiled with CUDA and cuda version need larger than or equal to 11.4"
+    "and device's compute capability must be 8.x or 90",
+)
 class TestBlockMultiHeadAttnEncDecPTQDequant(unittest.TestCase):
     def setUp(self):
         paddle.disable_static()
@@ -1641,6 +1654,13 @@ def test_all(self):
         )
 
 
+@unittest.skipIf(
+    not core.is_compiled_with_cuda()
+    or get_cuda_version() < 11040
+    or not is_sm_supported,
+    "core is not compiled with CUDA and cuda version need larger than or equal to 11.4"
+    "and device's compute capability must be 8.x or 90",
+)
 class TestBlockMultiHeadAttnEncDecPTQDequantQuantShiftSmooth(unittest.TestCase):
     def setUp(self):
         paddle.disable_static()
@@ -2013,6 +2033,13 @@ def test_all(self):
         )
 
 
+@unittest.skipIf(
+    not core.is_compiled_with_cuda()
+    or get_cuda_version() < 11040
+    or not is_sm_supported,
+    "core is not compiled with CUDA and cuda version need larger than or equal to 11.4"
+    "and device's compute capability must be 8.x or 90",
+)
 class TestBlockMultiHeadAttnEncDecQuant(unittest.TestCase):
     def setUp(self):
         paddle.disable_static()
@@ -2282,6 +2309,13 @@ def test_all(self):
         )
 
 
+@unittest.skipIf(
+    not core.is_compiled_with_cuda()
+    or get_cuda_version() < 11040
+    or not is_sm_supported,
+    "core is not compiled with CUDA and cuda version need larger than or equal to 11.4"
+    "and device's compute capability must be 8.x or 90",
+)
 class TestBlockMultiHeadAttnEncDecCacheKVDynamicQuant(unittest.TestCase):
     def setUp(self):
         paddle.disable_static()
@@ -2339,16 +2373,16 @@ def setUp(self):
         self.cache_k = paddle.zeros(shape=self.cache_shape, dtype='uint8')
         self.cache_v = paddle.zeros(shape=self.cache_shape, dtype='uint8')
         self.cache_k_quant_scales = paddle.zeros(
-            shape=[self.num_head], dtype='float32'
+            shape=[self.batch_size, self.num_head], dtype='float32'
         )
         self.cache_v_quant_scales = paddle.zeros(
-            shape=[self.num_head], dtype='float32'
+            shape=[self.batch_size, self.num_head], dtype='float32'
         )
         self.cache_k_dequant_scales = paddle.zeros(
-            shape=[self.num_head], dtype='float32'
+            shape=[self.batch_size, self.num_head], dtype='float32'
         )
         self.cache_v_dequant_scales = paddle.zeros(
-            shape=[self.num_head], dtype='float32'
+            shape=[self.batch_size, self.num_head], dtype='float32'
         )
 
         self.block_tables = paddle.zeros(
@@ -2510,6 +2544,7 @@ def test_all(self):
                 self.scale,
                 cache_k_dequant_scales=self.cache_k_dequant_scales,
                 cache_v_dequant_scales=self.cache_v_dequant_scales,
+                use_cachekv_int8="dynamic",
             )
             .transpose([0, 2, 1, 3])
             .reshape([self.batch_size, -1])
@@ -2555,6 +2590,13 @@ def test_all(self):
         )
 
 
+@unittest.skipIf(
+    not core.is_compiled_with_cuda()
+    or get_cuda_version() < 11040
+    or not is_sm_supported,
+    "core is not compiled with CUDA and cuda version need larger than or equal to 11.4"
+    "and device's compute capability must be 8.x or 90",
+)
 class TestBlockMultiHeadAttnEncDecCacheKVStaticQuant(unittest.TestCase):
     def setUp(self):
         paddle.disable_static()
@@ -2795,6 +2837,7 @@ def test_all(self):
                 self.scale,
                 cache_k_dequant_scales=self.cache_k_dequant_scales,
                 cache_v_dequant_scales=self.cache_v_dequant_scales,
+                use_cachekv_int8="static",
             )
             .transpose([0, 2, 1, 3])
             .reshape([self.batch_size, -1])