cvl-robot's diary

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

ルンバのワイヤレス制御用XBeeインターフェースを比較的綺麗に作りたいときのメモ

先日、ルンバのROI(ルンバオープンインターフェース)用USBシリアル変換ケーブルを綺麗に作りました、が、移動ロボットにケーブルがあると何かと不便です。
なので、今回はこれをその辺に転がっていたXbeeを使って無線化したいと思います。

1.購入する部品

品目メーカー・型番購入先参考価格
XBee(Coordinator用)digi・XBee S2Cなど秋月電子通商2,500円
XBee(Router用)digi・XBee S2Cなど秋月電子通商2,500円
XBeeUSB変換基板SparkFun・XBee Explorer USBスイッチサイエンス3,243円
XBee5Vマイコン用変換基板SparkFun・XBeeエクスプローラ5Vマイコン用スイッチサイエンス1,243円
コネクタminiDIN7pinオス千石電商155円
スーパー3端子レギュレータ5V500mARECOM・R-78E5.0-0.5秋月電子通商350円
小さなプラケース透明な奴秋月電子通商50円
AWG28ぐらいの細さで5cmぐらいの長さのケーブルが色違いで4本必要です。また、ビニールテープか熱収縮チューブの細い奴が2cm長ぐらい必要です。

2.はんだ付け

ルンバに取り付けるアダプタを先に作ります。
まず、ケースの加工をします。透明ケースの左肩部分にドリルで7mm程度の穴をあけます。保護のためにあらかじめ、セロテープなどを貼っておき、細いドリルから始めてじょじょに穴を広げていきます。とりあえずリーマの先が入るようになればOKです。miniDIN7ピンを取り出し、端子部分の部品だけを取り出します。箱の外側から、開けた穴に端子を合わせてぴったりハマり込むまでリーマで穴を広げます。

次に、部品を接続していきます。配線図は次の通り。
f:id:cvl-robot:20170830172555p:plain
スーパー三端子レギュレータの1番ピン(Vin)に赤色のケーブルを直にはんだ付けします。はんだ付け部分には熱収縮チューブなどをかぶせて絶縁してください。
スーパー三端子レギュレータの2番ピン(GND)、3番ピン(Vout)をXBeeエクスプローラ5Vマイコン用基板のGNDと5Vのスルーホールへはんだづけします。運良く基板上でピンは隣り合って並んでいますので、3端子レギュレータの端子を少し折り曲げてやれば綺麗に収まります。
あとはコネクタと基板の配線です。ケーブルは、ケースに開けた穴を通してからminiDIN端子にはんだ付けしてください。
a.miniDINの1番(または2番)の電源ピンと三端子レギュレータの1番ピンに接続した赤いケーブルをつなぎます。ここはルンバのバッテリーの電圧(14~16Vぐらい)が出ていますので、他に接触しないよう気を付けてください。
b.miniDINの6番(または7番)のGNDピンとXBeeエクスプローラ5Vマイコン用のGNDのスルーホールを黒いケーブルで接続します。
c.miniDINの4番のTxD端子をXBeeエクスプローラ5Vマイコン用のDINに繋ぎます。
d.miniDINの3番のRxD端子をXBeeエクスプローラ5Vマイコン用のDOUTに繋ぎます。
c,dについては、TxDが入力DIN、RxDが出力DOUTと互い違いにつながっていることに留意してください。

完成写真はこんな感じ。
f:id:cvl-robot:20170830173524p:plain
スーパー三端子レギュレータのこの品番のものはコンデンサを内蔵しているそうなので、外付け部品は不要です。
秋月の透明ケースは、ちょっと大きすぎたので100円ショップなどで、もっと収まりの良いケースを探した方が良さそうです。XBeeの基板はアンテナが見えるように表側に配置してください。

PC側は特に何もする必要がありません。USBケーブルでPCに接続してください。
たまたまですが、スーパー三端子レギュレータが入っていた透明のケースは、XBeeエクスプローラUSBの基板の幅にぴったり合います。
f:id:cvl-robot:20170830173917p:plain

3.Xbeeの設定

特に難しい設定は必要なく、透過モード(AT)で2つを繋げば良いようです。

取り付け先FirmwareIDDHDLBD
PC側ZigBee Coordinator AT[上下同じもの]例.631[RoombaXBeeのSH][RoombaXBeeのSL]115200
RoombaZigBee Router AT[上下同じもの]例.631[PC側XBeeのSH][PC側XBeeのSL]115200
ただ、LEDをBlinkさせようと高速で通信しようとするとエラーで落ちるので、設定に不備があるかも知れません。もっと詳しい人のページを調べてください。

4.テスト

テストは、http://www.irobotweb.com/~/media/MainSite/PDFs/About/STEM/Create/Create_2_Tethered_Driving_2.pdf
のアプリを使って、有線と同じように確認することができます。
f:id:cvl-robot:20170830175128p:plain

今日のオーディオアンプ

