Tất tần tật về Kernel

Cú pháp lệnh gọi kernel

Trong CUDA, lệnh gọi kernel có dạng tương tự như một hàm C thông thường, nhưng với một điểm khác biệt quan trọng: bạn cần thêm cấu hình thực thi của kernel trong dấu ngoặc ba dấu góc (<<< >>>).

kernel_function<<<grid_size, block_size>>>(parameters);

• grid_size: Xác định số lượng block trong grid (mạng lưới) mà bạn muốn kernel của mình thực thi. Tham số này cho phép bạn kiểm soát số lượng block trong mạng lưới mà kernel sẽ chạy. Bạn có thể xác định grid dưới dạng một giá trị đơn lẻ (1D), hai chiều (2D) hoặc ba chiều (3D), tùy thuộc vào bài toán của bạn.

  • Ví dụ: <<<dim3(grid_dim_x, grid_dim_y), block_size>>>

• block_size: Xác định số lượng thread trong mỗi block. Tham số này xác định số lượng thread trong mỗi block. Cũng giống như grid size, bạn có thể xác định block size dưới dạng 1D, 2D hoặc 3D. Số lượng thread tối đa cho mỗi block thường là 1024 đối với hầu hết các GPU.

  • Ví dụ: <<<grid_size, dim3(block_dim_x, block_dim_y, block_dim_z)>>>

Các tham số tùy chọn trong lệnh gọi kernel

Ngoài grid_size và block_size, bạn còn có thể thêm các tham số khác bên trong dấu ngoặc ba dấu góc, như:

• Shared Memory Size: Xác định lượng bộ nhớ chia sẻ mà mỗi block sử dụng. Bạn có thể truyền một giá trị kích thước bộ nhớ chia sẻ dưới dạng tham số thứ ba.

  • Ví dụ: <<<grid_size, block_size, shared_mem_size>>>

• Stream ID: Tham số thứ tư tùy chọn xác định stream mà kernel sẽ chạy trên đó. Điều này hữu ích khi bạn muốn thực thi song song nhiều kernel trên các stream khác nhau.

  • Ví dụ: <<<grid_size, block_size, shared_mem_size, stream_id>>>

Ví dụ về lệnh gọi kernel

my_kernel<<<dim3(16, 16), dim3(8, 8)>>>(arg1, arg2);

Trong ví dụ này:

• dim3(16, 16) xác định một grid 2D với kích thước 16x16 block.

• dim3(8, 8) xác định mỗi block có 8x8 thread.

• Các tham số arg1 và arg2 là các tham số được truyền cho kernel.

Điều chỉnh hiệu suất

• Kích thước Grid và Block: Kích thước grid và block cần được tinh chỉnh để tận dụng tối đa tài nguyên GPU. Điều này phụ thuộc vào bài toán cụ thể, cũng như số lượng tài nguyên như bộ nhớ chia sẻ, register, và khả năng đồng thời của GPU.

• Bộ nhớ chia sẻ và Stream: Sử dụng bộ nhớ chia sẻ một cách hiệu quả và lập kế hoạch cho việc thực thi song song trên các stream có thể cải thiện đáng kể hiệu suất của kernel.

Mã định nghĩa một kernel:

kernel_function<<<grid_size, block_size>>>(parameters);

Ví dụ một kernel được định nghĩa như sau:

kernel_function<<<3, 6>>>(parameters);

Ở đây có nghĩa là kernel này sẽ thực hiện trên 3 block, mỗi block có 6 thread.

Quay lại bài toán cộng hai vector:

Code C cho phép cộng này:

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
    }
}

Code Kernel:

// 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
    }
}

Sự khác nhau giữa code C và code Kernel?

Code kernel không có vòng for vì cùng CUDA sẽ chạy N luồng và thực hiện cùng một hàm này song song còn code C sẽ dùng 1 luồng và thực thi tuần tự

Các giới hạn của Kernel:

• Chỉ truy cập được vào bộ nhớ của device

• Return type phải là void

• Không hỗ trợ số lượng tham số không cố định

• Không hỗ trợ biến static

• Không hỗ trợ con trỏ hàm (function pointers)

