Bộ công cụ CUDA hiện đại trong thực tiễn: Hướng dẫn tối ưu hóa từng bước

NVIDIA CUDA vẫn là nền tảng của điện toán tăng tốc bằng GPU, hỗ trợ mọi thứ từ mô phỏng khoa học đến huấn luyện AI quy mô lớn.

Tuy nhiên, việc viết mã CUDA đúng đắn, dễ bảo trì và có hiệu suất cao có thể là một thách thức: các lỗi bộ nhớ ẩn mình ngay trước mắt, các nút thắt hiệu suất có thể không hiển thị rõ ràng nếu thiếu công cụ đo lường phù hợp, và các thuật toán GPU tự viết thường khó đạt được mức độ tối ưu như các thư viện đã được tối ưu hóa. May mắn thay, bộ công cụ CUDA hiện đại đã phát triển đáng kể, và nhiều thách thức trong số này hiện đã có các giải pháp đơn giản.

Trong bài viết blog này, chúng tôi sẽ hướng dẫn qua các công cụ mà NVIDIA cung cấp để gỡ lỗi, đo hiệu năng và cải thiện mã của bạn. Chỉ với những thay đổi nhỏ ở từng dòng mã mỗi lần, chúng tôi sẽ làm cho mã ví dụ an toàn hơn, dễ bảo trì hơn và nhanh hơn.

Trong sáu bước tăng dần, bài viết này sẽ bao gồm:

Cách dễ dàng tìm ra lỗi đánh chỉ mục bằng cách áp dụng API CCCL hiện đại và Compute Sanitizer
Cách cải thiện các bài kiểm tra hiệu năng Nsight Systems với NVTX
Cách sử dụng các thuật toán tối ưu của CUB ở cấp độ khối và thiết bị
Cách quản lý bộ nhớ GPU thông qua các bộ chứa được phân bổ
Cách tăng tốc việc truyền dữ liệu từ host sang thiết bị bằng các vùng bộ nhớ đã ghim (pinned memory)
Cách song song hóa công việc trên GPU bằng cách cung cấp cho mỗi thread một luồng (stream) riêng và các phép truyền dữ liệu không đồng bộ

Để hỗ trợ bài viết blog này, chúng tôi cung cấp mã nguồn và tùy chọn chạy trên Google Colab.

Điểm bắt đầu: Ví dụ về quy trình xử lý hình ảnh

Từ một luồng đầu vào gồm các hình ảnh màu đỏ, xanh lá cây và xanh dương, bắt đầu bằng việc chuyển dữ liệu từ CPU sang GPU. Sau đó, chuyển đổi các hình ảnh đó từ RGB sang thang độ xám.

Hình 1. Phần đầu tiên của quy trình xử lý hình ảnh. Đầu tiên, sao chép các hình ảnh RGB từ bộ nhớ máy chủ (host) sang thiết bị (device), sau đó chuyển đổi chúng từ RGB sang thang độ xám.

Sau đó, đối với mỗi khối 32×32 pixel trong ảnh, tính giá trị trung vị bằng cách sắp xếp các pixel và chọn giá trị ở giữa. Cuối cùng, sao chép giá trị trung vị của mỗi khối trở lại CPU.

Hình 2. Phần thứ hai của quy trình xử lý ảnh: Chia mỗi ảnh thang độ xám thành các khối 32×32 pixel, sắp xếp các pixel trong mỗi khối để chọn giá trị trung vị, sau đó sao chép các giá trị trung vị từ thiết bị (device) về máy chủ (host)

Ví dụ mã nguồn cơ bản

Dưới đây là toàn bộ mã nguồn bắt đầu. Mỗi bước trong bài viết này đều cải thiện mã nguồn đó.

