四畳半テクノポリス

地方国立大学のクズ大学生の吐き溜めです。やる気をください

Gstreamer+OpenCVの環境構築

Gstreamer+OpenCVの環境構築

 今回ロボコンの大会に出場することになり、友人のPC上に私と同じ環境を構築する必要が発生したため、大会で使用するプログラムのための環境構築について親愛なるMRI氏のために解説する。
 ロボットを遠隔操縦する場合ロボットに搭載されたカメラの映像をネットワーク経由でストリーミングできると便利である。MJPG-streamer等の導入が容易なシステムも存在するが、GstreamerにはOpenCVと容易に連結し画像処理を行うことができ、H264が利用できるためRasberryPiのハードウェアエンコーダを利用したH264による高速高画質なストリーミングが行えるなど多くの利点が存在する。

環境構築

①Gstreamerのインストール

 OpenCVをビルドする前にGstreamerをインストールしておく必要がある。これは公式に従いUbuntuであれば次のように実行する。

 sudo apt-get install libgstreamer1.0-0 gstreamer1.0-plugins-base gstreamer1.0-plugins-good gstreamer1.0-plugins-bad gstreamer1.0-plugins-ugly gstreamer1.0-libav gstreamer1.0-doc gstreamer1.0-tools 
OpenCVのダウンロード

 適当なディレクトリにOpenCVをダウンロードする。デバッグを行う際にエラーメッセージ内でOpenCVのライブラリを参照することがあるのでhomeにOpenCVというディレクトリを作っておくと良いかもしれない。今回はOpenCV3.4を利用することにする。

mkdir OpenCV
cd OpenCV
wget  https://jaist.dl.sourceforge.net/project/opencvlibrary/opencv-android/3.4.0/opencv-3.4.2-android-sdk.zip
unzip opencv-3.4.2.zip
OpenCVのcmake

 ダウンロードしたOpenCVのzipファイルを解凍し、解凍したディレクトリ内にbuildという名前のディレクトリを作成し中に移動する。
ここからが難関である。Gstreamerと競合するライブラリがあるとGstreamerが上手く動かなくなるので、競合するライブラリをすべて無効化しておく必要がある。

cd opencv-3.4.2
mkdir build
cmake -D CMAKE_BUILD_TYPE=RELEASE  -D  INSTALL_C_EXAMPLES=ON   -D WITH_1394=OFF -D WITH_GSTREAMER=ON -D WITH_FFMPEG=OFF -D WITH_QT=ON ..

重要なのは次の3つの項目である。

  1. WITH_1394=OFF
  2. WITH_FFMPEG=OFF
  3. WITH_GSTREAMER=ON

1つ目のWITH_1394はIEEE1394のことらしくカメラのインタフェースに関する項目であるがGstreamerと競合するようなので無効化する。2つ目のWITH_FFMPEGであるがこれはメディアフレームワークでありGstreamerと思いっきりキャラがかぶっているので無効化する。3つめのWITH_GSTREAMERであるが、これを有効化することでGSTREAMERが利用可能になる。

この設定でcmakeを行うと次のようなログが出力されるので確認してほしい。Video I/Oの項目でGstreamer以外にYESが無ければ成功である。もし上記の設定でcmakeを行い、YESの項目が存在すればここにcmakeのオプション一覧があるので、参照してYESになっている項目を無効化する。

--   Video I/O:
--     GStreamer:                   
--       base:                      YES (ver 1.8.3)
--       video:                     YES (ver 1.8.3)
--       app:                       YES (ver 1.8.3)
--       riff:                      YES (ver 1.8.3)
--       pbutils:                   YES (ver 1.8.3)
--     libv4l/libv4l2:              NO
--     v4l/v4l2:                    linux/videodev2.h
--     gPhoto2:                     NO
OpenCVのビルド

 OpenCVのビルドとインストールを行う。makeのスレッド数はコア数の倍程度が好ましいとされているが、RasberryPi3の場合は -j 4 程度にしておいたほうが無難である。

make -j 8 #8スレッド
sudo make install


以上で環境構築は終わりである。

使用例

 OpenCVにはこのGstreamerを利用するための拡張機能が存在し、VideoCaptureやVideoWriterを使ってGstreamerのパイプラインとOpenCVを連結することが可能である。

