squall/cutlass
1
0
Fork 0
Code Issues Pull Requests Actions 1 Packages Projects Releases Wiki Activity

fix cp.async L2 prefetch typo (#1187)

Browse Source
This commit is contained in:
reed 2023-11-29 05:58:04 +08:00 committed by GitHub
parent 8098336d51
commit eb01d5449d
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
1 changed files with 1 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
include/cute/arch/copy_sm80.hpp
View File

@ -86,7 +86,7 @@ struct SM80_CP_ASYNC_CACHEGLOBAL
#if defined(CUTE_ARCH_CP_ASYNC_SM80_ENABLED) #if defined(CUTE_ARCH_CP_ASYNC_SM80_ENABLED)
TS const* gmem_ptr = &gmem_src; TS const* gmem_ptr = &gmem_src;
uint32_t smem_int_ptr = cast_smem_ptr_to_uint(&smem_dst); uint32_t smem_int_ptr = cast_smem_ptr_to_uint(&smem_dst);
asm volatile("cp.async.cg.shared.global.L2::128BB [%0], [%1], %2;\n" asm volatile("cp.async.cg.shared.global.L2::128B [%0], [%1], %2;\n"
:: "r"(smem_int_ptr), :: "r"(smem_int_ptr),
"l"(gmem_ptr), "l"(gmem_ptr),
"n"(sizeof(TS))); "n"(sizeof(TS)));