cvl-robot's diary

研究ノート メモメモ https://github.com/dotchang/

openFrameworksのshaderでGLSL4.2のatomic counter bufferを使って画素の数をカウントしてみる(その3)

普段使用している開発環境ではnVidiaグラフィックカードを使用しているのですが、intel製グラフィックスチップの環境にこのコードを持っていくと動かないことが分かりました。
原因を調査したところ、shader側ではなくopengl側のbufferのbindの仕方が問題のようです。

software.intel.com

一部引用すると、

// create and configure the Atomic Counter Buffer
glGenBuffers(1, &acb);                                                       
glBindBufferBase(GL_ATOMIC_COUNTER_BUFFER, 0, acb);                          
glBufferData(GL_ATOMIC_COUNTER_BUFFER, nCounters * 4, NULL, GL_STATIC_COPY); 
glBindBufferBase(GL_ATOMIC_COUNTER_BUFFER, 0, acb);                          

glBindBufferではなく、glBindBufferBaseを使っていることが分かります。

書式を習って書き直すと、こんな感じですね。こちらの書き方なら、nVidiaでもintelでも動きました。

// アトミックカウンタバッファの作成
void create(int sz, int idx, int off)
{
	// パラメータの設定
	size = sz;        // バッファのサイズ
	index = idx;      // マウント先のインデックス
	offset = off;     // マウント先のメモリオフセット
	data.resize(sz);

	// バッファの作成
	// create and configure the Atomic Counter Buffer
	glGenBuffers(1, &buffer);                                                              
	glBindBufferBase(GL_ATOMIC_COUNTER_BUFFER, index, buffer);                                 
	glBufferData(GL_ATOMIC_COUNTER_BUFFER, size * sizeof(GLuint), NULL, GL_DYNAMIC_DRAW);
	glBindBufferBase(GL_ATOMIC_COUNTER_BUFFER, index, buffer);
}

# でもAMDだと動かないとかいうオチがありそう。

openFrameworksのshaderでGLSL4.2のatomic counter bufferを使って画素の数をカウントしてみる(その2)

大量のatomic counterを使って色分布のヒストグラムを作りたい、と考えたときに普通の発想であれば配列化してしまえば良い、と考えますが、資料[1]によるとshader内でそのような書き方はできない、とあります。
どういうことなのか、見てみましょう。

まず、openGL側の関数を引数を足して、少し汎用化します。

// from here: http://www.lighthouse3d.com/tutorials/opengl-atomic-counters/

void create_a_buffer_for_atomic_counters(GLuint& buffer, int size = 1, int index = 0)
{
	glGenBuffers(1, &buffer);
	// bind the buffer and define its initial storage capacity
	glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, buffer);
	glBufferData(GL_ATOMIC_COUNTER_BUFFER, sizeof(GLuint) * size, NULL, GL_DYNAMIC_DRAW);
	// unbind the buffer 
	glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, index);
}

void reset_the_atomic_counter_buffers(GLuint& buffer, int size = 1, int index = 0) {
	GLuint *userCounters;
	glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, buffer);
	// map the buffer, userCounters will point to the buffers data
	userCounters = (GLuint*)glMapBufferRange(GL_ATOMIC_COUNTER_BUFFER,
		0,
		sizeof(GLuint) * size,
		GL_MAP_WRITE_BIT | GL_MAP_INVALIDATE_BUFFER_BIT | GL_MAP_UNSYNCHRONIZED_BIT
	);
	// set the memory to zeros, resetting the values in the buffer
	memset(userCounters, 0, sizeof(GLuint) * size);
	// unmap the buffer
	glUnmapBuffer(GL_ATOMIC_COUNTER_BUFFER);

	glBindBufferBase(GL_ATOMIC_COUNTER_BUFFER, index, buffer); // added important
}

void simpler_reset_the_atomic_counter_buffers(GLuint& buffer, int size = 1, int index = 0, int offset = 0) {
	glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, buffer);

	GLuint *a = new GLuint[size];
	memset(a, 0, size * sizeof(GLuint));
	glBufferSubData(GL_ATOMIC_COUNTER_BUFFER, offset, sizeof(GLuint) * size, a);
	glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, index);

	glBindBufferBase(GL_ATOMIC_COUNTER_BUFFER, index, buffer); // added important
	delete[] a;
8}

void read_back_the_values_from_the_buffer(GLuint& acb, int* userVariables, int size = 1, int index = 0, int offset = 0)
{
	GLuint *userCounters;
	glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, acb);
	// again we map the buffer to userCounters, but this time for read-only access
	userCounters = (GLuint*)glMapBufferRange(GL_ATOMIC_COUNTER_BUFFER,
		(GLintptr)offset,
		sizeof(GLuint) * size,
		GL_MAP_READ_BIT
	);

	// copy the values to other variables because...
	for (int i = 0; i < size; i++) {
		userVariables[i] = (int)userCounters[i];
	}
	// ... as soon as we unmap the buffer
	// the pointer userCounters becomes invalid.
	glUnmapBuffer(GL_ATOMIC_COUNTER_BUFFER);
}