使用例としてgstreamerのテスト信号をキャプチャしOpenCVで表示するプログラムを掲載する。

gstreamer_test.cpp

#include <opencv2/opencv.hpp>

int main(int argc, char *argv[])
{
	cv::VideoCapture gstreamer;
	//gst-launchと同じコマンドを実行することができる。
	gstreamer.open("videotestsrc  ! appsink");                                 //カラーバー
	//	gstreamer.open("videotestsrc pattern=ball ! appsink");      //ボール
	//	gstreamer.open("videotestsrc pattern=snow ! appsink");   //じゃみじゃみ

	if (!gstreamer.isOpened()) 
	{
		printf("=ERR= fail to open\n");
		return -1;
	}

	while (1)
	{
		cv::Mat GstCap;
		gstreamer >> GstCap;
		cv::imshow("Gstreamer test", GstCap);

		int key = cv::waitKey(1);

		//’q’で終了
		if (key == 'q')
		{
			break;
		}
	}
	return 0;
}

コンパイル

g++ `pkg-config --cflags opencv`  gstreamer_test.cpp -std=c++11 -g `pkg-config --libs opencv` 

BINファイル-HEXファイル 変換ツール

mbedコンパイラのbinファイルをHexファイルに変換します

下のボタンを押しbinファイルを読み込むことでIntelHEXに変換され出力される。


 ダウンロード  

 LPC1114は安価で高性能で入門者にとても良いマイコンであり、mbedコンパイラを使って開発を行うことが出来る。

 mbedインターフェースを用いずFlashMagicを使って書き込みを行おうとした場合、どうしてもbinファイルをHEXファイルに変換する必要が発生する。このような場合HEXツールやBIN2HEX等のツールが用いられるが、前者は使い方がややこしく、後者は64ビット環境で動作位しない。このような環境依存の問題がややこしく感じたため、このツールを開発した。

なんか途中でファイル名の設定とかめんどくさくなったのでその辺適当

院試

自分の心の整理も兼ねて大学の院試の現状について書こうと思う、私が今年受験する大学は次の3つである。

 正直言って金沢大学はかなり厳しい気がする。現状の僕の学力では、専門はかろうじて解けるレベルだし、数学はかなり怪しいレベルである。金沢大学の先生が僕のことをかなり気に入ってくれているようなので中々複雑な心境である。

 NAISTに関して言えば院の学生とコネクションが出来たので、小論文を添削していただける事になりかなり有利に進められそうである。TOEICも585点でもそこまで悪い方では無いらしいし、小論文さえ良ければ大丈夫らしい、体感で「6割は行けるだろう」という感じがする。特に立地がのどかな上大阪日本橋にかなり近いということで非常に気に入っている。凄く行きたい。

 JAISTもとても良い大学であり高専時代はこの大学院へ専攻科から編入で進学しようと考えていたので受験することにした。興味のある研究分野もあるので受けることにした。行って後悔することは無いと思う。

 以上が僕の受ける大学である。できることであれば金沢大学へ進学したいところだが中々厳しそうである。現在の研究室の卒業研究を中途半端に進めて閉まっている現状なので、先生に土下座して院試の勉強とゼミだけさせてもらえるようにしたほうがいいかもしれない。

 とりあえずマクシミン戦略でNAISTJAISTを攻めて金沢へ広げていく感じで行こうと思う。

 万が一全部落ちてしまったら、卒業後研究生にでもなって、受験後に海外インターンシップでも行こうと考えている。というか考えてでもいないと心が潰れそうである。

 

 

「クジラの文化 竜の文明」大沢 昇

 以前神保町の古本市で三冊500円のまとめ売りで買った本である。

本のタイトルのクジラは日本のことであり、竜は中国のことである。

内容としては日米の文化の対比を歴史、文化、地理的要因、宗教、フォークロアなどの様々な側面から考察している本である。また中国と日本の類似性を語る上でアジア圏の文化と欧米文化との比較も登場する。

f:id:toriten1024:20170502173207j:plain

概要

