From 3284c88967e76e9702190edbf60acb29ec2ebff0 Mon Sep 17 00:00:00 2001
From: aamir <araihan@ece.ubc.ca>
Date: Tue, 31 Jul 2018 14:21:18 -0700
Subject: added more regression tests

---
 cutlass-example/cutlass_example.cu | 246 +++++++++++++++++++++++++++++++++++++
 cutlass-example/launch.sh          | 121 ++++++++++++++++--
 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 
-- 
cgit v1.3