5

カーネルで展開ループを 8 から 9 に増やすと、out of resourcesエラーで中断します。

リソース不足による CUDA の起動失敗を診断するにはどうすればよいですか?を読みました。パラメータの不一致とレジスタの使いすぎが問題になる可能性がありますが、ここではそうではないようです。

n私のカーネルは、点と重心の間の距離を計算し、m各点について最も近い重心を選択します。8 次元では機能しますが、9 次元では機能しません。dimensions=9距離計算のために 2 行を設定してコメントを外すと、pycuda._driver.LaunchError: cuLaunchGrid failed: launch out of resources.

この動作の原因は何だと思いますか? out of resources*の原因となるその他の iusses は何ですか?

Quadro FX580 を使用しています。これが最小限の（っぽい）例です。実際のコードで展開するには、テンプレートを使用します。

import numpy as np
from pycuda import driver, compiler, gpuarray, tools
import pycuda.autoinit


## preference
np.random.seed(20)
points = 512
dimensions = 8
nclusters = 1

## init data
data = np.random.randn(points,dimensions).astype(np.float32)
clusters = data[:nclusters]

## init cuda
kernel_code = """

      // the kernel definition 
    __device__ __constant__ float centroids[16384];

    __global__ void kmeans_kernel(float *idata,float *g_centroids,
    int * cluster, float *min_dist, int numClusters, int numDim) {
    int valindex = blockIdx.x * blockDim.x + threadIdx.x ;
    float increased_distance,distance, minDistance;
    minDistance = 10000000 ;
    int nearestCentroid = 0;
    for(int k=0;k<numClusters;k++){
      distance = 0.0;
      increased_distance = idata[valindex*numDim] -centroids[k*numDim];
      distance = distance +(increased_distance * increased_distance);
      increased_distance =  idata[valindex*numDim+1] -centroids[k*numDim+1];
      distance = distance +(increased_distance * increased_distance);
      increased_distance =  idata[valindex*numDim+2] -centroids[k*numDim+2];
      distance = distance +(increased_distance * increased_distance);
      increased_distance =  idata[valindex*numDim+3] -centroids[k*numDim+3];
      distance = distance +(increased_distance * increased_distance);
      increased_distance =  idata[valindex*numDim+4] -centroids[k*numDim+4];
      distance = distance +(increased_distance * increased_distance);
      increased_distance =  idata[valindex*numDim+5] -centroids[k*numDim+5];
      distance = distance +(increased_distance * increased_distance);
      increased_distance =  idata[valindex*numDim+6] -centroids[k*numDim+6];
      distance = distance +(increased_distance * increased_distance);
      increased_distance =  idata[valindex*numDim+7] -centroids[k*numDim+7];
      distance = distance +(increased_distance * increased_distance);
      //increased_distance =  idata[valindex*numDim+8] -centroids[k*numDim+8];
      //distance = distance +(increased_distance * increased_distance);

      if(distance <minDistance) {
        minDistance = distance ;
        nearestCentroid = k;
        } 
      }
      cluster[valindex]=nearestCentroid;
      min_dist[valindex]=sqrt(minDistance);
    } 
 """
mod = compiler.SourceModule(kernel_code)
centroids_adrs = mod.get_global('centroids')[0]    
kmeans_kernel = mod.get_function("kmeans_kernel")
clusters_gpu = gpuarray.to_gpu(clusters)
cluster = gpuarray.zeros(points, dtype=np.int32)
min_dist = gpuarray.zeros(points, dtype=np.float32)

driver.memcpy_htod(centroids_adrs,clusters)

distortion = gpuarray.zeros(points, dtype=np.float32)
block_size= 512

## start kernel
kmeans_kernel(
    driver.In(data),driver.In(clusters),cluster,min_dist,
    np.int32(nclusters),np.int32(dimensions),
    grid = (points/block_size,1),
    block = (block_size, 1, 1),
)
print cluster
print min_dist
4

1 に答える 1

8

block_size(512) が大きすぎるため、レジスタが不足しています。

ptxasカーネルがコメント行で 16 個のレジスタを使用していることを報告します。

$ nvcc test.cu -Xptxas --verbose
ptxas info    : Compiling entry function '_Z13kmeans_kernelPfS_PiS_ii' for 'sm_10'
ptxas info    : Used 16 registers, 24+16 bytes smem, 65536 bytes cmem[0]

行のコメントを外すと、レジスタの使用が 17 に増え、実行時にエラーが発生します。

$ nvcc test.cu -run -Xptxas --verbose
ptxas info    : Compiling entry function '_Z13kmeans_kernelPfS_PiS_ii' for 'sm_10'
ptxas info    : Used 17 registers, 24+16 bytes smem, 65536 bytes cmem[0]
error: too many resources requested for launch

カーネルの各スレッドで使用される物理レジスタの数によって、実行時に起動できるブロックのサイズが制限されます。SM 1.0 デバイスには、スレッドのブロックで使用できる 8K のレジスタがあります。これをカーネルのレジスタ要求と比較できます: 17 * 512 = 8704 > 8K. 16 個のレジスタで、元のコメント付きカーネルは次のようにきしむだけです: 16 * 512 = 8192 == 8K.

アーキテクチャが指定されていない場合nvcc、デフォルトで SM 1.0 デバイス用のカーネルをコンパイルします。PyCUDA も同じように動作する可能性があります。

問題を解決するには、減らすblock_size(たとえば 256) か、SM 2.0 デバイス用にカーネルをコンパイルするように PyCUDA を構成する方法を見つけることができます。QuadroFX 580 などの SM 2.0 デバイスは、元block_sizeの 512 には十分な 32K レジスタを提供します。

于 2011-10-01T02:16:33.397 に答える