Fix: Invoke f16f32 in WGMMA

ashvardanian · ashvardanian · commit 44234214fcd2 · 2025-02-10T23:17:47.000Z
diff --git a/less_slow.cu b/less_slow.cu
@@ -603,7 +603,7 @@ __global__ void tops_f16f32_sm90wgmma_64x256x16_loop128_cuda_kernel() {
     std::uint64_t b_descriptor = wgmma_descriptor((std::uint64_t)b_shared, 128 * 256 / 8, 128, 0, 0);
     wgmma_fence();
     for (int i = 0; i != 128; ++i) {
-        wgmma_bf16f32_64x256x16(c_registers, a_descriptor, b_descriptor);
+        wgmma_f16f32_64x256x16(c_registers, a_descriptor, b_descriptor);
         wgmma_commit_group();
     }
     wgmma_sync_group();

Original file line number	Diff line number	Diff line change
`@@ -603,7 +603,7 @@ __global__ void tops_f16f32_sm90wgmma_64x256x16_loop128_cuda_kernel() {`
`603`	`603`	`std::uint64_t b_descriptor = wgmma_descriptor((std::uint64_t)b_shared, 128 * 256 / 8, 128, 0, 0);`
`604`	`604`	`wgmma_fence();`
`605`	`605`	`for (int i = 0; i != 128; ++i) {`
`606`		`- wgmma_bf16f32_64x256x16(c_registers, a_descriptor, b_descriptor);`
	`606`	`+ wgmma_f16f32_64x256x16(c_registers, a_descriptor, b_descriptor);`
`607`	`607`	`wgmma_commit_group();`
`608`	`608`	`}`
`609`	`609`	`wgmma_sync_group();`