【发布时间】：2016-11-01 00:04:34
【问题描述】：
我正在尝试为每个线程块复制图像补丁和相关围裙到共享内存。
在我的数据被复制（我使用 矩阵）到共享内存后，我想要一个关系来映射共享内存中的掩码中心，我考虑用于卷积和掩码的中心在图像缓冲区中。
我想要这样，因为如果我尝试对图像进行卷积，共享内存中掩码的中心似乎与 全局内存 中存储的图像缓冲区的中心不对应。 p>
在下面的代码中，我编写了一个简单图像 黑白侵蚀算法 的示例，当我将 卷积 的结果放入输出图像时，似乎中心不对应。
我使用 512x512px 图像编写示例
我的设置是：
//block and grid size
dim3 block(16,16);
dim3 grid(512/(block.x),512/(block.y),1);
这是我的内核：
#define STREL_SIZE 5
#define TILE_W 16
#define TILE_H 16
#define R (STREL_SIZE/2)
//size of tile image + apron
#define BLOCK_W (TILE_W+(2*R))
#define BLOCK_H (TILE_H+(2*R))
__global__ void erode_multiple_img_SM_v2(unsigned char * buffer_in,
unsigned char * buffer_out,
int w,int h ){
// Data cache: threadIdx.x , threadIdx.y
__shared__ unsigned char data[TILE_W +STREL_SIZE ][TILE_H +STREL_SIZE ];
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
// global mem address of this thread
int gLoc = row*w +col;
int x, y; // image based coordinate
if((col<w)&&(row<h)) {
data[threadIdx.x][threadIdx.y]=buffer_in[gLoc];
if (threadIdx.y > (h-STREL_SIZE))
data[threadIdx.x][threadIdx.y + STREL_SIZE]=buffer_in[gLoc + STREL_SIZE];
if (threadIdx.x >(w-STREL_SIZE))
data[threadIdx.x + STREL_SIZE][threadIdx.y]=buffer_in[gLoc+STREL_SIZE];
if ((threadIdx.x >(w-STREL_SIZE)) && (threadIdx.y > (h-STREL_SIZE)))
data[threadIdx.x+STREL_SIZE][threadIdx.y+STREL_SIZE] = buffer_in[gLoc+2*STREL_SIZE];
//wait for all threads to finish read
__syncthreads();
unsigned char min_value = 255;
for(x=0;x<STREL_SIZE;x++){
for(y=0;y<STREL_SIZE;y++){
min_value = min( (data[threadIdx.x+x][threadIdx.y+y]) , min_value);
}
}
buffer_out[gLoc]= min_value;
}
}
我的输入图像：
我的内核输出是：
其中w是图片的宽度，等于512， 其中h是图片的高度，等于512。
我调用内核：
erode_multiple_img_SM<<<grid,block>>>(dimage_src,dimage_dst,512,512);
dimage_src 是输入图像一个数组缓冲区而不是矩阵，dimage_dst 是输出图像一个缓冲区。
每个缓冲区的大小为 nElem * nImg * sizeof(unsigned char) 其中 nElem=512*512 是缓冲区的大小，nImg 是我想要处理的图像数量等于 1。 我哪里错了？
代码更新：
__global__ void erode_multiple_img_SM_v2(unsigned char * buffer_in,
unsigned char * buffer_out,
int w,int h ){
// Data cache: threadIdx.x , threadIdx.y
__shared__ unsigned char data[TILE_W + STREL_SIZE-1 ][TILE_H + STREL_SIZE-1 ];
// global mem address of this thread
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
int gLoc = row*w +col;
// each threads loads four values from global memory into shared mem
int x, y; // image based coordinate
if((col<w)&&(row<h)) {
data[threadIdx.x][threadIdx.y] = buffer_in[gLoc];
if (threadIdx.y > (TILE_H-STREL_SIZE+1))
data[threadIdx.x][threadIdx.y + STREL_SIZE-1]=buffer_in[(row + STREL_SIZE-1)*w + col];
if (threadIdx.x > (TILE_W-STREL_SIZE+1))
data[threadIdx.x + STREL_SIZE-1][threadIdx.y] = buffer_in[row*w+col + STREL_SIZE-1];
if ((threadIdx.x > (TILE_W-STREL_SIZE+1)) && (threadIdx.y > (TILE_H-STREL_SIZE+1)))
data[threadIdx.x + STREL_SIZE-1][threadIdx.y + STREL_SIZE-1] = buffer_in[(row + STREL_SIZE-1)*w + col + STREL_SIZE-1];
//wait for all threads to finish read
__syncthreads();
unsigned char min_value = 255;
for(x=0;x<STREL_SIZE;x++){
for(y=0;y<STREL_SIZE;y++){
min_value = min( (data[threadIdx.x+x][threadIdx.y+y]) , min_value);
}
}
buffer_out[gLoc]= min_value;
}
}
我现在的输出是：
更新 2（版本 2 -working-）：
我已经实现了另一个版本的算法。为此，我遵循slide，我发现它非常有用且解释清楚，尤其是作者谈论卷积的部分幻灯片 27。
我将 block 和 grid 设置更改为：
dim3 block(20,20);
dim3 grid(512/(block.x)+ block.x,512/(block.y)+block.y);
内核调用却保持不变：
erode_multiple_img_SM<<<grid,block>>>(dimage_src,dimage_dst,512,512);
内核的参数在哪里：
- dimage_src： 大小为 height x width 的 unsigned char 缓冲区，其中包含输入图像。
- dimage_dst:**大小为 **height x width 的 unsigned char 缓冲区，其中包含我的内核生成的输出图像。
- 512：第三个参数是图片的宽度。
- 512：第四个参数是图片的高度。
记住我的图像样本是黑白，但是这个版本的侵蚀也可以使用灰度。
这里是我的工作内核：
#define STREL_W 5
#define STREL_H 5
#define STREL_SIZE 5
#define TILE_W 16
#define TILE_H 16
#define R (STREL_SIZE/2)
#define BLOCK_W (TILE_W+(2*R))
#define BLOCK_H (TILE_H+(2*R))
__global__ void erode_multiple_img_working(unsigned char * buffer_in,
unsigned char * buffer_out,
int w,int h ){
__shared__ unsigned char fast_acc_mat[BLOCK_w][BLOCK_H];
int ty = threadIdx.y;
int tx = threadIdx.x;
int row_o = blockIdx.y * TILE_W + ty;
int col_o = blockIdx.x * TILE_H + tx;
int row_i = row_o - R;
int col_i = col_o - R;
//in of img size
if((row_i >= 0) && (row_i < h) && (col_i >= 0) && (col_i < w) ){
fast_acc_mat[ty][tx] = buffer_in[ row_i * w + col_i];
}
else{
fast_acc_mat[ty][tx] = 0;
}
__syncthreads();
if( ty < TILE_H && tx < TILE_W ){
unsigned char min_val=255;
for(int i = 0; i < STREL_SIZE; i++) {
for(int j = 0; j < STREL_SIZE; j++) {
min_val = min( fast_acc_mat[i+ty][j+tx] , min_val );
}
}
if(row_o < h && col_o < w)
buffer_out[row_o * w + col_o] = min_val;
}
}
这是我侵蚀的图像（输出）：
我实现了一个方案，该方案展示了 Eric 描述的算法部分如何在共享内存中加载 TILE 的像素：
【问题讨论】：
-
我把代码放在 github 上好吗？
-
我不能把所有的代码都放在帖子里，因为它太大了。但如果您认为这有助于阅读问题，我也会调用内核？
-
代码属于问题，不在外部链接中。所以是clear about this（“最短代码...在问题本身）。如果代码对于问题来说太长了，那么你还没有创建正确的@ 987654329@.
-
不够好。你如何启动你的内核？
w?h?
标签： c++ image-processing cuda