C++ / SDL로 망델브로 집합 그려보기 – 7

	virtual void resize();

	inline std::complex<real_t> getPosition() const { return { pos_x, pos_y }; }
	void setPosition(real_t x, real_t y);
	virtual void move(int32_t rel_px, int32_t rel_py);
	...

	virtual void update(bool rerender_all = true, bool clear_surface = true);

#include "mandelbrot_tbb.h"
#include "mandelbrot_cuda.h"
...

	unique_ptr<Mandelbrot> mandelbrot = make_unique<MandelbrotCUDA>(renderer);

	enum class Acc {
		CPU      = 0,
		CPU_TBB  = 1,
		GPU_CUDA = 2
	};

#include "mandelbrot_tbb.h"
#include "mandelbrot_cuda.h"
#include "time.h"
...

	switch (settings.accelerator) {
	case Acc::CPU:
		mandelbrot = make_unique<Mandelbrot>(renderer);
		break;
	case Acc::CPU_TBB:
		mandelbrot = make_unique<MandelbrotTBB>(renderer);
		break;
	case Acc::GPU_CUDA:
		mandelbrot = make_unique<MandelbrotCUDA>(renderer);
		break;
	}
...

void GUI::showAcceleratorUI(std::unique_ptr<Mandelbrot>& mandelbrot)
{
	if (ImGui::CollapsingHeader("accelerator")) {
		static const char* items[] = { "CPU", "CPU - TBB", "GPU - CUDA" };
		
		ImGui::Text("accelerator :");
		if (ImGui::Combo(IMGUI_NO_LABEL, (int*)&settings.accelerator, items, 3))
			acceleratorChanged(mandelbrot);

		if (settings.accelerator == Acc::CPU_TBB) {
			auto* man_tbb = dynamic_cast<MandelbrotTBB*>(mandelbrot.get()); // will not fail

			auto value = (int)man_tbb->getMaxConcurrency();
			ImGui::Text("max concurrency:");
			if (ImGui::InputInt(IMGUI_NO_LABEL, &value, 1)) {
				value = SDL_clamp(value, 1, thread::hardware_concurrency());
				man_tbb->setMaxConcurrency(value);
			}
		} else if (settings.accelerator == Acc::GPU_CUDA) {
			static const char* sizes[] = { "1x1", "2x2", "4x4", "8x8", "16x16" };

			auto* man_cuda = dynamic_cast<MandelbrotCUDA*>(mandelbrot.get()); // will not fail

			auto block_size = (int)log2(man_cuda->getBlockSize());
			ImGui::Text("block size:");
			if (ImGui::Combo(IMGUI_NO_LABEL, &block_size, sizes, 5))
				man_cuda->setBlockSize(pow(block_size + 1, 2));
		}
	}
}

	inline uint32_t getBlockSize() const { return block_size; }
	inline void setBlockSize(uint32_t size) { block_size = size; }

	inline uint32_t getTotalSample() const { return sample_total; }
	void setTotalSample(uint32_t sample);

	inline uint32_t getSamplePerLaunch() const { return sample_per_launch; }
	inline void setSamplePerLaunch(uint32_t sample) { sample_per_launch = sample; }

	inline uint32_t getSampleCount() const { return sample_count; }
...

	uint32_t block_size;
	uint32_t sample_total;
	uint32_t sample_per_launch;
	uint32_t sample_count;

	struct PixelInfo {
		union {
			struct {
				uint32_t sample_count;
				uint32_t acc_r;
				uint32_t acc_g;
				uint32_t acc_b;
			};
			bool rendered = false;
		};
	};

template <class T>
__device__ inline uint32_t mandelbrot(T& cx, T& cy, uint32_t max_iter);
__device__ inline uint32_t make_color(uint8_t r, uint8_t g, uint8_t b, uint8_t a);
__device__ inline uint32_t lerp_color(uint32_t a, uint32_t b, double t);

__device__ inline void acc_color(uint32_t col, uint32_t& r, uint32_t& g, uint32_t& b) {
	struct Color { uint8_t b, g, r, a; } color = *(Color*)&col;

	r += color.r;
	g += color.g;
	b += color.b;
}

static __device__ __forceinline__ uint32_t lcg(uint32_t& prev)
{
	prev = (1664525u * prev + 1013904223u);
	return prev & 0x00FFFFFF;
}

static __device__ __forceinline__ float rnd(uint32_t& prev)
{
	return ((float)lcg(prev) / (float)0x01000000);
}
...

template <class T>
__device__ uint32_t get_color_idx(T iter) {
	return (uint32_t)(params.color_scale * 256 * iter / params.iter) % 256;
}

