summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--cutlass-example/Makefile11
-rw-r--r--cutlass-example/cutlass_example.cu96
-rwxr-xr-xcutlass-example/launch.sh39
-rw-r--r--src/cuda-sim/ptx_ir.h12
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;
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 06:14:04 +0000