Ошибка "Out of resources" при развертывании цикла - PullRequest
4 голосов
/ 28 сентября 2011

Когда я увеличиваю развертывание с 8 до 9 циклов в моем ядре, оно прерывается с ошибкой out of resources.

Я прочитал в Как диагностировать сбой запуска CUDA из-за нехватки ресурсов? , что несоответствие параметров и чрезмерное использование регистров могут быть проблемой, но здесь это не так. .

Мое ядро ​​вычисляет расстояние между n точками и m центроидами и выбирает для каждой точки ближайший центроид. Это работает для 8 измерений, но не для 9. Когда я устанавливаю dimensions=9 и раскомментирую две строки для вычисления расстояния, я получаю pycuda._driver.LaunchError: cuLaunchGrid failed: launch out of resources.

Как вы думаете, могло ли это быть причиной такого поведения? Какие еще иусы могут вызвать out of resources*?

Я использую Quadro FX580. Вот минимальный (ish) пример. Для раскрутки в реальном коде я использую шаблоны.

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

1 Ответ

8 голосов
/ 01 октября 2011

У вас закончились регистры, потому что ваш 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), либо найти способ настроить PyCUDA для компиляции вашего ядра для устройства SM 2.0. Устройства SM 2.0, такие как QuadroFX 580, имеют 32 тыс. Регистров, что более чем достаточно для исходных block_size из 512.

...