__global__ void mandelbrot_kernel(
	uint32_t* device_surface,
	PixelInfo* device_pixel_info,
	uint32_t   sample_total,
	uint32_t   sample,
	real_t     min_x,
	real_t     max_y,
	real_t     dp
) {
	if (params.stop_all) return;

	int w = blockDim.x * blockIdx.x + threadIdx.x;
	int h = blockDim.x * blockIdx.y + threadIdx.y;

	if (params.width <= w || params.height <= h) return;

	auto pixel_off  = h * params.width + w;
	auto& info      = *(device_pixel_info + pixel_off);

	if (info.sample_count >= sample_total) return;
 
	auto seed  = (uint32_t)clock();
	auto count = min(sample_total - info.sample_count, sample);

	uint32_t col;
	uint32_t r = info.acc_r;
	uint32_t g = info.acc_g;
	uint32_t b = info.acc_b;

	for (uint32_t i = 0; i < count; ++i) {
		if (params.stop_all) return;
		
		real_t cx = min_x + dp * (w + rnd(seed));
		real_t cy = max_y - dp * (h + rnd(seed));

		auto iterated = mandelbrot<real_t>(cx, cy, params.iter);

		if (iterated == params.iter) {
			col = 0xff000000;
		} else if (!params.color_smooth) {
			col = cuda_colormap[params.color_idx][get_color_idx(iterated)];
		} else {
			double log_zn  = log(cx * cx + cy * cy) / 2.;
			double nu      = log(log_zn / log(2.)) / log(2.);
			auto real_iter = iterated + 3.5 - nu;

			auto col1 = cuda_colormap[params.color_idx][get_color_idx(real_iter)];
			auto col2 = cuda_colormap[params.color_idx][get_color_idx(real_iter + 1.)];

			col = lerp_color(col1, col2, fmod(real_iter, 1.));
		}

		acc_color(col, r, g, b);
	}

	info.sample_count += count;
	info.acc_r         = r;
	info.acc_g         = g;
	info.acc_b         = b;

	auto& pixel = *(device_surface + pixel_off);
	pixel = make_color(r / info.sample_count, g / info.sample_count, b / info.sample_count);
}

	block_size        = 8;
	sample_total      = 1;
	sample_per_launch = 1;
	sample_count      = 0;
...

void MandelbrotCUDA::setTotalSample(uint32_t sample)
{
	auto clear = sample_count > (sample_total = sample);
	update(clear, clear);
}
...

void MandelbrotCUDA::drawSurface()
{
	dim3 grid((width - 1) / block_size + 1, (height - 1) / block_size + 1);
	dim3 block(block_size, block_size);

	real_t min_x = pos_x - 2. * scale * aspect;
	real_t max_y = pos_y + 2. * scale;
	real_t dp    = 4. * scale / height;

	while (sample_count < sample_total) {
		if (stop_all) break;

		mandelbrot_kernel<<<grid, block, 0, streams[1]>>> (
			device_surface,
			device_pixel_info,
			sample_total,
			sample_per_launch,
			min_x, max_y, dp);

		sample_count = min(sample_count + sample_per_launch, sample_total);
		cudaStreamSynchronize(streams[1]);
	}

	cudaStreamSynchronize(streams[1]);
}

void MandelbrotCUDA::update(bool rerender_all, bool clear_surface)
{
	Mandelbrot::update(rerender_all, clear_surface);
	
	size_t size = surface->w * surface->h;

	if (rerender_all) {
		cudaMemset(device_pixel_info, 0, size * sizeof(PixelInfo));
		sample_count = 0;
	}
	if (clear_surface)
		cudaMemset(device_surface, 0, size * sizeof(uint32_t));

	Constants constants = {
		width, height, iter, color_idx,color_scale, smooth, stop_all
	};

	cudaMemcpyToSymbol(params, &constants, sizeof(Constants));
}

			static const char* sizes[] = { "1x1", "2x2", "4x4", "8x8", "16x16" };

			auto* man_cuda = dynamic_cast<MandelbrotCUDA*>(mandelbrot.get()); // will not fail

			auto block_size = (int)log2(man_cuda->getBlockSize());
			ImGui::Text("block size:");
			if (ImGui::Combo(IMGUI_NO_LABEL, &block_size, sizes, 5))
				man_cuda->setBlockSize(pow(block_size + 1, 2));

			auto sample_total = (int)man_cuda->getTotalSample();
			ImGui::Text("total sample:");
			if (ImGui::InputInt(IMGUI_NO_LABEL, &sample_total))
				man_cuda->setTotalSample(sample_total = max(sample_total, 1));
			
			auto spl = (int)man_cuda->getSamplePerLaunch();
			ImGui::Text("sample per launch:");
			ImGui::InputInt(IMGUI_NO_LABEL, &spl);
			man_cuda->setSamplePerLaunch(spl = SDL_clamp(spl, 1, sample_total));

			ImGui::Text("sampled: %d", man_cuda->getSampleCount());