CUDA histogram2d不工作

-1

由於看起來缺乏像樣的CUDA二維直方圖（我可以找到...指針歡迎），我試圖用pyCUDA自己實現它。CUDA histogram2d不工作

這裏的直方圖應該是什麼樣子（使用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，轉置輸出等來跨越atomicAdd地址），但似乎我錯過了一些細微的東西，可能是指針算術的工作原理。任何幫助，將不勝感激。

來源

2016-08-19 scnerd

對於n個樣本，in_x和in_y具有形狀（n，）和in_w具有形狀（n，4）。有形狀（num_blocks，yres，xres，4）。我的目標是爲每個CUDA塊分配足夠的空間以便擁有自己的（yres，xres，4）區域來寫入（所以原子添加不會彼此阻塞），然後我對軸0進行求和以獲得最終直方圖 – scnerd

爲CUDA部分的輸出數組分配的數組使用Numpy的默認float64而不是float32，因此內存是預期的兩倍。下面是新的直方圖輸出：

我還是非常感激有助於解釋爲什麼這些直方圖都來自彼此不同的意見或答案。

來源

2016-08-19 13:39:35 scnerd

CUDA histogram2d不工作

回答

相關問題