diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml deleted file mode 100644 index 2f8120a..0000000 --- a/.github/workflows/ci.yml +++ /dev/null @@ -1,34 +0,0 @@ -name: CI - -on: - push: - branches: main - paths: ['**.cu','**.c','**.cpp', '**.h', '**CMakeLists.txt'] - pull_request: - branches: main - paths: ['**.cu','**.c','**.cpp', '**.h', '**CMakeLists.txt'] - -jobs: - build-and-test: - runs-on: ubuntu-latest - - steps: - - name: Checkout code - uses: actions/checkout@v4 - - - name: Setup python - uses: actions/setup-python@v5 - with: - python-version: '3.10' - - - name: Install dependencies - run: | - pip install pandas - - - name: Build project - run: | - make build - - - name: Run test suite - run: | - make test \ No newline at end of file diff --git a/.github/workflows/cpp-linter.yml b/.github/workflows/cpp-linter.yml index 3ec0fe6..8aec3f9 100644 --- a/.github/workflows/cpp-linter.yml +++ b/.github/workflows/cpp-linter.yml @@ -1,10 +1,10 @@ name: cpp-linter on: pull_request: - branches: main + branches: cuda-staging paths: ['**.cu','**.cpp','**.c', '**.h', '**CMakeLists.txt'] push: - branches: main + branches: cuda-staging paths: ['**.cu','**.cpp','**.c', '**.h', '**CMakeLists.txt'] permissions: diff --git a/include/matrix.h b/include/matrix.h deleted file mode 100644 index 7176922..0000000 --- a/include/matrix.h +++ /dev/null @@ -1,19 +0,0 @@ -#pragma once - -typedef struct { - int rows; - int cols; - float* data; // array -} matrix; -void initmalloc(matrix* d_mat, matrix* h_mat, int rows, int cols); -void dealloc(matrix* d_mat); - -matrix* new_matrix(int rows, int cols); - -__global__ void matrix_mul(matrix* a, matrix* b, matrix* result); - -__global__ void matrix_add(matrix* a, matrix* b); - -__global__ void relu(matrix* a); - -__global__ void softmax(matrix* a); diff --git a/src/main.cu b/src/main.cu index 5079299..1a12d40 100644 --- a/src/main.cu +++ b/src/main.cu @@ -1,4 +1,4 @@ -#include "../include/matrix.h" +#include "matrix.cuh" #include #include #include @@ -21,9 +21,11 @@ matrix* weights[NUM_LAYERS]; matrix* biases[NUM_LAYERS]; // device weights and biases; -matrix* d_weights; -matrix* d_biases; -matrix* d_input; +matrix* d_weights[7]; +matrix* d_biases[7]; +matrix** d_inputs; + +int* results; char letters[52] = {'A', 'a', 'B', 'b', 'C', 'c', 'D', 'd', 'E', 'e', 'F', 'f', 'G', 'g', 'H', 'h', 'I', 'i', 'J', 'j', 'K', 'k', 'L', 'l', 'M', 'm', 'N', 'n', 'O', 'o', 'P', 'p', 'Q', 'q', 'R', 'r', @@ -101,91 +103,70 @@ void read_tensor(matrix* a, const char* fileName) { } void propagate_fwd(matrix* weights, matrix* input_layer, matrix* output_layer, matrix* biases) { - matrix_mul<<<1, 1>>>(weights, input_layer, output_layer); + matrix_mul<<<1, 1>>>(weights->data, input_layer->data, output_layer->data, weights->rows, weights->cols); cudaDeviceSynchronize(); - matrix_add<<<1, 1>>>(output_layer, biases); + matrix_add<<<1, 1>>>(output_layer->data, biases->data, biases->rows); cudaDeviceSynchronize(); } -__global__ void get_max(matrix* a, int* d_int) { - int idx = 0; - float res = a->data[0]; - for (int i = 0; i < a->rows; i++) { - if (res < a->data[i]) { - res = a->data[i]; - idx = i; - } - } - *d_int = idx; -} - int infer(matrix* d_input) { - matrix* mdl_layers[NUM_LAYERS]; - matrix* d_mdl_layers; - - mdl_layers[0] = new_matrix(98, 1); - mdl_layers[1] = new_matrix(65, 1); - mdl_layers[2] = new_matrix(50, 1); - mdl_layers[3] = new_matrix(30, 1); - mdl_layers[4] = new_matrix(25, 1); - mdl_layers[5] = new_matrix(40, 1); - mdl_layers[6] = new_matrix(52, 1); - - CUDA_CHECK(cudaMalloc(&d_mdl_layers, NUM_LAYERS * sizeof(matrix))); - - initmalloc(&d_mdl_layers[0], mdl_layers[0], 98, 1); - initmalloc(&d_mdl_layers[1], mdl_layers[1], 65, 1); - initmalloc(&d_mdl_layers[2], mdl_layers[2], 50, 1); - initmalloc(&d_mdl_layers[3], mdl_layers[3], 30, 1); - initmalloc(&d_mdl_layers[4], mdl_layers[4], 25, 1); - initmalloc(&d_mdl_layers[5], mdl_layers[5], 40, 1); - initmalloc(&d_mdl_layers[6], mdl_layers[6], 52, 1); - - propagate_fwd(&d_weights[0], d_input, &d_mdl_layers[0], &d_biases[0]); - relu<<<1, 1>>>(&d_mdl_layers[0]); + matrix* outputs[2]; + outputs[0] = new_matrix_d(98, 1); + outputs[1] = new_matrix_d(65, 1); + + propagate_fwd(d_weights[0], d_input, outputs[0], d_biases[0]); + relu<<<1, 1>>>(outputs[0]->data, 98); cudaDeviceSynchronize(); - propagate_fwd(&d_weights[1], &d_mdl_layers[0], &d_mdl_layers[1], &d_biases[1]); - relu<<<1, 1>>>(&d_mdl_layers[1]); + propagate_fwd(d_weights[1], outputs[0], outputs[1], d_biases[1]); + cudaMemsetAsync(outputs[0], 0, 50 * sizeof(float)); + relu<<<1, 1>>>(outputs[1]->data, 65); cudaDeviceSynchronize(); - propagate_fwd(&d_weights[2], &d_mdl_layers[1], &d_mdl_layers[2], &d_biases[2]); - relu<<<1, 1>>>(&d_mdl_layers[2]); + propagate_fwd(d_weights[2], outputs[1], outputs[0], d_biases[2]); + cudaMemsetAsync(outputs[1], 0, 30 * sizeof(float)); + relu<<<1, 1>>>(outputs[0]->data, 50); cudaDeviceSynchronize(); - propagate_fwd(&d_weights[3], &d_mdl_layers[2], &d_mdl_layers[3], &d_biases[3]); - relu<<<1, 1>>>(&d_mdl_layers[3]); + propagate_fwd(d_weights[3], outputs[0], outputs[1], d_biases[3]); + cudaMemsetAsync(outputs[0], 0, 25 * sizeof(float)); + relu<<<1, 1>>>(outputs[1]->data, 30); cudaDeviceSynchronize(); - propagate_fwd(&d_weights[4], &d_mdl_layers[3], &d_mdl_layers[4], &d_biases[4]); - relu<<<1, 1>>>(&d_mdl_layers[4]); + propagate_fwd(d_weights[4], outputs[1], outputs[0], d_biases[4]); + cudaMemsetAsync(outputs[1], 0, 40 * sizeof(float)); + relu<<<1, 1>>>(outputs[0]->data, 25); cudaDeviceSynchronize(); - propagate_fwd(&d_weights[5], &d_mdl_layers[4], &d_mdl_layers[5], &d_biases[5]); - relu<<<1, 1>>>(&d_mdl_layers[5]); + propagate_fwd(d_weights[5], outputs[0], outputs[1], d_biases[5]); + cudaMemsetAsync(outputs[0], 0, 52 * sizeof(float)); + relu<<<1, 1>>>(outputs[1]->data, 40); cudaDeviceSynchronize(); - propagate_fwd(&d_weights[6], &d_mdl_layers[5], &d_mdl_layers[6], &d_biases[6]); - softmax<<<1, 1>>>(&d_mdl_layers[6]); + propagate_fwd(d_weights[6], outputs[1], outputs[0], d_biases[6]); + softmax<<<1, 1>>>(outputs[0]->data, 52); cudaDeviceSynchronize(); - int* d_int; - int h_int = 0; + int* d_res; + cudaMalloc(&d_res, sizeof(int)); - CUDA_CHECK(cudaMalloc((void**)&d_int, sizeof(int))); - get_max<<<1, 1>>>(&d_mdl_layers[6], d_int); + argmax<<<1, 1>>>(outputs[0]->data, 52, d_res); cudaDeviceSynchronize(); - CUDA_CHECK(cudaMemcpy(&h_int, d_int, sizeof(int), cudaMemcpyDeviceToHost)); - dealloc(&d_mdl_layers[0]); - dealloc(&d_mdl_layers[1]); - dealloc(&d_mdl_layers[2]); - dealloc(&d_mdl_layers[3]); - dealloc(&d_mdl_layers[4]); - dealloc(&d_mdl_layers[5]); - dealloc(&d_mdl_layers[6]); + cudaFree(outputs[0]->data); + free(outputs[0]); + cudaFree(outputs[1]->data); + free(outputs[1]); - return h_int; + int h_res; + cudaMemcpy(&h_res, d_res, sizeof(int), cudaMemcpyDeviceToHost); + return h_res; +} + +void process(int input_size) { + for (int i = 1; i <= input_size; i++) { + results[i] = infer(d_inputs[i]); + } } int main(int argc, char* argv[]) { @@ -215,29 +196,24 @@ int main(int argc, char* argv[]) { biases[4] = new_matrix(25, 1); biases[5] = new_matrix(40, 1); biases[6] = new_matrix(52, 1); - read_model(argv[1]); - // initialize d_weights struct matrix arr; - CUDA_CHECK(cudaMalloc(&d_weights, NUM_LAYERS * sizeof(matrix))); - CUDA_CHECK(cudaMalloc(&d_biases, NUM_LAYERS * sizeof(matrix))); - - initmalloc(&d_weights[0], weights[0], 98, 225); - initmalloc(&d_weights[1], weights[1], 65, 98); - initmalloc(&d_weights[2], weights[2], 50, 65); - initmalloc(&d_weights[3], weights[3], 30, 50); - initmalloc(&d_weights[4], weights[4], 25, 30); - initmalloc(&d_weights[5], weights[5], 40, 25); - initmalloc(&d_weights[6], weights[6], 52, 40); - initmalloc(&d_biases[0], biases[0], 98, 1); - initmalloc(&d_biases[1], biases[1], 65, 1); - initmalloc(&d_biases[2], biases[2], 50, 1); - initmalloc(&d_biases[3], biases[3], 30, 1); - initmalloc(&d_biases[4], biases[4], 25, 1); - initmalloc(&d_biases[5], biases[5], 40, 1); - initmalloc(&d_biases[6], biases[6], 52, 1); - - // Run program + d_weights[0] = copy_to_device(weights[0]); + d_weights[1] = copy_to_device(weights[1]); + d_weights[2] = copy_to_device(weights[2]); + d_weights[3] = copy_to_device(weights[3]); + d_weights[4] = copy_to_device(weights[4]); + d_weights[5] = copy_to_device(weights[5]); + d_weights[6] = copy_to_device(weights[6]); + + d_biases[0] = copy_to_device(biases[0]); + d_biases[1] = copy_to_device(biases[1]); + d_biases[2] = copy_to_device(biases[2]); + d_biases[3] = copy_to_device(biases[3]); + d_biases[4] = copy_to_device(biases[4]); + d_biases[5] = copy_to_device(biases[5]); + d_biases[6] = copy_to_device(biases[6]); + const char* directory_path = argv[2]; struct dirent* entry; DIR* dir = opendir(directory_path); @@ -253,9 +229,11 @@ int main(int argc, char* argv[]) { size++; } } - int* results = (int*)malloc((size + 1) * sizeof(int)); + + results = (int*)malloc((size + 1) * sizeof(int)); + d_inputs = (matrix**)malloc((size + 1) * sizeof(matrix*)); + dir = opendir(directory_path); - matrix* d_input; while ((entry = readdir(dir)) != NULL) { if (entry->d_type == DT_REG) { @@ -267,11 +245,7 @@ int main(int argc, char* argv[]) { strcat(file_name, "/"); strcat(file_name, entry->d_name); read_tensor(input, file_name); - CUDA_CHECK(cudaMalloc(&d_input, 255 * sizeof(matrix))); - initmalloc(d_input, input, 1, 225); - results[file_num] = infer(d_input); - dealloc(d_input); - + d_inputs[file_num] = copy_to_device(input); free(input); } } @@ -280,6 +254,9 @@ int main(int argc, char* argv[]) { free(file_num_str); closedir(dir); + // Process + process(size); + // Write to csv file FILE* csv_file = fopen("results.csv", "w+"); fprintf(csv_file, "image_number, guess\n"); diff --git a/src/matrix.cu b/src/matrix.cu index 3069f38..e8ed37d 100644 --- a/src/matrix.cu +++ b/src/matrix.cu @@ -1,4 +1,5 @@ -#include "../include/matrix.h" +#include "matrix.cuh" +#include "util.cuh" #include #include #include @@ -13,101 +14,62 @@ matrix* new_matrix(int rows, int cols) { return res; } -__global__ void ptref(matrix* d_mat, float* d_res, int* d_cols, int* d_rows) { - d_mat->data = d_res; - d_mat->cols = *d_cols; - d_mat->rows = *d_rows; -} - -// Allocate device memory for matrix dimensions and data -void initmalloc(matrix* d_mat, matrix* h_mat, int rows, int cols) { - int* d_cols; - int* d_rows; - float* d_res; - cudaMalloc(&d_cols, sizeof(int)); - cudaMalloc(&d_rows, sizeof(int)); - cudaMalloc(&d_res, rows * cols * sizeof(float)); - - cudaMemcpy(d_rows, &rows, sizeof(int), cudaMemcpyHostToDevice); - cudaMemcpy(d_cols, &cols, sizeof(int), cudaMemcpyHostToDevice); - - cudaMemcpy(d_res, h_mat->data, (rows * cols * sizeof(float)), cudaMemcpyHostToDevice); - - // Call kernel to initialize the matrix structure on the device - ptref<<<1, 1>>>(d_mat, d_res, d_cols, d_rows); - cudaDeviceSynchronize(); +matrix* new_matrix_d(int rows, int cols) { + matrix* res = (matrix*)malloc(sizeof(matrix)); + res->rows = rows; + res->cols = cols; + res->cols = cols; + cudaMalloc((void**)&(res->data), rows * cols * sizeof(float)); + return res; } -void dealloc(matrix* d_mat) { - cudaFree(&d_mat->data); - cudaFree(&d_mat->cols); - cudaFree(&d_mat->rows); - - cudaFree(d_mat); +matrix* copy_to_device(matrix* h_mat) { + matrix* res = new_matrix_d(h_mat->rows, h_mat->cols); + CUDA_CHECK(cudaMemcpy(res->data, h_mat->data, h_mat->rows * h_mat->cols * sizeof(float), cudaMemcpyHostToDevice)); + return res; } -// Loop unrolling optimisation with a factor of 8 which should be enough to saturate a Zen3 core -__global__ void matrix_mul(matrix* weights, matrix* inputs, matrix* __restrict__ result) { - - int res_rows = result->rows; - int w_width = weights->cols; - float* w_data = weights->data; - float* i_data = inputs->data; - - int u_limit = w_width - (UNROLL_FACTOR - 1); - - for (int cur_row = 0; cur_row < res_rows; cur_row++) { - float sum0 = 0; - float sum1 = 0; - float sum2 = 0; - float sum3 = 0; - float sum4 = 0; - float sum5 = 0; - float sum6 = 0; - float sum7 = 0; - - int row_offs = cur_row * w_width; - - int k = 0; - for (; k < u_limit; k += UNROLL_FACTOR) { - sum0 += w_data[row_offs + k] * i_data[k]; - sum1 += w_data[row_offs + k + 1] * i_data[k + 1]; - sum2 += w_data[row_offs + k + 2] * i_data[k + 2]; - sum3 += w_data[row_offs + k + 3] * i_data[k + 3]; - sum4 += w_data[row_offs + k + 4] * i_data[k + 4]; - sum5 += w_data[row_offs + k + 5] * i_data[k + 5]; - sum6 += w_data[row_offs + k + 6] * i_data[k + 6]; - sum7 += w_data[row_offs + k + 7] * i_data[k + 7]; - } - - for (; k < w_width; k++) { - sum0 += w_data[row_offs + k] * i_data[k]; +__global__ void matrix_mul(float* weight, float* input, float* result, int w_rows, int w_cols) { + for (int i = 0; i < w_rows; i++) { + float sum = 0; + for (int j = 0; j < w_cols; j++) { + sum += weight[i * w_cols + j] * input[j]; } - - (result->data)[cur_row] = sum0 + sum1 + sum2 + sum3 + sum4 + sum5 + sum6 + sum7; // + sum8 + sum9; + result[i] = sum; } } -__global__ void matrix_add(matrix* a, matrix* b) { - - for (int i = 0; i < a->rows; i++) { - (a->data)[i] += (b->data)[i]; +__global__ void matrix_add(float* a, float* b, int rows) { + for (int i = 0; i < rows; i++) { + a[i] += b[i]; } } -__global__ void relu(matrix* a) { - for (int i = 0; i < a->rows; i++) { - if ((a->data)[i] < (float)0) - (a->data)[i] = (float)0; +__global__ void relu(float* a, int rows) { + for (int i = 0; i < rows; i++) { + if ((a)[i] < (float)0) + (a)[i] = (float)0; } } -__global__ void softmax(matrix* a) { +__global__ void softmax(float* a, int rows) { float res = (float)0; - for (int i = 0; i < a->rows; i++) { - res += exp((a->data)[i]); + for (int i = 0; i < rows; i++) { + res += exp(a[i]); } - for (int i = 0; i < a->rows; i++) { - (a->data)[i] /= res; + for (int i = 0; i < rows; i++) { + a[i] /= res; } } + +__global__ void argmax(float* a, int rows, int* des) { + int res = a[0]; + int idx = 0; + for (int i = 0; i < rows; i++) { + if (res < a[i]) { + res = a[i]; + idx = i; + } + } + *des = idx; +} \ No newline at end of file diff --git a/src/matrix.cuh b/src/matrix.cuh new file mode 100644 index 0000000..83005cf --- /dev/null +++ b/src/matrix.cuh @@ -0,0 +1,23 @@ +#pragma once + +typedef struct { + int rows; + int cols; + float* data; // array +} matrix; + +matrix* new_matrix(int rows, int cols); + +matrix* copy_to_device(matrix* h_mat); + +matrix* new_matrix_d(int rows, int cols); + +__global__ void matrix_mul(float* a, float* b, float* c, int rows, int cols); + +__global__ void matrix_add(float* a, float* b, int rows); + +__global__ void relu(float* a, int rows); + +__global__ void softmax(float* a, int rows); + +__global__ void argmax(float* a, int rows, int* res); \ No newline at end of file diff --git a/src/util.cuh b/src/util.cuh new file mode 100644 index 0000000..a6f988a --- /dev/null +++ b/src/util.cuh @@ -0,0 +1,8 @@ +#pragma once + +#define CUDA_CHECK(call) \ + cudaError_t err = call; \ + if (err != cudaSuccess) { \ + fprintf(stderr, "CUDA error in %s (%s:%d): %s\n", __func__, __FILE__, __LINE__, cudaGetErrorString(err)); \ + exit(EXIT_FAILURE); \ + } \ No newline at end of file