void alternative_read_back_the_values_from_the_buffer(GLuint& buffer, int* userVariables, int size = 1, int offset = 0) {
	GLuint *userCounters = new GLuint[size];
	glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, buffer);
	glGetBufferSubData(GL_ATOMIC_COUNTER_BUFFER, offset, sizeof(GLuint) * size, userCounters);
	glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, 0);
	for (int i = 0; i < size; i++) {
		userVariables[i] = (int)userCounters[i];
	}
}

次に、shaderのプログラムにヒストグラムカウントを書き足します。このコードは、正しくありませんので実行する必要はありません。

	string fragmentShaderProgram = STRINGIFY(
		#version 450 compatibility\n
		#extension GL_ARB_shader_atomic_counters : enable\n
		#extension GL_EXT_gpu_shader4 : enable\n

		layout(binding = 0, offset = 0) uniform atomic_uint atRed;
	layout(binding = 0, offset = 4) uniform atomic_uint atGreen;
	layout(binding = 0, offset = 8) uniform atomic_uint atBlue;

	layout(binding = 1, offset = 0) uniform atomic_uint myIntReds[256]; // このような書き方はできない

	uniform sampler2DRect tex0;

	in vec4 gl_FragCoord;
	out vec4 colorOut;

	void main() {
		vec2 pos = vec2(gl_FragCoord.x, gl_FragCoord.y); // テクスチャ上の座標を取得する

		float r = texture2DRect(tex0, pos).r;
		float g = texture2DRect(tex0, pos).g;
		float b = texture2DRect(tex0, pos).b;
		float a = texture2DRect(tex0, pos).a;

		if (r > 0.5) atomicCounterIncrement(atRed);
		if (g > 0.5) atomicCounterIncrement(atGreen);
		if (b > 0.5) atomicCounterIncrement(atBlue);

		int myIntRed = int(255.0 * r);
		atomicCounterIncrement(myIntReds[myIntRed]);    // これが動かない
		
		colorOut = vec4(r, g, b, a);
	}
	);

テスト用にコードをヒストグラム読み出し部分を適当に少し足します。

	alternative_read_back_the_values_from_the_buffer(myIntReds, reds, 256, 0);
	simpler_reset_the_atomic_counter_buffers(myIntReds, 256, 1, 0);

	int total_r = 0;
	cout << "rgb count=" << cnt[0] << ", " << cnt[1] << ", " << cnt[2] << std::endl;
	for (int i = 0; i < 256; i++) {
		cout << reds[i] << " ";
		total_r += reds[i];
	}
	cout << endl;
	cout << "total_r = " << total_r << endl;

この実行結果は次のようになります。

rgb count=376021, 307367, 232583
7892 3783 3359 3061 2629 1981 1615 1382 1459 1467 1248 1096 1057 1115 1235 1385
1550 1586 1718 1687 1682 1835 1935 1732 2161 2389 2205 1987 1767 1834 2068 2172
2262 2346 2496 2424 2433 2307 2201 2100 2310 2432 2833 2850 2733 2635 2548 2506
2640 2545 2430 2519 2800 3017 2999 2913 2957 2863 2797 2919 2741 2629 2606 2882
3006 3142 3070 3010 2866 2852 2791 2738 2631 2614 2952 3162 3255 3218 3281 3240
3262 3121 2968 3045 2985 3048 2936 2904 2784 2931 3256 3305 3345 3404 3496 3762
3736 3701 3673 3798 3818 3964 4087 3994 4313 4753 5354 5516 5438 5546 5543 6006
6266 6478 6180 6295 5983 5980 5763 5492 5302 5149 5151 5020 4973 4917 5057 5040
5244 5282 5130 5041 5094 5015 5120 5144 5026 5061 4963 4996 5198 5190 5219 5381
5372 5433 5693 5315 5543 5858 5471 5574 5635 5475 5377 5445 5303 5166 4675 4436
4240 4088 3786 3372 3223 2987 3021 2831 2753 2667 2788 2772 2706 2740 2532 2458
2461 2324 2288 2182 2192 2248 2253 2434 2516 2378 2311 2094 2067 2043 2088 1979
1951 1883 1964 2052 2092 2037 2096 1982 1898 1855 1835 1778 1818 1824 1879 1811
1822 1838 1795 1671 1682 1636 1669 1744 1770 1766 1749 1734 1791 1809 1800 1831
1868 1929 1837 1759 1688 1613 1570 1564 1483 1470 1414 1391 1346 1362 1322 1414
1439 1345 1309 1281 1302 1272 1169 1211 1243 1213 1274 1538 2622 4229 8019 8966

