Skip to content

GitLab

Commit 505c5ff3 authored by zhuwenwen's avatar zhuwenwen
Browse files

Merge remote-tracking branch 'origin/0.6.2-zhagnshao' into v0.6.2-dev

parents 2b1b8c7d 76058dcd
Show whitespace changes
Inline Side-by-side
Showing with 2 additions and 2 deletions
csrc/attention/attention_kernels_opt_tc.cu
View file @ 505c5ff3
...@@ -235,7 +235,7 @@ __device__ void paged_attention_kernel_TC( ...@@ -235,7 +235,7 @@ __device__ void paged_attention_kernel_TC(
__shared__ half4x2 q_vecs[REUSE_KV_TIMES][16]; __shared__ half4x2 q_vecs[REUSE_KV_TIMES][16];
//if(thread_idx==0)printf("blockIdx.x==%d,q_boundary=%d,head_idx=%d,kv_head_idx=%d\n",blockIdx.x,q_boundary,head_idx,kv_head_idx); //if(thread_idx==0)printf("blockIdx.x==%d,q_boundary=%d,head_idx=%d,kv_head_idx=%d\n",blockIdx.x,q_boundary,head_idx,kv_head_idx);
for(int i=0;i<REUSE_KV_TIMES;i++){ for(int i=0;i<q_boundary;i++){
if(thread_idx<16){ if(thread_idx<16){
q_vecs[i][thread_idx]=*reinterpret_cast<const half4x2*>(q_ptr+i*HEAD_SIZE+thread_idx*8); q_vecs[i][thread_idx]=*reinterpret_cast<const half4x2*>(q_ptr+i*HEAD_SIZE+thread_idx*8);
} }
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment