孤独明镜

菩提本无树,明镜亦非台,本来无一物,何处惹尘埃。

嗨,我是Dragon,一名全栈开发者。现居上海,就职于一家游戏开发公司。做过Linux/QT、WEB(PHP/Python)、游戏开发,目前主攻图形/图像方向,对数学、文学、哲学、摄影等非常的喜爱。


使用CUDA计算Haar小波变换

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

__global__ void
_cuda_haar(float *A, float *B, int w, int h, int b_w, int b_h, int row_scan)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;

    int extern_w = b_w * 2;
    int extern_h = b_h * 2;

    if (j > extern_w)
        return;

    if (i > extern_h)
        return;

    if (row_scan)//行变换,因为数据为HSL格式,所以只提取L分量,即i * 3 + 2位置的数据,生成的数据只是一个L分量,并*255量化
    {
        if (j % 2 != 0)
        {
            if (j > w)
            {
                B[i * extern_w + b_w + j/2] = A[(i * w + j - 1) * 3 + 2] * 255;
            }else
            {
                B[i * extern_w + b_w + j/2] = (A[(i * w + j - 1) * 3 + 2] - A[(i * w + j) * 3 + 2]) * 255 * 0.5;
            }

        }else
        {
            if (j == w)
            {
                B[i * extern_w + j/2] = A[(i * w + j) * 3 + 2] * 255;
            }else
            {
                B[i * extern_w + j/2] = (A[(i * w + j) * 3 + 2] + A[(i * w + j + 1) * 3 + 2]) * 255 * 0.5;
            }
        }
    }else
    {
        if (i % 2 != 0)
        {
            if (i > h)
            {
                B[(i/2 + b_h) * extern_w + j] = A[(i - 1) * extern_w + j];
            }else
            {
                B[(i/2 + b_h) * extern_w + j] = (A[(i - 1) * extern_w + j] - A[i * extern_w + j]) * 0.5;
            }
        }else
        {
            if (i == h)
            {
                B[i/2 * extern_w + j] = A[i * extern_w + j];
            }else
            {
                B[i/2 * extern_w + j] = (A[i * extern_w + j] + A[(i + 1) * extern_w + j]) * 0.5;
            }
        }
    }
}

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

转载请注明地址:孤独明镜