• Có tính chất bất đồng bộ (asynchronous behavior)

1. Chỉ truy cập vào bộ nhớ device

Các kernel của CUDA chỉ có thể truy cập trực tiếp vào bộ nhớ nằm trên GPU (bộ nhớ device), không thể truy cập vào bộ nhớ nằm trên CPU (bộ nhớ host). Nói cách khác, khi một kernel đang chạy, nó chỉ có thể làm việc với dữ liệu đã được sao chép từ host (CPU) sang device (GPU). Việc phân biệt giữa bộ nhớ host và device là quan trọng trong CUDA vì CPU và GPU có không gian bộ nhớ riêng biệt. Dữ liệu phải được chuyển đổi giữa hai không gian này bằng các hàm như cudaMemcpy().

__global__ void kernel_function(float *device_array) {
    device_array[0] = 42.0f;  // Truy cập vào bộ nhớ device
}

2. Phải có kiểu trả về void

Các kernel CUDA luôn phải có kiểu trả về void, nghĩa là chúng không thể trả về giá trị trực tiếp. Thay vào đó, kết quả phải được truyền lại cho host thông qua bộ nhớ device (global memory) hoặc shared memory. Kernel có thể thay đổi dữ liệu được truyền vào dưới dạng con trỏ, nhưng không thể trả về giá trị như một hàm C++ thông thường.

Lý do: Vì các kernel CUDA chạy song song trên nhiều thread, việc trả về giá trị từ mỗi thread sẽ rất phức tạp. Thay vào đó, GPU thay đổi bộ nhớ mà tất cả các thread có thể truy cập hoặc ghi kết quả trực tiếp vào global memory.

__global__ void kernel_function(float *data) {
    int idx = threadIdx.x;
    data[idx] = idx * 2.0f;  // Kết quả được ghi vào bộ nhớ device, không phải trả về.
}

3. Không hỗ trợ số lượng tham số không cố định

• Các kernel CUDA không hỗ trợ danh sách tham số có độ dài thay đổi, giống như cách sử dụng với các hàm như printf (ví dụ: void foo(...);). Trong kernel, số lượng và kiểu dữ liệu của các tham số phải được xác định rõ ràng khi định nghĩa và gọi hàm.

• Lý do: Danh sách tham số có độ dài thay đổi tạo ra sự mơ hồ trong cơ chế truyền tham số, và việc quản lý các tham số thay đổi trong môi trường thực thi song song sẽ làm giảm hiệu suất của GPU.

// Điều này không được phép trong các kernel CUDA
__global__ void kernel_function(int num, ...) {
    // tham số thay đổi không được phép
}

4. Không hỗ trợ biến static

• Các kernel CUDA không hỗ trợ việc sử dụng biến static, loại biến này giữ lại giá trị của nó giữa các lần gọi hàm khác nhau. Biến static thường giữ giá trị của nó giữa các lần gọi hàm trên CPU, nhưng các kernel CUDA chạy nhiều thread song song, làm cho việc quản lý trạng thái của các biến này trở nên khó khăn khi chúng được sử dụng đồng thời bởi nhiều thread và block.

• Lý do: Bản chất song song của CUDA có nghĩa là nhiều thread và block có thể thay đổi cùng một biến static cùng lúc, dẫn đến các tình huống race condition hoặc hành vi không xác định. Việc quản lý các biến static một cách nhất quán giữa các thread là rất phức tạp, vì vậy điều này bị cấm.

// Điều này không được phép trong các kernel CUDA
__global__ void kernel_function() {
    static int counter = 0;  // Biến static không được phép
    counter++;
}

5. Không hỗ trợ con trỏ hàm (function pointers)

• Các kernel CUDA không hỗ trợ việc sử dụng con trỏ hàm. Trong C++ chuẩn, một con trỏ hàm có thể được truyền và gọi, nhưng trong CUDA, điều này bị cấm vì kiến trúc GPU không hỗ trợ đầy đủ các lệnh gọi hàm gián tiếp, mà con trỏ hàm dựa vào.

