Kết hợp quản lý bộ nhớ tùy chỉnh và Lực đẩy trong CUDA

Question

Trong dự án của tôi, tôi đã triển khai bộ cấp phát bộ nhớ tùy chỉnh để tránh các cuộc gọi không cần thiết đến cudaMalloc sau khi ứng dụng đã "khởi động". Hơn nữa, tôi sử dụng hạt nhân tùy chỉnh để điền mảng cơ bản, hoạt động số học giữa các mảng, v.v. và muốn đơn giản hóa mã của tôi bằng cách sử dụng Thrust và loại bỏ các hạt nhân này. Mỗi mảng trên thiết bị được tạo và truy cập thông qua các con trỏ thô (bây giờ) và tôi muốn sử dụng các phương thức device_vector và Thrust s trên các đối tượng này, nhưng tôi thấy mình chuyển đổi giữa các con trỏ thô và device_ptr<> mọi lúc, hơi lộn xộn lên mã.

Câu hỏi khá mơ hồ của tôi: Bạn sẽ tổ chức việc sử dụng quản lý bộ nhớ tùy chỉnh như thế nào, các phương pháp mảng và cuộc gọi đến hạt nhân tùy chỉnh theo cách dễ đọc nhất là Thrust?

Answer 1

Giống như tất cả các vùng chứa tiêu chuẩn C++, bạn có thể tùy chỉnh cách thrust::device_vector phân bổ bộ nhớ bằng cách cung cấp bộ nhớ bằng "allocator" của riêng bạn. Theo mặc định, phân bổ của thrust::device_vector là thrust::device_malloc_allocator, phân bổ lưu trữ (deallocates) với cudaMalloc (cudaFree) khi hệ thống phụ trợ của Thrust là CUDA.

Thỉnh thoảng, bạn có thể tùy chỉnh cách device_vector phân bổ bộ nhớ, chẳng hạn như trong trường hợp OP, người muốn phân bổ lưu trữ trong một phân bổ lớn được thực hiện khi khởi tạo chương trình. Điều này có thể tránh được chi phí mà có thể phát sinh bởi nhiều cuộc gọi riêng lẻ đến sơ đồ phân bổ cơ bản, trong trường hợp này là cudaMalloc.

Cách đơn giản để cung cấp device_vector trình phân bổ tùy chỉnh là được kế thừa từ device_malloc_allocator. Người ta có thể về nguyên tắc tác giả một phân bổ toàn bộ từ đầu, nhưng với cách tiếp cận kế thừa, chỉ cần cung cấp các chức năng thành viên allocate và deallocate. Khi phân bổ tùy chỉnh được xác định, nó có thể được cung cấp cho device_vector làm tham số mẫu thứ hai của nó.

Ví dụ này đang chứng minh làm thế nào để cung cấp một cấp phát tùy chỉnh mà in một thông điệp khi phân bổ và deallocation:

#include <thrust/device_malloc_allocator.h> 
#include <thrust/device_vector.h> 
#include <iostream> 

template<typename T> 
    struct my_allocator : thrust::device_malloc_allocator<T> 
{ 
    // shorthand for the name of the base class 
    typedef thrust::device_malloc_allocator<T> super_t; 

    // get access to some of the base class's typedefs 

    // note that because we inherited from device_malloc_allocator, 
    // pointer is actually thrust::device_ptr<T> 
    typedef typename super_t::pointer pointer; 

    typedef typename super_t::size_type size_type; 

    // customize allocate 
    pointer allocate(size_type n) 
    { 
    std::cout << "my_allocator::allocate(): Hello, world!" << std::endl; 

    // defer to the base class to allocate storage for n elements of type T 
    // in practice, you'd do something more interesting here 
    return super_t::allocate(n); 
    } 

    // customize deallocate 
    void deallocate(pointer p, size_type n) 
    { 
    std::cout << "my_allocator::deallocate(): Hello, world!" << std::endl; 

    // defer to the base class to deallocate n elements of type T at address p 
    // in practice, you'd do something more interesting here 
    super_t::deallocate(p,n); 
    } 
}; 

int main() 
{ 
    // create a device_vector which uses my_allocator 
    thrust::device_vector<int, my_allocator<int> > vec; 

    // create 10 ints 
    vec.resize(10, 13); 

    return 0; 
} 

Dưới đây là kết quả:

$ nvcc my_allocator_test.cu -arch=sm_20 -run 
my_allocator::allocate(): Hello, world! 
my_allocator::deallocate(): Hello, world! 

Trong ví dụ này, lưu ý rằng chúng ta nghe từ my_allocator::allocate() một lần khi số vec.resize(10,13). my_allocator::deallocate() được gọi một lần khi vec vượt quá phạm vi vì nó phá hủy các phần tử của nó.