一章 「顔」と「国の形」

 主に地理や歴史的な側面に関しての比較を行っている。国の成立の経緯や皇帝と天皇の違いなどについて詳しく解説している。最初は「日本鬼子」や「艦隊これくしょん」といったオタクカルチャーや日本から中国に流入した漢字の話から入り、それを起点に日本と中国の文化の根本的差異や類似点に関して話を広げていく。

 特に興味を引かれたのは皇帝に関する話である。私たちの日本の天皇は伝説上では神の血を引いていることになっており、血筋によって決められるものであるが、中国の皇帝は能力のある人間が選ばれるシステムであり、どのような身分の人間であっても運さえ良ければ皇帝になれたということである。そのせいか、天皇は位を子供に譲ることで退位するのがもっとも多いパターンであったようだが、皇帝が退位する要因としては一番が寿命や病死などによる退位、二番目が廃位、三番目が殺害であったという。

二章「水の文化」と「火の文明」

 機構的要因や生活様式から生じる、衣食住の文化に関する対比を行っている。タイトルの「水」と「火」という部分についててであるが、これはそれぞれの国においてのの浄化の象徴である。

 巨大な大陸にあり、大河が流れている中国における浄化の衝動は食べ物や水を消毒できる「火」であり、それが調理法などにも大きな影響を与えている。それに対し、細長い島国であり、中央の山脈から左右に短い川が流れている。かつて日本を訪れたオランダ人土木技術者も日本の川を見て「これは川でなく滝だ」といったという。このような短く水源から海までの距離の近い日本の川は淀むことがなく、常にきれいな水が手に入ったこのようなことから日本で「水に流す」という言葉があるように、日本における浄化の象徴は水であるという。

三章 どちらも現実主義だが

 国民性というか、国民の気質やものの考え方に関する対比を行っている。この章で私が面白いと感じたのは3節の”「縮みの志向」と「巨大願望」”である。簡単に言えば日本人は何でも小さく纏めてしまうことを好み、中国人は大きく巨大で壮大なものを好むというはなしである。

 このような傾向は建造物で顕著に現れており、中国の都市部へ行けば巨大な建物が沢山あるという、以前私は親戚に会うためにシンガポールを訪れた時に、多くの巨大なビルを目にし、度肝を抜かれた経験がある。日本にも都庁をはじめとする単純に大きい建造物は多数存在するのだが、それらとシンガポールの建物は空間のつくりがまったく異なっているのだ。巨大な吹き抜けやマリーナベイサンズのような一見無茶な建造物を多数みかけた、このような傾向はやはりシンガポールの経済の中枢を華僑が回していることが影響しているのだろう。

 また、本に登場するエピソードで、大学生として興味を惹かれるものに、こんな話があった。中国からきた留学生が日本の大学で研究テーマを決めるとき大きく壮大なテーマを選ぼうとする。すると日本人の指導教官からダメだしをもらい、もっと小さなテーマを選ぶよう促される。その後中国に帰りかえり研究テーマを尋ねられて答えると、「わざわざ留学までして、そんな小さなテーマを扱ったのか」と文句を言われるのだという。このような一見文化と関係なさそうなアカデミックな領域まで国民性が関与してくるのは意外であった。

四章 明るい競争社会の裏側

 移動手段や文学や芸術に関しての対比が行われている。タイトルの意味は良く分からなかった。

 移動手段に関しては国土の性質が強く影響しており、中国の辞書では日本で言う「衣食住」の項目が「衣食住行」となっているのだと云う。比較し狭い面積でかつ細長い日本では東海道などに見られるよう徒歩による文化が発達し、縦横それぞれにとても広い国土をもつ中国では乗り物による移動の文化が発達したのだという。

 文学に関しては宗教の影響が強く影響している。古くより日本では仏教神道が信仰されており、幽霊が創作の題材に取り扱われる。これには日本の神道が単なる体系性をもった多神教ではなく、地域の特性を色濃く持つ土着宗教に近い性質を持っていたり、仏教に関してもかなりのローカライズがなされている、ことが影響しているのかもしれない、それに対し「儒教思想の支配が強い中国では孔子が人知の及ばないところに関しては語るべきでない」といっているように幽霊などは民の時代までは、あまり文学の題材として扱われることはなかったという。中国は欧米や日本と比べ創世に関するハッキリした伝承がないためフォークロアに関しても自由な発想のものが多く、孔子以前の古い民話や伝説が失われてしまっていると思うと少し残念である。