total_r = 786432

1024×768=786432です。あれ?????動いていそうですね?
本来は答えが不定のはずで、偶然でしょうか。。。
上手く動かない、が、この記事の結論のつもりだったのですが、困ったな。


追記:もうチョイましにするため、クラス化しました。

class AtomicCounterBuffer
	{
	public:
		AtomicCounterBuffer()
			: size(1), index(0), offset(0)
		{
			data.reserve(size);
		}

		~AtomicCounterBuffer()
		{
			data.clear();
			glDeleteBuffers(1, &buffer);
		}

		// アトミックカウンタバッファの作成
		void create(int sz, int idx, int off)
		{
			// パラメータの設定
			size = sz;        // バッファのサイズ
			index = idx;      // マウント先のインデックス
			offset = off;     // マウント先のメモリオフセット
			data.resize(sz);

			// バッファの作成
			glGenBuffers(1, &buffer);
			// bind the buffer and define its initial storage capacity
			glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, buffer);
			glBufferData(GL_ATOMIC_COUNTER_BUFFER, sizeof(GLuint) * size, NULL, GL_DYNAMIC_DRAW);
			// unbind the buffer 
			glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, index);
		}

		// カウンタを全てゼロにリセット
		void reset() {
			GLuint *userCounters;
			glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, buffer);
			// map the buffer, userCounters will point to the buffers data
			userCounters = (GLuint*)glMapBufferRange(GL_ATOMIC_COUNTER_BUFFER,
				0,
				sizeof(GLuint) * size,
				GL_MAP_WRITE_BIT | GL_MAP_INVALIDATE_BUFFER_BIT | GL_MAP_UNSYNCHRONIZED_BIT
			);
			// set the memory to zeros, resetting the values in the buffer
			memset(userCounters, 0, sizeof(GLuint) * size);
			// unmap the buffer
			glUnmapBuffer(GL_ATOMIC_COUNTER_BUFFER);

			glBindBufferBase(GL_ATOMIC_COUNTER_BUFFER, index, buffer); // added important
		}

		// カウンタを全てゼロにリセット
		// reset関数と同じ
		void simpler_reset() {
			glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, buffer);

			GLuint *a = new GLuint[size];
			memset(a, 0, size * sizeof(GLuint));
			glBufferSubData(GL_ATOMIC_COUNTER_BUFFER, offset, sizeof(GLuint) * size, a);
			glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, index);

			glBindBufferBase(GL_ATOMIC_COUNTER_BUFFER, index, buffer); // added important
			delete[] a;
		}

		// カウンタ値を読み出してベクターに格納
		int* read_back()
		{
			GLuint *userCounters;
			glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, buffer);
			// again we map the buffer to userCounters, but this time for read-only access
			userCounters = (GLuint*)glMapBufferRange(GL_ATOMIC_COUNTER_BUFFER,
				(GLintptr)offset,
				sizeof(GLuint) * size,
				GL_MAP_READ_BIT
			);

			// copy the values to other variables because...
			for (int i = 0; i < size; i++) {
				data[i] = (int)userCounters[i];
			}
			// ... as soon as we unmap the buffer
			// the pointer userCounters becomes invalid.
			glUnmapBuffer(GL_ATOMIC_COUNTER_BUFFER);

			return data.data();
		}

		// カウンタ値を読み出してベクターに格納
		// read_back関数と同じ
		int* alternative_read_back() {
			GLuint *userCounters = new GLuint[size];
			glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, buffer);
			glGetBufferSubData(GL_ATOMIC_COUNTER_BUFFER, offset, sizeof(GLuint) * size, userCounters);
			glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, 0);
			for (int i = 0; i < size; i++) {
				data[i] = (int)userCounters[i];
			}
			delete[] userCounters;

			return data.data();
		}

		// dataの参照を直接返す
		std::vector<int>& values() {
			return data;
		}

	protected:
		GLuint buffer;
		std::vector<int> data;
		int size;
		int index;
		int offset;
	};


[1] http://www.lighthouse3d.com/tutorials/opengl-atomic-counters/

openFrameworksのshaderでGLSL4.2のatomic counter bufferを使って画素の数をカウントしてみる

GLSLは、OpenGLのシェーダを使うための記述言語です。並列計算はとっても速いので使えると格好良いです。が、GLSLはバージョン毎にすぐ記述方法やデータの受け渡し方が変わったりするので、混乱しがちです。
openFrameworksでは新しいものは面倒みないと割り切っているのか、#version 120もしくは#version 150までの対応が普通のようです。

