added regression and debugged the failing testcase

4 files changed, 153 insertions, 5 deletions

diff --git a/cutlass-example/Makefile b/cutlass-example/Makefile
index 8156847..f72b732 100644
--- a/cutlass-example/Makefile
+++ b/cutlass-example/Makefile
@@ -34,4 +34,13 @@ INC_DIR = ./
 LDFLAGS = --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -I$(INC_DIR)
 
 myprog: $(obj)
-	$(CC) -o $@ $^ $(LDFLAGS)
+	$(CC) $(CFLAGS) -o $@ $^ $(LDFLAGS)
+
+clean: 
+	rm -rf gemm_tt_*
+	rm -rf host_results_*
+	rm -rf gpgpusim_power_report__*
+	rm -rf _cuobjdump*
+	rm -rf gpgpu_inst_stats.txt
+	rm -rf myprog
+	echo Clean done
diff --git a/cutlass-example/cutlass_example.cu b/cutlass-example/cutlass_example.cu
index d3e8f89..6371164 100644
--- a/cutlass-example/cutlass_example.cu
+++ b/cutlass-example/cutlass_example.cu
@@ -6,11 +6,103 @@
 #include <gemm.h>
 
 int main(int argc, char* argv[]) {
-    typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
-                                        cutlass::MatrixLayout::kRowMajor, 
+
+#ifdef WMMA_GEMM_16x16x16_NT
+  typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+                                        cutlass::MatrixLayout::kRowMajor,
+                                        cutlass::Shape<32, 16, 16> >
+      WmmaGemmTraits;
+  run_gemm<WmmaGemmTraits>(16, 16, 16);
+#endif
+
+#ifdef WMMA_GEMM_16x16x32_NT
+  typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+                                        cutlass::MatrixLayout::kRowMajor,
+                                        cutlass::Shape<32, 16, 16> >
+      WmmaGemmTraits;
+  run_gemm<WmmaGemmTraits>(16, 16, 32);
+#endif
+
+#ifdef WMMA_GEMM_16x16x16_NN
+  typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+                                        cutlass::MatrixLayout::kColumnMajor,
+                                        cutlass::Shape<32, 16, 16> >
+      WmmaGemmTraits;
+  run_gemm<WmmaGemmTraits>(16, 16, 16);
+#endif
+
+
+#ifdef WMMA_GEMM_16x16x32_NN
+  typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+                                        cutlass::MatrixLayout::kColumnMajor,
+                                        cutlass::Shape<32, 16, 16> >
+      WmmaGemmTraits;
+  run_gemm<WmmaGemmTraits>(16, 16, 32);
+#endif
+
+#ifdef WMMA_GEMM_16x16x16_TT
+  typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+                                        cutlass::MatrixLayout::kRowMajor,
+                                        cutlass::Shape<32, 16, 16> >
+      WmmaGemmTraits;
+  run_gemm<WmmaGemmTraits>(16, 16, 16);
+#endif
+
+#ifdef WMMA_GEMM_16x16x32_TT
+  typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+                                        cutlass::MatrixLayout::kRowMajor,
+                                        cutlass::Shape<32, 16, 16> >
+      WmmaGemmTraits;
+  run_gemm<WmmaGemmTraits>(16, 16, 32);
+#endif
+
+#ifdef WMMA_GEMM_16x16x16_TN
+  typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+                                        cutlass::MatrixLayout::kColumnMajor,
+                                        cutlass::Shape<32, 16, 16> >
+      WmmaGemmTraits;
+  run_gemm<WmmaGemmTraits>(16, 16, 16);
+#endif
+
+#ifdef WMMA_GEMM_16x16x32_TN
+  typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+                                        cutlass::MatrixLayout::kColumnMajor,
+                                        cutlass::Shape<32, 16, 16> >
+      WmmaGemmTraits;
+  run_gemm<WmmaGemmTraits>(16, 16, 32);
+#endif
+
+#ifdef WMMA_16x16x16_GEMM_256x256x128_NT
+  typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+                                        cutlass::MatrixLayout::kRowMajor,
+                                        cutlass::Shape<32, 128, 128> >
+      WmmaGemmTraits;
+  run_gemm<WmmaGemmTraits>(256, 256, 128);
+#endif
+
+#ifdef WMMA_16x16x16_GEMM_256x256x128_NN
+  typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kColumnMajor,
+                                        cutlass::MatrixLayout::kColumnMajor,
+                                        cutlass::Shape<32, 128, 128> >
+      WmmaGemmTraits;
+  run_gemm<WmmaGemmTraits>(256, 256, 128);
+#endif
+
+#ifdef WMMA_16x16x16_GEMM_256x256x128_TT
+  typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+                                        cutlass::MatrixLayout::kRowMajor,
+                                        cutlass::Shape<32, 128, 128> >
+      WmmaGemmTraits;
+  run_gemm<WmmaGemmTraits>(256, 256, 128);
+#endif
+
+#ifdef WMMA_16x16x16_GEMM_256x256x128_TN
+  typedef cutlass::gemm::WmmaGemmTraits<cutlass::MatrixLayout::kRowMajor,
+                                        cutlass::MatrixLayout::kColumnMajor,
                                         cutlass::Shape<32, 128, 128> >
       WmmaGemmTraits;
   run_gemm<WmmaGemmTraits>(256, 256, 128);
+#endif
 
 }
 