#define CUDA_CHECK_ERROR(call) do { \
    cudaError_t err = call; \
    if (err != cudaSuccess) { \
        std::cerr << "CUDA error in " << __FILE__ << " at line " << __LINE__ << ": " \
                  << cudaGetErrorString(err) << std::endl; \
        std::exit(EXIT_FAILURE); \
    } \
} while (0)
// Alias for an image pixel
using pixel_t = uint8_t;
// Kernel converting the red, green and blue images into a single gray image
__global__ void computeRGBToGray(const pixel_t* d_image_r, const pixel_t* d_image_g, const pixel_t* d_image_b, pixel_t* d_image_gray, int width, int height) {
    // Compute the thread global index in the grid
    const int x = threadIdx.x + blockIdx.x * blockDim.x;
    const int y = threadIdx.y + blockIdx.y * blockDim.y;
    // Boundary check selecting only threads within the image boundary
    if (x < width && y < height) {
        // Compute the thread index in the image
        const int i = x + y * width;
        // Convert from rgb to grayscale and store the result in global memory
        d_image_gray[i] = static_cast(0.299f * d_image_r[i] + 0.587f * d_image_g[i] + 0.114f * d_image_b[i]);
    }
}
// Kernel computing the median of each tile in the grayscale image
template 
__global__ void computeMedian(pixel_t *d_image_gray, pixel_t *d_median, int width, int height) {
    // Compute the thread global index in the grid
    const int x = threadIdx.x + blockIdx.x * blockDim.x;
    const int y = threadIdx.y + blockIdx.y * blockDim.y;
    // Boundary check selecting only threads within the image boundary
    if (!(x < width && y < height))
        return;
    // Allocate the shared memory in which we will store the tile
    __shared__ pixel_t tile[TILE_WIDTH * TILE_WIDTH];
    // Compute the thread index in the image
    const int index = x + y * width;
    // Load the tile's grayscale value from global memory into shared memory
    tile[index] = d_image_gray[index];
    // Synchronize to make sure all threads have loaded their data
    __syncthreads();
    // Sort the tile array using a single threaded bubble sort
    if (threadIdx.x == 0 && threadIdx.y == 0) {
        for (int i = 0; i < TILE_WIDTH * TILE_WIDTH; ++i)
            for (int j = i + 1; j < TILE_WIDTH * TILE_WIDTH; ++j)
                if (tile[i] > tile[j])
                    cuda::std::swap(tile[i], tile[j]);
        // Each thread block stores the median, found in the middle index after sorting, in the global median array
        const int medianIndex = (TILE_WIDTH * TILE_WIDTH) / 2;
        d_median[blockIdx.x + blockIdx.y * gridDim.x] = tile[medianIndex];
    }
}
int main() {
    // Define all the example constants
    constexpr auto TILE_WIDTH = 32;
    constexpr auto HISTO_SIZE = 256;
    constexpr auto NB_TILE_X = 250;
    constexpr auto NB_TILE_Y = NB_TILE_X;
    constexpr auto IMAGE_LENGTH = TILE_WIDTH * NB_TILE_X;
    constexpr auto IMAGE_SIZE = IMAGE_LENGTH * IMAGE_LENGTH;
    constexpr auto NB_IMAGES = 3;
    constexpr auto INIT_VALUE = 4;
    // Allocate the CPU memory to store the images tiles medians and for the red, green, blue and grayscale images
    std::vector<std::vector> h_images_r(NB_IMAGES, std::vector(IMAGE_SIZE, 4));
    std::vector<std::vector> h_images_g(NB_IMAGES, std::vector(IMAGE_SIZE, 4));
    std::vector<std::vector> h_images_b(NB_IMAGES, std::vector(IMAGE_SIZE, 4));
    std::vector<std::vector> h_images_gray(NB_IMAGES, std::vector(IMAGE_SIZE, 0));
    std::vector<std::vector> h_medians(NB_IMAGES, std::vector(NB_TILE_X * NB_TILE_Y));
    // Run the image processing pipeline for each image, in parallel
    #pragma omp parallel for
    for (int i = 0; i < NB_IMAGES; ++i)
    {
        pixel_t *d_image_r, *d_image_g, *d_image_b, *d_image_gray, *d_median;
        // Allocate the GPU memory for each container
        CUDA_CHECK_ERROR(cudaMalloc(&d_image_r, IMAGE_SIZE * sizeof(pixel_t)));
        CUDA_CHECK_ERROR(cudaMalloc(&d_image_g, IMAGE_SIZE * sizeof(pixel_t)));
        CUDA_CHECK_ERROR(cudaMalloc(&d_image_b, IMAGE_SIZE * sizeof(pixel_t)));
        CUDA_CHECK_ERROR(cudaMalloc(&d_image_gray, IMAGE_SIZE * sizeof(pixel_t)));
        CUDA_CHECK_ERROR(cudaMalloc(&d_median, (NB_TILE_X * NB_TILE_Y) * sizeof(pixel_t)));
        // Copy the memory of each container from CPU to GPU
        CUDA_CHECK_ERROR(cudaMemcpy(d_image_r, h_images_r[i].data(), IMAGE_SIZE * sizeof(pixel_t), cudaMemcpyHostToDevice));
        CUDA_CHECK_ERROR(cudaMemcpy(d_image_g, h_images_g[i].data(), IMAGE_SIZE * sizeof(pixel_t), cudaMemcpyHostToDevice));
        CUDA_CHECK_ERROR(cudaMemcpy(d_image_b, h_images_b[i].data(), IMAGE_SIZE * sizeof(pixel_t), cudaMemcpyHostToDevice));
        // Launch a GPU kernel to convert the RGB images to grayscale
        dim3 blockSize(TILE_WIDTH, TILE_WIDTH);
        dim3 gridSize(cuda::ceil_div(IMAGE_LENGTH, blockSize.x), cuda::ceil_div(IMAGE_LENGTH, blockSize.y));
        computeRGBToGray<<<gridsize, blocksize="">>>(d_image_r, d_image_g, d_image_b, d_image_gray, IMAGE_LENGTH, IMAGE_LENGTH);
        CUDA_CHECK_ERROR(cudaGetLastError());
        // Launch the GPU kernel to compute the median of every tile in the image
        computeMedian<tile_width, histo_size=""><<<gridsize, blocksize="">>>(d_image_gray, d_median, IMAGE_LENGTH, IMAGE_LENGTH);
        CUDA_CHECK_ERROR(cudaGetLastError());
        // Copy the GPU median memory back to the CPU
        CUDA_CHECK_ERROR(cudaMemcpy(h_medians[i].data(), d_median, (NB_TILE_X * NB_TILE_Y) * sizeof(pixel_t), cudaMemcpyDeviceToHost));
        // Free the GPU memory
        CUDA_CHECK_ERROR(cudaFree(d_image_r));
        CUDA_CHECK_ERROR(cudaFree(d_image_g));
        CUDA_CHECK_ERROR(cudaFree(d_image_b));
        CUDA_CHECK_ERROR(cudaFree(d_image_gray));
        CUDA_CHECK_ERROR(cudaFree(d_median));
    }
    return 0;
}
</gridsize,></tile_width,></gridsize,></std::vector</std::vector</std::vector</std::vector</std::vector