DDFA方式で有名なDENON PMA-50の後継機PMA-60が発表になりましたね。欲しいです。

ルンバの制御用USBシリアルケーブルを綺麗に作りたいときのメモ

一年戦争の時のジオン軍よろしく、移動ロボットを闇雲に量産しようという計画を行っています。そういえば、研究室に転がっている(掃除用)ルンバも動かせるはずだと思いだし、ケーブルを作ってテストしてみました。

1.購入する部品

品目メーカー・型番購入先参考価格
USBシリアルケーブル(5V)FTDI・TTL-232R-5V秋月電子通商1950円
コネクタminiDIN7pinオス千石電商210円

2.はんだ付け

FTDIのUSBシリアル変換ケーブルのピン配置は、次のリンクにあります。
http://akizukidenshi.com/download/ds/ftdi/TTL-232R_pin.pdf
ルンバのオープンインターフェースのピン配置は、次のリンクの3ページ目に有ります。
http://www.irobot.lv/uploaded_files/File/iRobot_Roomba_500_Open_Interface_Spec.pdf
しかし、上下・裏表どっちの方向から見ているのか分かりにくいので、[1]のページで中盤で紹介されている写真を参考にします。
清書しなおすとこんな感じ。
f:id:cvl-robot:20170828170508p:plain
はんだ付けが必要なのはTxD, RxD, GNDの3本だけ。TxD, RxDはクロスするように、TxDにはRxD、RxDにはTxDを繋ぐよう互い違いに接続します。
GNDピンは6番,7番どちらでも構いません。両方繋いでも構いません。
Vpwrは、大した電流を引っ張れないので使わない方が賢明だと思います。
また、起動時のシリアルの通信速度115200bpsから19200bpsに変更するためにはBRC端子を使う必要がありますが、このケーブルでは変更できないのでdefultの115200bps固定になります。
あと、

超重要:はんだ付けの前に、miniDINケーブルのカバーをUSBシリアルケーブルに通しておくことを忘れないように!

完成品はこんな感じ。3.3V-5Vレベルシフトの回路等が不要なので、両端以外にデッパリの無い綺麗な仕上がりになります。
f:id:cvl-robot:20170828172421p:plain

3.テスト

ルンバの動かし方については、ROSのノードとして動かす方法や、自分で作ってみよう系の記事が山のように出てくるのでそちらを参考にしてもらうことにして、ここでは簡単に配線が正しく作れたかどうかを確認する方法を記します。
[2]の電通大の記事を読むと、Windows用アプリケーションが公式に配られているとあります。
http://www.irobotweb.com/~/media/MainSite/PDFs/About/STEM/Create/Create_2_Tethered_Driving_2.pdf
これをダウンロードして実行すれば良さそうです。
f:id:cvl-robot:20170828171619p:plain

ルンバのカバーのツメを丁寧にはがして外してテストします。
加工してしまうと保障が受けられなくなってしまうかもしれないので、以下は自分の責任で行ってください。
カバーのエッジとバリで指を切らないように注意してください。特に横に撫でるように指を動かすと、ケガします。
f:id:cvl-robot:20170828172738p:plain
動いたー。

取り外したカバーのコネクタ部分に、ドリルとリーマを使ってコネクタが挿せる大きさの穴をあけて、カバーを基に戻して完成。
f:id:cvl-robot:20170828173219p:plain
綺麗にできました。

テストに使用したルンバは、ビックカメラモデルのRoomba631です。

ルンバ631 (量販店モデル)

ルンバ631 (量販店モデル)

FPFH特徴量をbinaryファイルに出力

FastGlobalRegistrationのテストデータを自分で作りたいと思います。
github.com
PCLに慣れている人ならば問題ないのでしょうが、上記サイトにはFPFH特徴量を出力するサンプルコードの抜粋しか載っていません。
欠けている部分を他のプログラムから探してきて、動くようにしたいと思います。

1.FPFH特徴量のサンプルプログラム

PCD形式のデータからFPFH特徴量を計算するサンプルプログラムを探すと、下記リンクに奇麗なコードがありました。
PCL/OpenNI tutorial 4: 3D object recognition (descriptors) - robotica.unileon.es

2.PCL1.8のインストール

下記サイトで提供してくださっているinstallerを使って、PCLをインストールします。インストーラに素直に従えば大丈夫です。
Point Cloud Library 1.8.0 has been released – Summary?Blog
3rd partyフォルダの下にあるopenNIのインストールも行ってください。

3.CMakeLists.txtの準備

適当にプロジェクト用のフォルダを作ってください。ここでは、

E://workspace/pcl/FPFHbinary

とします。このフォルダにCMakeLists.txtというファイルを作り、下記の例のようにします。