ここでは、どうしても超高速に画素の数を数えたくなる用事ができたので、最近(#version 420)で追加されたらしいatomic counter bufferと言うのを、動くかどうか分からないまま恐る恐るテストしてみたいと思います。

main.cpp

#include "ofMain.h"
#include "ofApp.h"

//========================================================================
int main( ){
	//ofGLFWWindowSettings settings;
	ofGLWindowSettings settings;
	settings.setGLVersion(4, 5); //version of opengl corresponding to your GLSL version
	settings.width = 1280;
	settings.height = 720;
	ofCreateWindow(settings);

	// this kicks off the running of my app
	// can be OF_WINDOW or OF_FULLSCREEN
	// pass in width and height too:
	ofRunApp(new ofApp());
}

ofApp.h

#pragma once

#include "ofMain.h"

class ofApp : public ofBaseApp{

	public:
		void setup();
		void update();
		void draw();

		void keyPressed(int key);
		void keyReleased(int key);
		void mouseMoved(int x, int y );
		void mouseDragged(int x, int y, int button);
		void mousePressed(int x, int y, int button);
		void mouseReleased(int x, int y, int button);
		void mouseEntered(int x, int y);
		void mouseExited(int x, int y);
		void windowResized(int w, int h);
		void dragEvent(ofDragInfo dragInfo);
		void gotMessage(ofMessage msg);
		
		ofShader shader;
		ofPixels pix;
		ofTexture tex;
		ofFbo fbo;

                GLuint atomicsBuffer;
};

ofApp.cpp

#include "ofApp.h"
// #include "ofxTimeMeasurements.h"

// from here: http://www.lighthouse3d.com/tutorials/opengl-atomic-counters/

void create_a_buffer_for_atomic_counters(GLuint& acb)
{
	glGenBuffers(1, &acb);
	// bind the buffer and define its initial storage capacity
	glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, acb);
	glBufferData(GL_ATOMIC_COUNTER_BUFFER, sizeof(GLuint) * 3, NULL, GL_DYNAMIC_DRAW);
	// unbind the buffer 
	glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, 0);
}

void reset_the_atomic_counter_buffers(GLuint& acb) {
	GLuint *userCounters;
	glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, acb);
	// map the buffer, userCounters will point to the buffers data
	userCounters = (GLuint*)glMapBufferRange(GL_ATOMIC_COUNTER_BUFFER,
		0,
		sizeof(GLuint) * 3,
		GL_MAP_WRITE_BIT | GL_MAP_INVALIDATE_BUFFER_BIT | GL_MAP_UNSYNCHRONIZED_BIT
	);
	// set the memory to zeros, resetting the values in the buffer
	memset(userCounters, 0, sizeof(GLuint) * 3);
	// unmap the buffer
	glUnmapBuffer(GL_ATOMIC_COUNTER_BUFFER);

	glBindBufferBase(GL_ATOMIC_COUNTER_BUFFER, 0, acb); // added important
}

void simpler_reset_the_atomic_counter_buffers(GLuint& acb) {
	glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, acb);

	GLuint a[3] = { 0,0,0 };
	glBufferSubData(GL_ATOMIC_COUNTER_BUFFER, 0, sizeof(GLuint) * 3, a);
	glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, 0);

	glBindBufferBase(GL_ATOMIC_COUNTER_BUFFER, 0, acb); // added important
}

void read_back_the_values_from_the_buffer(GLuint& acb, int& redPixels, int& greenPixels, int& bluePixels)
{
	GLuint *userCounters;
	glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, acb);
	// again we map the buffer to userCounters, but this time for read-only access
	userCounters = (GLuint*)glMapBufferRange(GL_ATOMIC_COUNTER_BUFFER,
		0,
		sizeof(GLuint) * 3,
		GL_MAP_READ_BIT
	);

	// copy the values to other variables because...
	redPixels = (int)userCounters[0];
	greenPixels = (int)userCounters[1];
	bluePixels = (int)userCounters[2];
	// ... as soon as we unmap the buffer
	// the pointer userCounters becomes invalid.
	glUnmapBuffer(GL_ATOMIC_COUNTER_BUFFER);
}

void alternative_read_back_the_values_from_the_buffer(GLuint& acb, int& redPixels, int& greenPixels, int& bluePixels) {
	GLuint userCounters[3];
	glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, acb);
	glGetBufferSubData(GL_ATOMIC_COUNTER_BUFFER, 0, sizeof(GLuint) * 3, userCounters);
	glBindBuffer(GL_ATOMIC_COUNTER_BUFFER, 0);
	redPixels = userCounters[0];
	greenPixels = userCounters[1];
	bluePixels = userCounters[2];
}


// from here: http://tokyo.supersoftware.co.jp/code/720

#define STRINGIFY(A) #A

//--------------------------------------------------------------
void ofApp::setup() {
	string fragmentShaderProgram = STRINGIFY(
		#version 450 compatibility\n
		#extension GL_ARB_shader_atomic_counters : enable\n
		#extension GL_EXT_gpu_shader4 : enable\n

		layout(binding = 0, offset = 0) uniform atomic_uint atRed;
		layout(binding = 0, offset = 4) uniform atomic_uint atGreen;
		layout(binding = 0, offset = 8) uniform atomic_uint atBlue;

		uniform sampler2DRect tex0;

		in vec4 gl_FragCoord;
		out vec4 colorOut;

		void main() {
			vec2 pos = vec2(gl_FragCoord.x, gl_FragCoord.y); // テクスチャ上の座標を取得する
	
			float r = texture2DRect(tex0, pos).r;
			float g = texture2DRect(tex0, pos).g;
			float b = texture2DRect(tex0, pos).b;
			float a = texture2DRect(tex0, pos).a;

			if (r > 0.5) atomicCounterIncrement(atRed);
			if (g > 0.5) atomicCounterIncrement(atGreen);
			if (b > 0.5) atomicCounterIncrement(atBlue);

			colorOut = vec4(r, g, b, a);
		}
	);
	shader.setupShaderFromSource(GL_FRAGMENT_SHADER, fragmentShaderProgram);
	shader.linkProgram();

	// 画像の読み込み
	ofLoadImage(pix, "Koala.jpg"); // 汎用性のためにofPixelsで読み込み
	tex.allocate(pix, true);       // trueが要るのかどうか不明
	tex.setTextureWrap(GL_CLAMP_TO_BORDER, GL_CLAMP_TO_BORDER);
	fbo.allocate((int)tex.getWidth(), (int)tex.getHeight());

	// メインウィンドウの大きさ調整
	ofSetWindowShape((int)tex.getWidth(), (int)tex.getHeight());

        // アトミックカウンタバッファの初期化
	create_a_buffer_for_atomic_counters(atomicsBuffer);
	//reset_the_atomic_counter_buffers(atomicsBuffer);
	simpler_reset_the_atomic_counter_buffers(atomicsBuffer);
}

//--------------------------------------------------------------
void ofApp::update() {

}

//--------------------------------------------------------------
void ofApp::draw() {
	int r, g, b;
	ofBackground(0);

	//TS_START("measurement1");
	fbo.begin();
	shader.begin();
	shader.setUniformTexture("tex0", tex, 0); // テクスチャを渡す
	ofDrawRectangle(-tex.getWidth() / 2, -tex.getHeight() / 2, tex.getWidth(), tex.getHeight()); // 原点を調整してfboにレンダリング
	shader.end();
	fbo.end();

        // アトミックカウンタバッファから結果の受け取りとリセット
	//read_back_the_values_from_the_buffer(atomicsBuffer, r, g, b);
	alternative_read_back_the_values_from_the_buffer(atomicsBuffer, r, g, b);
	simpler_reset_the_atomic_counter_buffers(atomicsBuffer);
	//TS_STOP("measurement1");

	fbo.draw(0, 0);

	cout << "rgb count=" << r << ", " << g << ", " << b << std::endl;
}

//--------------------------------------------------------------
void ofApp::keyPressed(int key) {

}

//--------------------------------------------------------------
void ofApp::keyReleased(int key) {

}

//--------------------------------------------------------------
void ofApp::mouseMoved(int x, int y) {

}

//--------------------------------------------------------------
void ofApp::mouseDragged(int x, int y, int button) {

}

//--------------------------------------------------------------
void ofApp::mousePressed(int x, int y, int button) {

}

//--------------------------------------------------------------
void ofApp::mouseReleased(int x, int y, int button) {

}

//--------------------------------------------------------------
void ofApp::mouseEntered(int x, int y) {

}

//--------------------------------------------------------------
void ofApp::mouseExited(int x, int y) {

}

//--------------------------------------------------------------
void ofApp::windowResized(int w, int h) {

}

//--------------------------------------------------------------
void ofApp::gotMessage(ofMessage msg) {

}

//--------------------------------------------------------------
void ofApp::dragEvent(ofDragInfo dragInfo) {

}

atomic counterは、shaderの中でatomic_uint型という特殊な変数を使いますが、これを受け渡す方法はofShaderには(まだ)用意されていません。
[1]の記事を参考に、というかまんま借りてきて関数化して使うことにします。ところがそのままだと上手く配列を渡せていないようです。
[2]の記事を参考に、初期化の関数の最後に

glBindBufferBase(GL_ATOMIC_COUNTER_BUFFER, 0, acb);

の一行を加えてみますと、上手く動きました。

コアラの色をカウントしてみましょう。処理時間は1024×768pixelsで0.8msecでした。オーバーヘッドが大きいことを考えて、100回ループをまわしたところおよそ30msecでした。
f:id:cvl-robot:20170519213323p:plain

rgb count=376021, 307367, 232583

