CUDA中的HAAR小波变换

Posted 2023-03-13

【中文标题】CUDA中的HAAR小波变换

【英文标题】：HAAR wavelet transform in CUDA

【发布时间】：2012-05-30 08:44:27

【问题描述】：

我尝试在 CUDA 中为一维数组实现 HAAR 小波变换。

算法

我在输入数组中有 8 个索引

在这种情况下if(x_index>=o_width/2 || y_index>=o_height/2) 我将有 4 个线程，它们应该是 0、2、4、6，我计划用它们中的每一个处理输入中的两个索引。

我计算 avg.EG：如果我的线程 id 为 '0'-则 avg 为 (input[0]+input[1])/2，然后同时我得到将输入的差异 [ 0]-avg 等其他线程。

现在重要的是输出的位置。我为输出创建了一个单独的 thread_id，因为使用索引 0、2、4、6 会在将输出放置在正确的索引中造成困难。

我的 avgs 应该放在前 4 个索引中，即输出的 0、1、2、3，而 o_thread_id 应该是 0、1、2、3。 同样，为了在 4、5、6、7 处放置差异，我将 0、1、2、3 与 '4' 相加，如代码所示

问题

我的输出全为零！！！无论我改变什么，我都会得到它。

代码

__global__ void cal_haar(int input[],float output [],int i_widthstep,int o_widthstep,int o_width,int o_height)


    int x_index=blockIdx.x*blockDim.x+threadIdx.x;
    int y_index=blockIdx.y*blockDim.y+threadIdx.y;

    if(x_index>=o_width/2 || y_index>=o_height/2) return;

    int i_thread_id=y_index*i_widthstep+(2*x_index);
    int o_thread_id=y_index*o_widthstep+x_index;

    float avg=(input[i_thread_id]+input[i_thread_id+1])/2;
    float diff=input[i_thread_id]-avg;
    output[o_thread_id]=avg;
    output[o_thread_id+4]=diff;



void haar(int input[],float output [],int i_widthstep,int o_widthstep,int o_width,int o_height)


    int * d_input;
    float * d_output;

    cudaMalloc(&d_input,i_widthstep*o_height);
    cudaMalloc(&d_output,o_widthstep*o_height);

    cudaMemcpy(d_input,input,i_widthstep*o_height,cudaMemcpyHostToDevice);

    dim3 blocksize(16,16);
    dim3 gridsize;
    gridsize.x=(o_width+blocksize.x-1)/blocksize.x;
    gridsize.y=(o_height+blocksize.y-1)/blocksize.y;

    cal_haar<<<gridsize,blocksize>>>(d_input,d_output,i_widthstep,o_widthstep,o_width,o_height);


    cudaMemcpy(output,d_output,o_widthstep*o_height,cudaMemcpyDeviceToHost);

    cudaFree(d_input);
    cudaFree(d_output);

以下是我的主要功能：-

void main()

    int in_arr[8]=1,2,3,4,5,6,7,8;
    float out_arr[8];
    int i_widthstep=8*sizeof(int);
    int o_widthstep=8*sizeof(float);
    haar(in_arr,out_arr,i_widthstep,o_widthstep,8,1);

    for(int c=0;c<=7;c++)
    cout<<out_arr[c]<<endl;
    cvWaitKey();

你能告诉我哪里出了问题，它给了我零作为输出吗？ 谢谢。

【问题讨论】：

您是否尝试使用 * 而不是 []？

对不起，我不明白。你能提一下具体的代码行吗？

例如这个签名global void cal_haar(int input[],float output [],int i_widthstep,float o_widthstep,int o_width,int o_height)尝试使用像global void cal_haar(int* input,float* output,int i_widthstep,float o_widthstep,int o_width,int o_height)这样的指针。

在您的问题中包含其他相关信息。在评论中阅读和关注非常困难。

好的，所以 widthstep 以字节为单位，它负责 malloc 和 memcpy。但是在内核中，您使用它来计算浮点数组的索引，这意味着您将访问尚未分配的内存。并遵循@talonmies 的建议并添加一些错误检查。另请参阅 cuda-memcheck 工具以帮助您找到此类错误。

【参考方案1】：

您的代码的问题是以下情况：

if(x_index>=o_width/2 || y_index>=o_height/2) return;

给定o_height = 1，我们有o_height/2 = 0（o_height 是int，所以我们这里有整数除法和向下舍入），所以没有线程执行任何操作。为了实现你想要的，你可以在这里做浮点运算，或者使用(o_height+1)/2和(o_width+1)/2：它会使用“算术”舍入执行除法（你将有( x_index >= (8+1)/2 /*= 4*/ && y_index >= (1+1)/2 /*= 1*/ )）

此外，当您在 Y 维度中有 1 个以上的线程时，寻址会出现问题，因为那时您的 i_thread_id 和 o_thread_id 计算将是错误的（_withstep 的大小以字节为单位，但您将其用作数组索引）。

【讨论】：

我正在使用 i_widthstep 从全局 x_index 和 y_index 计算块内的线程 ID，为此我需要使用绝对以字节为单位的 widthstep。我的输出对于一维数组也是正确的，所以我认为这是这样做的方法。但是是的，你提到的整数除法存在问题。