CUDA Cumulative Product Optimization Guide

__global__ void cumprod_naive(float* x, float* y, int n) {
    if (threadIdx.x == 0) {
        y[0] = x[0];
        for (int i = 1; i < n; i++)
            y[i] = y[i-1] * x[i];  // May overflow!
    }
}

#include <cub/cub.cuh>

struct MultOp {
    __device__ float operator()(float a, float b) { return a * b; }
};

void cumprod_opt(float* x, float* y, int n) {
    size_t temp_bytes = 0;
    MultOp mult_op;
    float init = 1.0f;

    cub::DeviceScan::InclusiveScan(
        nullptr, temp_bytes, x, y, mult_op, n);

    void* d_temp;
    cudaMalloc(&d_temp, temp_bytes);

    cub::DeviceScan::InclusiveScan(
        d_temp, temp_bytes, x, y, mult_op, n);
}