• Lý do: Con trỏ hàm tạo ra mức độ gián tiếp mà sẽ không hiệu quả để quản lý trên GPU. Việc hỗ trợ con trỏ hàm trên nhiều thread song song trong các đơn vị thực thi khác nhau trên GPU sẽ làm giảm hiệu suất và tạo ra các thách thức về đồng bộ hóa.

// Điều này không được phép trong các kernel CUDA
__global__ void kernel_function(void (*func_ptr)(int)) {
    func_ptr(5);  // Con trỏ hàm không được phép trong các kernel CUDA
}

6. Có tính chất bất đồng bộ (asynchronous behavior)

• Các kernel CUDA có tính chất bất đồng bộ, nghĩa là khi một kernel được khởi chạy, CPU không chờ kernel hoàn tất thực thi. Thay vào đó, kernel được khởi chạy trên GPU, và quyền điều khiển lập tức được trả lại cho CPU để tiếp tục thực thi các lệnh khác. CPU và GPU có thể thực thi đồng thời trừ khi được đồng bộ hóa rõ ràng.

• Lý do: Điều này cho phép sự chồng lấp tốt hơn giữa việc tính toán và truyền dữ liệu, dẫn đến hiệu suất được cải thiện. Ví dụ, trong khi GPU đang thực thi một kernel, CPU có thể thực hiện các tác vụ khác hoặc chuẩn bị dữ liệu cho thao tác tiếp theo trên GPU.

kernel_function<<<16, 256>>>(device_data);  // Kernel được khởi chạy bất đồng bộ
cudaDeviceSynchronize();  // Dùng để chờ kernel hoàn tất

Nếu không có cudaDeviceSynchronize(), CPU có thể tiếp tục làm việc trong khi GPU xử lý kernel ở nền. Tuy nhiên, nếu bạn cần đảm bảo kernel đã hoàn tất trước khi tiếp tục, bạn có thể gọi cudaDeviceSynchronize() để buộc CPU chờ GPU thực thi xong.

Tóm gọn lại: Các hàm kernel CUDA bị giới hạn trong việc tương tác với bộ nhớ device và phải tuân theo các hạn chế như kiểu trả về void, không có tham số thay đổi, không có biến static và không có con trỏ hàm. Ngoài ra, các kernel CUDA được khởi chạy bất đồng bộ, cho phép CPU tiếp tục thực thi trong khi GPU xử lý kernel song song. Những hạn chế này nhằm tối ưu hóa hiệu suất song song trên GPU.

Code đầy đủ

void checkResult(float *hostRef, float *gpuRef, const int N) { 
    double epsilon = 1.0E-8;
    int match = 1;
    for (int i = 0; i < N; i++) {
        if (abs(hostRef[i] - gpuRef[i]) > epsilon) { 
            match = 0;
            printf("Arrays do not match!\n");
            printf("host %5.2f gpu %5.2f at current %d\n", hostRef[i], gpuRef[i], i); 
            break;
        } 
    }
    if (match) 
        printf("Arrays match.\n\n");
    return; 
}

Hàm checkResult được sử dụng để so sánh hai mảng (hoặc vector) của các số dấu phẩy động, thường để kiểm tra kết quả từ tính toán trên host CPU (hostRef) so với tính toán trên GPU (gpuRef). Dưới đây là chi tiết về cách hàm này hoạt động:

Các tham số:

• float *hostRef: Con trỏ trỏ đến một mảng lưu trữ kết quả tính toán trên host (CPU).

• float *gpuRef: Con trỏ trỏ đến một mảng lưu trữ kết quả tính toán trên GPU.

• const int N: Số phần tử trong các mảng.

Chức năng:

1. Thiết lập ngưỡng dung sai (epsilon):
   Giá trị nhỏ epsilon (1.0E-8) được định nghĩa để chấp nhận những khác biệt nhỏ do độ chính xác của dấu phẩy động. Các so sánh giữa hai mảng sẽ cho phép sự khác biệt lên đến giá trị nhỏ này.

