diff options
Diffstat (limited to 'cutlass-example')
| -rw-r--r-- | cutlass-example/Makefile | 11 | ||||
| -rw-r--r-- | cutlass-example/cutlass_example.cu | 96 | ||||
| -rwxr-xr-x | cutlass-example/launch.sh | 39 |
3 files changed, 143 insertions, 3 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 |