- html - 出于某种原因,IE8 对我的 Sass 文件中继承的 html5 CSS 不友好?
- JMeter 在响应断言中使用 span 标签的问题
- html - 在 :hover and :active? 上具有不同效果的 CSS 动画
- html - 相对于居中的 html 内容固定的 CSS 重复背景?
我有一个图像特征矩阵 A 是 n*m*31 矩阵 填充用于过滤,我有 B 作为对象过滤器 k*l*31。我想获得一个输出矩阵C是p * r * 31,图像A的大小没有填充。我尝试编写一个 CUDA 代码来在 A 上运行过滤器 B 并获得 C。
我假设 A 上的每个过滤操作和过滤器 B 都被一个线程 block 占用,因此每个线程 block 内都会有 k*l 操作。并且每个移位的过滤操作将在不同的线程 block 上完成。对于 A(0,0),过滤将在 thread_block(0,0) 上,对于 A(0,1),将在 thread_block(1,0) 上,依此类推。我也有第三个维度为 31。第三个维度的每个空间都将自行计算。因此,通过对矩阵进行正确的 3d 索引,我可能能够以非常并行的形式进行所有操作。
所以操作是
A n*m*31 X B k*l*31 = C p*r*31
如何有效地为操作进行内核索引?
最佳答案
这是我编写的一些代码,大致按照您的描述进行。
它创建一个 3D 数据集(称为 cell
,但它与您的数组 A
相当),用随机数据填充它,然后计算结果 3D基于 A
中的数据的输出数组(称为 node
,但它与您的数组 C
相当)。 A
的数据大小大于 C
的大小(您称之为“填充”)以允许传递函数 B
A
的边界元素。在我的例子中,函数 B 只是在 A
的 3D 立方体积中找到与 B 的大小相关联的最小值(我称之为 WSIZE
,创建一个 3D 区域WSIZE
x WSIZE
x WSIZE
) 并将结果存储在 C
中。
此特定代码试图通过将输入 A 的某个区域复制到每个 block 的共享内存中来利用数据重用。每个 block 计算多个输出点(即它多次计算 B 以填充 C 的区域),以利用 B 的相邻计算的数据重用机会。
这可能有助于您入门。您显然必须用您想要的 B 函数替换 B (我的最小查找代码)。此外,在我的情况下,您需要将 B 域从立方域修改为与您的 B 尺寸相对应的任何类型的矩形棱镜。这也会影响共享内存操作,因此您可能希望在第一次迭代中放弃共享内存,只是为了让事情在功能上正确,然后添加共享内存优化,看看您可以获得什么好处。
#include <stdio.h>
#include <stdlib.h>
// these are just for timing measurments
#include <time.h>
// Computes minimum in a 3D volume, at each output point
// To compile it with nvcc execute: nvcc -O2 -o grid3d grid3d.cu
//define the window size (cubic volume) and the data set size
#define WSIZE 6
#define DATAXSIZE 100
#define DATAYSIZE 100
#define DATAZSIZE 20
//define the chunk sizes that each threadblock will work on
#define BLKXSIZE 8
#define BLKYSIZE 8
#define BLKZSIZE 8
// for cuda error checking
#define cudaCheckErrors(msg) \
do { \
cudaError_t __err = cudaGetLastError(); \
if (__err != cudaSuccess) { \
fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
msg, cudaGetErrorString(__err), \
__FILE__, __LINE__); \
fprintf(stderr, "*** FAILED - ABORTING\n"); \
return 1; \
} \
} while (0)
// device function to compute 3D volume minimum at each output point
__global__ void cmp_win(int knode[][DATAYSIZE][DATAXSIZE], const int kcell[][DATAYSIZE+(WSIZE-1)][DATAXSIZE+(WSIZE-1)])
{
__shared__ int smem[(BLKZSIZE + (WSIZE-1))][(BLKYSIZE + (WSIZE-1))][(BLKXSIZE + (WSIZE-1))];
int tempnode, i, j, k;
int idx = blockIdx.x*blockDim.x + threadIdx.x;
int idy = blockIdx.y*blockDim.y + threadIdx.y;
int idz = blockIdx.z*blockDim.z + threadIdx.z;
if ((idx < (DATAXSIZE+WSIZE-1)) && (idy < (DATAYSIZE+WSIZE-1)) && (idz < (DATAZSIZE+WSIZE-1))){
smem[threadIdx.z][threadIdx.y][threadIdx.x]=kcell[idz][idy][idx];
if ((threadIdx.z > (BLKZSIZE - WSIZE)) && (idz < DATAZSIZE))
smem[threadIdx.z + (WSIZE-1)][threadIdx.y][threadIdx.x] = kcell[idz + (WSIZE-1)][idy][idx];
if ((threadIdx.y > (BLKYSIZE - WSIZE)) && (idy < DATAYSIZE))
smem[threadIdx.z][threadIdx.y + (WSIZE-1)][threadIdx.x] = kcell[idz][idy+(WSIZE-1)][idx];
if ((threadIdx.x > (BLKXSIZE - WSIZE)) && (idx < DATAXSIZE))
smem[threadIdx.z][threadIdx.y][threadIdx.x + (WSIZE-1)] = kcell[idz][idy][idx+(WSIZE-1)];
if ((threadIdx.z > (BLKZSIZE - WSIZE)) && (threadIdx.y > (BLKYSIZE - WSIZE)) && (idz < DATAZSIZE) && (idy < DATAYSIZE))
smem[threadIdx.z + (WSIZE-1)][threadIdx.y + (WSIZE-1)][threadIdx.x] = kcell[idz+(WSIZE-1)][idy+(WSIZE-1)][idx];
if ((threadIdx.z > (BLKZSIZE - WSIZE)) && (threadIdx.x > (BLKXSIZE - WSIZE)) && (idz < DATAZSIZE) && (idx < DATAXSIZE))
smem[threadIdx.z + (WSIZE-1)][threadIdx.y][threadIdx.x + (WSIZE-1)] = kcell[idz+(WSIZE-1)][idy][idx+(WSIZE-1)];
if ((threadIdx.y > (BLKYSIZE - WSIZE)) && (threadIdx.x > (BLKXSIZE - WSIZE)) && (idy < DATAYSIZE) && (idx < DATAXSIZE))
smem[threadIdx.z][threadIdx.y + (WSIZE-1)][threadIdx.x + (WSIZE-1)] = kcell[idz][idy+(WSIZE-1)][idx+(WSIZE-1)];
if ((threadIdx.z > (BLKZSIZE - WSIZE)) && (threadIdx.y > (BLKYSIZE - WSIZE)) && (threadIdx.x > (BLKXSIZE - WSIZE)) && (idz < DATAZSIZE) && (idy < DATAYSIZE) && (idx < DATAXSIZE))
smem[threadIdx.z+(WSIZE-1)][threadIdx.y+(WSIZE-1)][threadIdx.x+(WSIZE-1)] = kcell[idz+(WSIZE-1)][idy+(WSIZE-1)][idx+(WSIZE-1)];
}
__syncthreads();
if ((idx < DATAXSIZE) && (idy < DATAYSIZE) && (idz < DATAZSIZE)){
tempnode = knode[idz][idy][idx];
for (i=0; i<WSIZE; i++)
for (j=0; j<WSIZE; j++)
for (k=0; k<WSIZE; k++)
if (smem[threadIdx.z + i][threadIdx.y + j][threadIdx.x + k] < tempnode)
tempnode = smem[threadIdx.z + i][threadIdx.y + j][threadIdx.x + k];
knode[idz][idy][idx] = tempnode;
}
}
int main(int argc, char *argv[])
{
typedef int cRarray[DATAYSIZE+WSIZE-1][DATAXSIZE+WSIZE-1];
typedef int nRarray[DATAYSIZE][DATAXSIZE];
int i, j, k, u, v, w, temphnode;
const dim3 blockSize(BLKXSIZE, BLKYSIZE, BLKZSIZE);
const dim3 gridSize(((DATAXSIZE+BLKXSIZE-1)/BLKXSIZE), ((DATAYSIZE+BLKYSIZE-1)/BLKYSIZE), ((DATAZSIZE+BLKZSIZE-1)/BLKZSIZE));
// these are just for timing
clock_t t0, t1, t2, t3;
double t1sum=0.0f;
double t2sum=0.0f;
double t3sum=0.0f;
// overall data set sizes
const int nx = DATAXSIZE;
const int ny = DATAYSIZE;
const int nz = DATAZSIZE;
// window (cubic minimization volume) dimensions
const int wx = WSIZE;
const int wy = WSIZE;
const int wz = WSIZE;
// pointers for data set storage via malloc
nRarray *hnode; // storage for result computed on host
nRarray *node, *d_node; // storage for result computed on device
cRarray *cell, *d_cell; // storage for input
// start timing
t0 = clock();
// allocate storage for data set
if ((cell = (cRarray *)malloc(((nx+(wx-1))*(ny+(wy-1))*(nz+(wz-1)))*sizeof(int))) == 0) {fprintf(stderr,"malloc Fail \n"); return 1;}
if ((node = (nRarray *)malloc((nx*ny*nz)*sizeof(int))) == 0) {fprintf(stderr,"malloc Fail \n"); return 1; }
if ((hnode = (nRarray *)malloc((nx*ny*nz)*sizeof(int))) == 0) {fprintf(stderr, "malloc Fail \n"); return 1; }
// synthesize data
for(i=0; i<(nz+(wz-1)); i++)
for(j=0; j<(ny+(wy-1)); j++)
for (k=0; k<(nx+(wx-1)); k++){
cell[i][j][k] = rand(); // unless we use a seed this will produce the same sequence all the time
if ((i<nz) && (j<ny) && (k<nx)) {
node[i][j][k] = RAND_MAX;
hnode[i][j][k] = RAND_MAX;
}
}
t1 = clock();
t1sum = ((double)(t1-t0))/CLOCKS_PER_SEC;
printf("Init took %3.2f seconds. Begin compute\n", t1sum);
// allocate GPU device buffers
cudaMalloc((void **) &d_cell, (((nx+(wx-1))*(ny+(wy-1))*(nz+(wz-1)))*sizeof(int)));
cudaCheckErrors("Failed to allocate device buffer");
cudaMalloc((void **) &d_node, ((nx*ny*nz)*sizeof(int)));
cudaCheckErrors("Failed to allocate device buffer2");
// copy data to GPU
cudaMemcpy(d_node, node, ((nx*ny*nz)*sizeof(int)), cudaMemcpyHostToDevice);
cudaCheckErrors("CUDA memcpy failure");
cudaMemcpy(d_cell, cell, (((nx+(wx-1))*(ny+(wy-1))*(nz+(wz-1)))*sizeof(int)), cudaMemcpyHostToDevice);
cudaCheckErrors("CUDA memcpy2 failure");
cmp_win<<<gridSize,blockSize>>>(d_node, d_cell);
cudaCheckErrors("Kernel launch failure");
// copy output data back to host
cudaMemcpy(node, d_node, ((nx*ny*nz)*sizeof(int)), cudaMemcpyDeviceToHost);
cudaCheckErrors("CUDA memcpy3 failure");
t2 = clock();
t2sum = ((double)(t2-t1))/CLOCKS_PER_SEC;
printf(" Device compute took %3.2f seconds. Beginning host compute.\n", t2sum);
// now compute the same result on the host
for (u=0; u<nz; u++)
for (v=0; v<ny; v++)
for (w=0; w<nx; w++){
temphnode = hnode[u][v][w];
for (i=0; i<wz; i++)
for (j=0; j<wy; j++)
for (k=0; k<wx; k++)
if (temphnode > cell[i+u][j+v][k+w]) temphnode = cell[i+u][j+v][k+w];
hnode[u][v][w] = temphnode;
}
t3 = clock();
t3sum = ((double)(t3-t2))/CLOCKS_PER_SEC;
printf(" Host compute took %3.2f seconds. Comparing results.\n", t3sum);
// and compare for accuracy
for (i=0; i<nz; i++)
for (j=0; j<ny; j++)
for (k=0; k<nx; k++)
if (hnode[i][j][k] != node[i][j][k]) {
printf("Mismatch at x= %d, y= %d, z= %d Host= %d, Device = %d\n", i, j, k, hnode[i][j][k], node[i][j][k]);
return 1;
}
printf("Results match!\n");
free(cell);
free(node);
cudaFree(d_cell);
cudaCheckErrors("cudaFree fail");
cudaFree(d_node);
cudaCheckErrors("cudaFree fail");
return 0;
}
关于image-processing - 用于图像过滤的 3d CUDA 内核索引?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/14920931/
我正在尝试学习 Knockout 并尝试创建一个照片 uploader 。我已成功将一些图像存储在数组中。现在我想回帖。在我的 knockout 码(Javascript)中,我这样做: 我在 Jav
我正在使用 php 编写脚本。我的典型问题是如何在 mysql 中添加一个有很多替代文本和图像的问题。想象一下有机化学中具有苯结构的描述。 最有效的方法是什么?据我所知,如果我有一个图像,我可以在数据
我在两个图像之间有一个按钮,我想将按钮居中到图像高度。有人可以帮帮我吗? Entrar
下面的代码示例可以在这里查看 - http://dev.touch-akl.com/celebtrations/ 我一直在尝试做的是在 Canvas 上绘制 2 个图像(发光,然后耀斑。这些图像的链接
请检查此https://jsfiddle.net/rhbwpn19/4/ 图像预览对于第一篇帖子工作正常,但对于其他帖子则不然。 我应该在这里改变什么? function readURL(input)
我对 Canvas 有疑问。我可以用单个图像绘制 Canvas ,但我不能用单独的图像绘制每个 Canvas 。- 如果数据只有一个图像,它工作正常,但数据有多个图像,它不工作你能帮帮我吗? va
我的问题很简单。如何获取 UIImage 的扩展类型?我只能将图像作为 UIImage 而不是它的名称。图像可以是静态的,也可以从手机图库甚至文件路径中获取。如果有人可以为此提供一点帮助,将不胜感激。
我有一个包含 67 个独立路径的 SVG 图像。 是否有任何库/教程可以为每个路径创建单独的光栅图像(例如 PNG),并可能根据路径 ID 命名它们? 最佳答案 谢谢大家。我最终使用了两个答案的组合。
我想将鼠标悬停在一张图片(音乐专辑)上,然后播放一张唱片,所以我希望它向右移动并旋转一点,当它悬停时我希望它恢复正常动画片。它已经可以向右移动,但我无法让它随之旋转。我喜欢让它尽可能简单,因为我不是编
Retina iOS 设备不显示@2X 图像,它显示 1X 图像。 我正在使用 Xcode 4.2.1 Build 4D502,该应用程序的目标是 iOS 5。 我创建了一个测试应用(主/细节)并添加
我正在尝试从头开始以 Angular 实现图像 slider ,并尝试复制 w3school基于图像 slider 。 下面我尝试用 Angular 实现,谁能指导我如何使用 Angular 实现?
我正在尝试获取图像的图像数据,其中 w= 图像宽度,h = 图像高度 for (int i = x; i imageData[pos]>0) //Taking data (here is the pr
我的网页最初通过在 javascript 中动态创建图像填充了大约 1000 个缩略图。由于权限问题,我迁移到 suPHP。现在不用标准 标签本身 我正在通过这个 php 脚本进行检索 $file
我正在尝试将 python opencv 图像转换为 QPixmap。 我按照指示显示Page Link我的代码附在下面 img = cv2.imread('test.png')[:,:,::1]/2
我试图在这个 Repository 中找出语义分割数据集的 NYU-v2 . 我很难理解图像标签是如何存储的。 例如,给定以下图像: 对应的标签图片为: 现在,如果我在 OpenCV 中打开标签图像,
import java.util.Random; class svg{ public static void main(String[] args){ String f="\"
我有一张 8x8 的图片。 (位图 - 可以更改) 我想做的是能够绘制一个形状,给定一个 Path 和 Paint 对象到我的 SurfaceView 上。 目前我所能做的就是用纯色填充形状。我怎样才
要在页面上显示图像,你需要使用源属性(src)。src 指 source 。源属性的值是图像的 URL 地址。 定义图像的语法是: 在浏览器无法载入图像时,替换文本属性告诉读者她们失去的信息。此
**MMEditing是基于PyTorch的图像&视频编辑开源工具箱,支持图像和视频超分辨率(super-resolution)、图像修复(inpainting)、图像抠图(matting)、
我正在尝试通过资源文件将图像插入到我的程序中,如下所示: green.png other files 当我尝试使用 QImage 或 QPixm
我是一名优秀的程序员,十分优秀!