2. Duyệt qua các mảng:
   Vòng lặp for duyệt qua các phần tử của mảng từ i = 0 đến i < N. Trong mỗi lần lặp:

   • Nó kiểm tra sự chênh lệch tuyệt đối giữa hostRef[i] và gpuRef[i].

   • Nếu sự khác biệt lớn hơn epsilon, nó sẽ đặt biến match về 0, nghĩa là các mảng không khớp nhau.

3. Hiển thị chi tiết khác biệt:
   Nếu có sự khác biệt (sự khác biệt lớn hơn epsilon):

   • Nó sẽ in ra thông báo "Arrays do not match!".

   • Nó cũng in ra giá trị của hostRef[i] và gpuRef[i] cho phần tử mà sự khác biệt xảy ra, cũng như chỉ số i.

   • Hàm sau đó sẽ thoát khỏi vòng lặp vì đã phát hiện sự khác biệt.

4. Kết quả cuối cùng:
   Sau vòng lặp:

   • Nếu match vẫn bằng 1 (nghĩa là không có sự khác biệt), nó sẽ in ra "Arrays match." để chỉ ra rằng tất cả các phần tử trong cả hai mảng đều nằm trong phạm vi dung sai cho phép.

5. Trả về:
   Hàm trả về kiểu void (không trả về giá trị gì), vì nó chỉ in ra kết quả là các mảng có khớp nhau hay không.

Mục đích:

Hàm này thường được sử dụng trong tính toán GPU để kiểm tra xem kết quả từ GPU (ví dụ: sử dụng CUDA hoặc OpenCL) có nhất quán với những gì được tính toán trên CPU hay không. Do các khác biệt trong tính toán dấu phẩy động giữa CPU và GPU, một giá trị epsilon nhỏ thường được sử dụng để xử lý sai số làm tròn.

Ví dụ cụ thể

Code này sẽ mô phỏng lập trình cộng hai vector giữa C và dùng CUDA. Sau đó kiểm tra hai kết quả này có giống nhau không.

Trong code có một số hàm kiểm tra xem CUDA có hoạt động đúng không ví dụ như là

err = cudaMalloc((void **)&d_a, size); // Cấp phát bộ nhớ cho mảng d_a trên GPU

Chi tiết

#include <stdio.h>  // Thư viện cho đầu vào và đầu ra tiêu chuẩn
#include <stdlib.h> // Thư viện cho các hàm cấp phát bộ nhớ động
#include <math.h>   // Thư viện cho các hàm toán học

// Hàm checkResult để kiểm tra kết quả giữa CPU và GPU
void checkResult(float *hostRef, float *gpuRef, const int N) { 
    double epsilon = 1.0E-8; // Sai số chấp nhận được
    int match = 1; // Biến lưu trữ trạng thái có khớp hay không
    for (int i = 0; i < N; i++) { // Lặp qua tất cả các phần tử
        if (abs(hostRef[i] - gpuRef[i]) > epsilon) { // Kiểm tra sự khác biệt giữa giá trị CPU và GPU
            match = 0; // Nếu có khác biệt lớn hơn epsilon thì không khớp
            printf("Arrays do not match!\n"); // In ra thông báo không khớp
            printf("host %5.2f gpu %5.2f at current %d\n", hostRef[i], gpuRef[i], i); // In ra giá trị khác biệt
            break; // Dừng kiểm tra nếu tìm thấy sự khác biệt
        } 
    }
    if (match)  // Nếu không có khác biệt nào
        printf("Arrays match.\n\n"); // In ra thông báo khớp
    return;  // Trả về
}

// Kernel ví dụ mô phỏng một phép tính trên GPU (đơn giản hóa)
__global__ void addVectors(float *a, float *b, float *result, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x; // Tính chỉ số của thread hiện tại
    if (i < N) {  // Đảm bảo thread không vượt quá giới hạn của mảng
        result[i] = a[i] + b[i]; // Thực hiện phép cộng
    }
}

// Hàm khởi tạo mảng với một số giá trị ngẫu nhiên
void initializeArray(float *arr, int N) {
    for (int i = 0; i < N; i++) { // Lặp qua tất cả các phần tử
        arr[i] = rand() % 100 / 10.0f;  // Gán giá trị ngẫu nhiên từ 0 đến 10 cho mảng
    }
}

