2011-08-03 4 views
0

C/CUDA를 사용하여 이미지 스택에 3D 회전 루틴을 구현하려고합니다. 대부분 계산 시간이 단축됩니다. ImageJ 소스 코드를 코드의 기초로 사용했기 때문에 회전은 원점에 대해 자유롭지 않고 축을 따라 자유롭게 적용됩니다. 내가 겪은 재미있는 문제가 있습니다. 작은 문제없이 Y 축에 대한 객체의 회전을 구현했습니다. 그러나 X 축과 비슷한 코드로 회전하려고하면 문제가 발생합니다. 나는 X 회전에서 상당한 스트라이핑은이 예로서, 거기 것으로 나타났습니다 : 내가하고 있던 Y 회전에서 발생되지 않은C/CUDA에서 3D 회전과 관련된 문제가 발생했습니다.

http://i.imgur.com/dkecs.png

.

나는 각 축에 대해 회전을 수행하기 위해 실행되는 CUDA 커널을 제공했습니다 (rotateY는 작동하는 것이고 rotationX는 스트라이핑을 제공하는 것입니다). 나는 그들이 구현에서 매우 유사하다면, 내가 왜 다른 것에 문제를 일으키는 지에 대한 제안을 누군가가 줄 수 있는지 궁금해했다.

편집 : 나는 atomicMin()이 제대로 작동하지 않는 문제를 줄였습니다. 모든 오프셋이 올바르게 설정되어 있어도 z 버퍼가 올바르게 변경되지 않습니다. 왜 이것이 작동하지 않을지 알고있는 사람은 알고있는 것이 좋습니다.

__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; 
    } 
} 
} 
+0

난 당신이 http://rsbweb.nih.gov/ij/developer/source/ij/plugin/Projector.java.html – whoplisp

+0

에 기능 doProjections를 참조 가정 사용 해봤 CUDA-GDB 또는 NSight를 사용하여 atomicMin 문제를 디버깅 하시겠습니까? – harrism

답변

1

rotationX의 함수 프로토 타입에 매개 변수 projectionwidth가 없습니다. 이것은 현재 오류에 대한 최선의 후보입니다.

enter image description here

+0

X 회전의 projectionwidth가 width와 같으므로 차이가 없습니다. 얇은 스택 (이 경우 20 개 이미지)에서도 Y 회전의 프로젝션 폭은 너비와 동일합니다. 제가 그 일을 할 때 문제가되지 않았습니다. – Vic

관련 문제