ちゃんと数字があっているのか確認のため100×100サイズに一点だけ赤くした画像
f:id:cvl-robot:20170519211800p:plain
で試してみると、結果は

rgb count=1, 0, 0

でした。良かった、合っていますね。

速度の比較のために、CPUでの画像操作を想定したダミーコードの実行時間を測定したところ、1回の走査で8.9msecで、100回で788.9msecした。20倍ぐらいShaderの方が速そうですね。

	unsigned char *image = new unsigned char[1024 * 768 * 4];
	TS_START("measurement2");
#ifdef _OPENMP
#pragma omp parallel for
#endif
	for (int i = 0; i < 1024; i++) {
		for (int j = 0; j < 768; j++) {
			image[(j * 1024 + i) * 4 + 0] = 255;
			image[(j * 1024 + i) * 4 + 1] = 255;
			image[(j * 1024 + i) * 4 + 2] = 255;
			image[(j * 1024 + i) * 4 + 3] = 255;
		}
	}
	TS_STOP("measurement2");
	delete[] image;

追記:公平を期すためにOpenMPによる簡単な並列化をしたところ、一回の走査で2.5msec、100回で80-100msecでした。
CPUは、DUAL XEON E5-2620の24論理スレッド構成です。

参考にした記事はこちらです。
[1] http://www.lighthouse3d.com/tutorials/opengl-atomic-counters/
[2] https://shikihuiku.wordpress.com/2012/08/15/atomicoperationinpixelshader/
[3] http://wlog.flatlib.jp/item/1635
[4] https://rauwendaal.net/category/glsl/
GLSLをStringで渡す方法はこちらを参考にしました。
[5] http://tokyo.supersoftware.co.jp/code/720

openFrameworksでOpenGLの処理を速くしたいときの確認項目

この方法で、(にわかに信じられませんが)場合によっては10倍ぐらい速くなることがあります。

普段、あまりいじることの無いmain.cppの最初で呼び出している

	ofSetupOpenGL(1024,768,OF_WINDOW);			// <-------- setup the GL context

を、

	ofGLFWWindowSettings settings;
	//ofGLWindowSettings settings;
	settings.setGLVersion(4, 5); //version of opengl corresponding to your GLSL version
	settings.setSize(1280, 720);
	//settings.width = 1280; // old version of OF
	//settings.height = 720;
	ofCreateWindow(settings);

に変更します。GLVersionやWindowサイズの数値は、使用している環境に合わせて調整してください。

ofShaderを使うときだけ変更する必要があるのかと思いきや、ofTextureやofFboを使うときにも効果があります。テスト環境はnVidia Quadro K4000, Windows7です。

# 常識なのかな、、、今まで気がつかなかった。

今日のNew Balance

靴はNewBalanceが良いよ、と時々聞きますが、NewBalanceにも型番が一杯あって何を選べばいいのか分かりません。次の3モデルの中から、適当に選べば良いそうです。

M1400

MRL996

ML574

M1400は最上位モデルです。一番良い靴がほしいときにはこれを選びます。ちなみに(ひどい偏見ですが)年をとって小銭を持ったオタクの人は、総じてニューバランスのM1400NVを履くそうです。
MRL996は街歩き向けの軽量でクッションが薄めのタイプです。
ML574はMRL996に比べるとクッションが厚めでやや重いらしいです、が、ほとんど変わりません。安いです。定番のオーソドックスな色にはタブにClassicと書かれてます。デザイン的に落ち着いているので、これを選ぶと無難です。

ofFboにofTextureから大きな画像の一部を斜めに切り出す方法

大きな画像から、任意の位置と角度で、適当な矩形を切り出す方法をメモしておきます。

注記:普通にやるならこちらです。
cvl-robot.hateblo.jp


f:id:cvl-robot:20170515212006p:plain
画像サイズ1024×768。描画なしで800fps程度。描画ありだと、verticalSyncに引っ張られて60fps。
普通にOpenCVで書くのとどっちが速いだろう???

カーソル操作は自動車(戦車?)のように移動します。
← 左に1度旋回
→ 右に1度旋回
↑ 前進
↓ 後退

マウスは、ドラッグした位置に切り抜き中心が移動します。スクロール上で左に旋回、下で右に旋回します。 

ofApp.cpp

#include "ofApp.h"

//--------------------------------------------------------------
void ofApp::setup(){
	// 描画の有無とフレームレートの設定
	flag_draw = true;
	ofSetVerticalSync(flag_draw);
	ofSetFrameRate(0);

	// 切り抜き画像サイズ
	crop_size.x = 100.f;
	crop_size.y = 100.f;

	// 画像の読み込み
	ofLoadImage(pix, "Koala.jpg"); // 汎用性のためにPixelsで読み込み
	tex.allocate(pix);
	tex.setTextureWrap(GL_CLAMP_TO_BORDER, GL_CLAMP_TO_BORDER);

	// 切り抜き中心位置の初期化
	pos.x = tex.getWidth() / 2.f;
	pos.y = tex.getHeight() / 2.f;

	// fboの準備
	fbo.allocate((int)crop_size.x, (int)crop_size.y);

	// メインウィンドウの大きさ調整
	ofSetWindowShape((int)ceil(tex.getWidth() + crop_size.x), (int)tex.getHeight());
	flag_updated = true; // 更新フラグ
}