cmake_minimum_required(VERSION 2.6 FATAL_ERROR)
project(FPFH_BINARY)
set(PCL_DIR "E:/Program Files/PCL 1.8.0/cmake/PCLConfig.cmake")
find_package(PCL 1.7 REQUIRED)
include_directories(${PCL_INCLUDE_DIRS})
link_directories(${PCL_LIBRARY_DIRS})
add_definitions(${PCL_DEFINITIONS})
add_executable(fpfh_binary fpfh_binary.cpp)
target_link_libraries(fpfh_binary ${PCL_COMMON_LIBRARIES} ${PCL_IO_LIBRARIES} ${PCL_FEATURES_LIBRARIES} ${PCL_SEARCH_LIBRARIES})

適当にプロジェクト名、およびファイル名はfpfh_binaryと命名することにします。
target_link_librariesに、common, io, features, searchを追加します。

4.fpfh_binary.cppの準備

githubのサンプルコードと、チュートリアルのサンプルコードを合体させて動くようにします。
githubの方はPointNormal型を入力データとしているので、チュートリアルの方のPointXYZとNormalを別々に扱う形式に変更します。

#include <pcl/io/pcd_io.h>
#include <pcl/features/normal_3d.h>
#include <pcl/features/fpfh_omp.h>

int
main(int argc, char** argv)
{
	// Object for storing the point cloud.
	pcl::PointCloud<pcl::PointXYZ>::Ptr cloud(new pcl::PointCloud<pcl::PointXYZ>);
	// Object for storing the normals.
	pcl::PointCloud<pcl::Normal>::Ptr normals(new pcl::PointCloud<pcl::Normal>);
	// Object for storing the FPFH descriptors for each point.
	// pcl::PointCloud<pcl::FPFHSignature33>::Ptr descriptors(new pcl::PointCloud<pcl::FPFHSignature33>());

	// Read a PCD file from disk.
	if (pcl::io::loadPCDFile<pcl::PointXYZ>(argv[1], *cloud) != 0)
	{
		return -1;
	}

	// Note: you would usually perform downsampling now. It has been omitted here
	// for simplicity, but be aware that computation can take a long time.

	// Estimate the normals.
	pcl::NormalEstimation<pcl::PointXYZ, pcl::Normal> normalEstimation;
	normalEstimation.setInputCloud(cloud);
	normalEstimation.setRadiusSearch(0.03);
	pcl::search::KdTree<pcl::PointXYZ>::Ptr kdtree(new pcl::search::KdTree<pcl::PointXYZ>);
	normalEstimation.setSearchMethod(kdtree);
	normalEstimation.compute(*normals);

	// Assume a point cloud with normal is given as
	// pcl::PointCloud<pcl::PointNormal>::Ptr object

	pcl::FPFHEstimationOMP<pcl::PointXYZ, pcl::Normal, pcl::FPFHSignature33> fest;
	pcl::PointCloud<pcl::FPFHSignature33>::Ptr object_features(new pcl::PointCloud<pcl::FPFHSignature33>());
	fest.setRadiusSearch(0.03); 
	fest.setInputCloud(cloud);
	fest.setInputNormals(normals);
	fest.compute(*object_features);

	// FILE* fid = fopen("features.bin", "wb");
	FILE* fid = fopen(argv[2], "wb");
	int nV = cloud->size(), nDim = 33;
	fwrite(&nV, sizeof(int), 1, fid);
	fwrite(&nDim, sizeof(int), 1, fid);
	for (int v = 0; v < nV; v++) {
		const pcl::PointXYZ &pt = cloud->points[v];
		float xyz[3] = { pt.x, pt.y, pt.z };
		fwrite(xyz, sizeof(float), 3, fid);
		const pcl::FPFHSignature33 &feature = object_features->points[v];
		fwrite(feature.histogram, sizeof(float), 33, fid);
	}
	fclose(fid);
}

5. テストデータの準備

適当にpcdデータを用意します。FPFH特徴量の計算は比較的速いのですが、FastGlobalRegistrationでは大きなデータはとんでもなく長い時間がかかるので、スタンフォードバニー辺りでテストしておくと良いです。
sourceforge.net

FPFHの出力は、コマンドプロンプトで次の例のようにします。第一引数が入力pcdで、第二引数が出力です。

fpfh_binary.exe scan_000.pcd scan_000.bin

ファイルサイズは入力の100倍ぐらいの大きさになります。

FastGlobalRegistrationのMatlabのMexファイルのコンパイル

Matlabコマンドプロンプトから次のオプションを付けてコンパイルしてMEXファイルを作ります。

mex read_features.cpp -I../External/Eigen -I../External
mex fast_global_registration.cpp ../FastGlobalRegistration/app.cpp -I../External/Eigen -I../External -I../FastGlobalRegistration

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

今日のアニメ化待ち漫画

邪神ちゃんドロップキック、アニメになったら絶対かわいい。

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.width = 1280;
	settings.height = 720;
	ofCreateWindow(settings);

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

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

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

今日のNew Balance

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

M1400

MRL996

ML574

[ニューバランス] スニーカー  NB ML574 VNNAVY(VN) 26.5 D

[ニューバランス] スニーカー NB ML574 VNNAVY(VN) 26.5 D

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