Add test for matmul

.gitignore

@@ -4,4 +4,5 @@
 *.prof
 test/*.bin
 test/test_add
-test/test_fma
+test/test_fma
+test/test_matmul

test/Makefile

@@ -4,21 +4,29 @@ CFLAGS = -DTEST
 
 TEST_ADD_CU = test_add.cu
 TEST_FMA_CU = test_fma.cu
-KERNEL_CU = ../matmulf8_kernel.cu
+TEST_MATMUL_CU = test_matmul.cu
+KERNEL_CU = ../matmulf8_kernel.cu ../matmulf8.cu ../load_core.cu
 OUTPUT_ADD = test_add
 OUTPUT_FMA = test_fma
+OUTPUT_MATMUL = test_matmul
 BIN_FILE = test_add.bin test_fma.bin
 
-OUTPUT_EXE = $(OUTPUT_ADD) $(OUTPUT_FMA)
+OUTPUT_EXE = $(OUTPUT_ADD) $(OUTPUT_FMA) $(OUTPUT_MATMUL)
 
-all: $(OUTPUT_EXE) $(BIN_FILE)
+test: $(OUTPUT_EXE) $(BIN_FILE)
+	./$(OUTPUT_ADD)
+	./$(OUTPUT_FMA)
+	./$(OUTPUT_MATMUL)
 
 $(OUTPUT_ADD): $(TEST_ADD_CU) $(KERNEL_CU)
 	$(NVCC) $(CFLAGS) $(TEST_ADD_CU) $(KERNEL_CU) -o $(OUTPUT_ADD)
 
 $(OUTPUT_FMA): $(TEST_FMA_CU) $(KERNEL_CU)
 	$(NVCC) $(CFLAGS) $(TEST_FMA_CU) $(KERNEL_CU) -o $(OUTPUT_FMA)
 
+$(OUTPUT_MATMUL): $(TEST_MATMUL_CU) $(KERNEL_CU)
+	$(NVCC) $(CFLAGS) $(TEST_MATMUL_CU) $(KERNEL_CU) -o $(OUTPUT_MATMUL)
+
 $(BIN_FILE): gen_test_bin.py
 	python gen_test_bin.py
 

test/gen_test_bin.py

@@ -34,4 +34,5 @@ def gen_fma_test_bin():
 
 if __name__ == '__main__':
     gen_add_test_bin()
-    gen_fma_test_bin()
+    gen_fma_test_bin()
+    print(float8_e5m2(128.0).tobytes().hex())

test/test_matmul.cu

@@ -0,0 +1,75 @@
+#include "../load_core.cuh"
+#include "../matmulf8.cuh"
+#include <cstdio>
+#include <random>
+#include "test_matmul.h"
+
+void test_cross_line(u_int8_t *A, u_int8_t *B, u_int8_t *C, int n, int m, int p,
+                     int target_x, int target_y,
+                     int *acore, int *mcore){
+    int success = 1;
+    for(int i = 0; i < n * m; i++) {
+        A[i] = 0;
+    }
+    for(int i = 0; i < m * p; i++) {
+        B[i] = 0;
+    }
+    for (int i = 0; i < m; i++) {
+        A[target_x * m + i] = 0x3c;
+    }
+    for (int i = 0; i < m; i++) {
+        // B is transposed
+        B[target_y * m + i] = 0x3c;
+    }
+    for (int i = 0; i < n * p; i++) {
+        C[i] = 0xff;
+    }
+
+    float t = matmul((int *)A, (int *)B, (int *)C, n, m, p, acore, mcore);
+
+    for (int i = 0; i < n; i++) {
+        for (int j = 0; j < p; j++) {
+            if (i == target_x && j == target_y){
+                // 0x58 is float8_e5m2 of 128.0
+                if(C[i * p + j] != 0x58){
+                    printf("Target(%d, %d) expected 0x58(128.0), got: %x\n", i, j, C[i * p + j]);
+                    success = 0;
+                }
+            }else{
+                if(C[i * p + j] != 0){
+                    printf("(%d, %d) expected 0, got: %x", i, j, C[i * p + j]);
+                    success = 0;
+                }
+            }
+        }
+    }
+    if(!success){
+        printf("Test failed for (%d, %d)\n", target_x, target_y);
+    }else{
+        printf("Test passed for (%d, %d)\n", target_x, target_y);
+    }
+}
+
+int main(){
+    cudaSetDevice(0);
+
+    int* acore = load_core("../apdcore.bin");
+    int* mcore = load_core("../mltcore.bin");
+    int n = 128, m = 128, p = 128;
+    u_int8_t *A, *B, *C;
+
+    cudaMallocHost(&A, n * m);
+    cudaMallocHost(&B, m * p);
+    cudaMallocHost(&C, n * p);
+    std::mt19937 gen(std::random_device{}());
+    std::uniform_int_distribution<int> dis(0, 128);
+
+    for(int t = 0; t < 10; t++){
+        int x = dis(gen), y = dis(gen);
+        test_cross_line(A, B, C, n, m, p, x, y, acore, mcore);
+    }
+
+    cudaFreeHost(A);
+    cudaFreeHost(B);
+    cudaFreeHost(C);
+}

test/test_matmul.h

@@ -0,0 +1,2 @@
+#pragma once
+void test_cross_line(int m, u_int8_t *A, u_int8_t *B, int n, int p, u_int8_t *C, int *acore, int *mcore);