请教有关纹理内存的问题!!!

  • 1 replies
  • 1481 views
请教有关纹理内存的问题!!!
« 于: 四月 20, 2021, 09:42:31 am »
大家好,想请教各位关于纹理拾取的问题:
我现在的情况是,我想用较少的线程数去处理一个已经被纹理绑定的二维cuda数组,假如总的线程数是256,那么纹理拾取的时候应该怎么设置x,y呢?就是texRef2D(x,y),x和y怎么和线程联系起来,并且最终可以拾取到对应位置的数组元素?

下面是我的代码:
__global__ void flowdirection(int* output, int width, int height) {//计算流向和的核函数


    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int tid = x + y * blockDim.x * gridDim.x;//线程的绝对地址

    int DB[8] = { 1,2,4,8,16,32,64,128 };
    float flag[8] = { 0 };
    int count = 0;


    if (x < width && y < height) {

        if (tex2D(texRef, x, y) == 999) {
            return;
        }
        else {

            flag[0] = tex2D(texRef, x, y) - tex2D(texRef, x + 1, y);//1   
            flag[1] = (tex2D(texRef, x, y) - tex2D(texRef, x + 1, y + 1)) / sqrt(2.0);//2
            flag[2] = tex2D(texRef, x, y) - tex2D(texRef, x, y + 1);//4
            flag[3] = (tex2D(texRef, x, y) - tex2D(texRef, x - 1, y + 1)) / sqrt(2.0);//8
            flag[4] = tex2D(texRef, x, y) - tex2D(texRef, x - 1, y);//16
            flag[5] = (tex2D(texRef, x, y) - tex2D(texRef, x - 1, y - 1)) / sqrt(2.0);//32;
            flag[6] = tex2D(texRef, x, y) - tex2D(texRef, x, y - 1);//64
            flag[7] = (tex2D(texRef, x, y) - tex2D(texRef, x + 1, y - 1)) / sqrt(2.0);//128


            count = maxIndex(flag, 8);

            if (DB[count] > 0) {//可以解决中心像元为最低点、9个数据全为“无数据”的情况

                output[tid] = DB[count];

            }
        }




    }

    __syncthreads();//什么时候使用线程同步



}

Re: 请教有关纹理内存的问题!!!
« 回复 #1 于: 四月 22, 2021, 02:07:22 pm »
大家好,想请教各位关于纹理拾取的问题:
我现在的情况是,我想用较少的线程数去处理一个已经被纹理绑定的二维cuda数组,假如总的线程数是256,那么纹理拾取的时候应该怎么设置x,y呢?就是texRef2D(x,y),x和y怎么和线程联系起来,并且最终可以拾取到对应位置的数组元素?

下面是我的代码:
__global__ void flowdirection(int* output, int width, int height) {//计算流向和的核函数


    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int tid = x + y * blockDim.x * gridDim.x;//线程的绝对地址

    int DB[8] = { 1,2,4,8,16,32,64,128 };
    float flag[8] = { 0 };
    int count = 0;


    if (x < width && y < height) {

        if (tex2D(texRef, x, y) == 999) {
            return;
        }
        else {

            flag[0] = tex2D(texRef, x, y) - tex2D(texRef, x + 1, y);//1   
            flag[1] = (tex2D(texRef, x, y) - tex2D(texRef, x + 1, y + 1)) / sqrt(2.0);//2
            flag[2] = tex2D(texRef, x, y) - tex2D(texRef, x, y + 1);//4
            flag[3] = (tex2D(texRef, x, y) - tex2D(texRef, x - 1, y + 1)) / sqrt(2.0);//8
            flag[4] = tex2D(texRef, x, y) - tex2D(texRef, x - 1, y);//16
            flag[5] = (tex2D(texRef, x, y) - tex2D(texRef, x - 1, y - 1)) / sqrt(2.0);//32;
            flag[6] = tex2D(texRef, x, y) - tex2D(texRef, x, y - 1);//64
            flag[7] = (tex2D(texRef, x, y) - tex2D(texRef, x + 1, y - 1)) / sqrt(2.0);//128


            count = maxIndex(flag, 8);

            if (DB[count] > 0) {//可以解决中心像元为最低点、9个数据全为“无数据”的情况

                output[tid] = DB[count];

            }
        }




    }

    __syncthreads();//什么时候使用线程同步



}

你好, 关于你的几个问题:
(1)如果将线程的x/y坐标, 映射到问题的数据集上的x/y坐标(例如你这里的纹理的), 这个和具体问题有关, 而你已经给出写法了:
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
该x/y的映射是否合理, 需要你对着自己的问题看一下(例如你这里将线程分别按照x和y方向平铺开).
而并不存在一个固定的从blockIdx/threadIdx到(x,y,z)的映射, 这个是灵活的, 可以随意选择的.

(2)如何使用纹理拾取, 将固定线程所需要的位置的数据, 读取出来. 这个在手册上有相当多的章节进行描述, 可以自行看一下. 而你已经给出了你在使用tex2D()了. 注意也可以考虑归一化坐标, 这样线程规模和原始的数据集规模可以自动缩放, 不需要严格对应.(手册也提到了)

(3)何时在使用__syncthreads(), 以及它有什么用, 请自行阅读手册. 这里不进行解释. (以及, 你给出的代码中并不需要使用__syncthreads()的, 除非你的 maxIndex()中有对output指向的缓冲区的特殊使用).

(4)这个回复没有对你的maxIndex()做出任何评价. (当然, 你也没问).