五章 「クジラの文化」と「竜の文明」

 タイトル回収の章である。日本がクジラ、中国が竜に例えた理由やそれに対する今後の展望にかんして書かれている。

 竜は欧米の伝説ではお姫様が竜にさらわれたり、黙示録の獣が竜であるように悪の化身として扱われるが、中国では神聖なものとして崇められ、中国人は自分たちは竜の血を引くと自称するという。筆者は中国の竜の特徴は様々な生き物の特徴の複合であることを挙げている。竜の頭は麒麟にており、鹿に似た角と、鯉に似た髭を持つなどのキメラ的特長を持つ、このことから他民族が入り混じり多数の文化が組み合わさった中国の文化との類似性なども含め中国を竜の文明としている。

 それに対し、日本の象徴としてクジラを挙げている、日本の文化として紀元前よりクジラを食べる文化を持ち、鯨の肉から、髭、内臓に至るまで余すことなく使い、クジラを弔う「青海島鯨墓」があるなど、鯨が日本文化に密着していること、中国と近すぎず離れすぎず独自に進化したことなどを挙げ日本をクジラの文化としてる。

感想

 まず購入したときの値段に対してかなり楽しめる本だったので満足していることを書いておきたい、中古なので筆者にお金がまわらない事を少し申し訳なく思う。

 私は以前よりエンジニアとして東洋のシリコンバレーと呼ばれる「シンセン」対し強い憧れを持っていたので中国に興味はあったのだが、技術的な側面でしか見ていなかった中国に対し様々な側面で見識を深めることができた。

 エンジニアとして興味を惹かれたものに「落」という概念があった。中国には「盗」と「落」とう二つの概念があり、「盗」は盗みで罪であるが、自分の仕事場のものを少し持ち帰って使うことは「落」にあたり、罪ではないということである。聞く話によるとシンセンのジャンク屋に行くとiPhoneのCPUなどが平然と売られているという、iPhoneの製造は中国のFoxConの工場で行われているというが、こういった部品が流出している事態も「落」に当たるのかと考えてしまった。

大学のサークルで技術書典2に出展してきました。

技術書典

 コミックマーケットでは以前より技術島と呼ばれる一角があって、技術関連の同人誌が並べらており、メイカーズムーブメント以降その数は増えているものと思われる。しかし、世界最大級の同人誌即売会であるコミックマーケット、技術書の占める割合は極一部に過ぎず、技術書を購入しようと思って行っても、エロ本の中をかき分け無いと、だどりつけない状況にある。

 4月9日に開かれた、技術書典はコミックマーケットとは異なり。技術書が主役、いや、技術書と同人ハード/ソフトのみの同人誌即売会である。よって、お子様連れや中学生でも安心して同人技術書を購入する事ができる。

大学のサークルで出展

 私が所属するTUT-CC(豊橋技術科学大学コンピュータクラブ)は以前よりコミックマーケットに参加しており、毎回各部員が自主的に学んできたことに対して、思い思いの記事を書いてまとめた会報的な部誌を頒布している。実を言うと私がこのサークルに入ることにしたのは、豊橋技科大に合格した年の冬のコミックマーケットTUT-CCが出展しているのを見つけた事が決め手となっており、今回部誌に投稿出来たことは非常に光栄に感じている。

 今回は新任の部長さんが取りまとめをやってくださったのだが、私はひどい乱文を校閲させてしまったり、gitの部誌のリポジトリで自分の記事をセルフマージしてしまったり、モノクロ画像しか使えないのに図が白黒だったので気づかずにカラー画像を挿入して印刷を遅らせてしまったり、散々迷惑をかけてしまって大変申し訳無かった。(本当にありがとうございました。)そんなこんなでなんとか印刷は間に合い技術書典に出展することは出来た。

 ちなみに私は以前掲載したステレオマッチングに関する記事を詳細に解説したものを投稿した。図もかなり増やしてわかりやすく説明したつもりだ。

 開場直後は弊サークルの同人誌を買っていく人は少なかったものの、技術書専門の同人誌即売会ということで結局身内が集まり、買っていった人の半分くらいは、部長さんの知り合いと、私の高専時代の後輩(都立高専なんで会場の近所に住んでいる)、それから豊橋技科大のOBだったと思う。残りは立ち読みして興味を持ってくれた他大学の学生や社会人で、私の記事に興味を持って買って行ってくれた人もいて、とてもうれしく感じた。

 結果としては既刊と新刊ともにすべて売り切ってしまった。コミックマーケットでは同じ分量を吸って半分売れ残っていたようなので、技術書典恐るべしである。

