使用CUDA计算Haar小波变换

在《Haar小波变换的快速实现》一文里我们提到了Haar小波变换的计算，在这里我们使用CUDA实现文中提到的计算方式。

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)//行变换，因为数据为HSL格式，所以只提取L分量，即i * 3 + 2位置的数据，生成的数据只是一个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

}

代码有点乱，本来行变换和列变换是一次完成的，中间使用一个临时全局存储空间进行存储，但发现做列变换时从全局存储空间提出的数据有一部分是“脏”的，使用_threadfence做了同步处理。问题暂时没有找到，就使用这个替代方法了：先做行变换，然后将数据拷回内存，再拷回显存，第二次调用此内核函数，做列变换。
变换结果如下：

作者：renjihe 出处： http://www.renjihe.com
本站所有文章均是原创，版权归本站所有。欢迎转载，但未经同意必须保留此段声明，且在文章页面明显位置给出原文连接，否则保留追究法律责任的权利.