出外到了異國，總會好奇那邊的人吃些什麼東西，儘管事先做了survey，但關於俄羅斯的食記資料並不多，所以行前也只查了一下食物的中俄名稱而已，就這樣帶著好奇心(的胃)到了俄國。
首先是研討會所提供的中餐，在科學院的餐廳中用餐，服務生慢慢把一道道的食物端上來，桌上則已擺著麵包和吐司，吐司細細長長的很好入口，其中黑黑的吐司在飛機餐也有，偏硬的口感。
這碗滿江紅的湯，應該是所謂的羅宋湯了，據說羅宋是Russian的音譯，所以羅宋湯意思是俄國的湯。湯內是滿滿的蔬菜，有沒有肉忘了XD，口感是酸辣湯但去掉辣的感覺，還蠻好喝的。
主菜登場了，是焗烤蘑菇雞胸肉佐軟爛薯條與小黃瓜，上面那層軟嫩焗烤配上富咬勁的雞肉，組合起來的口感非常絕配。
排好餐點再點名一次，剛才漏點了一道沙拉，這沙拉本身沒什麼特別，不過好像有加怪怪起士的樣子。
morris手裡的杯中物，就是伏特加啦，如水般清澈的外觀下，藏著濃厚酒香及一喝就能讓人了解自己食道長度的功能。當時是在研討會的歡迎晚會中用餐，大會很特別地選了一個賽馬場辦buffet，而且會場中央只有桌子，大家拿了食物還得站著吃，可能這樣方便大家走動與交流吧。喝這杯伏特加之前，旁人說要一口乾掉才好，於是就給它一飲而盡啦，瞬間喉嚨以下胃以上都熱了起來，也算是嚐到俄國名產了，再一杯嗎? 不了XD。
說到咱們到俄羅斯最常吃的食物，可能會跌破一堆人眼鏡。居然是...俄式養老乃瀧，咳...，就是日本料理啦。原因無它，離旅館最近的餐廳就是它咩，話說旅館附近吃的店家也超少的，選擇不多。
門面其實不怎麼日式，還好有這個菜單看版幫忙加分。
這邊出賣一下幫助我們許多的鄭秘書，讓他露個臉，這一餐是他帶我們過來的，用超流利俄文點了一桌菜，不禁要想，以後幾天沒有他在我們要怎麼點餐啊!! 哈。這家的日本料理雖不算非常好，但吃著較熟悉的食物，也多了一份安心感(講得俄國食物很可怕似的)。
進到店裡會發現，哇! 吸煙區比非吸煙區大好多啊~~ 這邊比較禮遇吸煙客喔。 後來我們用英文跟服務生說要非吸煙的座位，服務生點頭似乎懂了，開始帶位，結果還是讓我們坐在吸煙區裡，呃.......(無言)
研討會的結束晚會是在Tsaritsino park中舉行，大家坐定在十人大桌前，像在吃傳統中式合菜似的，但不同的是菜都已經上好啦，好有效率唷。而且也不用擔心菜太遠很難夾，因為服務生會幫每個人分菜，一道一道慢慢來，挺高級的感覺。
這一桌菜大部分都很可口，唯獨要特別點名這個長得像瑞士捲的菜色，裡頭包的是魚耶，真有創意，可惜morris對此並不買單啊。聽大家說在這麼內陸的地方，魚的價格是比牛豬啦高得多，但是老實說這魚有些許腥味，不過也許是morris對此特別敏感吧XD。
吃完大餐吃野餐，這一餐是在參觀MAKS的行程中吃的，會場中人潮佔滿了室內座內，我們只好在外頭吃，在攤販買了一份小朋友大量出走的烤肉，印象中花了700多台幣吧XD，也就幾塊肉，一個麵包，一杯啤酒。嗚，搶錢啊。
在回程的飛機上享用俄航的飛機餐，這次又受到高級的禮遇啦，有高貴的魚肉耶。這時候倒希望俄航能虧待客人一下，提供個牛肉吧...。
2011/03/30
2011/03/17
[Stranger in 莫斯科] 路過英法泰的轉機之行
<啟程>
問morris莫斯科有多遠? 回答你: 17個小時那麼遠 XD
目前國內並沒有直航莫斯科的航班，所以得轉機，而我們又訂到轉最久的那一班...
出發了，搭乘長榮BR87 (這班次有點囧) 的飛機，前往巴黎戴高樂機場。
託某人的福，有幸進入長榮的貴賓室坐一坐，當個貴賓，吃吃喝喝，好不舒服。
接近搭機時間了，便前去候機室等候，這時裡面已人聲鼎沸了，找位子坐了下來，才發現咱們一群被法國人包圍了。
可能是morris第一次看到那麼多法國人吧，其中有一個小朋友吸引我的目光，正用著可愛的童音跟身旁的媽媽講話，此時是中文發音，過了不久，他突然跟對面的男人講起了法文! 原來小朋友是中法混血啊。
稍後搭機的過程，已經不記得了，總之13.5個小時就這麼過去了，順利降落戴高樂機場。不過機場離巴黎市似乎有段距離，所以看不到地標巴黎鐵塔了。
戴高樂有2個航站，其中2航站還分2A、2B、2C到2G等7個區域，之間的連結透過是接駁車來轉運，前往接駁車站的路上，經過了一個長長的隧道，腳下是電動的手扶梯，少走了不少路。
在候車室拍了航站分布圖及行車時間告示牌，還遇到台灣來的旅行團，他們要去奧地利，也在這轉機。
在接駁車上看到這台算是民航界的超跑吧，已退休的協和號超音速客機。
戴高樂機場內的免稅店，都是些超級名牌，但是身價不到那個level，覺得這邊很難逛。 後來到了紀念品店挑了幾張明信片，結果連明信片都很高貴，這薄薄卡片難道是LV做的嗎? 結帳時，店員親切的用法文問候"bonjour"，真好聽! 問她可不可以付美金，她回答了一串morris有聽沒有懂的英文，找完錢之後，才知道她的意思是找錢時會找歐元啦。
穿過免稅店後，來到候機室，大型的圓拱屋頂，居然是用木質做裝飾，很有質感，個人很喜歡這設計。
來到外國，除了拍風景照，也要記得拍人像照喔!
停留不久，咱們就要告別法國了，旅行的下一棒交給俄羅斯航空，準備前往莫斯科SHEREMETYEVO機場，飛行時間是3.5小時。
<回程>
說起來這是一篇沒有踏上俄羅斯的俄羅斯遊記! 哈! 這次不用直述的方法講了，讓我們直接跳過在俄羅斯的日子，到踏上歸途的航班上。回程由莫斯科出發，要在倫敦與曼谷各停留一次，才回到台灣。
搭乘俄航飛往倫敦希斯洛機場的班機，原來會經過倫敦市的上空，一路上可以看到泰悟士河，甚至是倫敦鐵橋! 可惜morris一開始顧慮到快降落了，不敢拿相機出來拍，直到看到有人在猛拍了，才跟著拍，所以倫敦鐵橋就沒圖了。
因為轉機時間相當短，我們馬不停蹄趕往登機門，雖然身處這陌生又號稱全球前幾大繁忙的機場，但指標意外地非常清楚，這一點真的要誇獎一下，讓我們沒迷路順利找到了目的地。
人家說時間嘛，擠一下就有了 XD 所以我們還逛了英國知名百貨公司Harrods在機場的免稅店。
東西不便宜，隨便一個馬克杯就要10幾英鎊起跳，台幣大概500多。
英式下午茶專區，morris買了一組有水果香味的茶葉，後來發現，俗~有夠俗，比在台灣買便宜多啦。
在曼谷機場，飛機下來加個油，我們則下來放風一下，這個曼谷機場，屋頂也是圓拱設計，跟戴高樂還有幾份相似呢。不過感覺就稍微冷調一點，少了那麼點木質的溫暖。在此我們又進了長榮貴賓室，喝喝泰式奶茶，讚!
最後就順利回到台灣囉。
~END~
2011/03/14
[Stranger in 莫斯科] 前往俄羅斯之二三事
這一篇講的是前往俄羅斯首府莫斯科之前的準備工作，不過這些已經是二年前的事了......。
時間: 2009年的7月
morris接到指示要前往莫斯科，當下不禁暗想"Are you kidding ?" "很遠耶......" "甘有可能 ?" "那裡不是很冷嗎 ?"，種種不可置信的話語都浮現了，但離出發日已近，回過神來已經在處理相關手續了，而當中又發生了一些插曲，讓旅行差點無法如期成行，前往俄羅斯有這麼難嗎?
首先是錢的問題，俄羅斯使用的貨幣是盧布(RUB) (俄文：Рубль)，而盧布在台灣的銀行是換不到的! 因此只能去當地換了，他們可收美金與歐元，於是我們就準備好美金，當地在機場和"路邊攤"可換，別懷疑真的很像路邊攤，而且路邊攤換的匯率比機場還漂亮呢! 建議在機場換少許夠用即可，到市區的路邊攤再換多一點。
再來整理一下需要辦的東西，除了護照已有之外，要辦的有:
1.訂旅館
2.邀請函
3.簽證
4.機票
訂俄羅斯的飯店時，發現一個特點，就是怎麼找都找不到單一旅館的網頁，只能連到一個統合的訂旅館公司，再從中選擇你要訂的飯店，例如。最後訂的飯店是Hotel Varshava(華沙)。
邀請函比較特別，它是辦理簽證所需的一份文件，根據外交部領事事務局說法，邀請函是由旅館或旅行社所提供，上面會註明你所待的城市、預計待的時間等等。原本，想透過所參加的研討會主辦單位來取得，幾番的書信往來之後確定無望，只好向訂旅館公司申請，花錢了事，30美金!
申請信寄過去，想說這樣就OK了吧，結果日子一天天過去，沒消沒息，整整過了一個禮拜才收到回覆，寫著:
"I am sorry for the delay, I was out of the office."
看到差點昏倒，你們服務業是沒有職務代理人嗎? 不過外國人還真直接，實話實說。
邀請函(tourist visa support voucher)長這樣
收到邀請函、旅館訂房證明、研討會參加證明之後，就可以開始申請簽證了，這部分由旅行社代勞，取得了俄羅斯簽證，而簽證有效日期呢跟我們的申請日期一模模一樣樣，表示你多待一天都不行喔! 這時候突然覺得這個前共黨老大哥真的很嚴，不禁懷疑是否不歡迎人家去呢。
機票部分，也是由旅行社處理，可以選擇的航空公司有長榮、日航、韓航。這三家以韓航的飛行時間最短，日航則是要在東京過夜再飛，而長榮的飛行時間最長，將近17個小時，因此最早被我們排除，而把目標放在日、韓的飛機，但是事情又開始不順了，日韓的機位都額滿，一直後補不上，最後還是要坐長榮囉，準備在機上度過3/4天吧!
2010/11/15
使用CUDA實作image convolution
Image convolution是影像處理上常用的方法之一，可應用在影像模糊、影像強化、擷取邊緣...等等。在數學上，convolution代表兩個函數(ex: f and g)的重疊運算，結果會產生第3個函數(h) :
f * g = h
應用在影像處理時，函數f為原影像，另一個函數g則為filter，而h可視為原影像被convolution後的結果。 convolution在2維discrete domain的運算過程，可用下圖來表示:
要計算source image上每一個像素的convolution結果，需要以該像素為中心取出與filter相同大小的區域(window)，然後該區域的每一個像素與filter相對應的位置做點對點的相乘，最後計算相乘結果的總和，即為該像素convolution的結果。
在影像上計算每一像素的convolution時，會遇到邊界像素的window超出影像範圍的問題，超出的部分就稱為apron，如下圖所示:
而apron的尺寸會等於filter半徑的寬度。當apron所處的位置已超出影像範圍，必需特別給定此apron的值，才能與filter進行運算，而最常見的方法是補0，或者重複邊界像素的值。 進行2維convolution時，一般需要2維影像與2維的filter進行運算，然而某些2維的filter可由2個1維的filter來組成，當此filter可分離時，2維的計算可等效於連續使用2個1維的filter進行運算，如此一來，最大的好處是減少了計算量，例如當filter尺寸為n X m時，原本每個像素需要進行n*m次乘法，分離filter後，只需要n+m次乘法。
測試平台
- OS: Win XP x64
- CPU: Core 2 Extreme X9650 3.0GHz (4 cores)
- VGA: Geforce GTX280
- nVidia CUDA 3.1
使用CPU計算Image convolution
先使用CPU計算的程式來測試，使用3張不同尺寸的圖片，filter使用高斯濾波器，設定filter kernel的尺寸為17*17(半徑為8)，並分離成2個一維的filter來計算。設定超出影像範圍的apron其值為0。效能測試中，皆進行50次運算，再計算平均花費時間。
CPU:
Parallelization with OpenMP:
CUDA程式實作
在nvidia CUDA SDK中，有convolution的範例程式convolutionSeparable及convolutionTexture，這兩支程式做了最佳化而有很好的效能，但它們並沒有實際接受影像輸入，而是用亂數產生float type的資料，模擬一個channel的影像。接下來，我們先不考慮最佳化，而用最直接的方法實作，再一步步改進程式，最後比較各種方法的差異。
The most naive approach
最基本的方法是將影像資料送到device上的global memory，再讓每個thread在計算convolution的kernel中存取它。整張影像將被分成許多block，如下圖所示，而每一個thread負責計算一個像素的結果，因此在這個應用中，將沒有任何idle的thread。設定每一個block的size為16 X 16。
效能上比平行化的CPU程式快1.29倍。
convolution kernel的code如下:
1: __global__ void my2DConvKernel(float *d_Result, float *d_Data, int dataW, int dataH)2: {3: // original image based coordinate4: int y = blockIdx.y * blockDim.y + threadIdx.y;5: int x = blockIdx.x * blockDim.x + threadIdx.x;6:7: int BiasY = y - KERNEL_RADIUS;8: int BiasX = x - KERNEL_RADIUS;9:10: float sum = 0;11: for(int j = 0; j < KERNEL_LENGTH; ++j)12: {13: //out of image range14: if (BiasY + j < 0 || BiasY + j >= dataH)15: continue;16:17: for(int i = 0; i < KERNEL_LENGTH; ++i)18: {19: //out of image range20: if (BiasX + i < 0 || BiasX + i >= dataW)21: continue;22:23: sum += d_Data[(BiasY + j) * dataW + BiasX + i] *24: c_Kernel[KERNEL_LENGTH * j + i];25: }26: }27:28: d_Result[y * dataW + x] = sum;29: }30:
其中24行的c_Kernel為filter，在呼叫kernel前就先把它傳至constant memory中，雖然constant memory是唯讀，但具有快取，適合將唯讀的資料放於此，加速存取。
2D convolution using shared memory
上一個程式，只把影像資料存到global memory，每讀一次原圖的值都要存取一次，然而存取global memory所需要的時間(latency)是很長的，它也沒有快取，很容易讓運算的瓶頸卡在資料傳輸上。改善的方法是利用shared memory，存取的速度上會快很多，甚至與存取暫存器的速度相同，但它有大小限制，每一個multiprocessor只有16KB的shared memory。接下來，我們把一個block送進shared memory中，其中含要處理的像素以及apron區域，如下圖所示:
在此仍然讓一個thread只處理一個像素的結果，在load資料至shared memory時，每個thread負責一個像素的傳輸，而在計算階段，只有負責傳輸白色區域的thread需要計算出結果，負責傳輸apron的thread會idle，因此雖然利用到了shared memory，但卻減少能同時計算的thread數量，有得也有失。然而在此測試中，filter半徑為8時，若白色區域尺寸維持16 X 16，會使block尺寸達到32 X 32，如此一來會超過每個block容許512個thread的上限(顯卡硬體限制，因不同顯卡而異)，所以在此設定白色區域尺寸為4 X 4。
效能上比上一個方法"慢"1.96倍。
convolution kernel的code如下:
1: __global__ void NaiveSharedMemoryGPUKernel(float *d_Result, float *d_Data, int dataW, int dataH)2: {3: __shared__ float data[KERNEL_RADIUS * 2 + ACTIVE_T_W][KERNEL_RADIUS * 2 + ACTIVE_T_H];4:5: // global mem address for this thread6: const int gLoc = blockIdx.y * ACTIVE_T_H * dataW +7: blockIdx.x * ACTIVE_T_W +8: threadIdx.y * dataW + threadIdx.x -9: KERNEL_RADIUS * dataW - KERNEL_RADIUS;10:11: // original image based coordinate12: int x = blockIdx.x * ACTIVE_T_W - KERNEL_RADIUS + threadIdx.x;13: int y = blockIdx.y * ACTIVE_T_H - KERNEL_RADIUS + threadIdx.y;14:15: if (x < 0 || x > dataW - 1 || y < 0 || y > dataH - 1 )16: data[threadIdx.x][threadIdx.y] = 0;17: else18: data[threadIdx.x][threadIdx.y] = d_Data[gLoc];19:20: __syncthreads();21:22: float sum = 0;23: if(threadIdx.x >= KERNEL_RADIUS &&24: threadIdx.x < KERNEL_RADIUS + ACTIVE_T_W &&25: threadIdx.y >= KERNEL_RADIUS &&26: threadIdx.y < KERNEL_RADIUS + ACTIVE_T_H)27: {28: // row wise29: for (int i = -KERNEL_RADIUS; i <= KERNEL_RADIUS; ++i)30: {31: // col wise32: for (int j = -KERNEL_RADIUS; j <= KERNEL_RADIUS; ++j)33: {34: sum += data[threadIdx.x + i][threadIdx.y + j] *35: c_Kernel[KERNEL_RADIUS + i] *36: c_Kernel[KERNEL_RADIUS + j];37: }38: }39:40: d_Result[gLoc] = sum;41: }42: }
以上的ACTIVE_T_W與ACTIVE_T_H為白色區的長與寬，皆設定為4。第3行利用__shared__來宣告一個使用shared memory的變數，data。第20行的__syncthreads()為CUDA函式，讓同一個block內的thread在此同步，才能繼續執行，在此因為要把data的資料都填完才能做後續的計算，因此需要做同步的動作。
Separable convolution using shared memory
此方法先將2D的filter分解成2個一維的filter，再接連對影像做X方向及Y方向的convolution。當只考慮X方向的convolution時，每個block只需包含X方向的apron，如下圖:
如此一來，不用考慮Y方向的apron，即可增加每個block中用來計算的thread。另外，也可減少影像上同一區域被重複讀進shared memory的次數，原本2D convolution中，同一區域最多會被讀取9次，separable convolution方法中最多只會被讀取6次。
效能上比上一個方法快4.63倍。
convolution kernel的code如下:
1: __global__ void mySeparableRowKernel(float *d_Result, float *d_Data, int dataW, int dataH)2: {3: __shared__ float data[KERNEL_RADIUS * 2 + ACTIVE_T_W][ACTIVE_T_H];4:5: // global mem address for this thread6: const int gLoc = blockIdx.y * ACTIVE_T_H * dataW +7: blockIdx.x * ACTIVE_T_W * dataW +8: threadIdx.x -9: KERNEL_RADIUS + threadIdx.y;10:11: // original image based coordinate12: int x = blockIdx.x * ACTIVE_T_W - KERNEL_RADIUS + threadIdx.x;13:14: if (x < 0 || x > dataW - 1)15: data[threadIdx.x][threadIdx.y] = 0;16: else17: data[threadIdx.x][threadIdx.y] = d_Data[gLoc];18:19: __syncthreads();20:21: float sum = 0;22: if(threadIdx.x >= KERNEL_RADIUS &&23: threadIdx.x < KERNEL_RADIUS + ACTIVE_T_W)24: {25: // row wise26: for (int i = -KERNEL_RADIUS; i <= KERNEL_RADIUS; ++i)27: {28: sum += data[threadIdx.x + i][threadIdx.y] *29: c_Kernel[KERNEL_RADIUS + i];30: }31:32: d_Result[gLoc] = sum;33: }34: }
1: __global__ void mySeparableColKernel(float *d_Result, float *d_Data, int dataW, int dataH)2: {3: __shared__ float data[ACTIVE_T_W][KERNEL_RADIUS * 2 + ACTIVE_T_H];4:5: // global mem address for this thread6: const int gLoc = blockIdx.y * ACTIVE_T_H * dataW +7: blockIdx.x * ACTIVE_T_W +8: threadIdx.y * dataW +9: threadIdx.x -10: KERNEL_RADIUS * dataW;11:12: // original image based coordinate13: int y = blockIdx.y * ACTIVE_T_H - KERNEL_RADIUS + threadIdx.y;14:15: if (y < 0 || y > dataH - 1)16: data[threadIdx.x][threadIdx.y] = 0;17: else18: data[threadIdx.x][threadIdx.y] = d_Data[gLoc];19:20: __syncthreads();21:22: float sum = 0;23: if(threadIdx.y >= KERNEL_RADIUS &&24: threadIdx.y < KERNEL_RADIUS + ACTIVE_T_H)25: {26: // col wise27: for (int j = -KERNEL_RADIUS; j <= KERNEL_RADIUS; ++j)28: {29: sum += data[threadIdx.x][threadIdx.y + j] *30: c_Kernel[KERNEL_RADIUS + j];31: }32:33: d_Result[gLoc] = sum;34: }35: }
nvidia convolution application: convolutionSeparable
nvidia SDK中的範例程式convolutionSeparable，使用了一些最佳化技巧來提升程式執行效率，在此介紹一下。首先在記憶體存取上，存取global memory資料時有符合"coalescence"，滿足的條件有2:
- 每一個half-wrap的第一個thread有aligned在適當的記憶體位置 (該位置必需是每個thread存取的資料大小的16倍)
- "連續"地讀取記憶體
為了滿足第一項條件，以row的convolution來說，每個block除了包含需計算結果的區域及apron區外，還多讀取了一些額外的資料，如下圖的綠色區域，其目的就是要讓讀取記憶體的初始位置能align在適當的位置。
從範例程式中得知，這個位置的所在是從白色區域往前算的第16個資料的倍數，當appron寬度小於16時，就往前多載入16個資料，若寬度大於16，則載入16的倍數的資料，如此一來讀取的範圍既可涵蓋appron又滿足位置的aligned。 接下來談第二項條件，連續讀取記憶體的部分，先瞭解GPU的運作方式，當一個thread去存取記憶體並等待資料時，GPU會先切換到下個thread並執行，因此存取記憶體的順序應該是:
所以下圖讓thread 0存取紅色部分，thread 1與2各別存取黃色與綠色部分，並沒有滿足連續存取記憶體。
而是要如下圖所示，才能滿足連續存取記憶體。
nvidia的程式透過memory coalescence的最佳化，並將某些迴圈以loop unrolling的方式展開，來提高執行效率，結果如下:
效能上比上一個方法快1.24倍。
測試結果
下圖為本篇所有測試的效能比較圖。
最後結果，nvidia的方法比沒有平行化的CPU程式還快上12.63倍，也比平行化的CPU程式快上3.78倍。其中2D convolution使用shared memory那一項，因為每個block只有4X4個thread能執行計算，比其他的測試少，因此效能甚至比直接存取global memory還不好。但整體而言，使用shared memory仍對效能有所幫助。
這次的測試程式(包含nvidia code)，其實只能在一些理想條件下運作，例如影像只能有單channel、影像的長與寬必需是block尺寸的倍數、filter的尺寸不能太大以免單一block超過512個thread…等等。所以若要能處理各類型的影像與不同大小filter進行convolution，這些程式都還需要再加以改進。
2010/11/08
迷迭香與非洲菫
朋友突然說想要種非洲菫，因為這種花有清淨空氣的功用，而且又適合室內栽培，好像很好養，morris想說也來種看看，就說好了一起去買。
除了非洲菫，其實morris一直很想種迷迭香，除了香氣蠻特殊的，更重要的是可以做迷迭香烤雞啊! 於是今天就將這兩種花草一起入手了。
morris在大潤發前的花園買了迷迭香，現場有兩個品種，一種是直立迷迭香，另一種是匍匐迷迭香，後來買了小盆的直立迷迭香。
計畫把迷迭香擺在陽台，因為它需要日照，另外藉助它可驅蟲的功效，當做擋小蟲的防護罩。
小小一株迷迭香
大潤發花園裡賣的非洲菫，這一株是花開最好看的一株。
後來想說貨比三家，去逛了體育館旁的假日花市，那邊可以看到其他的花色，甚至連花瓣的形式也不同。
最後挑了這株花色白中透著粉紅的非洲菫。
這種花蠻適合懶人種的，老闆說一週澆一次水就好，希望它在本人照顧下能好好活著，並扮演好空氣清靜機的角色 XD
訂閱：
文章 (Atom)