戦利品

 同人誌即売会ということでもちろん他のサークルの頒布している同人誌も購入してきた。あまりお金を持って行かなかったのと、買い物に行ける状態になったのが客足が減ってからだったので多くの本を購入することは出来なかった。買えたのはこの四冊である。

f:id:toriten1024:20170412024250j:plain

 立ち読みしてお金の許す範囲で購入したので基本どれも面白いものを購入したつもりだが、おまんじゅう本舗刊のSTM32マイコン入門は入門書として非常に良く出来ており、スイッチサイエンスの同人誌コーナーで販売してもいいんじゃないかというクオリティである。

 

今回初めて同人誌即売会に出展側として参加したが、お客さんとして参加するのとは違い緊張感を持った中々面白い体験だった。いつかは自分で同人誌を書いて参加してみたとい思う。

BootCampが「no bootable device」になったときにwindowsの復旧出来た話

パーティション編集に失敗

 先日macbootにubuntuをインストールしようと考え、OSXの下方にubuntu用の領域を作成したところ、bootcampでインストールされていたwindows7を起動しても「no bootable device -- insert boot disk and press any key」としか表示されなくなってしまった。おそらくパーティション編集時にマスタブートレコードを傷つけてしまい、ブートローダのアドレスなどがおかしくなってしまったものと考えられる。

optionキーを押しながら起動すればmacに入ることはできるのでそこから再インストールすることも可能であったが、windows7にはフルパッケージのCygwinFPGAの開発環境のようなインストールのややこしいソフトがいくつかインストールされており、できることなら再インストールして環境を再構築することは避けたいと考えた。そこでwindous7のリカバリーメディアによるマスターブートレコードの修復を試みた。

よってこの記事はwidowsの領域を破壊していないが、マスターブートレコードや ブートマネージャに問題が発生し、mac上のwindowsが起動できなくなってしまった人の役に立つものだと思う。

リカバリメディアの作成

要はwindows7のインストールメディアを用意すればいいのである。OSXを起動してインストールメディアを作成する。インストールメディアのisoファイルはbootcampでインストールした人間であれば、おそらくダウンロードしたものがOSX上に残っていると思うが、残っていない場合はプロダクトキーを使って、次のサイトからダウンロードすることができる。

https://www.microsoft.com/ja-jp/software-download/windows7

このisoファイルをbootcampアシスタントではなくmacのディスクユーティリティより「ディスクの作成」を使って、DVDに書き込む。

windowsの復旧

ディスクを入れ起動時に「Option」を押すとbootメニューの中に円盤のアイコンが現れると思う、これがwindowsのインストールメディアであるので、選択し、起動する。

言語選択の画面が表示されるので日本語を選択し、次に進むと今すぐインストールの左下に「コンピュータを修復する」という項目があるはずである。

それをクリックすると「システム回復オプション」というウィンドウが表示されるので、上下2つのチェックボックスのうち上の方を選択し、次に進む。

するといくつか修復の項目が表示されると思うので、「スタートアップ修復」を選択すると修復が開始する。

と、まぁ 普通のスタートアップ修復である。

 

OSというのはどうしても大量のソフトウェアをインストールするうちの構築に数日かかるような環境が構築されてしまいがちなので、macを使っているのであれはwindwsのパーティションも含め普段からTimeMachineでバックアップを取っていいたほうが良いであろう。

CUDAによるジュリア集合の描画

 定期試験も来週に迫っているわけであるが、凝りもせず、こんな駄文をインターネットに垂れ流している。というのも、今週にもあった定期試験を対した努力もせず、パス出来たからである。人間というのは、楽な課題だけ選んで挑戦していれば、努力せずともそこそこの成果を得られてしまう、それだけならいいのであるが、結果を得ることで自分が優秀な人間であると勘違いしてしまうことがある、これは恐ろしいことである。知らず、知らずのうちに能力がどんどん下がっていってしまうからである。

 僕は以前も書いたように豊橋技科大のコンピュータクラブに所属しプログラムを書いているわけであるが、ゲームのような一般人から見てもわかりやすいものを作っているわけでも無く、文化祭の際にこれといって展示できるようなものが無かったた。以前からCUDAとOpenCVによる画像処理を行っていたため、CUDAのカーネルを記述することが出来た。そこでCUDAを用いてマンデルブロ集合やジュリア集合と言ったフラクタルを描画するプログラムを作成した。