//--------------------------------------------------------------
void ofApp::update(){
	// cout << "FrameRate: " << ofToString(ofGetFrameRate(), 0) << "\t\r";
}

//--------------------------------------------------------------
// Texture tの中心位置p,角度aからFbo fにfのサイズで画像を切り出す関数
void ofApp::rotationalCrop(ofFbo& f, ofTexture& t, ofVec2f p, float a)
{
	float w = f.getWidth();
	float h = f.getHeight();
	// float diag = sqrtf(powf(w, 2.f) + powf(h, 2.f));	// 対角線の長さ

	// 外接する長方形の長辺
	float rad = ofDegToRad(a);
	float i = fabs(w*cos(rad)) + fabs(h*sin(rad));
	float j = fabs(w*sin(rad)) + fabs(h*cos(rad));
	float r = ceil((i>j)?i:j);

	f.begin();
	ofClear(0.f);
	ofPushMatrix();
	// Fboの中心で回転
	ofTranslate(0.5f*w, 0.5f*h);
	ofRotate(a);
	ofTranslate(-0.5f*w, -0.5f*h);
	// 長辺長の正方形を位置pを中心にして描画
	t.drawSubsection(-0.5f*(r-w), -0.5f*(r-h), r, r, p.x - 0.5f*r, p.y - 0.5f*r, r, r);
	ofPopMatrix();
	f.end();
}

//--------------------------------------------------------------
void ofApp::draw(){
	if (tex.isAllocated() && fbo.isAllocated() && flag_updated) {
		rotationalCrop(fbo, tex, pos, angle);
		flag_updated = false;
	}

	if (!flag_draw) return;

	if (tex.isAllocated()) {
		// テクスチャの表示
		tex.draw(0, 0);

		// カーソルの表示
		ofPushStyle();
		ofNoFill();
		ofPushMatrix();
		ofTranslate(pos.x, pos.y);
		ofRotate(-angle);
		ofDrawRectangle(-crop_size.x/2.f, -crop_size.y/2.f, crop_size.x, crop_size.y);
		ofPopMatrix();
		ofPopStyle();
	}

	if (fbo.isAllocated()) {
		// 切り抜き結果の表示
		ofPushMatrix();
		ofTranslate(tex.getWidth(), 0);
		fbo.draw(0, 0);
		ofPopMatrix();
	}

	// フレームレートの表示
	// ofDrawBitmapString(ofToString(ofGetFrameRate(), 0), 20, 20); // bug
}

//--------------------------------------------------------------
void ofApp::keyPressed(int key){
	// 回転
	if (key == OF_KEY_LEFT) {
		angle += 1.f;
		flag_updated = true;
	}
	else if (key == OF_KEY_RIGHT) {
		angle -= 1.f;
		flag_updated = true;
	}

	// 前進・後退
	if (key == OF_KEY_UP) {
		float r = 2.f;
		float rad = ofDegToRad(angle + 90.f);
		ofVec2f d(r*cos(rad), -r*sin(rad));
		pos = ofVec2f(ofClamp(pos.x + d.x, 0.f, tex.getWidth()), ofClamp(pos.y + d.y, 0.f, tex.getHeight()));
		flag_updated = true;
	}
	else if (key == OF_KEY_DOWN) {
		float r = 2.f;
		float rad = ofDegToRad(angle + 90.f);
		ofVec2f d(-r*cos(rad), r*sin(rad));
		pos = ofVec2f(ofClamp(pos.x + d.x, 0.f, tex.getWidth()), ofClamp(pos.y + d.y, 0.f, tex.getHeight()));
		flag_updated = true;
	}

	// 切り出したFboをファイルに保存
	// from here: http://www.atnr.net/how-to-save-fbo-screenshot-on-of/
	if (key == 's') {
		ofImage img;
		ofPixels pixels;

		fbo.readToPixels(pixels);
		img.setFromPixels(pixels);
		char fileNameStr[255];
		string date = ofGetTimestampString(); // タイムスタンプをファイル名にする
		sprintf(fileNameStr, "%s.png", date.c_str());
		img.save(fileNameStr, OF_IMAGE_QUALITY_BEST);
	}
}

//--------------------------------------------------------------
void ofApp::keyReleased(int key){
	
}

//--------------------------------------------------------------
void ofApp::mouseMoved(int x, int y ){
	
}

