Add TMA unittest app (#58)

* Add tma unittest

* add regular load to TMA benchmark

* make the regular load to have same access pattern as TMA load

* avoid compiler optimization

* move cuda mempcy to be before kernel launch

* add iteration count for tma ubench

* minor formatting

* move tma to ubench folder

* make setup script works with zsh

* fix the issue that ubench all return 1 even without issue

* add a sample test kernel for mbarrier PTX mapping to SASS

* update gitignore

* add gmma kernels for latency measurement

* increase iter to 1024

* add missed kernels

* add maxflops for gmma

* update block size

* update prints for MaxFlops_gmma

* fix a bug

* fix include after updating it

* fix for cpp and c source

* fix compile

* fix for pattern matching

* fix compilation for mbarrier

* Fix makefile for tma app

* generate SASS and PTX for TMA and GMMA workloads

* update makefile to force PTX to be embedded in final fat bin

* change naming

* comment out parboil as it is using python2

* Add GPU ubench to clean target

* Use dynamic linking by default for GPU apps

* Add test binaries for GMMA instruction

* Checkout CUTLASS during ci

* Use type to specify gmma ubench iteration count and update test code

* Fix typos

* Update Makefiles and setup_environment to use C++17 standard

* missed one rename

* Remove unused clean target and tma build steps from Makefile

---------

Co-authored-by: JRPAN <25518778+JRPan@users.noreply.github.com>

Author: William-An

.github/workflows/test-build.yml

@@ -30,6 +30,7 @@ jobs:
         run: |
              git config --global --add safe.directory /__w/gpu-app-collection/gpu-app-collection
              git submodule update --init -- src/cuda/cuda-samples
+             git submodule update --init -- src/cuda/cutlass-bench
              /bin/bash test-build.sh ci
 
       - name: Print Successful Apps

.gitignore

@@ -15,4 +15,13 @@ src/cuda/rodinia/3.1/cuda/particlefilter/particlefilter_naive
 src/cuda/rodinia/3.1/cuda/pathfinder/pathfinder
 4.2
 .venv/
-__pycache__/
+__pycache__/
+compile_commands.json
+.cache/
+tmp/
+
+# Ignoring files without extension (but keep Makefile and files with extensions)
+src/cuda/GPU_Microbenchmark/ubench/**/*
+!src/cuda/GPU_Microbenchmark/ubench/**/*/
+!src/cuda/GPU_Microbenchmark/ubench/**/*.*
+!src/cuda/GPU_Microbenchmark/ubench/**/Makefile

src/Makefile

@@ -109,8 +109,6 @@ GPU_Microbenchmark:
 	mkdir -p $(BINDIR)/$(BINSUBDIR)/
 	$(SETENV) $(MAKE) $(MAKE_ARGS) -C cuda/GPU_Microbenchmark
 	mv cuda/GPU_Microbenchmark/bin/* $(BINDIR)/$(BINSUBDIR)/
-clean_GPU_Microbenchmark:
-	find cuda/GPU_Microbenchmark/ubench -type f -executable -delete
 
 
 Deepbench_nvidia:
@@ -540,22 +538,23 @@ clean_shoc:
 	cd cuda/shoc-master/; $(MAKE) clean; $(MAKE) distclean
 
 clean_parboil:
-	$(SETENV) cd cuda/parboil; ./parboil clean cutcp cuda
-	$(SETENV) cd cuda/parboil; ./parboil clean bfs cuda
-	$(SETENV) cd cuda/parboil; ./parboil clean histo cuda
-	$(SETENV) cd cuda/parboil; ./parboil clean lbm cuda
-	$(SETENV) cd cuda/parboil; ./parboil clean mri-gridding cuda
-	$(SETENV) cd cuda/parboil; ./parboil clean mri-q cuda
-	$(SETENV) cd cuda/parboil; ./parboil clean sad cuda
-	$(SETENV) cd cuda/parboil; ./parboil clean sgemm cuda
-	$(SETENV) cd cuda/parboil; ./parboil clean spmv cuda
-	$(SETENV) cd cuda/parboil; ./parboil clean stencil cuda
-	$(SETENV) cd cuda/parboil; ./parboil clean tpacf cuda
+# Commented out as parboil uses Python2
+# $(SETENV) cd cuda/parboil; ./parboil clean cutcp cuda
+# $(SETENV) cd cuda/parboil; ./parboil clean bfs cuda
+# $(SETENV) cd cuda/parboil; ./parboil clean histo cuda
+# $(SETENV) cd cuda/parboil; ./parboil clean lbm cuda
+# $(SETENV) cd cuda/parboil; ./parboil clean mri-gridding cuda
+# $(SETENV) cd cuda/parboil; ./parboil clean mri-q cuda
+# $(SETENV) cd cuda/parboil; ./parboil clean sad cuda
+# $(SETENV) cd cuda/parboil; ./parboil clean sgemm cuda
+# $(SETENV) cd cuda/parboil; ./parboil clean spmv cuda
+# $(SETENV) cd cuda/parboil; ./parboil clean stencil cuda
+# $(SETENV) cd cuda/parboil; ./parboil clean tpacf cuda
 
 clean_parboil_hw_power:
-	$(SETENV) cd cuda/parboil; ./parboil clean mri-q cuda_k1
-	$(SETENV) cd cuda/parboil; ./parboil clean sad cuda_k1
-	$(SETENV) cd cuda/parboil; ./parboil clean sgemm cuda_k1
+# $(SETENV) cd cuda/parboil; ./parboil clean mri-q cuda_k1
+# $(SETENV) cd cuda/parboil; ./parboil clean sad cuda_k1
+# $(SETENV) cd cuda/parboil; ./parboil clean sgemm cuda_k1
 
 clean_lonestargpu-2.0:
 	$(setenv) $(MAKE) $(make_args) noinline=$(noinline) -C cuda/lonestargpu-2.0 clean
@@ -698,4 +697,7 @@ clean_cuda_samples:
 	$(MAKE) clean -C ./cuda/cuda-samples/build
 
 clean_huggingface:
-	rm -rf $(BINDIR)/$(BINSUBDIR)/huggingface
+	rm -rf $(BINDIR)/$(BINSUBDIR)/huggingface
+
+clean_GPU_Microbenchmark:
+	$(MAKE) clean -C ./cuda/GPU_Microbenchmark

src/cuda/GPU_Microbenchmark/.gitignore

@@ -1,3 +1,5 @@
 bin/
 *.o
 *.out
+*.a
+*.ptx

src/cuda/GPU_Microbenchmark/common/common.mk

@@ -9,16 +9,42 @@ CC := nvcc
 
 LIB :=
 
-release:
-	$(CC) $(NVCC_FLAGS) $(CUOPTS) $(SRC) -o $(EXE) -I$(INCLUDE) -L$(LIB) -lcudart
+# Generate object file list from SRC (for parallel compilation)
+CUDA_SRC_FILES := $(filter %.cu, $(SRC))
+CPP_SRC_FILES := $(filter %.cpp, $(SRC))
+C_SRC_FILES := $(filter %.c, $(SRC))
+
+# To preserve PTX in multi-step compilation, we have to compile the CUDA source files to .a files
+CUDA_LIB_FILES := $(CUDA_SRC_FILES:.cu=.a)
+
+# Host side source files
+CPP_OBJECT_FILES := $(CPP_SRC_FILES:.cpp=.o)
+C_OBJECT_FILES := $(C_SRC_FILES:.c=.o)
+OBJECT_FILES := $(CPP_OBJECT_FILES) $(C_OBJECT_FILES)
+
+# If multiple source files are provided, compile them separately and link
+# To preserve PTX in final binary: First create static library, then link to executable
+# This avoids nvlink stripping PTX during device linking
+release: $(CUDA_LIB_FILES) $(OBJECT_FILES)
+	$(CC) $(NVCC_FLAGS) $^ -o $(EXE) -L$(LIB) -lcudart --cudart shared
 	mv $(EXE) $(BIN_DIR)
 
+# Pattern rule for compiling individual .cu files to .o files
+%.a: %.cu
+	$(CC) $(NVCC_FLAGS) $(INCLUDE) $(CUOPTS) --lib $< -o $@
+
+%.o: %.cpp
+	$(CC) $(NVCC_FLAGS) $(INCLUDE) $(CUOPTS) -dc $< -o $@
+
+%.o: %.c
+	$(CC) $(NVCC_FLAGS) $(INCLUDE) $(CUOPTS) -dc $< -o $@
+
 tuner:
-	$(CC) $(NVCC_FLAGS) $(CUOPTS) -DTUNER $(SRC) -o $(EXE) -I$(INCLUDE) -L$(LIB) -lcudart
+	$(CC) $(NVCC_FLAGS) $(CUOPTS) -DTUNER $(SRC) -o $(EXE) $(INCLUDE) -L$(LIB) -lcudart --cudart shared
 	mv $(EXE) $(BIN_DIR)
 
 clean:
-	rm -f *.o; rm -f $(EXE)
+	rm -f *.o *.a *.ptx *.sass $(OBJECTS) $(CUDA_LIB_FILES); rm -f $(EXE) $(LIB_FILE)
 
 run:
 	./$(EXE)
@@ -36,7 +62,7 @@ nvsight:
 	nv-nsight-cu-cli --metrics gpc__cycles_elapsed.avg,sm__cycles_elapsed.sum,smsp__inst_executed.sum,sm__warps_active.avg.pct_of_peak_sustained_active,l1tex__t_sectors_pipe_lsu_mem_global_op_ld_lookup_hit.sum,l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,l1tex__t_sectors_pipe_lsu_mem_global_op_st_lookup_hit.sum,l1tex__t_sectors_pipe_lsu_mem_global_op_st.sum,lts__t_sectors_srcunit_tex_op_read.sum,lts__t_sectors_srcunit_tex_op_write.sum,lts__t_sectors_srcunit_tex_op_read_lookup_hit.sum,lts__t_sectors_srcunit_tex_op_write_lookup_hit.sum,lts__t_sector_op_read_hit_rate.pct,lts__t_sector_op_write_hit_rate.pct,lts__t_sectors_srcunit_tex_op_read.sum.per_second,dram__sectors_read.sum,dram__sectors_write.sum,dram__bytes_read.sum  --csv --page raw ./$(EXE) | tee nsight.csv
 
 ptx:
-	cuobjdump -ptx ./$(EXE)  tee ptx.txt
+	cuobjdump -ptx ./$(EXE) | tee $(EXE).ptx
 
 sass:
-	cuobjdump -sass ./$(EXE)  tee sass.txt
+	cuobjdump -sass ./$(EXE) | tee $(EXE).sass

src/cuda/GPU_Microbenchmark/hw_def/common/common.h

@@ -22,7 +22,7 @@ enum dram_model { GDDR5 = 1, GDDR5X = 2, GDDR6 = 3, HBM = 4 };
 
 // source:
 // https://stackoverflow.com/questions/466204/rounding-up-to-next-power-of-2
-unsigned round_up_2n(unsigned v) {
+inline unsigned round_up_2n(unsigned v) {
   v--;
   v |= v >> 1;
   v |= v >> 2;
@@ -34,9 +34,9 @@ unsigned round_up_2n(unsigned v) {
   return v;
 }
 
-unsigned round_up_2n(float n) { return round_up_2n((unsigned)ceil(n)); }
+inline unsigned round_up_2n(float n) { return round_up_2n((unsigned)ceil(n)); }
 
-bool isPowerOfTwo(int n) {
+inline bool isPowerOfTwo(int n) {
   if (n == 0)
     return false;
 
@@ -51,12 +51,12 @@ static const unsigned dram_model_burst_length[] = {0, 8, 8, 16, 2};
 static const unsigned dram_model_freq_ratio[] = {0, 4, 4, 4, 2};
 // atom size =
 // dram_model_channel_width*dram_model_mem_per_ctrlr*dram_model_burst_length
-unsigned get_atom_size_inByte(enum dram_model model) {
+inline unsigned get_atom_size_inByte(enum dram_model model) {
   return (dram_model_bus_width[model] / 8) * dram_model_mem_per_ctrlr[model] *
          dram_model_burst_length[model];
 }
 // CCD = dram_model_burst_length/dram_model_freq_ratio
-unsigned get_adjusted_CCD(enum dram_model model) {
+inline unsigned get_adjusted_CCD(enum dram_model model) {
   assert(dram_model_burst_length[model] % dram_model_freq_ratio[model] == 0);
   return dram_model_burst_length[model] / dram_model_freq_ratio[model];
 }

src/cuda/GPU_Microbenchmark/hw_def/common/gpuConfig.h

@@ -45,7 +45,7 @@ struct GpuConfig
     unsigned FBP_COUNT = 0;           // Frame Buffer Partitions
     unsigned L2_BANKS = 0;            // L2 Cache Banks (LTCs)
 };
-GpuConfig config;
+inline GpuConfig config;
 // Parses short flags like --sm 80 into a GpuConfig object
 inline void parseGpuConfigArgs(int argc, char *argv[])
 {
@@ -158,7 +158,7 @@ inline void gpuAssert(cudaError_t code, const char *file, int line,
     }
 }
 
-cudaDeviceProp deviceProp;
+inline cudaDeviceProp deviceProp;
 
 // NVIDIA RM API defines
 #define NV_IOCTL_MAGIC 'F'
@@ -237,7 +237,7 @@ inline unsigned queryGrInfo(uint32_t info_index)
     return result;
 }
 
-unsigned intilizeDeviceProp(unsigned deviceID, int argc, char *argv[])
+inline unsigned initializeDeviceProp(unsigned deviceID, int argc, char *argv[])
 {
     // Check if running in GPGPU-Sim by looking for gpgpusim.config
     std::ifstream configFile("gpgpusim.config");

src/cuda/GPU_Microbenchmark/ubench/atomics/Atomic_add_bw/atomic_add_bw.cu

@@ -51,7 +51,7 @@ __global__ void atomic_bw(uint64_t *startClk, uint64_t *stopClk, T *data1,
 int main(int argc, char *argv[])
 {
 
-  intilizeDeviceProp(0, argc, argv);
+  initializeDeviceProp(0, argc, argv);
 
   // Parse command line arguments for --fast flag
   uint32_t repeat_times = 2048; // default

src/cuda/GPU_Microbenchmark/ubench/atomics/Atomic_add_bw_conflict/atomic_add_bw_conflict.cu

@@ -40,7 +40,7 @@ __global__ void atomic_bw(uint32_t *startClk, uint32_t *stopClk, T *data1,
 int main(int argc, char *argv[])
 {
 
-  intilizeDeviceProp(0, argc, argv);
+  initializeDeviceProp(0, argc, argv);
   config.BLOCKS_NUM = config.SM_NUMBER * 2;
   config.TOTAL_THREADS = config.THREADS_PER_BLOCK * config.BLOCKS_NUM;
 

src/cuda/GPU_Microbenchmark/ubench/atomics/Atomic_add_lat/atomic_add_lat.cu

@@ -42,7 +42,7 @@ __global__ void atmoic_latency(uint32_t *startClk, uint32_t *stopClk, T *data1,
 int main(int argc, char *argv[])
 {
 
-  intilizeDeviceProp(0, argc, argv);
+  initializeDeviceProp(0, argc, argv);
 
   config.THREADS_PER_BLOCK = 1;
   config.THREADS_PER_SM = 1;