Skip to content

GitLab

Unverified Commit e9d3aa04 authored by Simon Mo's avatar Simon Mo Committed by GitHub
Browse files

Revert "[Kernel] Marlin_24: Ensure the mma.sp instruction is using the... 

Revert "[Kernel] Marlin_24: Ensure the mma.sp instruction is using the ::ordered_metadata modifier (introduced with PTX 8.5)" (#5149)
parent a22dea54
Show whitespace changes
Inline Side-by-side
Showing with 4 additions and 8 deletions
csrc/quantization/marlin/sparse/common/mma.h
View file @ e9d3aa04
...@@ -32,8 +32,7 @@ __device__ inline void mma_sp(const FragB& a_frag0, const FragB& a_frag1, ...@@ -32,8 +32,7 @@ __device__ inline void mma_sp(const FragB& a_frag0, const FragB& a_frag1,
float* c = reinterpret_cast<float*>(&frag_c); float* c = reinterpret_cast<float*>(&frag_c);
if (psel == 0) { if (psel == 0) {
asm volatile( asm volatile(
"mma.sp::ordered_metadata.sync.aligned.m16n8k32.row.col.f32.f16.f16." "mma.sp.sync.aligned.m16n8k32.row.col.f32.f16.f16.f32 "
"f32 "
"{%0, %1, %2, %3}, {%4, %5, %6, %7}, {%8, %9, %10,%11}, " "{%0, %1, %2, %3}, {%4, %5, %6, %7}, {%8, %9, %10,%11}, "
"{%12,%13,%14,%15}, %16, 0x0;\n" "{%12,%13,%14,%15}, %16, 0x0;\n"
: "=f"(c[0]), "=f"(c[1]), "=f"(c[2]), "=f"(c[3]) : "=f"(c[0]), "=f"(c[1]), "=f"(c[2]), "=f"(c[3])
...@@ -41,8 +40,7 @@ __device__ inline void mma_sp(const FragB& a_frag0, const FragB& a_frag1, ...@@ -41,8 +40,7 @@ __device__ inline void mma_sp(const FragB& a_frag0, const FragB& a_frag1,
"r"(b[4]), "r"(b[6]), "f"(c[0]), "f"(c[1]), "f"(c[2]), "f"(c[3]), "r"(b[4]), "r"(b[6]), "f"(c[0]), "f"(c[1]), "f"(c[2]), "f"(c[3]),
"r"(e[0])); "r"(e[0]));
asm volatile( asm volatile(
"mma.sp::ordered_metadata.sync.aligned.m16n8k32.row.col.f32.f16.f16." "mma.sp.sync.aligned.m16n8k32.row.col.f32.f16.f16.f32 "
"f32 "
"{%0, %1, %2, %3}, {%4, %5, %6, %7}, {%8, %9, %10,%11}, " "{%0, %1, %2, %3}, {%4, %5, %6, %7}, {%8, %9, %10,%11}, "
"{%12,%13,%14,%15}, %16, 0x0;\n" "{%12,%13,%14,%15}, %16, 0x0;\n"
: "=f"(c[4]), "=f"(c[5]), "=f"(c[6]), "=f"(c[7]) : "=f"(c[4]), "=f"(c[5]), "=f"(c[6]), "=f"(c[7])
...@@ -51,8 +49,7 @@ __device__ inline void mma_sp(const FragB& a_frag0, const FragB& a_frag1, ...@@ -51,8 +49,7 @@ __device__ inline void mma_sp(const FragB& a_frag0, const FragB& a_frag1,
"r"(e[0])); "r"(e[0]));
} else { } else {
asm volatile( asm volatile(
"mma.sp::ordered_metadata.sync.aligned.m16n8k32.row.col.f32.f16.f16." "mma.sp.sync.aligned.m16n8k32.row.col.f32.f16.f16.f32 "
"f32 "
"{%0, %1, %2, %3}, {%4, %5, %6, %7}, {%8, %9, %10,%11}, " "{%0, %1, %2, %3}, {%4, %5, %6, %7}, {%8, %9, %10,%11}, "
"{%12,%13,%14,%15}, %16, 0x1;\n" "{%12,%13,%14,%15}, %16, 0x1;\n"
: "=f"(c[0]), "=f"(c[1]), "=f"(c[2]), "=f"(c[3]) : "=f"(c[0]), "=f"(c[1]), "=f"(c[2]), "=f"(c[3])
...@@ -60,8 +57,7 @@ __device__ inline void mma_sp(const FragB& a_frag0, const FragB& a_frag1, ...@@ -60,8 +57,7 @@ __device__ inline void mma_sp(const FragB& a_frag0, const FragB& a_frag1,
"r"(b[4]), "r"(b[6]), "f"(c[0]), "f"(c[1]), "f"(c[2]), "f"(c[3]), "r"(b[4]), "r"(b[6]), "f"(c[0]), "f"(c[1]), "f"(c[2]), "f"(c[3]),
"r"(e[0])); "r"(e[0]));
asm volatile( asm volatile(
"mma.sp::ordered_metadata.sync.aligned.m16n8k32.row.col.f32.f16.f16." "mma.sp.sync.aligned.m16n8k32.row.col.f32.f16.f16.f32 "
"f32 "
"{%0, %1, %2, %3}, {%4, %5, %6, %7}, {%8, %9, %10,%11}, " "{%0, %1, %2, %3}, {%4, %5, %6, %7}, {%8, %9, %10,%11}, "
"{%12,%13,%14,%15}, %16, 0x1;\n" "{%12,%13,%14,%15}, %16, 0x1;\n"
: "=f"(c[4]), "=f"(c[5]), "=f"(c[6]), "=f"(c[7]) : "=f"(c[4]), "=f"(c[5]), "=f"(c[6]), "=f"(c[7])
......
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