Увеличение нагрузки на регистры без снижения occupancy просаживает производительность

Переношу работу с нейронными сетями на GPU, пока написал простую болванку для оценки производительности сети из 3-х слоев с количеством нейронов в кажом слое 256-128-5

  1. __device__ float sigmoid_symmetric(float sum)
  2. {
  3.         return (2.0f/(1.0f + __expf(-2.0f * sum)) - 1.0f);
  4. }
  5.  
  6. template <int totalNeurons>
  7. __global__ void
  8. //__launch_bounds__(128, 8)
  9. nn_test(float *weights, float *inputs, float *outputs, int inputPitch, int outputPitch)
  10. {
  11.         const int num_layers = 3;
  12.         const int layer_size[num_layers] = {256, 128, 5};
  13.         const int layer_index[num_layers+1] = {0, 256, 256 + 128, 256 + 128 + 5};
  14.         //const float steepness = 1;
  15.         const float max_sum = 150;///steepness;
  16.         //const float max_sum_neg = -150/steepness;
  17.  
  18.         __shared__ float val [totalNeurons];
  19.  
  20.     unsigned int tid = threadIdx.x;
  21.         float *locInputs = inputs + blockIdx.x * inputPitch;
  22.         float *locOutputs = outputs + blockIdx.x * outputPitch;
  23.  
  24.         //input layer
  25.         for(int i = tid; i < layer_size[0]; i += blockDim.x)
  26.         {
  27.                 //last neuron is the bias
  28.                 val[i] = (i != layer_size[0]-1) ? locInputs[i] : 1;
  29.         }
  30.        
  31.         //hidden layer(s)
  32.         float *weightsIt = weights;
  33.         for(int layer = 1; layer < num_layers; layer++)
  34.         {
  35.                 __syncthreads();
  36.                 if(layer > 1)
  37.                         weightsIt += layer_size[layer] * layer_size[layer - 1];
  38.  
  39.                 if(tid < layer_size[layer])
  40.                 {
  41.                         int num_connections = layer_size[layer - 1];
  42.                         float *weightsNeuron = weightsIt + tid;
  43.                         int neuronIndex = layer_index[layer] + tid;
  44.                         int neuronIndexPrev = layer_index[layer-1];
  45.                
  46.                         float neuron_sum = 0;
  47.                         for(int con = 0; con < num_connections; con ++)
  48.                         {
  49.                                 neuron_sum += weightsNeuron[con * num_connections] * val[neuronIndexPrev + con];
  50.                         }
  51.                         //neuron_sum *= 1;//steepness;
  52.                         //if(neuron_sum > max_sum)
  53.                         //      neuron_sum = max_sum;
  54.                         //else if(neuron_sum < -max_sum)
  55.                         //      neuron_sum = -max_sum;
  56.                         neuron_sum = fmaxf(-max_sum, fminf(neuron_sum, max_sum));
  57.  
  58.                         val[neuronIndex] = 0;//sigmoid_symmetric(neuron_sum);
  59.                 }
  60.         }
  61.  
  62.         //output layer
  63.         __syncthreads();
  64.         int outputIndex = layer_index[num_layers - 1];
  65.         if(tid < layer_size[num_layers - 1])
  66.         {
  67.                 locOutputs[tid] = val[outputIndex + tid];
  68.         }
  69. }

С coalesced чтением разобрался, это увеличило скорость вдвое, но возникли другие приколы.
Карта GTX460, SM 2.1, CUDA 4.0 RC2

запускаю блоками по 128 потоков, на мультипроцессор влезает по 8 блоков, итого 1024 потока на мультипроцессор, occupancy 67%

В таком виде каждый поток использует по 14 регистров.
Если neuron_sum = fmaxf(-max_sum, fminf(neuron_sum, max_sum)); заменить на

  1.                         if(neuron_sum > max_sum)
  2.                                 neuron_sum = max_sum;
  3.                         else if(neuron_sum < -max_sum)
  4.                                 neuron_sum = -max_sum;

или раскомментировать вызов sigmoid_symmetric, количество используемых регистров резко возрастает до 22 и скорость падает в 5 (пять!) раз. Уже весь мозг сломал, не могу понять почему.
На мультипроцессор при этом расходуется 22K регистров из 32K, occupancy не падает как по калькулятору, так и по отчету профайлера. Тогда где тут собака порылась?

Forums: 

А если все то же самое

А если все то же самое (включить ваш sigmoid_symmetric), но компилировать (nvcc) с -use-fast-math ?

Без разницы.Я уже

Без разницы.

Я уже разобрался, как и в соседней теме, при правках кода компилятор просто выкидывал нафиг внутренний цикл, отсюда и регистров меньше, и чудеса скорости. Потому тему можно считать исчерпанной.

Кернел получился compute bound, потому или оставить как есть (если большими пачками данные подсовывать), или усложнять схему и каждый нейрон обрабатывать в 2 или 4 потока (для пачек поменьше)

Забавно, что в самом первом варианте с не coalesced чтением процент попадания в L1 кэш был более 90%, в текущем всего 22% из-за чтения из разных мест, зато без конфликта банков.

P.S. Пытался сам себе в тему ответить, чтоб народ зря не будоражить, так антиспам не пропустил.

Антиспам должен в

Антиспам должен в сомнительных случаях капчу спросить, а не просто не пропустить?

Капча будет в самом низу

Спрашивал, вводил, с 3-х

Спрашивал, вводил, с 3-х попыток не прокатило.

Есть подозрение, что ему не понравилось, что первый же коммент от автора темы