毫无意义地使用 CUDA 共享内存

Question

我有两个版本的相同算法。它最初是卷积，但我对其进行了修改以减少它以检查我的瓶颈在哪里（请注意，每个循环只有一次访问全局内存）：

__global__
void convolve (unsigned char * Md, float * Kd, unsigned char * Rd, int width, int height, int kernel_size, int tile_width, int channels){

int row = blockIdx.y*tile_width + threadIdx.y;
int col = blockIdx.x*tile_width + threadIdx.x;

int sum = 0;
int pixel;
int local_pixel;
int working_pixel;

int row_offset = (kernel_size/2)*(width+kernel_size-1);
int col_offset = kernel_size/2;

for(int color=0; color<channels; color++){

    pixel = color*width*height + row*width + col;
    local_pixel = color*(width+kernel_size-1)*(height+kernel_size-1) + row*(width+kernel_size-1) + col + row_offset + col_offset;
    if(row < height  &&  col < width){
        for(int x=(-1)*kernel_size/2; x<=kernel_size/2; x++)
            for(int y=(-1)*kernel_size/2; y<=kernel_size/2; y++){
                working_pixel = local_pixel + x + y*(width+kernel_size-1);
                sum += (int)((float)Md[working_pixel]);// * ((float)Kd[x+kernel_size/2 + (y+kernel_size/2)*kernel_size]);

            }
        Rd[pixel] = (int) sum;
        sum = 0;
    }
}
}

这是共享内存版本（每个循环一次访问共享内存）

__global__
void convolve (unsigned char * Md, float * Kd, unsigned char * Rd, int width, int height, int kernel_size, int tile_width, int channels){

__shared__ unsigned char Mds[256 + 16*4 +4];

int row = blockIdx.y*tile_width + threadIdx.y;
int col = blockIdx.x*tile_width + threadIdx.x;

if(row < height  &&  col < width){

    int sum = 0;
    int pixel;  //the pixel to copy from Md (the input image)
    int local_pixel;    //the pixel in shared memory

    int start_pixel;    //the offset to copy the borders

    int mds_width = tile_width+kernel_size-1;
    int md_width = width+kernel_size-1;
    int md_height = height+kernel_size-1;

    for(int color=0; color<channels; color++){

        pixel = color*md_width*md_height + row*md_width + col  +  (kernel_size/2)*md_width + kernel_size/2; //position (including borders) + offset
        local_pixel = threadIdx.y*mds_width + threadIdx.x  +  (kernel_size/2)*mds_width + kernel_size/2;    //position + offset


        //Loading the pixels
        Mds[local_pixel] = Md[pixel];//bringing the central pixel itself (position + offset)


        __syncthreads();

        //Convolving
        for(int x=(-1)*kernel_size/2; x<=kernel_size/2; x++)
            for(int y=(-1)*kernel_size/2; y<=kernel_size/2; y++)
                sum += (int)((float)Mds[local_pixel + x + y*mds_width]); // * ((float)Kd[x+kernel_size/2 + (y+kernel_size/2)*kernel_size]);
        Rd[color*width*height + row*width + col] = (int) sum;
        sum = 0;

        __syncthreads();

    }
}
}

执行参数是

convolve<<<dimGrid,dimBlock>>>(Md,Kd,Rd,width,new_height,kernel_size,block_size,colors);

dimGrid = (1376,768)
dimBlock = (16,16)
Md is the read only image
Kd is the filter (3x3)
width = 22016
height = 12288
kernel_size = 3
block_size=16
colors=3

第一个算法得到 1249.59 毫秒，第二个算法得到 1178.2 毫秒，我觉得这很荒谬。 我认为寄存器的数量应该不是问题。用 ptxas 编译我得到：

ptxas info: 560 bytes gmem, 52 bytes cmem[14]
ptxas info: Compiling entry function '_Z8convolvePhPfS_iiiii' for 'sm_10'
ptxas info: Used 16 registers, 384 bytes smem, 4 bytes cmem[1]

