在《Haar小波變換的快速實(shí)現(xiàn)》一文里我們提到了Haar小波變換的計(jì)算，在這里我們使用CUDA實(shí)現(xiàn)文中提到的計(jì)算方式。

01	__global__?void

02	_cuda_haar(float?*A,?float?*B,?int?w,?int?h,?int?b_w,?int?b_h,?int?row_scan)

03

{

04	????int?i = blockIdx.x * blockDim.x + threadIdx.x;

05	????int?j = blockIdx.y * blockDim.y + threadIdx.y;

06

?

07	????int?extern_w = b_w * 2;

08	????int?extern_h = b_h * 2;

09

?

10	????if?(j > extern_w)

11	????????return;

12

?

13	????if?(i > extern_h)

14	????????return;

15

?

16	????if?(row_scan)//行變換，因?yàn)閿?shù)據(jù)為HSL格式，所以只提取L分量，即i * 3 + 2位置的數(shù)據(jù)，生成的數(shù)據(jù)只是一個(gè)L分量，并＊255量化

17

????{

18	????????if?(j % 2 != 0)

19

????????{

20	????????????if?(j > w)

21	????????????{

22	????????????????B[i * extern_w + b_w + j/2] = A[(i * w + j - 1) * 3 + 2] * 255;

23	????????????}else

24	????????????{

25	????????????????B[i * extern_w + b_w + j/2] = (A[(i * w + j - 1) * 3 + 2] - A[(i * w + j) * 3 + 2]) * 255 * 0.5;

26	????????????}

27

?

28	????????}else

29

????????{

30	????????????if?(j == w)

31	????????????{

32	????????????????B[i * extern_w + j/2] = A[(i * w + j) * 3 + 2] * 255;

33	????????????}else

34	????????????{

35	????????????????B[i * extern_w + j/2] = (A[(i * w + j) * 3 + 2] + A[(i * w + j + 1) * 3 + 2]) * 255 * 0.5;

36	????????????}

37

????????}

38

????}else

39

????{

40	????????if?(i % 2 != 0)

41

????????{

42	????????????if?(i > h)

43	????????????{

44	????????????????B[(i/2 + b_h) * extern_w + j] = A[(i - 1) * extern_w + j];

45	????????????}else

46	????????????{

47	????????????????B[(i/2 + b_h) * extern_w + j] = (A[(i - 1) * extern_w + j] - A[i * extern_w + j]) * 0.5;

48	????????????}

49	????????}else

50

????????{

51	????????????if?(i == h)

52	????????????{

53	????????????????B[i/2 * extern_w + j] = A[i * extern_w + j];

54	????????????}else

55	????????????{

56	????????????????B[i/2 * extern_w + j] = (A[i * extern_w + j] + A[(i + 1) * extern_w + j]) * 0.5;

57	????????????}

58

????????}

59

????}

60

}

代碼有點(diǎn)亂，本來行變換和列變換是一次完成的，中間使用一個(gè)臨時(shí)全局存儲(chǔ)空間進(jìn)行存儲(chǔ)，但發(fā)現(xiàn)做列變換時(shí)從全局存儲(chǔ)空間提出的數(shù)據(jù)有一部分是“臟”的，使用_threadfence做了同步處理。問題暫時(shí)沒有找到，就使用這個(gè)替代方法了：先做行變換，然后將數(shù)據(jù)拷回內(nèi)存，再拷回顯存，第二次調(diào)用此內(nèi)核函數(shù)，做列變換。
變換結(jié)果如下：

作者：renjihe 出處：http://www.renjihe.com
本站所有文章均是原創(chuàng)，版權(quán)歸本站所有。歡迎轉(zhuǎn)載，但未經(jīng)同意必須保留此段聲明，且在文章頁(yè)面明顯位置給出原文連接，否則保留追究法律責(zé)任的權(quán)利.

總結(jié)

以上是生活随笔為你收集整理的使用CUDA计算Haar小波变换的全部?jī)?nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯(cuò)，歡迎將生活随笔推薦給好友。