Đoạn mã này bắt đầu bằng cách định nghĩa hai kernel:

computeRGBToGray tải các giá trị của ba ảnh đầu vào đỏ, lục và lam để chuyển đổi và ghi chúng vào ảnh đầu ra thang độ xám.
computeMedian tính giá trị trung vị của mỗi khối trong ảnh xám đầu vào. Mỗi khối luồng tải khối dữ liệu từ bộ nhớ toàn cục vào bộ nhớ chia sẻ. Sau đó, một luồng đơn lẻ được sử dụng để sắp xếp mảng và ghi giá trị tìm thấy tại chỉ số giữa, tương ứng với giá trị trung vị, vào mảng đầu ra toàn cục chứa các giá trị trung vị.

Trong hàm main, sau khi định nghĩa các hằng số được sử dụng cho ví dụ, bộ nhớ CPU được cấp phát cho từng ảnh và cho các giá trị trung vị.

Quy trình xử lý hình ảnh sau đó được chạy cho từng trong ba hình ảnh, song song, bằng cách sử dụng OpenMP. Quy trình bắt đầu bằng việc phân bổ bộ nhớ cần thiết trên GPU trước khi chuyển dữ liệu từ CPU sang GPU. Hai kernel để chuyển đổi từ RGB sang thang độ xám và tính toán trung vị được khởi chạy sau đó. Cuối cùng, chúng tôi sao chép lại kết quả trung vị về CPU trước khi giải phóng bộ nhớ.

Mã này có một số khiếm khuyết sẽ được giải quyết, từng bước một.

