CUDA histogram2d가 작동하지 않습니다.

Question

CUDA에 대한 괜찮은 2D 히스토그램이 부족하여 (필자가 ... 포인터를 환영합니다), pyCUDA로 직접 구현하려고합니다. 여기

은 히스토그램 (NumPy와 사용)과 같아야 작업은 다음과 같습니다

code = ''' 
__global__ void histogram2d(const float *in_x, const float *in_y, const float *in_w, float *out) {{ 
    int start = blockIdx.x * blockDim.x + threadIdx.x; 

    float *block_out = &out[{xres} * {yres} * {num_chans} * blockIdx.x]; 

    for(int i = 0; i < {length}; i++) {{ 
     float x = in_x[start + i]; 
     float y = in_y[start + i]; 
     int w_idx = (start + i) * {num_chans}; 

     int xbin = (int) (((x - {xmin})/{xptp}) * {xres}); 
     int ybin = (int) (((y - {ymin})/{yptp}) * {yres}); 

     if (0 <= xbin && xbin < {xres} && 0 <= ybin && ybin < {yres}) {{ 
      for(int c = 0; c < {num_chans}; c++) {{ 
       atomicAdd(&block_out[(ybin * {xres} + xbin) * {num_chans} + c], in_w[w_idx + c]); 
      }} 
     }} 
    }} 
}} 
'''.format(**args) 

------ 

__global__ void histogram2d(const float *in_x, const float *in_y, const float *in_w, float *out) { 
    int start = blockIdx.x * blockDim.x + threadIdx.x; 

    float *block_out = &out[50 * 50 * 4 * blockIdx.x]; 

    for(int i = 0; i < 100; i++) { 
     float x = in_x[start + i]; 
     float y = in_y[start + i]; 
     int w_idx = (start + i) * 4; 

     int xbin = (int) (((x - -10.0)/20.0) * 50); 
     int ybin = (int) (((y - -10.0)/20.0) * 50); 

     if (0 <= xbin && xbin < 50 && 0 <= ybin && ybin < 50) { 
      for(int c = 0; c < 4; c++) { 
       atomicAdd(&block_out[(ybin * 50 + xbin) * 4 + c], in_w[w_idx + c]); 
      } 
     } 
    } 
} 

가있는 것으로 보인다 :

여기

내가 지금까지있어 무엇 색인 생성에 문제가 있지만 이전에 CUDA를 많이 수행하지 않았기 때문에 그것이 무엇인지 알 수는 없습니다. 여기에 내가 해당 파이썬이있을 거라고 생각 내용은 다음과 같습니다

def slow_hist(in_x, in_y, in_w, out, blockx, blockdimx, threadx): 
    start = blockx * blockdimx + threadx 

    block_out_addr = args['xres'] * args['yres'], args['num_chans'] * blockx 

    for i in range(args['length']): 
     x = in_x[start + i] 
     y = in_y[start + i] 
     w_idx = (start + i) * args['num_chans'] 

     xbin = int(((x - args['xmin'])/args['xptp']) * args['xres']) 
     ybin = int(((y - args['ymin'])/args['yptp']) * args['yres']) 

     if 0 <= xbin < args['xres'] and 0 <= ybin < args['yres']: 
      for c in range(args['num_chans']): 
       out[(ybin * args['xres'] + xbin) * args['num_chans'] + c] += in_w[w_idx + c] 

모든 코드가 이러한 이미지, at the Github page of this notebook (이 셀은 하단에) 포함하여 볼 수 있습니다.

이 CUDA 코드에서 잘못된 것은 무엇입니까? 나는 (포인터를 1, 4, 8, 16, transposing 등으로 atomicAdd 주소를 striding) 작은 조정을 많이 시도했지만 포인터 계산이 어떻게 작동하고 있는지 미묘한 부분을 놓친 것 같습니다. 어떤 도움을 주시면 감사하겠습니다.

Answer 1

CUDA 섹션의 출력 배열에 할당 된 배열은 float32 대신 Numpy의 기본 float64를 사용하므로 메모리가 예상보다 두 배 커졌습니다. 여기에 새로운 히스토그램 출력입니다 :

는 여전히 크게이 히스토그램은 모두 서로 너무 다른 이유를 설명하는 데 도움이 의견이나 답변을 주시면 감사하겠습니다.