CUB сегментированное сокращение не дает результатов

Question

Я пытаюсь использовать примитив суммы сегментированных сокращений CUB, и я застрял на этом.

Вот мой код:

int main() {


     const int N = 7;
     const int num_segments  = 3;
     int d_offsets[]= {0,3,3,7};


    int *h_data       = (int *)malloc(N * sizeof(int));
    int *h_result = (int *)malloc(num_segments * sizeof(int));


    for (int i=0; i<N; i++) {
        h_data[i] = 3;

    }


    int *d_data;
    cudaMalloc((int**)&d_data, N * sizeof(int));
    cudaMemcpy(d_data, h_data, N * sizeof(int), cudaMemcpyHostToDevice);


    int           *d_result;
    cudaMalloc((int**)&d_result, num_segments * sizeof(int));

    void            *d_temp_storage = NULL;
    size_t          temp_storage_bytes = 0;


    cudaMalloc((void**)&d_temp_storage, temp_storage_bytes);


    cub::DeviceSegmentedReduce::Sum(d_temp_storage, temp_storage_bytes, d_data, d_result,
        num_segments, d_offsets, d_offsets + 1);


    cudaMemcpy(h_result, d_result, num_segments*sizeof(int), cudaMemcpyDeviceToHost);




    printf("Results:\n");

   for (int i=0; i<num_segments; i++) {
        printf("CUB: %d\n", h_result[i]);

    }


}

Но в результате я получил это:

Results:
CUB: 0
CUB: 0
CUB: 0

Я не могу понять, в чем именно проблема. В реальном примере у меня очень большой массив с сегментами, равными 400. Можно ли оптимизировать код так, чтобы мне не нужно было объявлять и выделять память для d_offsets,

Answer 1

Вы действительно серьезно не пытались отладить свой код:

• Вы пропустили выделение памяти для d_results (вы это исправили)
• Вы пытались передать адрес памяти хоста для адреса памяти устройства в d_offsets, Конечно, это приводит к ошибке во время выполнения CUDA - но
• Вы не проверяли ошибки во время выполнения.
• Вы вызывали функцию CUB только один раз - хотя вы должны запустить ее дважды, чтобы она действительно что-то делала: один раз с nullptr в качестве пространства для царапин, чтобы получить размер пространства для царапин, а затем снова с реальным пространством для царапин, чтобы сделать работу. Это надоедливый API, но так оно и работает.

Вам неуместно тратить время сообщества SO на отладку кода, если вы сами не потратили на это время.

Тем не менее, есть кое-что, что вы могли бы сделать, чтобы избежать необходимости проверять ошибки, по крайней мере, использовать какую-то библиотеку, которая делает это за вас (например, с помощью ошибки). Если вы сделали это - например, используя мои обертки CUDA Runtime API (извините за самоподключение) и правильно распределили память для всего, что вам нужно, вы бы получили что-то вроде этого:

#include <cub/cub.cuh>
#include <cuda/api_wrappers.h>
#include <vector>
#include <cstdlib>

int main() {

    const int N = 7;
    const int num_segments  = 3;
    auto h_offsets = std::vector<int> {0,3,3,7};

    auto h_data = std::vector<int>(N);
    auto h_results = std::vector<int>(num_segments);

    std::fill(h_data.begin(), h_data.end(), 3);

    auto current_device = cuda::device::current::get();
    auto d_offsets = cuda::memory::device::make_unique<int[]>(
        current_device, h_offsets.size());
    auto d_data = cuda::memory::device::make_unique<int[]>(
        current_device, N);
    cuda::memory::copy(
        d_offsets.get(), &h_offsets[0], h_offsets.size() * sizeof(int));
    cuda::memory::copy(
        d_data.get(),  &h_data[0], h_data.size() * sizeof(int));
    auto d_results = cuda::memory::device::make_unique<int[]>(
        current_device, num_segments);

    auto d_start_offsets = d_offsets.get();
    auto d_end_offsets = d_start_offsets + 1; // aliasing, see CUB documentation

    size_t temp_storage_bytes = 0;

    // This call merely obtains a value for temp_storage_bytes, passed here
    // as a non-const reference; other arguments are unused
    cub::DeviceSegmentedReduce::Sum(
        nullptr, temp_storage_bytes, d_data.get(), d_results.get(),
        num_segments, d_start_offsets, d_end_offsets);

    auto d_temp_storage = cuda::memory::device::make_unique<char[]>(
        current_device, temp_storage_bytes);

    cub::DeviceSegmentedReduce::Sum(
        d_temp_storage.get(), temp_storage_bytes, d_data.get(), 
        d_results.get(), num_segments, d_start_offsets, d_end_offsets);

    cuda::memory::copy(
        &h_results[0], d_results.get(), num_segments * sizeof(int));

    std::cout << "Results:\n";

    for (int i=0; i<num_segments; i++) {
        std::cout << "Segment " << i << " data sums up to " << h_results[i] << "\n";
    }

    return EXIT_SUCCESS;
}

который работает:

Results:
Segment 0 data sums up to 9
Segment 1 data sums up to 0
Segment 2 data sums up to 12

Дополнительные советы:

• Всегда изучайте предупреждения компилятора.
• использование cuda-memcheck чтобы избежать утечек памяти / инициализации на неправильном устройстве / стороне хоста и т. д.
• Если вы используете API CUDA Runtime напрямую, вы должны проверять каждый вызов на наличие ошибок.