summaryrefslogtreecommitdiff
path: root/cutlass-example/cutlass_example.cu
diff options
context:
space:
mode:
Diffstat (limited to 'cutlass-example/cutlass_example.cu')
-rw-r--r--cutlass-example/cutlass_example.cu246
1 files changed, 246 insertions, 0 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
}
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 09:14:10 +0000