summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--cutlass-example/cutlass_example.cu246
-rwxr-xr-xcutlass-example/launch.sh121
2 files changed, 354 insertions, 13 deletions
diff --git a/cutlass-example/cutlass_example.cu b/cutlass-example/cutlass_example.cu
index 6371164..0e5722c 100644
--- a/cutlass-example/cutlass_example.cu
+++ b/cutlass-example/cutlass_example.cu
@@ -72,6 +72,64 @@ int main(int argc, char* argv[]) {
run_gemm<WmmaGemmTraits>(16, 16, 32);
#endif
+#ifdef WMMA_GEMM_32x32x32_NT
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+ cutlass::MatrixLayout::kRowMajor,
+ cutlass::Shape<32, 16, 16> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(32, 32, 32);
+#endif
+#ifdef WMMA_GEMM_32x32x32_NN
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+ cutlass::MatrixLayout::kColumnMajor,
+ cutlass::Shape<32, 16, 16> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(32, 32, 32);
+#endif
+#ifdef WMMA_GEMM_32x32x32_TT
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+ cutlass::MatrixLayout::kRowMajor,
+ cutlass::Shape<32, 16, 16> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(32, 32, 32);
+#endif
+#ifdef WMMA_GEMM_32x32x32_TN
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+ cutlass::MatrixLayout::kColumnMajor,
+ cutlass::Shape<32, 16, 16> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(32, 32, 32);
+#endif
+
+#ifdef WMMA_GEMM_128x128x128_NT
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+ cutlass::MatrixLayout::kRowMajor,
+ cutlass::Shape<32, 16, 16> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(128, 128, 128);
+#endif
+#ifdef WMMA_GEMM_128x128x128_NN
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+ cutlass::MatrixLayout::kColumnMajor,
+ cutlass::Shape<32, 16, 16> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(128, 128, 128);
+#endif
+#ifdef WMMA_GEMM_128x128x128_TT
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+ cutlass::MatrixLayout::kRowMajor,
+ cutlass::Shape<32, 16, 16> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(128, 128, 128);
+#endif
+#ifdef WMMA_GEMM_128x128x128_TN
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+ cutlass::MatrixLayout::kColumnMajor,
+ cutlass::Shape<32, 16, 16> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(128, 128, 128);
+#endif
+
#ifdef WMMA_16x16x16_GEMM_256x256x128_NT
typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
cutlass::MatrixLayout::kRowMajor,
@@ -104,6 +162,194 @@ int main(int argc, char* argv[]) {
run_gemm<WmmaGemmTraits>(256, 256, 128);
#endif
+#ifdef WMMA_GEMM_256x256x256_NT
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+ cutlass::MatrixLayout::kRowMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(256, 256, 256);
+#endif
+
+#ifdef WMMA_GEMM_256x256x256_NN
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+ cutlass::MatrixLayout::kColumnMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(256, 256, 256);
+#endif
+
+#ifdef WMMA_GEMM_256x256x256_TT
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+ cutlass::MatrixLayout::kRowMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(256, 256, 256);
+#endif
+
+#ifdef WMMA_GEMM_256x256x256_TN
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+ cutlass::MatrixLayout::kColumnMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(256, 256, 256);
+#endif
+#ifdef WMMA_GEMM_512x512x512_NT
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+ cutlass::MatrixLayout::kRowMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(512, 512, 512);
+#endif
+
+#ifdef WMMA_GEMM_512x512x512_NN
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+ cutlass::MatrixLayout::kColumnMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(512, 512, 512);
+#endif
+
+#ifdef WMMA_GEMM_512x512x512_TT
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+ cutlass::MatrixLayout::kRowMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(512, 512, 512);
+#endif
+
+#ifdef WMMA_GEMM_512x512x512_TN
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+ cutlass::MatrixLayout::kColumnMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(512, 512, 512);
+#endif
+
+#ifdef WMMA_GEMM_768x768x768_NT
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+ cutlass::MatrixLayout::kRowMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(768, 768, 768);
+#endif
+#ifdef WMMA_GEMM_768x768x768_NN
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+ cutlass::MatrixLayout::kColumnMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(768, 768, 768);
+#endif
+
+#ifdef WMMA_GEMM_768x768x768_TT
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+ cutlass::MatrixLayout::kRowMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(768, 768, 768);
+#endif
+
+#ifdef WMMA_GEMM_768x768x768_TN
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+ cutlass::MatrixLayout::kColumnMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(768, 768, 768);
+#endif
+
+
+#ifdef WMMA_GEMM_1024x1024x1024_NT
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+ cutlass::MatrixLayout::kRowMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(1024, 1024, 1024);
+#endif
+
+#ifdef WMMA_GEMM_1024x1024x1024_NN
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+ cutlass::MatrixLayout::kColumnMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(1024, 1024, 1024);
+#endif
+
+#ifdef WMMA_GEMM_1024x1024x1024_TT
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+ cutlass::MatrixLayout::kRowMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(1024, 1024, 1024);
+#endif
+
+#ifdef WMMA_GEMM_1024x1024x1024_TN
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+ cutlass::MatrixLayout::kColumnMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(1024, 1024, 1024);
+#endif
+#ifdef WMMA_GEMM_2048x2048x2048_NT
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+ cutlass::MatrixLayout::kRowMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(2048, 2048, 2048);
+#endif
+
+#ifdef WMMA_GEMM_2048x2048x2048_NN
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+ cutlass::MatrixLayout::kColumnMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(2048, 2048, 2048);
+#endif
+
+#ifdef WMMA_GEMM_2048x2048x2048_TT
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+ cutlass::MatrixLayout::kRowMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(2048, 2048, 2048);
+#endif
+
+#ifdef WMMA_GEMM_2048x2048x2048_TN
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+ cutlass::MatrixLayout::kColumnMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(2048, 2048, 2048);
+#endif
+#ifdef WMMA_GEMM_4096x4096x4096_NT
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+ cutlass::MatrixLayout::kRowMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(4096, 4096, 4096);
+#endif
+
+#ifdef WMMA_GEMM_4096x4096x4096_NN
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+ cutlass::MatrixLayout::kColumnMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(4096, 4096, 4096);
+#endif
+
+#ifdef WMMA_GEMM_4096x4096x4096_TT
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+ cutlass::MatrixLayout::kRowMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(4096, 4096, 4096);
+#endif
+
+#ifdef WMMA_GEMM_4096x4096x4096_TN
+ typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+ cutlass::MatrixLayout::kColumnMajor,
+ cutlass::Shape<64, 128, 128> >
+ WmmaGemmTraits;
+ run_gemm<WmmaGemmTraits>(4096, 4096, 4096);
+#endif
}
diff --git a/cutlass-example/launch.sh b/cutlass-example/launch.sh
index 2cae877..e48f53e 100755
--- a/cutlass-example/launch.sh
+++ b/cutlass-example/launch.sh
@@ -1,39 +1,134 @@
#!/bin/bash
make clean
make CFLAGS=-DWMMA_GEMM_16x16x16_NT
-./myprog>LOG_WMMA_GEMM_16x16x16_NT
+./myprog > LOG_WMMA_GEMM_16x16x16_NT
make clean
make CFLAGS=-DWMMA_GEMM_16x16x32_NT
-./myprog>LOG_WMMA_GEMM_16x16x32_NT
+./myprog > LOG_WMMA_GEMM_16x16x32_NT
make clean
make CFLAGS=-DWMMA_GEMM_16x16x16_NN
-./myprog>LOG_WMMA_GEMM_16x16x16_NN
+./myprog > LOG_WMMA_GEMM_16x16x16_NN
make clean
make CFLAGS=-DWMMA_GEMM_16x16x32_NN
-./myprog>LOG_WMMA_GEMM_16x16x32_NN
+./myprog > LOG_WMMA_GEMM_16x16x32_NN
make clean
make CFLAGS=-DWMMA_GEMM_16x16x16_TT
-./myprog>LOG_WMMA_GEMM_16x16x16_TT
+./myprog > LOG_WMMA_GEMM_16x16x16_TT
make clean
make CFLAGS=-DWMMA_GEMM_16x16x32_TT
-./myprog>LOG_WMMA_GEMM_16x16x32_TT
+./myprog > LOG_WMMA_GEMM_16x16x32_TT
make clean
make CFLAGS=-DWMMA_GEMM_16x16x16_TN
-./myprog>LOG_WMMA_GEMM_16x16x16_TN
+./myprog > LOG_WMMA_GEMM_16x16x16_TN
make clean
make CFLAGS=-DWMMA_GEMM_16x16x32_TN
-./myprog>LOG_WMMA_GEMM_16x16x32_TN
+./myprog > LOG_WMMA_GEMM_16x16x32_TN
make clean
make CFLAGS=-DWMMA_16x16x16_GEMM_256x256x128_NT
-./myprog>LOG_WMMA_16x16x16_GEMM_256x256x128_NT
+./myprog > LOG_WMMA_16x16x16_GEMM_256x256x128_NT
make clean
make CFLAGS=-DWMMA_16x16x16_GEMM_256x256x128_NN
-./myprog>LOG_WMMA_16x16x16_GEMM_256x256x128_NN
+./myprog > LOG_WMMA_16x16x16_GEMM_256x256x128_NN
make clean
make CFLAGS=-DWMMA_16x16x16_GEMM_256x256x128_TT
-./myprog>LOG_WMMA_16x16x16_GEMM_256x256x128_TT
+./myprog > LOG_WMMA_16x16x16_GEMM_256x256x128_TT
make clean
make CFLAGS=-DWMMA_16x16x16_GEMM_256x256x128_TN
-./myprog>LOG_WMMA_16x16x16_GEMM_256x256x128_TN
+./myprog > LOG_WMMA_16x16x16_GEMM_256x256x128_TN
make clean
-grep "Result Verifier" LOG_* | wc -l
+make CFLAGS=-DWMMA_GEMM_256x256x256_NT
+./myprog > LOG_WMMA_GEMM_256x256x256_NT
+make clean
+make CFLAGS=-DWMMA_GEMM_256x256x256_NN
+./myprog > LOG_WMMA_GEMM_256x256x256_NN
+make clean
+make CFLAGS=-DWMMA_GEMM_256x256x256_TT
+./myprog > LOG_WMMA_GEMM_256x256x256_TT
+make clean
+make CFLAGS=-DWMMA_GEMM_256x256x256_TN
+./myprog > LOG_WMMA_GEMM_256x256x256_TN
+make clean
+make CFLAGS=-DWMMA_GEMM_512x512x512_NT
+./myprog > LOG_WMMA_GEMM_512x512x512_NT
+make clean
+make CFLAGS=-DWMMA_GEMM_512x512x512_NN
+./myprog > LOG_WMMA_GEMM_512x512x512_NN
+make clean
+make CFLAGS=-DWMMA_GEMM_512x512x512_TT
+./myprog > LOG_WMMA_GEMM_512x512x512_TT
+make clean
+make CFLAGS=-DWMMA_GEMM_512x512x512_TN
+./myprog > LOG_WMMA_GEMM_512x512x512_TN
+make clean
+make CFLAGS=-DWMMA_GEMM_1024x1024x1024_NT
+./myprog > LOG_WMMA_GEMM_1024x1024x1024_NT
+make clean
+make CFLAGS=-DWMMA_GEMM_1024x1024x1024_NN
+./myprog > LOG_WMMA_GEMM_1024x1024x1024_NN
+make clean
+make CFLAGS=-DWMMA_GEMM_1024x1024x1024_TT
+./myprog > LOG_WMMA_GEMM_1024x1024x1024_TT
+make clean
+make CFLAGS=-DWMMA_GEMM_1024x1024x1024_TN
+./myprog > LOG_WMMA_GEMM_1024x1024x1024_TN
+make clean
+make CFLAGS=-DWMMA_GEMM_2048x2048x2048_NT
+./myprog > LOG_WMMA_GEMM_2048x2048x2048_NT
+make clean
+make CFLAGS=-DWMMA_GEMM_2048x2048x2048_NN
+./myprog > LOG_WMMA_GEMM_2048x2048x2048_NN
+make clean
+make CFLAGS=-DWMMA_GEMM_2048x2048x2048_TT
+./myprog > LOG_WMMA_GEMM_2048x2048x2048_TT
+make clean
+make CFLAGS=-DWMMA_GEMM_2048x2048x2048_TN
+./myprog > LOG_WMMA_GEMM_2048x2048x2048_TN
+make clean
+make CFLAGS=-DWMMA_GEMM_4096x4096x4096_NT
+./myprog > LOG_WMMA_GEMM_4096x4096x4096_NT
+make clean
+make CFLAGS=-DWMMA_GEMM_4096x4096x4096_NN
+./myprog > LOG_WMMA_GEMM_4096x4096x4096_NN
+make clean
+make CFLAGS=-DWMMA_GEMM_4096x4096x4096_TT
+./myprog > LOG_WMMA_GEMM_4096x4096x4096_TT
+make clean
+make CFLAGS=-DWMMA_GEMM_4096x4096x4096_TN
+./myprog > LOG_WMMA_GEMM_4096x4096x4096_TN
+make clean
+make CFLAGS=-DWMMA_GEMM_32x32x32_NT
+./myprog > LOG_WMMA_GEMM_32x32x32_NT
+make clean
+make CFLAGS=-DWMMA_GEMM_32x32x32_NN
+./myprog > LOG_WMMA_GEMM_32x32x32_NN
+make clean
+make CFLAGS=-DWMMA_GEMM_32x32x32_TT
+./myprog > LOG_WMMA_GEMM_32x32x32_TT
+make clean
+make CFLAGS=-DWMMA_GEMM_32x32x32_TN
+./myprog > LOG_WMMA_GEMM_32x32x32_TN
+make clean
+make CFLAGS=-DWMMA_GEMM_128x128x128_NT
+./myprog > LOG_WMMA_GEMM_128x128x128_NT
+make clean
+make CFLAGS=-DWMMA_GEMM_128x128x128_NN
+./myprog > LOG_WMMA_GEMM_128x128x128_NN
+make clean
+make CFLAGS=-DWMMA_GEMM_128x128x128_TT
+./myprog > LOG_WMMA_GEMM_128x128x128_TT
+make clean
+make CFLAGS=-DWMMA_GEMM_128x128x128_TN
+./myprog > LOG_WMMA_GEMM_128x128x128_TN
+make clean
+make CFLAGS=-DWMMA_GEMM_768x768x768_NT
+./myprog > LOG_WMMA_GEMM_768x768x768_NT
+make clean
+make CFLAGS=-DWMMA_GEMM_768x768x768_NN
+./myprog > LOG_WMMA_GEMM_768x768x768_NN
+make clean
+make CFLAGS=-DWMMA_GEMM_768x768x768_TT
+./myprog > LOG_WMMA_GEMM_768x768x768_TT
+make clean
+make CFLAGS=-DWMMA_GEMM_768x768x768_TN
+./myprog > LOG_WMMA_GEMM_768x768x768_TN
+make clean
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 03:59:42 +0000