Viết một chương trình CUDA

Kiến trúc bao gồm

• Host: CPU và bộ nhớ của CPU

• Device: GPU và bộ nhớ của GPU

Flow xử lý của CUDA:

• Copy dữ liệu từ bộ nhớ CPU vào bộ nhớ GPU

• Chạy kernel dể tính toán trên data lưu trữ trên bộ nhớ GPU.

• Copy data ngược trở lại từ bộ nhớ GPU về CPU

Trong CUDA, một kernel là một hàm chạy trên GPU và được thực thi song song bởi nhiều luồng (threads). Kernel được viết bằng cú pháp giống C và được gọi từ máy chủ (CPU) để thực thi trên thiết bị (GPU). Mỗi luồng thực thi kernel sẽ thực hiện các lệnh giống nhau nhưng trên dữ liệu khác nhau.

GPU phân bố nhớ thành hai loại

• Global Memory: Tương tự như bộ nhớ CPU

• Shared Memory: Tương tự như CPU Cache, tuy nhiên có thể quản lý trực tiếp với Cuda C kernel

Cộng hai mảng bất kỳ với CUDA.

Code

Đoạn mã này thực hiện các công việc sau:

1. Khởi tạo hai mảng với dữ liệu ngẫu nhiên.

2. Cộng từng phần tử tương ứng của hai mảng và lưu kết quả vào mảng thứ ba.

3. In kết quả của phép cộng và giải phóng bộ nhớ đã cấp phát.

Code bằng C++

#include <stdio.h>     // Thư viện chuẩn để sử dụng các hàm nhập/xuất
#include <stdlib.h>    // Thư viện chuẩn để cấp phát bộ nhớ và các hàm tiện ích khác
#include <string.h>    // Thư viện để thao tác với chuỗi
#include <time.h>      // Thư viện để làm việc với thời gian

// Hàm cộng hai mảng trên máy chủ
void sumArraysOnHost(float *A, float *B, float *C, const int N)
{
    // Vòng lặp để cộng từng phần tử của hai mảng
    for (int i = 0; i < N; i++)
    {
        C[i] = A[i] + B[i];  // Cộng phần tử tương ứng của mảng A và B và lưu vào mảng C
    }
}

// Hàm khởi tạo dữ liệu ngẫu nhiên cho mảng
void initData(float *ip, int size)
{
    // Vòng lặp để khởi tạo từng phần tử trong mảng
    for (int i = 0; i < size; i++)
    {
        ip[i] = (float)(rand() & 0xFF) / 10.0f;  // Gán giá trị ngẫu nhiên cho mỗi phần tử
    }
}

// Hàm chính của chương trình
int main(int argc, char **argv)
{
    int nElem = 1024;  // Số lượng phần tử trong mỗi mảng
    size_t nBytes = nElem * sizeof(float);  // Số byte cần thiết để lưu trữ mảng

    // Cấp phát bộ nhớ cho các mảng h_A, h_B và h_C
    float *h_A = (float *)malloc(nBytes);  // Cấp phát bộ nhớ cho mảng A
    float *h_B = (float *)malloc(nBytes);  // Cấp phát bộ nhớ cho mảng B
    float *h_C = (float *)malloc(nBytes);  // Cấp phát bộ nhớ cho mảng C

    initData(h_A, nElem);  // Khởi tạo dữ liệu cho mảng A
    initData(h_B, nElem);  // Khởi tạo dữ liệu cho mảng B

    sumArraysOnHost(h_A, h_B, h_C, nElem);  // Gọi hàm để cộng hai mảng

    // In kết quả của từng phép cộng
    for (int i = 0; i < nElem; i++)
    {
        printf("%f + %f = %f\n", h_A[i], h_B[i], h_C[i]);  // In ra từng phép cộng
    }

    free(h_A);  // Giải phóng bộ nhớ của mảng A
    free(h_B);  // Giải phóng bộ nhớ của mảng B
    free(h_C);  // Giải phóng bộ nhớ của mảng C

    return 0;  // Kết thúc chương trình
}

Kết quả in ra:

10.300000 + 4.300000 = 14.600000
19.799999 + 14.000000 = 33.799999
10.500000 + 18.200001 = 28.700001
11.500000 + 13.500000 = 25.000000
8.100000 + 2.700000 = 10.800000
25.500000 + 10.000000 = 35.500000
7.400000 + 24.500000 = 31.900000
23.600000 + 9.700000 = 33.299999
4.100000 + 17.100000 = 21.200001
20.500000 + 2.800000 = 23.299999
18.600000 + 23.100000 = 41.700001
17.100000 + 14.400000 = 31.500000
24.200001 + 9.100000 = 33.300003
25.100000 + 14.400000 = 39.500000
22.700001 + 3.000000 = 25.700001
7.000000 + 22.900000 = 29.900000
12.400000 + 0.200000 = 12.599999
19.400000 + 16.799999 = 36.199997
8.400000 + 1.700000 = 10.099999
...

Code bằng CUDA

Đoạn mã này thực hiện các công việc sau:

1. Khởi tạo dữ liệu ngẫu nhiên cho hai mảng trên host và sao chép chúng lên GPU.

2. Thực thi hàm kernel trên GPU để cộng từng phần tử của hai mảng và lưu kết quả vào mảng thứ ba.

3. Sao chép kết quả từ GPU về host, in kết quả và giải phóng bộ nhớ.

%%writefile add_cuda.cu

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>

// Hàm kernel để cộng hai mảng trên GPU
__global__ void sumArraysOnGPU(float *A, float *B, float *C, const int N)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;  // Xác định chỉ số phần tử
    if (i < N)
    {
        C[i] = A[i] + B[i];  // Cộng phần tử tương ứng của mảng A và B và lưu vào C
    }
}

// Hàm khởi tạo dữ liệu ngẫu nhiên cho mảng
void initData(float *ip, int size)
{
    for (int i = 0; i < size; i++)
    {
        ip[i] = (float)(rand() & 0xFF) / 10.0f;  // Gán giá trị ngẫu nhiên cho mỗi phần tử
    }
}

int main(int argc, char **argv)
{
    int nElem = 1024;  // Số lượng phần tử trong mảng
    size_t nBytes = nElem * sizeof(float);  // Số byte cần thiết để lưu trữ mảng

    // Cấp phát bộ nhớ cho các mảng trên host
    float *h_A = (float *)malloc(nBytes);
    float *h_B = (float *)malloc(nBytes);
    float *h_C = (float *)malloc(nBytes);

    // Khởi tạo dữ liệu cho các mảng trên host
    initData(h_A, nElem);
    initData(h_B, nElem);

    // Cấp phát bộ nhớ cho các mảng trên GPU
    float *d_A, *d_B, *d_C;
    cudaMalloc((float**)&d_A, nBytes);
    cudaMalloc((float**)&d_B, nBytes);
    cudaMalloc((float**)&d_C, nBytes);

    // Sao chép dữ liệu từ host sang device (GPU)
    cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);

    // Xác định kích thước khối và lưới
    int blockSize = 256;  // Số lượng thread mỗi khối
    int gridSize = (nElem + blockSize - 1) / blockSize;  // Số lượng khối

    // Gọi hàm kernel trên GPU
    sumArraysOnGPU<<<gridSize, blockSize>>>(d_A, d_B, d_C, nElem);

    // Sao chép kết quả từ device về host
    cudaMemcpy(h_C, d_C, nBytes, cudaMemcpyDeviceToHost);

    // In kết quả
    for (int i = 0; i < nElem; i++)
    {
        printf("%f + %f = %f\n", h_A[i], h_B[i], h_C[i]);
    }

    // Giải phóng bộ nhớ trên GPU và host
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    free(h_A);
    free(h_B);
    free(h_C);

    return 0;
}

Chạy code

!nvcc ./add_cuda.cu -o sum_cu

!./sum_cu

Trong đoạn code trên bạn sẽ thấy một số khái niệm liên quan tới luồng (thread), khối (block), lưới (grid), bạn hãy tham khảo 3 khái niệm này ở phía dưới nhé.

Các từ khóa quan trọng: luồng (thread), khối (block), lưới (grid). Mình sẽ dùng từ Tiếng Anh trong các phần tiếp.

Có hai phân cấp luồng bao gồm

• Khối của các luồng (blocks of threads)

• Lưới của các khối (grids of blocks)

Tất cả các thread sinh ra bởi một kernel gọi là lưới (grid), tất cả các threads trong một grid chia sẻ chung global memory.

Một grid được tạo bởi rất nhiều thread blocks.

Một block là một nhóm những thread, các thread này có thể phối hợp:

• Block-local synchorization

• Block-local shard memory

Thread từ block khác nhau không thể phối hợp.

Thread sử dụng hai tọa độ để phân biệt với thread khác:

• blockIdx: Vị trí của block trong grid

• threadIdx: Vị trí của thread trong block