void printResults(float *hostRef, float *gpuRef, const int N) {
    // In ra 10 phần tử đầu tiên của kết quả từ CPU và GPU để so sánh
    printf("First 10 elements of host and GPU results:\n");
    for (int i = 0; i < 10; i++) { // Lặp qua 10 phần tử đầu tiên
        printf("Index %d: host %5.2f gpu %5.2f\n", i, hostRef[i], gpuRef[i]); // In ra giá trị từ CPU và GPU
    }
}

int main() {
    int N = 1000;  // Kích thước của các mảng
    size_t size = N * sizeof(float); // Tính kích thước bộ nhớ cần cho mảng

    // Cấp phát bộ nhớ trên host (CPU)
    float *h_a = (float *)malloc(size);
    float *h_b = (float *)malloc(size);
    float *h_resultHost = (float *)malloc(size);
    float *h_resultGPU = (float *)malloc(size);

    // Khởi tạo mảng trên host
    initializeArray(h_a, N);
    initializeArray(h_b, N);

    // Thực hiện phép cộng trên host
    for (int i = 0; i < N; i++) {
        h_resultHost[i] = h_a[i] + h_b[i]; // Cộng từng phần tử của hai mảng
    }

    // Cấp phát bộ nhớ trên thiết bị (GPU)
    float *d_a, *d_b, *d_result;
    cudaError_t err;

    err = cudaMalloc((void **)&d_a, size); // Cấp phát bộ nhớ cho mảng d_a trên GPU
    if (err != cudaSuccess) { // Kiểm tra lỗi
        printf("CUDA malloc error (a): %s\n", cudaGetErrorString(err)); // In ra lỗi nếu có
    }

    err = cudaMalloc((void **)&d_b, size); // Cấp phát bộ nhớ cho mảng d_b trên GPU
    if (err != cudaSuccess) { // Kiểm tra lỗi
        printf("CUDA malloc error (b): %s\n", cudaGetErrorString(err)); // In ra lỗi nếu có
    }

    err = cudaMalloc((void **)&d_result, size); // Cấp phát bộ nhớ cho mảng d_result trên GPU
    if (err != cudaSuccess) { // Kiểm tra lỗi
        printf("CUDA malloc error (result): %s\n", cudaGetErrorString(err)); // In ra lỗi nếu có
    }

    // Sao chép dữ liệu từ host sang device
    err = cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice); // Sao chép mảng h_a sang d_a
    if (err != cudaSuccess) { // Kiểm tra lỗi
        printf("CUDA memcpy Host to Device error (a): %s\n", cudaGetErrorString(err)); // In ra lỗi nếu có
    }

    err = cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice); // Sao chép mảng h_b sang d_b
    if (err != cudaSuccess) { // Kiểm tra lỗi
        printf("CUDA memcpy Host to Device error (b): %s\n", cudaGetErrorString(err)); // In ra lỗi nếu có
    }

    // Định nghĩa số lượng blocks và threads
    int threadsPerBlock = 256; // Số lượng threads mỗi block
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock; // Số lượng blocks cần dùng

    // Khởi chạy kernel để thực hiện phép cộng trên GPU
    addVectors<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_result, N);
    err = cudaGetLastError();  // Kiểm tra lỗi khi khởi chạy kernel
    if (err != cudaSuccess) { // Nếu có lỗi
        printf("CUDA kernel launch error: %s\n", cudaGetErrorString(err)); // In ra thông báo lỗi
    }

    // Sao chép kết quả từ thiết bị về host
    err = cudaMemcpy(h_resultGPU, d_result, size, cudaMemcpyDeviceToHost); // Sao chép kết quả từ d_result sang h_resultGPU
    if (err != cudaSuccess) { // Kiểm tra lỗi
        printf("CUDA memcpy Device to Host error (result): %s\n", cudaGetErrorString(err)); // In ra lỗi nếu có
    }

    // In ra 10 phần tử đầu tiên của kết quả từ CPU và GPU để so sánh
    printResults(h_resultHost, h_resultGPU, N);

    // So sánh kết quả từ CPU và GPU
    checkResult(h_resultHost, h_resultGPU, N);

    // Giải phóng bộ nhớ trên thiết bị và host
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_result);
    free(h_a);
    free(h_b);
    free(h_resultHost);
    free(h_resultGPU);

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