1.フラクタル

 こんなブログを訪ねてくる人であれば大抵フラクタルについて理解はしているであろうが、一応簡単に説明しておく、僕はその分野の数学を専攻した事が無いので詳しい概念を理解しては居ないのだが、簡単に言うと無限の階層を持つ自己相似形の集まりである。

 分からない人は、そんなことを言っても分かりづらいと思うので、身近な例を挙げるとブロッコリーやカリフラワーの構造であろう。サラダに入っている小さく切られ、茹でられたブロッコリーと、スーパーで売られているブロッコリーは相似形である。更に、サラダに入っている茹でられたブロッコリーの一部を手でちぎると、そのかけらはサラダに入っているブロッコリーの切れ端とも、スーパーで売られているブロッコリーとも相似形である。簡単にいうと一部と全体の形が同じなのである。

 ブロッコリーはちぎっていくとつぼみ一つ一つになってしまうため、階層は有限であるが、数式で記述されるフラクタル図形は計算時間とアウトプットの解像度が許す限り無限の階層を持つことができる。

 

f:id:toriten1024:20170218041849p:plain

2.ジュリア集合とマンデルブロ集合

 僕は実際この集合の数学的な意味に関しては理解出来ていない、ただ映しだされる模様が美しいので好きなのである。次の画像がジュリア集合を描画した一例である。これは僕が作成したプログラムで描画したものである。

f:id:toriten1024:20170218042804j:plain

ジュリア集合は次の式で表される。

z_{k+1} = Z_{k^n+C}

 この計算を繰り返していくと、Zの値が変化し、発散する場所が出てくる。それに必要な計算回数をプロットしていくと以上の様な幾何学模様が現れるのである。zというのは複素数zでありz = z+yiという定義からX座標Y座標さえあれば計算することができる。

 このへんは数式で考えると頭が痛くなってくるので後々プログラムで説明するが、この計算の初期パラメータCを変化させることで、ジュリア集合は様々な形に変形し、美しい幾何学模様を描くのである。

3.実装

 今回作成したプログラムは前節せ説明したパラメータCを変化させジュリア集合の形が変化していくのを描画するだけのものである。様々に変化するジュリア集合の美しい模様を観察することができる。

 CUDAのカーネルは通常のC++と別に.cuファイルを作る必要があるため実装は2つのファイルに分けて行った。一つ目がOpenGLで描画を行うC++ファイル、2つ目がGPGPU計算を行うCUファイルである。

man.cu

#include "man.cuh"
#include "config.h"
#include <cuda_runtime.h>
#include <stdlib.h>
#include <stdio.h>
#include <malloc.h>
#include <device_launch_parameters.h>
#define BLOCK 32

__global__ void cal_pix(float* inMatA,float ci);

