Это обычно случай установки размера блока для оптимальной производительности и размера сетки в соответствии с общим объемом работы.У большинства ядер есть количество точек перекоса на Mp, где они работают лучше всего, и вы должны сделать несколько тестов / профилировок, чтобы увидеть, где это.Вам, вероятно, по-прежнему нужна логика переполнения в ядре, потому что размеры проблем редко бывают кратными размерам блоков.
РЕДАКТИРОВАТЬ: привести конкретный пример того, как это можно сделать для простого ядра (в данном случаепользовательская операция типа dscal BLAS уровня 1, выполняемая как часть факторизации Холецкого упакованных матриц симметричных полос):
// Fused square root and dscal operation
__global__
void cdivkernel(const int n, double *a)
{
__shared__ double oneondiagv;
int imin = threadIdx.x + blockDim.x * blockIdx.x;
int istride = blockDim.x * gridDim.x;
if (threadIdx.x == 0) {
oneondiagv = rsqrt( a[0] );
}
__syncthreads();
for(int i=imin; i<n; i+=istride) {
a[i] *= oneondiagv;
}
}
Для запуска этого ядра параметры выполнения рассчитываются следующим образом:
- Мы допускаем до 4 деформаций на блок (т. Е. 128 потоков).Обычно вы исправляете это в оптимальном количестве, но в этом случае ядро часто вызывается для очень маленьких векторов, поэтому иметь переменный размер блока имеет смысл.
- Затем мы вычисляем количество блоков в соответствии с общимобъем работы, всего до 112 блоков, что эквивалентно 8 блокам на MP на 14 MP Fermi Telsa.Ядро будет выполнять итерацию, если объем работы превышает размер сетки.
Результирующая функция-обертка, содержащая вычисления параметров выполнения и запуск ядра, выглядит следующим образом:
// Fused the diagonal element root and dscal operation into
// a single "cdiv" operation
void fusedDscal(const int n, double *a)
{
// The semibandwidth (column length) determines
// how many warps are required per column of the
// matrix.
const int warpSize = 32;
const int maxGridSize = 112; // this is 8 blocks per MP for a Telsa C2050
int warpCount = (n / warpSize) + (((n % warpSize) == 0) ? 0 : 1);
int warpPerBlock = max(1, min(4, warpCount));
// For the cdiv kernel, the block size is allowed to grow to
// four warps per block, and the block count becomes the warp count over four
// or the GPU "fill" whichever is smaller
int threadCount = warpSize * warpPerBlock;
int blockCount = min( maxGridSize, max(1, warpCount/warpPerBlock) );
dim3 BlockDim = dim3(threadCount, 1, 1);
dim3 GridDim = dim3(blockCount, 1, 1);
cdivkernel<<< GridDim,BlockDim >>>(n,a);
errchk( cudaPeekAtLastError() );
}
Возможно, это даетнекоторые советы о том, как разработать «универсальную» схему для установки параметров выполнения в зависимости от размера входных данных.