<commit_before>// This file is distributed under the MIT license.
// See the LICENSE file for details.

#pragma once

#ifndef VSNRAY_DETAIL_TILED_SCHED_H
#define VSNRAY_DETAIL_TILED_SCHED_H 1

#include "../math/detail/math.h"
#include "basic_sched.h"
#include "parallel_for.h"
#include "range.h"
#include "thread_pool.h"

namespace visionaray
{

struct int2
{
    int x;
    int y;
};

static thread_local int2 threadIdx;
static thread_local int2 blockIdx;
static thread_local int2 blockDim;
static thread_local int2 gridDim;

struct tiled_sched_backend
{
    explicit tiled_sched_backend(unsigned num_threads)
        : pool_(num_threads)
    {
    }

    void reset(unsigned num_threads)
    {
        pool_.reset(num_threads);
    }

    template <typename Func>
    void for_each_packet(
            tiled_range2d<int> const& tr,
            int packet_width,
            int packet_height,
            Func const& func
            )
    {
        blockDim.x = tr.rows().tile_size();
        blockDim.y = tr.cols().tile_size();

        gridDim.x = div_up(tr.rows().length(), tr.rows().tile_size());
        gridDim.y = div_up(tr.cols().length(), tr.cols().tile_size());

        visionaray::parallel_for(
            pool_,
            tr,
            [=](range2d<int> const& r)
            {
                blockIdx.x = r.rows().begin() / r.rows().length();
                blockIdx.y = r.cols().begin() / r.cols().length();

                for (int y = r.cols().begin(); y < r.cols().end(); y += packet_height)
                {
                    for (int x = r.rows().begin(); x < r.rows().end(); x += packet_width)
                    {
                        threadIdx.x = x % r.rows().length();
                        threadIdx.y = y % r.cols().length();

                        func(x, y);
                    }
                }
            });
    }

    thread_pool pool_;
};

template <typename R>
using tiled_sched = basic_sched<tiled_sched_backend, R>;

} // visionaray

#endif // VSNRAY_DETAIL_TILED_SCHED<commit_msg>Add launchDim and launchIdx to tiled scheduler<commit_after>// This file is distributed under the MIT license.
// See the LICENSE file for details.

#pragma once

#ifndef VSNRAY_DETAIL_TILED_SCHED_H
#define VSNRAY_DETAIL_TILED_SCHED_H 1

#include "../math/detail/math.h"
#include "basic_sched.h"
#include "parallel_for.h"
#include "range.h"
#include "thread_pool.h"

namespace visionaray
{

struct int2
{
    int x;
    int y;
};

static thread_local int2 threadIdx;
static thread_local int2 blockIdx;
static thread_local int2 launchIdx;
static thread_local int2 blockDim;
static thread_local int2 gridDim;
static thread_local int2 launchDim;

struct tiled_sched_backend
{
    explicit tiled_sched_backend(unsigned num_threads)
        : pool_(num_threads)
    {
    }

    void reset(unsigned num_threads)
    {
        pool_.reset(num_threads);
    }

    template <typename Func>
    void for_each_packet(
            tiled_range2d<int> const& tr,
            int packet_width,
            int packet_height,
            Func const& func
            )
    {
        blockDim.x = tr.rows().tile_size();
        blockDim.y = tr.cols().tile_size();

        gridDim.x = div_up(tr.rows().length(), tr.rows().tile_size());
        gridDim.y = div_up(tr.cols().length(), tr.cols().tile_size());

        launchDim.x = tr.rows().length();
        launchDim.y = tr.cols().length();

        visionaray::parallel_for(
            pool_,
            tr,
            [=](range2d<int> const& r)
            {
                blockIdx.x = r.rows().begin() / r.rows().length();
                blockIdx.y = r.cols().begin() / r.cols().length();

                for (int y = r.cols().begin(); y < r.cols().end(); y += packet_height)
                {
                    for (int x = r.rows().begin(); x < r.rows().end(); x += packet_width)
                    {
                        threadIdx.x = x % r.rows().length();
                        threadIdx.y = y % r.cols().length();

                        launchIdx.x = x;
                        launchIdx.y = y;

                        func(x, y);
                    }
                }
            });
    }

    thread_pool pool_;
};

template <typename R>
using tiled_sched = basic_sched<tiled_sched_backend, R>;

} // visionaray

#endif // VSNRAY_DETAIL_TILED_SCHED
{'fraction_non_alphanumeric': 0.12663454920853406, 'fraction_numerical': 0.0041293874741913286, 'mean_word_length': 1.90473017988008, 'pattern_counts': {'":': 0, '<': 17, '<?xml version=': 0, '>': 13, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 0, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '2124802', 'n_tokens_mistral': 1401, 'n_tokens_neox': 1358, 'n_words': 299}