1. Compute Sanitizer và API CCCL: Tìm lỗi dễ dàng và viết mã an toàn hơn

Hãy bắt đầu bằng cách chạy mã.

code_steps$ ./build/0_base_error_example
CUDA error in 0_base_error_example.cu at line 105: an illegal memory access was encountered

Mặc dù mã có một số kiểm tra lỗi, khi bạn nhận được thông báo lỗi như “truy cập bộ nhớ bất hợp pháp”, bạn nên bắt đầu bằng cách sử dụng compute-sanitizer để điều tra thêm.

Sử dụng Compute Sanitizer, bộ công cụ kiểm tra tính đúng đắn chức năng của NVIDIA, chúng ta có thể trực tiếp xác định một lỗi khó phát hiện:

$ compute-sanitizer ./build/0_base_error_example 
========= COMPUTE-SANITIZER
========= Invalid __shared__ write of size 1 bytes
=========     at void computeMedian<(int)32, (int)256>(unsigned char *, unsigned char *, int, int)+0x170 in 0_base_error_example.cu:55
=========     by thread (0,3,0) in block (20,0,0)
=========     Access at 0x6440 is out of bounds

Chạy đoạn mã trên trực tiếp cho thấy mã gặp phải lỗi ghi vượt quá phạm vi bộ nhớ chia sẻ tại dòng 55 của 0_base_error_example.cu.

tile[index] = d_image_gray[index];

Dòng này tải dữ liệu vào bộ nhớ chia sẻ một cách không chính xác bằng cách sử dụng chỉ mục toàn cục. Vì bộ nhớ chia sẻ được xác định ở cấp khối luồng, chúng ta cần thay đổi cách đánh chỉ mục. Để tránh các lỗi đánh chỉ mục, một API mới đã được giới thiệu trong CCCL để phân biệt giữa đánh chỉ mục toàn cục và đánh chỉ mục cấp khối. Để sử dụng nó, trước tiên bạn cần khởi chạy kernel của mình bằng API mới cuda::launch:

auto config = cuda::make_config(cuda::block_dims(...), cuda::grid_dims(...));
cuda::launch(stream, config, kernel_name<decltype(config)>, input)
</decltype(config)>

Sau đó, sử dụng API chỉ mục mới bên trong kernel:

template 
__global__ void kernel_name(Configuration config, ...) {
    // Retrieve and expand each global index
    const auto [x, y, z] = cuda::gpu_thread.index(cuda::grid, config);

    // Retrieve the block index structure (containing block_idx.x, .y, .z)
    const auto block_idx = cuda::gpu_thread.index(cuda::block, config);
}

Nếu không sử dụng compute-sanitizer hoặc API mới, lỗi này cũng có thể được phát hiện trực tiếp bằng cách sử dụng cuda::std::span hoặc biến thể n chiều của nó là cuda::std::mdspan thay vì các con trỏ thô. cuda::std::span và cuda::std::mdspan là các chế độ xem không sở hữu trên bộ nhớ liên tục và hữu ích để trừu tượng hóa chính xác vùng chứa. Truy cập dữ liệu thông qua các span an toàn hơn so với việc sử dụng các con trỏ thô, một phần vì ở chế độ gỡ lỗi, các truy cập vượt quá giới hạn sẽ kích hoạt một assertion.

Hạt nhân nên được cập nhật như sau:

// Alias for a 2-dimensional mdspan
template 
using span_2d = cuda::std::mdspan<t, cuda::std::dims<2="">>;

template 
__global__ void computeMedian(..., span_2d d_image_gray, ...)
</t,>

Nếu bạn chạy mã với những thay đổi này, bạn sẽ nhận được kết quả tương tự như sau:

$ ./build/1_span 
libcudacxx/include/cuda/std/__mdspan/mdspan.h:436: operator(): block: [16,0,0], thread: [0,30,0] Assertion `mdspan: operator() out of bounds access` failed.

Các truy cập bộ nhớ trong bộ nhớ chia sẻ cũng cần được bảo vệ bằng cách sử dụng cuda::shared_memory_mdspan như trong đoạn mã sau:

__shared__ pixel_t shared[TILE_WIDTH * TILE_WIDTH];
cuda::shared_memory_mdspan tile_2d(shared, TILE_WIDTH, TILE_WIDTH);