void draw_gpu( float *ptr_h){
	int matrixSize = sizeof( float) * V_MAX * (H_MAX*3) ;
	float* hMatA;
	hMatA = (float*)malloc(matrixSize);
	int col, row;
//初期化 for(col = 0; col < V_MAX; col++){ for(row = 0; row < H_MAX; row++){ hMatA[(col * H_MAX + row)*3 + 0] = 0; hMatA[(col * H_MAX + row)*3 + 1] = 0; hMatA[(col * H_MAX + row)*3 + 2] = 0; } } float* dMatA; cudaMalloc((void**)&dMatA,matrixSize); cudaMemcpy(dMatA,hMatA,matrixSize,cudaMemcpyHostToDevice); dim3 block(BLOCK, BLOCK); dim3 grid(V_MAX / BLOCK,H_MAX / BLOCK); static double cnt = 0.0; cnt = cnt +0.01; double cir = sin(cnt) * 0.5; //カーネルの呼び出し
cal_pix<<< grid, block >>>(dMatA,cir); cudaThreadSynchronize(); cudaMemcpy(ptr_h,dMatA,matrixSize,cudaMemcpyDeviceToHost); free(hMatA); cudaFree(hMatA); cudaThreadExit(); } __global__ void cal_pix(float* inMatA,float ci){ int col = blockIdx.x * blockDim.x + threadIdx.x; int row = blockIdx.y * blockDim.y + threadIdx.y; float x,y,xx,yy,zr,zi; double l; x = (3.0)*((float)col/V_MAX) - 1.5; y = (3.0)*((float)row/H_MAX) - 1.5; xx = x; yy = y; for(l = 0; l < 16; l+=0.05){ zr = xx*xx - yy*yy +ci; zi = 2*xx*yy+0.64; xx = zr; yy = zi; if((zr*zr + zi*zi) > 4) break; } inMatA[(col * H_MAX + row)*3 + 0] = (l > 15.5)? 0: l/16; inMatA[(col * H_MAX + row)*3 + 1] = (l > 15.5 || l < 1.0)? 0: (16 - l )/16; inMatA[(col * H_MAX + row)*3 + 2] = ( (int)(l*20) & 1)? l/16: (16 - l )/16 ; }

 CUDAカーネルはc++ファイルから呼び出す事が出来ないので間にdraw_gpu関数を挟む、この関数はGPU上に渡す画像の配列データとの初期化を行い、cal_pix関数へ渡す。この関数は渡された情報を基に自分の担当するピクセルがどのような条件で発散するかを調べ配列データに格納する。

gpu.cpp

#include <GL/glut.h>
#include "config.h"
#include "man.cuh"
#include <math.h>
#include <stdio.h>
float vram[(V_MAX * H_MAX)*3];

void idle(void)
{
  glutPostRedisplay();
}
void display(void)
{
  static float cnt = 0;
  cnt = cnt + 0.01;
  int i=0,j=0;
  glClear(GL_COLOR_BUFFER_BIT);
  glBegin(GL_POINTS);
  int pos;
  draw_gpu(vram); 
  for(i = 0; i < V_MAX ; i++){
	for(j = 0;j < H_MAX; j++){
	  pos = ((V_MAX * i) + j) * 3;
	  glColor3d(vram[pos+0] , vram[pos+1], vram[pos+2]); /* 赤 */
          glVertex2d(4.0*((float)i/V_MAX) - 2.0 ,4.0*((float)j/H_MAX) - 2.0);
	}
  }
	  
  glEnd();
  glutSwapBuffers();
}

void resize(int w, int h)
{
  glViewport(0, 0, w, h);
  glLoadIdentity();

  glOrtho(-w / (V_MAX/2), w / (V_MAX/2), -h / (H_MAX/2), h / (H_MAX/2), -1.0, 1.0);
}

void init(void)
{
  
}

int main(int argc, char *argv[])
{
  glutInitWindowSize(V_MAX, H_MAX);
  glutInit(&argc, argv);
  glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE);
  glutCreateWindow(argv[0]);
  glutDisplayFunc(display);
  glutReshapeFunc(resize);
  init();
  glutIdleFunc(idle);
  glutMainLoop();
  return 0;
}

 main.cppファイルはcudaを呼び出し描画するだけの簡単な処理を行う関数である。描画はOpenGLによって行っており、表示にはGULTを使用している。

 コンパイルは次のMakefileで行った

 

gpud: man.o gpu.o
	nvcc  -O2 -o gpu gpu.o man.o -lglut -lGLU -lGL -lm
gpu.o: gpu.cpp man.cuh config.h
	nvcc -arch=sm_20 -c  gpu.cpp -lglut -lGLU -lGL -lm
man.o: man.cu man.cuh config.h
	nvcc -arch=sm_20 -c man.cu
clean:
	rm man.o gpu.o

 

 CUDAを使って計算してみて驚いたことは、CUDAで計算してCPUに書き戻すまでのオーバーヘッドが思ったよりも大きかったことである。結果として発散までの計算回数が少ない場所ではCPUでジュリア集合の計算を行った場合に比べ、CPUで計算したほうが高速な部分が見受けられた。これには少し落胆してしまった。(とは言っても僕が使用しているGPUがGTX720というintelHDに毛が生えたレベルなのが原因かもしれないが)