Những giá trị này kernel có thể sử dụng trực tiếp. Khi chạy kernel, tọa độ blockIdx và threadIdx được gán cho mỗi thread của CUDA runtime. Dựa trên tọa độ này, bạn có thể gửi data cho thread tương ứng để thực thi.

Mỗi biến tọa độ này dạng uint3, CUDA loại vector có sẵn. Bạn có thể lấy 3 giá trị tương ứng từ blockIdx bao gồm blockIdx.x, blockIdx.y, blockIdx.z

Tương tự với threadIdx ta sẽ có threadIdx.x, threadIdx.y, threadIdx.z

Code

CUDA tổ chức grid và block 2 chiều (2D). 2D grid sẽ chứa 2D block.

• blockDim (Chiều block bằng số lượng thread)

• gridDim (Chiều grid bằng số lượng block)

Những biến này có dạng dim3, vector số.

Ta có thể truy cập được chi tiết blockDim qua blockDim.x, blockDim.y, blockDim.z

Thông thường, một grid được tổ chức dưới dạng mảng 2D của block, và block được tổ chức dưới dạng mảng 3D của thread.

Cả grid và block sử dụng loại dim3 với 3 số nguyên không âm, mặc định là 1.

Thực hành phần này bạn có thể làm như sau:

• Định nghĩa số lượng phần tử của dữ liệu

int nElem = 6

• Sau đó cỡ block và grid dựa trên block và chiều của dữ liệu.

# Kiểu dữ liệu dim3
# 1D block bao gồm 3 thread.
# Đây là 1D vì chỉ có một tham số sử dụng 

dim3 block(3)

# 1D grid với số lượng block được định nghĩa bằng việc chia số lượng phần tử cho số lượng block
# Đây là 1D vì chỉ có một tham số sử dụng
# block.x phản ánh số lượng thread trong block, vì là 1D nên lấy số lượng thread bằng block.x
# Tại sao lại cần nElem + block.x - 1: Thủ thuật này đảm bảo rằng số lượng blocks được làm tròn lên khi nElem không chia hết cho block.x. 
# Ví dụ, nếu bạn có 1025 phần tử và 256 threads trong mỗi block, bạn sẽ cần 5 blocks (không phải 4), vì 1025 phần tử cần thêm một block để xử lý 1 phần tử dư.
# Chia cho block.x: Chia nElem + block.x - 1 cho block.x sẽ cho ra tổng số blocks cần thiết để xử lý tất cả các phần tử. 
# Mỗi block xử lý block.x threads, do đó phép chia này phân phối các phần tử đều cho các blocks phù hợp.

dim3 grid((nElem + block.x - 1)/block.x)

Toàn bộ code như sau:

%%writefile dim.cu

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void checkIndex(void) {
    printf("threadIdx: (%d, %d, %d) blockIdx: (%d, %d, %d) blockDim: (%d, %d, %d) "
           "gridDim: (%d, %d, %d)\n",
           threadIdx.x, threadIdx.y, threadIdx.z,
           blockIdx.x, blockIdx.y, blockIdx.z,
           blockDim.x, blockDim.y, blockDim.z,
           gridDim.x, gridDim.y, gridDim.z);
}

int main(int argc, char **argv) {
    // Define total data elements
    int nElem = 6;

    // Define grid and block structure
    dim3 block(3);  // 3 threads in the x-dimension
    dim3 grid((nElem + block.x - 1) / block.x);  // Calculate the grid size

    // Check grid and block dimensions from host side
    printf("Grid dimensions: (%d, %d, %d)\n", grid.x, grid.y, grid.z);
    printf("Block dimensions: (%d, %d, %d)\n", block.x, block.y, block.z);

    // Check grid and block dimensions from device side
    checkIndex<<<grid, block>>>();

    // Reset device before you leave
    cudaDeviceReset();

    return 0;
}

Kết quả in ra:

Grid dimensions: (2, 1, 1)
Block dimensions: (3, 1, 1)
threadIdx: (0, 0, 0) blockIdx: (0, 0, 0) blockDim: (3, 1, 1) gridDim: (2, 1, 1)
threadIdx: (1, 0, 0) blockIdx: (0, 0, 0) blockDim: (3, 1, 1) gridDim: (2, 1, 1)
threadIdx: (2, 0, 0) blockIdx: (0, 0, 0) blockDim: (3, 1, 1) gridDim: (2, 1, 1)
threadIdx: (0, 0, 0) blockIdx: (1, 0, 0) blockDim: (3, 1, 1) gridDim: (2, 1, 1)
threadIdx: (1, 0, 0) blockIdx: (1, 0, 0) blockDim: (3, 1, 1) gridDim: (2, 1, 1)
threadIdx: (2, 0, 0) blockIdx: (1, 0, 0) blockDim: (3, 1, 1) gridDim: (2, 1, 1)