//--------------------------------------------------------------
void ofApp::mouseDragged(int x, int y, int button){
	pos = ofVec2f(ofClamp(x, 0, tex.getWidth()), ofClamp(y, 0, tex.getHeight()));
	flag_updated = true;
}

//--------------------------------------------------------------
void ofApp::mousePressed(int x, int y, int button){

}

//--------------------------------------------------------------
void ofApp::mouseReleased(int x, int y, int button){

}

//--------------------------------------------------------------
void ofApp::mouseEntered(int x, int y){

}

//--------------------------------------------------------------
void ofApp::mouseExited(int x, int y){

}

//--------------------------------------------------------------
void ofApp::windowResized(int w, int h){

}

//--------------------------------------------------------------
void ofApp::gotMessage(ofMessage msg){

}

//--------------------------------------------------------------
void ofApp::dragEvent(ofDragInfo dragInfo){ 

}

//--------------------------------------------------------------
void ofApp::mouseScrolled(int x, int y, float scrollX, float scrollY) {
	angle += scrollY;
	flag_updated = true;
}

ofApp.h

#pragma once

#include "ofMain.h"

#include "ofPixels.h"
#include "ofTexture.h"
#include "ofFbo.h"

class ofApp : public ofBaseApp {
public:
	void setup();
	void update();
	void draw();

	void keyPressed(int key);
	void keyReleased(int key);
	void mouseMoved(int x, int y);
	void mouseDragged(int x, int y, int button);
	void mousePressed(int x, int y, int button);
	void mouseReleased(int x, int y, int button);
	void mouseEntered(int x, int y);
	void mouseExited(int x, int y);
	void windowResized(int w, int h);
	void dragEvent(ofDragInfo dragInfo);
	void gotMessage(ofMessage msg);
	
	void mouseScrolled(int x, int y, float scrollX, float scrollY);
	void rotationalCrop(ofFbo& f, ofTexture& t, ofVec2f p, float a);

protected:
	ofPixels pix;
	ofTexture tex;
	ofFbo fbo;

	ofVec2f crop_size;
	ofVec2f pos;
	float angle;

private:
	bool flag_updated;
	bool flag_draw;
};

64ビットコンパイルすると、時々落ちるようになってしまったので、Framerateの表示をコメントアウトしました。
f:id:cvl-robot:20170518171139p:plain

Lidarを使った2D SLAM実装のためのお勉強

Atsushi SakaiさんのブログMyEnigmaで紹介されているmatlabのコードをいくつかc++(openFrameworks環境)に移植してみました。
myenigma.hatenablog.com
とても勉強になり大変に有難いです。移植の目的は、matlabに慣れたかったのと、現実的な実行速度を確認したかった為です。matlabpythonよりc++寄りで移植しやすいですね。

1. ofxGridMapSample.h


f:id:cvl-robot:20170511220657p:plain

2. ofxParticleFilterLocalization.h


f:id:cvl-robot:20170511220819p:plain

3. ofxICPSample.h


f:id:cvl-robot:20170511220928p:plain

main()関数やsetup(), update(), draw()の呼び出しは、openFrameworksのテンプレートから適当に作ってください。
また追加のaddonとしてofMatrix2x2をプロジェクトに加えてください。
GitHub - naokiring/ofMatrix2x2: 2x2 matrix class for openFrameworks

Globally Optimalな点群位置合わせ手法の調査

2010年台ごろに入ってから、3次元点群位置合わせの手法として"Globally Optimal*1"な物の提案が増えてきました。理論はおいて置いて、どのくらい実用になるかをテストするために、実装が公開されているものを中心にリストにまとめておきたいと思います。

1. Go-ICP: A Globally Optimal Solution to 3D ICP Point-Set Registration

Jiaolong Yang, et. al. ICCV2013, PAMI
iitlab.bit.edu.cn
解説

www.slideshare.net

2. Super4PCS: Fast Global Pointcloud Registration via Smart Indexing

Nicolas Mellado, Dror Aiger, Niloy J. Mitra CGF2014
SGP | UCL
github.com
つかってみた、試してみた解説
「 Super4PCS 」という点群結合のライブラリを試してみました - Natural Software
JVR ― 自腹でバーチャルリアリティ ― Super4PCS 速度面で実用的ではなさそう。

3. Fast Global Registration

Qian-Yi Zhou, Jaesik Park, Vladlen Koltun, ECCV2016
Fast Global Registration - Vladlen Koltun
github.com
入力が3次元点群+FPFHなどの3次元特徴量ベクトルなので、3次元特徴量の計算にも時間コストがかかることに留意。
速度的には、まだAlignment.exe(Indexed Imageを使った方法)で位置合わせした方が速そう。
読むべき論文多すぎて、サーベイがおっつかないよ。。。

*1:Globally Optimalの正しい和訳は知りませんが、端的な意訳は"初期位置あわせの要らない"だと思います。