Тепловая карта по палитре с использованием CUDA

Question

Я использую CUDA для рендеринга (используя raytracer) сцены на экран и хочу узнать, какие самые горячие точки на экране: я измеряю разницу между значениями, возвращаемыми clock64, чтобы узнать время выполнения для каждого пикселя наscreen:

float start = clock64();
frame[y * w + x] = TraceRay(x, y, w, h);
counters[y * w + x] = clock64() - start;`

В настоящее время я делаю что-то вроде тепловой карты:

auto p = thrust::cuda::par.on(stream);
thrust::device_ptr< const float > c = thrust::device_pointer_cast(counters);
auto m = thrust::minmax_element(p, c, c + w * h);
thrust::device_ptr< Color > f = thrust::device_pointer_cast(frame);
#ifndef __CUDACC_EXTENDED_LAMBDA__
#error "nvcc --expt-extended-lambda"
#endif
auto l = [=] __device__ (float c) -> Color
{
    auto color = (c - *m.first) / float(*m.second - *m.first);
    return {color, 0.0f, 1.0f - color, 1.0f};
};
thrust::transform(p, c, c + w * h, f, l);

Но распределение полученных значений слишком широкое, и большинство из них синие и некоторые маленькиеПодмножество наиболее интересных пикселей стало почти красным (самые горячие точки, которые сложнее всего рассчитать).Ценного градиента почти нет.

Также я хочу определить тепловую карту с помощью палитры с определением цветов в опорных точках (процентили для значений counters).Как это можно сделать в gnuplot:

Я думаю, что смогу ранг всех значений в counters ипримените к ним палитру:

• Прежде всего мне нужно отсортировать все значения из counters, ранее расширенные ими по (пустому) color полю и по координатам (x, y) (или просто по индексу в исходном линейном массиве).В компараторе сортировки должно участвовать только значение counters.

• После этого должна быть применена палитра.Простое присвоение полю color значения кусочно-линейной интерполяции цветов, заданных из палитры, сопоставленной с интервалом [0;1], а затем сопоставленной с линейным индексом отсортированного массива.

• В конце концов я могу нарисовать все color s, используя x и y координаты (или отсортировать их обратно по полю линейного индекса).

Есть ли местодля CUDA в реализации алгоритма?Как я знаю, существует радикальная сортировка, но применима ли она к структурам с key, не охватывающими всю структуру?

Answer 1

Следующее решение не является идеальным (одного из сортов можно полностью избежать в пользу перестановки; память можно повторно использовать вместо перераспределения в каждом следующем кадре; палитра может быть расширена), но вполне работоспособна. Он рисует в оттенках красного только верхние 5% большинства тяжелых блоков. Все остальные изображены в оттенках серого и голубого.

__global__
void drawHeatmap(unsigned int w, unsigned int h, const Color * heatmap, cudaSurfaceObject_t frame)
{
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
    if ((x >= w) || (y >= h)) {
        return;
    }
    surf2Dwrite< Color >(heatmap[w * y + x], frame, sizeof(Color) * x, y);
}

void CudaRaytracer::buildHeatmap(cudaStream_t stream,
                                 unsigned int w, unsigned int h,
                                 float * counters,
                                 cudaSurfaceObject_t frame)
{
    assert(counters);
    auto p = thrust::cuda::par.on(stream);
    thrust::device_ptr< float > c = thrust::device_pointer_cast(counters);
    const auto size = w * h;
    thrust::device_vector< unsigned int > indices(size);
    thrust::sequence(p, indices.begin(), indices.end());
    thrust::sort_by_key(p, c, c + size, indices.begin());
#ifndef __CUDACC_EXTENDED_LAMBDA__
#error "nvcc --expt-extended-lambda"
#endif
    auto make_palette = [=] __device__ (unsigned int index) -> Color
    {
        constexpr unsigned int palette_size = 3;
        static const float reference_points[palette_size] = {0.0f, 0.95f, 1.0f};
        float pos = index / float(size);
        unsigned int i = 0;
        for (; i < palette_size; ++i) {
            if (pos < reference_points[i]) {
                break;
            }
        }
        __syncwarp();
        static const Color palette[palette_size] = {{0.0f, 1.0f, 0.0f, 1.0f}, {0.0f, 0.0f, 1.0f, 1.0f}, {1.0f, 0.0f, 0.0f, 1.0f}};
        float weight = (pos - reference_points[i - 1]) / (reference_points[i] - reference_points[i - 1]);
        return weight * (palette[i - 1] - palette[i]) + palette[i];
    };
    thrust::device_vector< Color > heatmap(size);
    auto index = thrust::make_counting_iterator(0u);
    thrust::transform(p, index, index + size, heatmap.begin(), make_palette);
    thrust::sort_by_key(p, indices.begin(), indices.end(), heatmap.begin());
    auto gridSize = deriveGridSize(w, h);
    drawHeatmap<<< gridSize, blockSize, 0, stream >>>(w, h, heatmap.data().get(), frame);
    cudaStreamSynchronize(stream);
}

В GeForce RTX 2060 рендеринг составляет около 30 мс.