未验证 提交 974f8f32 编写于 作者: F feng_shuai 提交者: GitHub

fix:add no support for cuda_arch<700 (#47509)

上级 bfc45e34
...@@ -836,7 +836,13 @@ inline void MatMulWithHeadQK(const phi::GPUContext &context, ...@@ -836,7 +836,13 @@ inline void MatMulWithHeadQK(const phi::GPUContext &context,
FINAL_MASK); FINAL_MASK);
} else { } else {
if (bias_is_mask) { if (bias_is_mask) {
#ifndef __HIPCC__ #if defined(__HIPCC__) || (defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 700)
PADDLE_ENFORCE_EQ(bias_is_mask,
false,
platform::errors::InvalidArgument(
"QK_bias is mask can't be supported on rocm or "
"cuda_arch<700"));
#else
constexpr int ITEMS_PER_THREAD = 1; constexpr int ITEMS_PER_THREAD = 1;
bool is_half2 = true; bool is_half2 = true;
...@@ -853,11 +859,6 @@ inline void MatMulWithHeadQK(const phi::GPUContext &context, ...@@ -853,11 +859,6 @@ inline void MatMulWithHeadQK(const phi::GPUContext &context,
batch_size, batch_size,
head_num, head_num,
seq_len); seq_len);
#else
PADDLE_ENFORCE_EQ(bias_is_mask,
false,
platform::errors::InvalidArgument(
"rocm can't support that QK_bias is mask"));
#endif #endif
} else { } else {
SoftmaxKernelWithEltadd2<__half2><<<grid, block, 0, stream>>>( SoftmaxKernelWithEltadd2<__half2><<<grid, block, 0, stream>>>(
......
Markdown is supported
0% .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑!
想要评论请 注册