ブロックサイズに応じて動的にクラッシュするCUDAカーネル

Question

スパース行列、密ベクトル乗算を行いたい。 Matrixが圧縮行ストレージCRSのエントリを圧縮するための唯一の格納形式を想定します。私のカーネルは、次のようになります

：

__global__ void 
krnlSpMVmul1(
     float *data_mat, 
     int num_nonzeroes, 
     unsigned int *row_ptr, 
     float *data_vec, 
     float *data_result) 
{ 
    extern __shared__ float local_result[]; 
    local_result[threadIdx.x] = 0; 

    float vector_elem = data_vec[blockIdx.x]; 

    unsigned int start_index = row_ptr[blockIdx.x]; 
    unsigned int end_index = row_ptr[blockIdx.x + 1]; 

    for (int index = (start_index + threadIdx.x); (index < end_index) && (index < num_nonzeroes); index += blockDim.x) 
     local_result[threadIdx.x] += (data_mat[index] * vector_elem); 

    __syncthreads(); 

    // Reduction 

    // Writing accumulated sum into result vector 
} 

あなたがカーネルをできるだけ単純なことになって、それもいくつかのことを行いますされて見ることができるように間違っ（例えばvector_elemは常にだけではありません正しい値）。私はそれらのことを認識しています。

私の問題： ブロックサイズが32または64のスレッドを使用しているとします。私の行列の行に16以上の非ゼロ（例えば17）があるとすぐに最初の16の乗算が行われ、共有メモリに保存されます。私は、17番目の乗算の結果であるlocal_result[16]の値がちょうどゼロであることを知っています。 16または128スレッドのブロックサイズを使用すると、説明された問題が修正されます。

私はCUDAをかなり新しくしているので、最も簡単なことを見落としているかもしれませんが、私はそれ以上見ることができません。

大変助かりました！ talonmiesコメントへの

---

編集：

私は直接計算の後local_result[16]にあった値を印刷しました。それは0でした。

低減部：

int k = blockDim.x/2; 
while (k != 0) 
{ 
    if (threadIdx.x < k) 
     local_result[threadIdx.x] += local_result[threadIdx.x + k]; 
    else 
     return; 

    __syncthreads(); 

    k /= 2; 
} 

とどのように私は戻ってグローバルメモリへの結果の書き込み：私が得たすべてのthats

data_result[blockIdx.x] = local_result[0]; 

をそれにもかかわらず、ここで不足しているコードがあります。

今、私はすべてが非ゼロである17要素の1つの行からなる行列でシナリオをテストしています。

float data_mat[17] = { val0, .., val16 } 
unsigned int row_ptr[2] = { 0, 17 } 
float data_vec[17] = { val0 } // all values are the same 
float data_result[1] = { 0 } 

そして、私のラッパー関数の抜粋のthats：：バッファが擬似コードで次のようになり

float *dev_data_mat; 
unsigned int *dev_row_ptr; 
float *dev_data_vec; 
float *dev_data_result; 

// Allocate memory on the device 
HANDLE_ERROR(cudaMalloc((void**) &dev_data_mat, num_nonzeroes * sizeof(float))); 
HANDLE_ERROR(cudaMalloc((void**) &dev_row_ptr, num_row_ptr * sizeof(unsigned int))); 
HANDLE_ERROR(cudaMalloc((void**) &dev_data_vec, dim_x * sizeof(float))); 
HANDLE_ERROR(cudaMalloc((void**) &dev_data_result, dim_y * sizeof(float))); 

// Copy each buffer into the allocated memory 
HANDLE_ERROR(cudaMemcpy(
     dev_data_mat, 
     data_mat, 
     num_nonzeroes * sizeof(float), 
     cudaMemcpyHostToDevice)); 
HANDLE_ERROR(cudaMemcpy(
     dev_row_ptr, 
     row_ptr, 
     num_row_ptr * sizeof(unsigned int), 
     cudaMemcpyHostToDevice)); 
HANDLE_ERROR(cudaMemcpy(
     dev_data_vec, 
     data_vec, 
     dim_x * sizeof(float), 
     cudaMemcpyHostToDevice)); 
HANDLE_ERROR(cudaMemcpy(
     dev_data_result, 
     data_result, 
     dim_y * sizeof(float), 
     cudaMemcpyHostToDevice)); 

// Calc grid dimension and block dimension 
dim3 grid_dim(dim_y); 
dim3 block_dim(BLOCK_SIZE); 

// Start kernel 
krnlSpMVmul1<<<grid_dim, block_dim, BLOCK_SIZE>>>(
     dev_data_mat, 
     num_nonzeroes, 
     dev_row_ptr, 
     dev_data_vec, 
     dev_data_result); 

私は、これは簡単であると思いますが、それは任意の関心がある場合は、物事を説明します。

もう1つ：私は、BLOCK_SIZEを128に、33個の非ゼロを持つと、カーネルにも障害が発生することに気付きました。やはり最後の値は計算されていません。

Answer 1

動的に割り当てられた共有メモリサイズが正しくありません。今、あなたはこれをやっている：

krnlSpMVmul1<<<grid_dim, block_dim, BLOCK_SIZE>>>(.....) 

共有メモリサイズがバイトに与えられるべきです。ブロックごとに64個のスレッドを使用すると、16個の浮動小数点ワードに十分な共有メモリが割り当てられ、行ごとの魔法の17個のエントリでエラーが発生する理由を説明します。共有バッファオーバーフローが発生し、 GPUを実行し、カーネルを中止します。

krnlSpMVmul1<<<grid_dim, block_dim, BLOCK_SIZE * sizeof(float)>>>(.....) 

あなたは正しい動的共有メモリー・サイズを与えるとの問題を解消する必要があります

あなたはこのような何かをやっている必要があります。