解釈

Question

は、次のPythonコードを考えてみましょう。これを実行

from numpy import float64 
from pycuda import compiler, gpuarray 
import pycuda.autoinit 

# N > 960 is crucial! 
N = 961 
code = """ 
__global__ void kern(double *v) 
{ 
    double a = v[0]*v[2]; 
    double lmax = fmax(0.0, a), lmin = fmax(0.0, -a); 
    double smax = sqrt(lmax), smin = sqrt(lmin); 

    if(smax > 0.2) { 
     smax = fmin(smax, 0.2)/smax ; 
     smin = (smin > 0.0) ? fmin(smin, 0.2)/smin : 0.0; 
     smin = lmin + smin*a; 

     v[0] = v[0]*smin + smax*lmax; 
     v[2] = v[2]*smin + smax*lmax; 
    } 
} 
""" 
kernel_func = compiler.SourceModule(code).get_function("kern") 
kernel_func(gpuarray.zeros(3, float64), block=(N,1,1)) 

ができます：

Traceback (most recent call last): 
    File "test.py", line 25, in <module> 
    kernel_func(gpuarray.zeros(3, float64), block=(N,1,1)) 
    File "/usr/lib/python3.5/site-packages/pycuda/driver.py", line 402, in function_call 
    func._launch_kernel(grid, block, arg_buf, shared, None) 
pycuda._driver.LaunchError: cuLaunchKernel failed: too many resources requested for launch 

マイセットアップ：pycudaでのPython v3.5.2 == 2016年1月2日Ubuntu 16.04.1（64-bit）、kernel 4.4.0、nvcc V7.5.17ではnumpy == 1.11.1です。グラフィックスカードはNvidia GeForce GTX 480です。

これはマシンで再現できますか？このエラーメッセージの原因は何ですか？

備考：すべてのカーネルがv [0]とv [2]を変更しようとするため、原則として競合状態が存在することがわかりました。しかし、カーネルはifブロックの内側に届かないはずです！さらに、競合状態がなければエラーを再現できますが、はるかに複雑です。

Answer 1

ブロック当たりの制限を打つことはほぼ確実です。

relevant documentationを読むと、デバイスのブロックあたりの制限は32k 32ビットです。ブロックサイズが960スレッド（30ワープ）を超えると、カーネルの起動に必要なレジスタが多すぎるため、起動に失敗します。 NVIDIAはExcelのスプレッドシートとadviceを提供しています。スレッドごとにカーネルのレジスタ要件とカーネルがデバイスで起動するために使用できる制限ブロックサイズを決定する方法について説明します。