Kết quả chạy code như sau:

First 10 elements of host and GPU results:
Index 0: host 17.60 gpu 17.60
Index 1: host 12.00 gpu 12.00
Index 2: host 11.60 gpu 11.60
Index 3: host  4.90 gpu  4.90
Index 4: host 11.40 gpu 11.40
Index 5: host  9.40 gpu  9.40
Index 6: host 17.10 gpu 17.10
Index 7: host 14.90 gpu 14.90
Index 8: host 10.30 gpu 10.30
Index 9: host  8.20 gpu  8.20
Arrays match.

Vì các lệnh gọi CUDA là không đồng bộ (asynchronous), nên có thể khó xác định hàm nào gây ra lỗi. Một cách hữu ích để đơn giản hóa việc kiểm tra lỗi là định nghĩa một macro kiểm tra các lệnh gọi API CUDA.

Bạn có thể định nghĩa macro kiểm tra lỗi như sau:

#define CHECK(call) {
    const cudaError_t error = call;
    if (error != cudaSuccess) {
        printf("Error: %s:%d, ", __FILE__, __LINE__);
        printf("code:%d, reason: %s\n", error, cudaGetErrorString(error));
        exit(1);
    }
}

Ví dụ sử dụng 1: Kiểm tra chép bộ nhớ

Bạn có thể sử dụng macro này để bao bọc các lệnh cudaMemcpy như sau:

CHECK(cudaMemcpy(d_C, gpuRef, nBytes, cudaMemcpyHostToDevice));

Nếu việc sao chép bộ nhớ, hoặc một hoạt động không đồng bộ trước đó, gây ra lỗi, hàm sẽ:

- Báo cáo mã lỗi

- In ra thông báo lỗi dễ hiểu cho người dùng

- Dừng chương trình

Ví dụ sử dụng 2: Kernel Launch

Bạn cũng có thể sử dụng macro này sau khi khởi chạy một kernel để kiểm tra lỗi liên quan đến kernel. Ví dụ:

kernel_function<<<grid, block>>>(argument_list);
CHECK(cudaDeviceSynchronize());

Lệnh CHECK(cudaDeviceSynchronize()) đảm bảo rằng thread host bị chặn lại cho đến khi thiết bị (device) hoàn thành tất cả các tác vụ được yêu cầu trước đó. Nó cũng kiểm tra lỗi xảy ra trong quá trình khởi chạy kernel.

Code đầy đủ

Để đo lường thời gian tính toán của một kernel bất kỳ, ta sẽ dùng bộ đếm thời gian của CPU và GPU.

Trong bài lần này chúng ta sẽ tính toán trên một vector có hơn 16 triệu phần tử.

int nElem = 1 << 24;

Hàm tính toán thời gian trên CPU như sau:

double cpuSecond() {
    struct timeval tp;
    gettimeofday(&tp,NULL);
    return ((double)tp.tv_sec + (double)tp.tv_usec*1.e-6);
}

Đo lường thời gian CPU chạy như sau:

    // Add vectors at host side for result checks
    iStart = cpuSecond();
    sumArraysOnHost(h_A, h_B, hostRef, nElem);
    iElaps = cpuSecond() - iStart;

Đo lường thời gian GPU chạy như sau:

    iStart = cpuSecond();
    sumArraysOnGPU<<<grid, block>>>(d_A, d_B, d_C, nElem);
    cudaDeviceSynchronize();
    iElaps = cpuSecond() - iStart;
    printf("sumArraysOnGPU <<<%d,%d>>> Time elapsed %f sec\n", grid.x, block.x, iElaps);

Hàm cudaDeviceSynchronize() để chờ tất cả các thread của GPU hoàn thành.

Thời gian tính toán trên GPU:

./time Starting...
Using Device 0: NVIDIA L4
Vector size 16777216
sumArraysOnGPU <<<16384,1024>>> Time elapsed 0.130786 sec
Arrays match.