Bây giờ bạn có thể chạy và mọi thứ sẽ thực thi đúng cách mà không gặp lỗi.

Sử dụng API khởi chạy mới và cơ chế lập chỉ mục của nó, các khoảng cách vượt qua con trỏ thô, cùng với compute-sanitizer, các truy cập ngoài vùng biên giới hạn hoặc không xảy ra hoặc bị bắt ngay lập tức. Để biết thêm thông tin về compute-sanitizer, hãy xem Gỡ lỗi CUDA hiệu quả: Cách săn lùng lỗi với NVIDIA Compute Sanitizer.

2. Nsight Systems và NVTX: Đánh dấu hiệu năng mã của bạn một cách chính xác

Bây giờ mã nguồn đã không còn lỗi, nó sẵn sàng để được đo hiệu năng bằng NVIDIA Nsight Systems. Công cụ này cho phép bạn trực quan hóa dòng thời gian của chương trình: biết khi nào mỗi hàm được gọi và trong bao lâu.

Để dễ dàng trực quan hóa dòng thời gian hơn, chúng tôi bao bọc mọi phần mã nguồn thú vị bằng NVTX:

void image_compute(...)
{
  // NVTX range for the scope of the whole function
  nvtx3::scoped_range fun_scope("Image compute");

 // NVTX range that is pushed and then popped for a specific code section
  nvtxRangePushA("Kernel median");

  // Launch the GPU kernel to compute the median of every tile in the image
  ...

  // Pop the range at the end of the specific code section
  nvtxRangePop();
}

Điều này cho ra kết quả sau:

Hình 3. Dòng thời gian ban đầu của Nsight Systems với các chú thích NVTX. Các kernel tính toán trung vị chiếm ưu thế trong hồ sơ, chiếm gần như toàn bộ hoạt động của GPU và phần lớn thời gian xử lý hình ảnh 6,8 giây

Trong phần phần cứng GPU (CUDA HW) của kết quả bộ phân tích hiệu năng trong Hình 3 ở trên, nó được báo cáo rằng GPU chủ yếu bận rộn với các kernel (98,5% thời gian GPU) trong khi các thao tác bộ nhớ chỉ chiếm 1,5% thời gian GPU.

Trong số hai kernel, kernel tính toán trung vị chiếm phần lớn thời gian chạy với 2,1 giây cho mỗi hình ảnh (xem hộp màu vàng ở bên phải, nơi hiển thị thống kê cho computeMedian, và thời gian đã trôi qua là 2,142 giây).

Từ phần CPU (luồng), chúng ta thấy rằng việc tính toán hình ảnh mất tổng cộng 6,8 giây, với phần lớn thời gian được dành cho việc tính toán các giá trị trung vị cho ba ảnh xám.

Bây giờ chúng ta đã biết thao tác đầu tiên cần tối ưu hóa để tạo ra tác động lớn nhất. Để biết thêm thông tin về Nsight Systems, hãy xem Tối ưu hóa việc truyền dữ liệu bộ nhớ CUDA với NVIDIA Nsight Systems. Để biết thêm thông tin về NVTX, hãy xem Mẹo CUDA: Tạo các đường thời gian hồ sơ ứng dụng tùy chỉnh với NVTX.

3. CUB: Biểu diễn các thuật toán trực tiếp trên GPU

Khi xử lý các thuật toán phổ biến, việc viết các kernel tùy chỉnh dễ gây ra lỗi và rất có thể dẫn đến một cách triển khai kém hiệu quả. Bất cứ khi nào có thể, cả đối với các mẫu trên thiết bị và các nguyên thủy trong kernel, đều khuyến nghị sử dụng CUB.

CUB là thư viện thuật toán song song của NVIDIA, được phân phối thông qua CCCL. Nó cung cấp các hàm tối ưu hóa cao ở nhiều mức độ chi tiết: toàn bộ thiết bị (cub::Device*), cấp khối (cub::Block*), và cấp warp (cub::Warp*).

Đối với bước chuyển đổi từ RGB sang grayscale, chúng ta có thể thay thế kernel tùy chỉnh bằng cub::DeviceTransform::Transform. Nó áp dụng một hàm do người dùng cung cấp cho một bộ các iterator đầu vào và ghi kết quả vào một iterator đầu ra, thực thi trên GPU:

// Use CUB to convert the RGB images to grayscale
cub::DeviceTransform::Transform(
    cuda::std::make_tuple(d_image_r, d_image_g, d_image_b),  // inputs
    d_image_gray,                                            // output
    IMAGE_SIZE,                                              // size
    [] __host__ __device__ (pixel_t r, pixel_t g, pixel_t b) // functor
    {
        return static_cast(0.299f * r + 0.587f * g + 0.114f * b);
    },
    stream);

Đối với median, việc lập trình một phép sắp xếp song song cấp khối bằng tay là phức tạp và chậm. Thay vào đó, chúng tôi trực tiếp tận dụng phép sắp xếp radix cấp khối của CUB bên trong kernel:

// Declare and allocate the storage for CUB BlockRadixSort
using BlockRadixSort = cub::BlockRadixSort<...>;
__shared__ typename BlockRadixSort::TempStorage temp_storage;

// Load the tile's grayscale value from global memory
pixel_t thread_keys[1];
thread_keys[0] = d_image_gray(y, x);

// Perform the thread-block-level radix sort
BlockRadixSort(temp_storage).Sort(thread_keys);

// Select the thread found at the middle index
// Write its value which is, after sorting, the median, in the global median array
if (block_idx.x == TILE_WIDTH / 2 && block_idx.y == TILE_WIDTH / 2)
    d_median(grid_block_idx.y, grid_block_idx.x) = thread_keys[0];

Sau khi thay đổi này, chúng tôi tiếp tục chạy benchmark bằng Nsight Systems:

Hình 4. Dấu thời gian của Nsight Systems sau khi thay thế các thuật toán tùy chỉnh bằng CUB. Việc tính toán trung vị được giảm xuống còn micro giây và toàn bộ quy trình xử lý mất khoảng 635 mili giây, khiến việc phân bổ bộ nhớ trở thành chi phí chiếm ưu thế nhất

Thời gian cần thiết để tính toán trung vị hiện chỉ còn 773 micro giây (một lần nữa, hãy xem thời gian trôi qua trong hình ảnh nổi màu vàng của computeMedian), nhanh hơn 2717 lần. Tổng thời gian để tính toán cả ba hình ảnh hiện nay là 635 mili giây, nhanh hơn 10 lần.

Nếu chúng ta đánh giá lại nút thắt cổ chai hiện tại: thời gian dành cho việc phân bổ bộ nhớ chiếm khoảng 83% tổng thời gian chạy tính toán hình ảnh.

Điều này có thể được cải thiện đáng kể.

4. Các bộ chứa bộ nhớ gộp: Quản lý bộ nhớ tiện lợi và nhanh hơn

Việc cấp phát bộ nhớ GPU bằng cudaMalloc có thể gây ra những tác động tiêu cực ngoài ý muốn: rò rỉ bộ nhớ do quên gọi cudaFree và các thao tác bộ nhớ tốn kém ở những phần quan trọng trong mã của bạn.

Thay vào đó, chúng tôi khuyên bạn nên sử dụng các bộ chứa bộ nhớ bất đồng bộ của CCCL, cuda::device_buffer. Giống như std::vector trong C++, bộ nhớ sẽ tự động được giải phóng khi bộ chứa vượt ra khỏi phạm vi.

Ngoài ra, một bộ cấp phát bộ nhớ (memory pool) hỗ trợ cho bộ đệm, do đó các thao tác cấp phát và giải phóng bộ nhớ lặp lại sẽ không phải trả toàn bộ chi phí của cudaMalloc / cudaFree mỗi lần.

Để sử dụng các bộ chứa bộ nhớ GPU, chúng tôi cập nhật mã như sau:

// Resource to handle the GPU memory allocations
cuda::device_memory_pool_ref device_resource = cuda::device_default_memory_pool(cuda::device_ref{0});

// Explained at a later stage, unimportant for now
cuda::stream stream{cuda::device_ref{0}};

// Allocate the GPU memory using uninitialized containers
cuda::device_buffer d_image_r = cuda::make_buffer(stream, device_resource, IMAGE_SIZE, cuda::no_init);
...