Bạn có thể thấy mặc định các chiều không có sẽ là 1, ví dụ

• Grid dạng 1D có 2 block sẽ có chiều (2, 1, 1)

• Block dạng 1D có 3 thread sẽ có chiều (3, 1, 1)

• Thread có tọa độ 3D trong block

Một số chú ý quan trọng:

Đoạn code:

// Check grid and block dimensions from host side
printf("Grid dimensions: (%d, %d, %d)\n", grid.x, grid.y, grid.z);
printf("Block dimensions: (%d, %d, %d)\n", block.x, block.y, block.z);

là chạy trên host tức CPU, ta truy cập được chiều của grid và block thông qua

các giá trị grid.x, grid.y, grid.z và block.x, block.y, block.z

còn đoạn code:

printf("threadIdx: (%d, %d, %d) blockIdx: (%d, %d, %d) blockDim: (%d, %d, %d) "
           "gridDim: (%d, %d, %d)\n",
           threadIdx.x, threadIdx.y, threadIdx.z,
           blockIdx.x, blockIdx.y, blockIdx.z,
           blockDim.x, blockDim.y, blockDim.z,
           gridDim.x, gridDim.y, gridDim.z);

là chạy trên device tức GPU và ta truy cập chiều của grid và block thông qua các giá trị: blockDim.x, blockDim.y, blockDim.z và gridDim.x, gridDim.y, gridDim.z.

2 yếu tố quan trọng để xác định chiều của một block

• Thuộc tính của kernel

• Giới hạn của GPUs

Đặc điểm hiệu suất của kernel

• Workload Distribution (Phân phối khối lượng công việc): Kích thước block ảnh hưởng trực tiếp đến cách phân phối công việc giữa các thread của GPU. Lý tưởng nhất, mỗi block nên chứa đủ số lượng thread để tận dụng hết khả năng phần cứng, trong khi tránh việc sử dụng quá mức hoặc thiếu hụt tài nguyên GPU.

• Occupancy (Tỷ lệ sử dụng): Occupancy đề cập đến tỷ lệ giữa các warp đang hoạt động (các nhóm 32 thread) so với số lượng warp tối đa mà GPU có thể hỗ trợ. Occupancy cao thường dẫn đến hiệu suất tốt hơn, nhưng cần lưu ý rằng occupancy tối đa không phải lúc nào cũng đảm bảo hiệu suất tối ưu nếu băng thông bộ nhớ hoặc các tài nguyên khác trở thành điểm nghẽn.

• Memory Access Patterns (Cách truy cập bộ nhớ): Kích thước block ảnh hưởng đến cách các thread truy cập bộ nhớ. Bạn cần đảm bảo rằng kích thước block được đặt sao cho các thread có thể truy cập bộ nhớ một cách liên tục (coalesce memory accesses), giảm độ trễ của bộ nhớ toàn cục. Các truy cập không liên kết hoặc kích thước block không phù hợp có thể dẫn đến truy cập bộ nhớ không hiệu quả, làm giảm hiệu suất.

• Shared Memory Usage (Sử dụng bộ nhớ chia sẻ): Số lượng thread trên mỗi block cũng quyết định cách bộ nhớ chia sẻ được phân bổ giữa chúng. Nhiều thread trên mỗi block có thể giúp giảm thiểu các truy cập bộ nhớ toàn cục bằng cách tăng cường sử dụng bộ nhớ chia sẻ, nhưng điều này bị giới hạn bởi kích thước bộ nhớ chia sẻ có sẵn cho mỗi block.

Giới hạn tài nguyên của GPU

• Maximum Threads per Block (Số thread tối đa trên mỗi block): Mỗi kiến trúc GPU có giới hạn cố định về số lượng thread trên mỗi block (thường là 1024 thread trên nhiều GPU). Bạn cần đảm bảo rằng kích thước block của bạn không vượt quá giới hạn này.

• Registers per Thread (Số lượng register trên mỗi thread): Mỗi thread sử dụng một số lượng register nhất định, đây là tài nguyên có giới hạn. Nếu một block thread yêu cầu nhiều register hơn số lượng có sẵn trên GPU, số block có thể hoạt động sẽ giảm, có thể làm giảm occupancy.

