我试图实现用于用C图像堆栈一个3D旋转例程/ CUDA(主要是为了加快计算次数)。我用ImageJ的源代码作为代码的基础,所以旋转不会自由地围绕原点,而是沿着轴。虽然我遇到了一个有趣的问题。我实现了一个关于Y轴的对象的旋转,没有什么问题。但是,当我尝试围绕X轴旋转时,使用非常类似的代码,会出现问题。我注意到,在X轴旋转,有显著条纹,像这样的例子:故障在C 3D旋转/ CUDA
这不是在Y旋转我在做的发生。
我已经提供了正在运行办各轴旋转的CUDA内核(是的rotationY的作品之一,是的rotationX的一个,使该分段)。我在想,如果有人可以提供任何建议,为什么我会得到一个,而不是其他问题,只要它们在实施非常相似。
编辑:我已经缩小的问题降到atomicMin()无法正常工作。即使所有偏移量设置正确,zbuffer也不会正确更改。如果有人知道为什么这可能不起作用,这将是很好的知道。
__global__ void rotationY(int *input, int *projArray, int costheta, int sintheta, int width, int height, int depth, int xcenter, int zcenter,
int projectionwidth, int projectionsize, int *zbuffer, int adjCue, int depthCueSurf, int zmax, int zdiff){
int i=threadIdx.x + blockDim.x*blockIdx.x;
int zcostheta;
int zsintheta;
int offset;
int k, z, point, xnew, znew;
int y=i/width;
int x=i-y*width-xcenter;
int xcostheta = x*costheta;
int xsintheta = x*sintheta;
int offsetinit = y*projectionwidth;
zbuffer[i]=32767;
__syncthreads();
for(k=1; k<=depth; k++){
z = (int)(k-1+.5) - zcenter;
zcostheta = z*costheta;
zsintheta = z*sintheta;
point = i + (k-1)*width*height;
if(input[point]>0){
xnew = (xcostheta + zsintheta)/8192 + xcenter;
znew = (zcostheta - xsintheta)/8192 + zcenter;
offset = offsetinit + xnew;
if (offset<0 || offset>=projectionsize) offset = 0;
atomicMin(&zbuffer[offset],znew);
}
__syncthreads();
if(input[point]>0){
if(znew<=zbuffer[offset]) projArray[offset] = adjCue*input[point]/100+depthCueSurf*input[point]*(zmax-znew)/zdiff;
}
}
}
__global__ void rotationX(int *input, int *projArray, int costheta, int sintheta, int width, int height, int depth, int ycenter, int zcenter,
int projectionsize, int *zbuffer, int adjCue, int depthCueSurf, int zmax, int zdiff) {
int i=threadIdx.x + blockDim.x*blockIdx.x;
int zcostheta;
int zsintheta;
int offset;
int k, z, point, ynew, znew;
int y=i/width;
int x=i-y*width;
y=y-ycenter;
int ycostheta = y*costheta;
int ysintheta = y*sintheta;
zbuffer[i]=32767;
__syncthreads();
for(k=1; k<=depth; k++){
z = (int)(k-1+.5) - zcenter;
zcostheta = z*costheta;
zsintheta = z*sintheta;
point = i + (k-1)*width*height;
if(input[point]>0){
ynew = (ycostheta - zsintheta)/8192 + ycenter;
znew = (ysintheta + zcostheta)/8192 + zcenter;
offset = x + ynew*width;
if (offset<0 || offset>=projectionsize) offset = 0;
atomicMin(&zbuffer[offset], znew);
}
__syncthreads();
if(input[point]>0){
if(znew<=zbuffer[offset]) projArray[offset] = adjCue*input[point]/100+depthCueSurf*input[point]*(zmax-znew)/zdiff;
}
}
}
我想你指的是功能doProjections上http://rsbweb.nih.gov/ij/developer/source/ij/plugin/Projector.java.html – whoplisp
您是否尝试过使用CUDA-GDB或NSight调试atomicMin问题? – harrism