Sau khi thay đổi này, chúng tôi phân tích lại timeline:

Hình 5. Dấu thời gian của Nsight Systems sau khi áp dụng các bộ chứa bộ nhớ thiết bị được chia sẻ. Chi phí phân bổ bộ nhớ gần như bị loại bỏ, khiến việc truyền dữ liệu từ bộ nhớ máy chủ sang thiết bị trở thành nút thắt cổ chai tiếp theo có thể quan sát được

Thời gian dành cho việc phân bổ bộ nhớ hiện gần như không đáng kể; thời gian cần thiết để tính toán một hình ảnh đã được cải thiện 2,6 lần.

Thời gian xử lý trên GPU hiện bị chi phối bởi bộ nhớ. Hầu hết thời gian để tính toán tất cả các hình ảnh đều dành cho việc sao chép màu đỏ, xanh lá cây và xanh dương, cho ba hình ảnh, từ CPU sang GPU.

Có thể tăng tốc đáng kể các phép truyền bộ nhớ từ thiết bị chủ sang thiết bị.

5. Bộ nhớ ghim: Chuyển đổi bộ nhớ từ host sang device nhanh hơn

Các vùng phân bổ dữ liệu của CPU mặc định là có thể hoán đổi (pageable), điều mà GPU không thể truy cập trực tiếp. Trình điều khiển CUDA trước tiên phải phân bổ một mảng host tạm thời được khóa trang, hay còn gọi là ghim (pinned), sao chép dữ liệu từ host vào mảng đã ghim, và sau đó chuyển dữ liệu từ mảng đã ghim sang bộ nhớ device.

Khi biết trước rằng bộ nhớ CPU sẽ được sao chép sang GPU, nên trực tiếp cấp phát bằng bộ nhớ ghim (pinned memory).

CCCL cung cấp một container chủ (host container) bộ nhớ ghim, cuda::host_buffer, có thể được tạo thông qua nhà máy cuda::make_pinned_buffer:

// Allocate the CPU memory to store the image tiles, medians, and for the red, green, blue, and grayscale images
// Those CPU containers, contrary to std::vector, are allocated using pinned memory
std::vector<cuda::host_buffer> h_images_r(NB_IMAGES, cuda::make_pinned_buffer(stream, IMAGE_SIZE, ...));
...
</cuda::host_buffer

Sau những thay đổi đó, chúng ta có thể chạy lại bài kiểm tra hiệu năng:

Hình 6. Dải thời gian của Nsight Systems sau khi phân bổ dữ liệu máy chủ trong bộ nhớ ghim; các phép chuyển đổi từ máy chủ sang thiết bị ngắn hơn nhiều và cả ba hình ảnh đều hoàn thành trong khoảng 25 mili giây, mặc dù các thao tác của chúng vẫn được thực thi tuần tự trên luồng mặc định

Thời gian thực hiện các phép chuyển bộ nhớ từ host sang device đã được giảm đáng kể; hiện chỉ mất 25 ms để tính toán tất cả các hình ảnh, nhanh hơn 10 lần.

Để biết thêm thông tin về bộ nhớ ghim (pinned memory), hãy xem Cách tối ưu hóa các phép chuyển dữ liệu trong CUDA C/C++.

Một hành vi gây bất ngờ có thể đã thu hút sự chú ý của một số độc giả ngay từ đầu:

Mặc dù chúng ta đang sử dụng các luồng CPU khác nhau, tất cả các thao tác (bộ nhớ và kernel) đều đang được thực thi tuần tự trên GPU.

Hãy cùng khắc phục nó.

6. Streams: Song song hóa các thao tác trên GPU

Theo mặc định, tất cả các thao tác (kernel, phân bổ bộ nhớ hoặc truyền dữ liệu) đều được khởi chạy trên những gì chúng tôi gọi là luồng mặc định: nó có thể được xem như một hàng đợi các tác vụ mà GPU cần thực hiện theo thứ tự.

Trong ví dụ này, chúng ta cần một luồng cho mỗi hình ảnh/luồng xử lý. CCCL cung cấp cuda::stream, một phiên bản tự quản lý, sở hữu của các luồng CUDA. Nó có thể được khởi tạo đơn giản bên trong vòng lặp song song, vì vậy mỗi luồng OpenMP sẽ có hàng đợi công việc GPU riêng của mình.