Ngoài ra bạn có thể dùng công cụ có sẵn để đo như ncu

!ncu ./time 

Kết quả:

/content/./time Starting...
==PROF== Connected to process 3177 (/content/time)
Using Device 0: NVIDIA L4
Vector size 16777216
==PROF== Profiling "sumArraysOnGPU" - 0: 0%....50%....100% - 9 passes
sumArraysOnGPU <<<16384,1024>>> Time elapsed 0.415159 sec
Arrays match.

==PROF== Disconnected from process 3177
[3177] time@127.0.0.1
  sumArraysOnGPU(float *, float *, float *, int) (16384, 1, 1)x(1024, 1, 1), Context 1, Stream 7, Device 0, CC 8.9
    Section: GPU Speed Of Light Throughput
    ----------------------- ------------- ------------
    Metric Name               Metric Unit Metric Value
    ----------------------- ------------- ------------
    DRAM Frequency          cycle/nsecond         6.17
    SM Frequency            cycle/usecond       787.10
    Elapsed Cycles                  cycle      553,307
    Memory Throughput                   %        95.11
    DRAM Throughput                     %        95.11
    Duration                      usecond       701.54
    L1/TEX Cache Throughput             %        16.52
    L2 Cache Throughput                 %        30.54
    SM Active Cycles                cycle   498,586.10
    Compute (SM) Throughput             %        13.10
    ----------------------- ------------- ------------

    INF   The kernel is utilizing greater than 80.0% of the available compute or memory performance of the device. To   
          further improve performance, work will likely need to be shifted from the most utilized to another unit.      
          Start by analyzing DRAM in the Memory Workload Analysis section.                                              

    Section: Launch Statistics
    -------------------------------- --------------- ---------------
    Metric Name                          Metric Unit    Metric Value
    -------------------------------- --------------- ---------------
    Block Size                                                 1,024
    Function Cache Configuration                     CachePreferNone
    Grid Size                                                 16,384
    Registers Per Thread             register/thread              16
    Shared Memory Configuration Size           Kbyte            8.19
    Driver Shared Memory Per Block       Kbyte/block            1.02
    Dynamic Shared Memory Per Block       byte/block               0
    Static Shared Memory Per Block        byte/block               0
    Threads                                   thread      16,777,216
    Waves Per SM                                              282.48
    -------------------------------- --------------- ---------------

    Section: Occupancy
    ------------------------------- ----------- ------------
    Metric Name                     Metric Unit Metric Value
    ------------------------------- ----------- ------------
    Block Limit SM                        block           24
    Block Limit Registers                 block            4
    Block Limit Shared Mem                block            8
    Block Limit Warps                     block            1
    Theoretical Active Warps per SM        warp           32
    Theoretical Occupancy                     %        66.67
    Achieved Occupancy                        %        51.51
    Achieved Active Warps Per SM           warp        24.72
    ------------------------------- ----------- ------------

Xem phần đã được đề cập tại đây.

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

Ví dụ đề bài yêu cầu ta tính toán trên một ma trận 2D có chiều (12, 6) định dạng (nx,ny) như bên dưới.

Bản chất ma trận này sẽ được lưu trữ dưới dạng mảng 72 phần từ chỉ mục từ 0 đến 71.

Để tính toán được ma trận này ta sẽ phân bổ 72 phần tử của mảng cho 72 luồng để tính toán. 72 luồng này được bố trí dưới dạng grid 2D và block 2D như sau:

• Ta đặt sẵn block 2D bao gồm (4, 2) thread

Công thức tính ra chiều grid:

grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);

// Tương đương

grid(3, 3)

• Grid 2D bao gồm (3, 3) block

Nhiệm vụ phải xác định được chỉ mục của thread tương ứng với chỉ mục của phần tử.

Nhìn vào bên trên bạn có thể thấy một số thông tin như sau:

• Thread tính toán trên phần tử thứ 67 của mảng có tọa độ trên grid là (7,5), tọa độ local trong block là (3,1)

• Thread tính toán trên phần tử thứ 26 của mảng có tọa độ trên grid là (2, 2), tọa độ local trong block là (2,0)