而我的设备信息是：

Name: GeForce GTX 660 Ti
Minor Compute Capability: 0
Major Compute Capability: 3
Warp Size: 32
Max Treads per Block: 1024
Max Threads Dimension: (1024,1024,64)
Max Grid Size: (2147483647,65535,65535)
Number of SM: 7
Max Threads Per SM: 2048
Regs per Block (SM): 65536
Total global Memory: 2146762752
Shared Memory per Block: 49152

有没有人对这种糟糕的改进有任何暗示？我不知道还有人可以问..

编辑： 我今天使用的是不同的 nvidia 卡，因为我无法访问实验室。它还具有计算能力 3.0。 我把两个 if 语句都放在了循环之外。 我正在使用 -arch compute_30 -code sm_30 进行编译 我移除了所有的铸件。 全局矩阵现在声明为 const unsigned char * restrict Md 这次我使用了一个 9x9 的过滤器，它使每个像素在被引入共享内存后可以重复使用 81 次。

我从终端得到 3138.41 ms（全球版本）和 3120.96 ms（共享版本）。 在视觉分析器中需要更长的时间。这就是我得到的（截图） http://cl.ly/image/1X372l242S2u

和我一样迷失..

请在这里找到这个易于编译和执行的算法：

http://cl.ly/213l2X3S1v3a

./convolution 8000 4000 159 9 edge_detection_9.txt 0 为全局内存版本 ./convolution 8000 4000 159 9 edge_detection_9.txt 1 为共享内存版本

Answer 1

引起我注意的第一件事：

ptxas info: Compiling entry function '_Z8convolvePhPfS_iiiii' for 'sm_10'

您的卡具有计算能力 3.0，因此您应该使用 sm_30 进行编译。 sm_10 缺少 sm_30 的许多特性，共享内存更小，寄存器更少。

---

接下来我要做的是将 if 语句放在 for 循环之外的两个内核中，以便进行适当的内核比较。

---

接下来，我将增加 kernel_size 以突出共享内存的影响。您的内核中只有 9 次访问（如果我计算正确），这意味着：

• 在第一个内核中，您将全局内存中的元素直接读入寄存器并使用它们
• 在第二个内核中，每个线程从全局内存中读取一个元素，每个线程在共享缓存中进行 9 次访问。
• 由于您没有大量重用共享缓存中的元素，因此访问全局内存所付出的代价太大。

此外，sum += (int)((float)Mds[local_pixel + x + y*mds_width]); 会在共享缓存中产生一些库冲突，从而降低其吞吐量。

如果您的 kernel_size 始终为 3，您还可以通过展开它们并使用固定索引来替换 for 循环，以帮助编译器。

---

我还担心 从 uchar 转换为 float 到 int 的惩罚。我知道这些操作代价高昂，进一步降低了共享缓存使用增益。为什么你要投射例如(int) sum;，因为 Rd 是无符号字符？为什么不将 Rd 声明为 int* ？

---

我发现您的内核中也需要 Kd。正如现在声明的那样，它存储在全局内存中。如果它只是一个 3x3 过滤器，您可以对其进行硬编码或将其加载到循环外的线程局部变量中，这有可能被存储到寄存器中。

如果这不起作用，您可以尝试将系数存储到共享内存中。实际上复制了系数。每个线程都有，以避免银行冲突。共享内存在 Kepler 上有 32 个端口，因此可以在 warp 中的所有线程上同时提取系数。

---

总结 我认为您的共享缓存内核付出了访问全局内存、共享内存库冲突、使用 sm_10 和多种类型转换的代价，从而大大降低了共享缓存的收益。一般建议使用 CUDA Visual Profiler 来验证这些点。

---

另外，我会尝试通过将 Md 声明为 const __restrict__ 来使用 纹理缓存。与全局内存访问相比，这可能会显示一些加速，因为这是一个多端口缓存，具有旨在减少银行冲突的特殊映射。实际上，我希望这甚至比共享内存的情况更好。