• Shared Memory per Block (Bộ nhớ chia sẻ trên mỗi block): Bộ nhớ chia sẻ cũng bị giới hạn cho mỗi block. Nếu block yêu cầu nhiều bộ nhớ chia sẻ hơn số lượng có sẵn trên GPU, kernel có thể không chạy hoặc chạy với số lượng block hoạt động ít hơn, giảm hiệu suất.

• Warp Scheduling (Lập lịch warp): GPU thực thi các thread theo nhóm gọi là warp (thường là 32 thread). Kích thước block nên là bội số của kích thước warp để tránh sự phân kỳ của thread và việc sử dụng tài nguyên không hiệu quả. Ví dụ, kích thước block 128 hoặc 256 thread thường hoạt động tốt cho nhiều bài toán.

Các yếu tố khác

• Kernel-Specific Optimizations (Tối ưu hóa đặc thù của kernel): Một số kernel hưởng lợi từ kích thước block lớn hơn nếu chúng có thể sử dụng hiệu quả nhiều bộ nhớ chia sẻ hoặc register, trong khi những kernel khác có thể ưa thích các block nhỏ hơn để tăng số lượng block có thể chạy đồng thời trên GPU (cải thiện occupancy).

• Hardware Limitations (Giới hạn phần cứng): GPU hiện đại (ví dụ: từ NVIDIA và AMD) có nhiều kiến trúc khác nhau (như Volta, Ampere hoặc RDNA), mỗi kiến trúc có các giới hạn phần cứng khác nhau (ví dụ: giới hạn về warp, thread mỗi SM, bộ nhớ chia sẻ mỗi SM). Bạn cần tính đến phần cứng cụ thể mà bạn đang nhắm tới.

Tóm tắt các yếu tố:

1. Cân bằng Occupancy – Kích thước block nên tối đa hóa occupancy, nhưng cũng cần phù hợp với việc sử dụng bộ nhớ và các giới hạn tài nguyên khác.

2. Memory Coalescing (Gộp bộ nhớ) – Chọn kích thước block sao cho các thread có thể truy cập bộ nhớ hiệu quả, giảm độ trễ.

3. Giới hạn tài nguyên – Đảm bảo kích thước block nằm trong giới hạn về register, bộ nhớ chia sẻ và số thread mà GPU hỗ trợ.

Nhiệm vụ khi lập trình Kernel là mapping được chính xác luồng nào thực hiện trên dữ liệu nào.

Logic tính toán của CUDA sẽ sử dụng Kernel để cùng chạy logic dưới đây trên các luồng, tức là cùng một thời điểm các luồng này sẽ chạy cùng một logic:

// Hàm kernel để cộng hai mảng trên GPU
__global__ void sumArraysOnGPU(float *A, float *B, float *C, const int N)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;  // Xác định chỉ số phần tử
    if (i < N)
    {
        C[i] = A[i] + B[i];  // Cộng phần tử tương ứng của mảng A và B và lưu vào C
    }
}

Trong ví dụ dưới bạn có thể thấy ta sử dụng:

Sẽ dùng grid với 3 block (gridDim = 3), mỗi block có 2 thread (blockDim = 2) tức là 6 thread để thực hiện tính toán trên một mảng có 5 phần tử:

• Diễn giải như sau

  • Block 1 chỉ mục 0 trong grid:

    • Thread 1 trong block này thực hiện tính toán với phần tử chỉ mục 0 trong mảng

    • Thread 2 trong block này thực hiện tính toán với phần tử chỉ mục 1 trong mảng

  • Block 2 chỉ mục 1 trong grid:

    • Thread 1 trong block này thực hiện tính toán với phần tử chỉ mục 2 trong mảng

    • Thread 2 trong block này thực hiện tính toán với phần tử chỉ mục 3 trong mảng

  • Block 3 chỉ mục 2 trong grid:

    • Thread 1 trong block này thực hiện tính toán với phần tử chỉ mục 4 trong mảng

    • Thread 2 trong block này không tính toán vì chỉ mục 5 ngoài chiều dài của mảng

Đầy đủ code tại đây.

Code CUDA cộng hai vector:

https://colab.research.google.com/drive/15e23Dk_b6RbUH27NRu2nMirxbjCKC2Gt?usp=sharing

Code CUDA kiểm tra chiều:

https://colab.research.google.com/drive/1VLk6FWeTJK5mkr8TLUiZ0JGECLTLsDfg?usp=sharing