Để tận dụng hiệu quả các stream, chúng ta cũng cần sử dụng API bất đồng bộ: mỗi thao tác GPU được CPU khởi chạy không nên chờ đợi cho đến khi hoàn thành. Để làm đầy GPU, mỗi luồng CPU nên khởi chạy càng nhiều thao tác càng tốt, càng nhanh càng tốt, mà không chờ chúng hoàn thành trước đó. Các kernel và lệnh gọi thiết bị CUB đã là bất đồng bộ theo mặc định và được khởi chạy trên stream đã truyền. Để khởi chạy các bản sao bất đồng bộ giữa host và device, chúng ta sử dụng API mới cuda::copy_bytes của CCCL.

Trong bất kỳ mã CUDA hiện đại nào, khuyến nghị không bao giờ dựa vào stream mặc định mà luôn sử dụng các stream.

Chúng tôi cập nhật mã nguồn như sau:

Một init_stream chuyên dụng được sử dụng cho các lần phân bổ ban đầu của các bộ nhớ đệm host đã ghim. Mỗi vòng lặp của vòng lặp for song song giờ đây sở hữu một cuda::stream riêng cho đường ống tính toán:

// Stream used for initial host buffer allocations
cuda::stream init_stream{cuda::device_ref{0}};

...

// Allocate the host pinned buffers on init_stream:
std::vector<cuda::host_buffer> h_images_r(NB_IMAGES, cuda::make_pinned_buffer(init_stream, IMAGE_SIZE, ...));

...

// Sync before launching operations on another stream:
init_stream.sync();

#pragma omp parallel for
for (int i = 0; i < NB_IMAGES; ++i)
{
    // One different stream per thread
    cuda::stream stream{cuda::device_ref{0}};

    ...

    // GPU buffer allocations using the stream owned by each thread
    cuda::device_buffer d_image_r = cuda::make_buffer(stream, device_resource, IMAGE_SIZE, cuda::no_init);

    ...

    // Copy the memory of each container from CPU to GPU asynchronously using the stream owned by each thread
    cuda::copy_bytes(stream, h_images_r[i], d_image_r);
    
    ...

    // Use CUB to convert the RGB images to grayscale asynchronously using the per thread stream
    cub::DeviceTransform::Transform(..., stream.get());

    // Launch the GPU kernel to compute the median of every tile in the image using the per thread stream
    cuda::launch(stream, ...);

    // Copy the GPU median memory back to the CPU
    cuda::copy_bytes(stream, d_median, h_medians[i]);
    
    ...


    // To make sure the copy bytes is finished before accessing results on the host
    stream.sync();
}
</cuda::host_buffer

Sau những thay đổi đó, chúng ta có thể xem xét lần cuối cùng về dòng thời gian:

Hình 7. Dòng thời gian Nsight Systems cuối cùng sau khi gán một luồng CUDA cho mỗi hình ảnh; các phép truyền bộ nhớ và các kernel chồng chéo trên ba luồng, giảm tổng thời gian xử lý từ 6,8 giây xuống còn khoảng 23 mili giây

Bây giờ chúng ta đã có sự chồng lấp hoàn toàn giữa các kernel và các bản sao bộ nhớ.

Thời gian cuối cùng để tính toán cả ba hình ảnh sau tất cả các cải tiến của chúng ta là 23 mili giây, bắt đầu từ 6,8 giây.

Đến lượt bạn

Sử dụng Bộ công cụ cho nhà phát triển CUDA, chúng tôi đã giúp mã an toàn hơn, dễ bảo trì hơn và nhanh hơn. Không sử dụng các tối ưu hóa cấp thấp, nhưng mã vẫn nhanh hơn 300 lần.

Hãy tự thử nghiệm mã nguồn này và chạy trên Google Colab nếu bạn muốn.

Chúng tôi cũng đã xây dựng một lớp học đầy đủ để tìm hiểu cách sử dụng các công cụ này một cách chi tiết. Khóa học này có sẵn miễn phí trên YouTube cùng với các liên kết để thực hành trên Google Colab.

____
Bài viết liên quan