• Tương tự như vậy. Chạy code để xem chi tiết.

Cách xác định tọa độ của thread trên grid

Để xác định được tọa độ của thread trong grid hay cặp (7,5) ta sử dụng công thức sau:

Để xác định được ix hay giá trị 7 của thread trong grid ta lấy vị trí x của block trong grid nhân với chiều x của block cộng với vị trí local x của thread trong block

int ix = threadIdx.x + blockIdx.x * blockDim.x;

Để xác định được iy hay giá trị 5 của thread trong Grid ta lấy vị trí y của block trong grid nhân với chiều y của block cộng với vị trí local y của thread trong block

int iy = threadIdx.y + blockIdx.y * blockDim.y;

Cách xác định chỉ mục phần tử cần tính toán

Để xác định được con số 67 ta cần đoạn code sau:

Ta lấy tọa độ iy của thread nhân với số cột nx cộng với tọa độ ix là ra được vị trí của luồng/dữ liệu tương ứng trên mảng khi mảng được duỗi từ ma trận.

int idx = iy * nx + ix;

Toàn bộ code xác định luồng nào thực hiện trên phần tử nào dưới đây:

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

#define CHECK(call)                                                      \
{                                                                        \
    const cudaError_t error = call;                                       \
    if (error != cudaSuccess)                                             \
    {                                                                    \
        printf("Error: %s:%d, ", __FILE__, __LINE__);                     \
        printf("code:%d, reason: %s\n", error, cudaGetErrorString(error));\
        exit(-10 * error);                                                \
    }                                                                    \
}

void initialInt(int *ip, int size)
{
    for (int i = 0; i < size; i++)
    {
        ip[i] = i;
    }
}

void printMatrix(int *C, const int nx, const int ny)
{
    int *ic = C;
    printf("\nMatrix: (%d, %d)\n", nx, ny);

    for (int iy = 0; iy < ny; iy++)
    {
        for (int ix = 0; ix < nx; ix++)
        {
            printf("%3d ", ic[ix]);
        }
        ic += nx;
        printf("\n");
    }
    printf("\n");
}

__global__ void printThreadIndex(int *A, const int nx, const int ny)
{
    int ix = threadIdx.x + blockIdx.x * blockDim.x;
    int iy = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int idx = iy * nx + ix;

    if (ix < nx && iy < ny)
    {
        printf("thread_id (%d,%d) block_id (%d,%d) coordinate (%d,%d) "
               "global index %2d ival %2d\n",
               threadIdx.x, threadIdx.y, blockIdx.x, blockIdx.y, ix, iy, idx, A[idx]);
    }
}

int main(int argc, char **argv)
{
    printf("%s Starting...\n", argv[0]);

    // Get device information
    int dev = 0;
    cudaDeviceProp deviceProp;
    CHECK(cudaGetDeviceProperties(&deviceProp, dev));
    printf("Using Device %d: %s\n", dev, deviceProp.name);
    CHECK(cudaSetDevice(dev));

    // Set matrix dimension
    int nx = 12;
    int ny = 6;
    int nxy = nx * ny;
    int nBytes = nxy * sizeof(int);

    // Malloc host memory
    int *h_A = (int *)malloc(nBytes);

    // Initialize host matrix with integers
    initialInt(h_A, nxy);
    printMatrix(h_A, nx, ny);

    // Malloc device memory
    int *d_MatA;
    CHECK(cudaMalloc((void **)&d_MatA, nBytes));

    // Transfer data from host to device
    CHECK(cudaMemcpy(d_MatA, h_A, nBytes, cudaMemcpyHostToDevice));

    // Set up execution configuration
    dim3 block(4, 2);
    dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);

    printf("grid (%d,%d) block (%d,%d)\n", grid.x, grid.y, block.x, block.y);

    // Invoke the kernel
    printThreadIndex<<<grid, block>>>(d_MatA, nx, ny);
    CHECK(cudaDeviceSynchronize());

    // Free host and device memory
    CHECK(cudaFree(d_MatA));
    free(h_A);

    // Reset device
    CHECK(cudaDeviceReset());

    return 0;
}