diff --git a/cutlass-example/launch.sh b/cutlass-example/launch.sh
new file mode 100755
index 0000000..2cae877
--- /dev/null
+++ b/cutlass-example/launch.sh
@@ -0,0 +1,39 @@
+#!/bin/bash
+make clean
+make CFLAGS=-DWMMA_GEMM_16x16x16_NT
+./myprog>LOG_WMMA_GEMM_16x16x16_NT
+make clean
+make CFLAGS=-DWMMA_GEMM_16x16x32_NT
+./myprog>LOG_WMMA_GEMM_16x16x32_NT
+make clean
+make CFLAGS=-DWMMA_GEMM_16x16x16_NN
+./myprog>LOG_WMMA_GEMM_16x16x16_NN
+make clean
+make CFLAGS=-DWMMA_GEMM_16x16x32_NN
+./myprog>LOG_WMMA_GEMM_16x16x32_NN
+make clean
+make CFLAGS=-DWMMA_GEMM_16x16x16_TT
+./myprog>LOG_WMMA_GEMM_16x16x16_TT
+make clean
+make CFLAGS=-DWMMA_GEMM_16x16x32_TT
+./myprog>LOG_WMMA_GEMM_16x16x32_TT
+make clean
+make CFLAGS=-DWMMA_GEMM_16x16x16_TN
+./myprog>LOG_WMMA_GEMM_16x16x16_TN
+make clean
+make CFLAGS=-DWMMA_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
+make clean
+make CFLAGS=-DWMMA_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
+make clean
+make CFLAGS=-DWMMA_16x16x16_GEMM_256x256x128_TN
+./myprog>LOG_WMMA_16x16x16_GEMM_256x256x128_TN
+make clean
+grep "Result Verifier" LOG_* | wc -l
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index 62d7c7c..cff312b 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -434,11 +434,15 @@ public:
       m_uid = get_uid();
       m_valid = true;
       m_type = memory_t;
-      m_value.m_vector_symbolic = new const symbol*[4];
+      m_value.m_vector_symbolic = new const symbol*[8];
       m_value.m_vector_symbolic[0] = addr1;
       m_value.m_vector_symbolic[1] = addr2;
       m_value.m_vector_symbolic[2] = NULL;
       m_value.m_vector_symbolic[3] = NULL;
+      m_value.m_vector_symbolic[4] = NULL;
+      m_value.m_vector_symbolic[5] = NULL;
+      m_value.m_vector_symbolic[6] = NULL;
+      m_value.m_vector_symbolic[7] = NULL;
       m_addr_offset = 0;
       m_vector = false;
       m_neg_pred = false;
@@ -572,11 +576,15 @@ public:
       m_valid = true;
       m_vector = true;
       m_type = vector_t;
-      m_value.m_vector_symbolic = new const symbol*[4];
+      m_value.m_vector_symbolic = new const symbol*[8];
       m_value.m_vector_symbolic[0] = s1;
       m_value.m_vector_symbolic[1] = s2;
       m_value.m_vector_symbolic[2] = s3;
       m_value.m_vector_symbolic[3] = s4;
+      m_value.m_vector_symbolic[4] = NULL;
+      m_value.m_vector_symbolic[5] = NULL;
+      m_value.m_vector_symbolic[6] = NULL;
+      m_value.m_vector_symbolic[7] = NULL;
       m_addr_offset = 0;
       m_neg_pred = false;
       m_is_return_var = false;