Похоже, что максимально допустимое количество резидентных блоков на SM. Но в то время как другие «жесткие» ограничения легко найти (например, через «cudaGetDeviceProperties»), максимальное количество резидентных блоков, похоже, не документировано широко.
В следующем примере кода я настраиваю ядро с одним потоком на блок. Чтобы проверить гипотезу о том, что этот графический процессор (P100) имеет максимум 32 резидентных блока на SM, я создаю сетку из 56 * 32 блоков (56 = количество SM на P100). Каждому ядру требуется 1 секунда для обработки (с помощью процедуры «сна»), поэтому, если я правильно настроил ядро, код должен занять 1 секунду. Результаты сроков подтверждают это. Конфигурирование с 32 * 56 + 1 блоками занимает 2 секунды, предполагая, что 32 блока на SM являются максимально допустимыми для SM.
Интересно, почему этот предел не стал более доступным? Например, он не отображается `cudaGetDeviceProperties '. Где я могу найти этот предел для разных графических процессоров? Или, может быть, это не реальный предел, а вытекает из других жестких ограничений?
Я использую CUDA 10.1
#include <stdio.h>
#include <sys/time.h>
double cpuSecond() {
struct timeval tp;
gettimeofday(&tp,NULL);
return (double) tp.tv_sec + (double)tp.tv_usec*1e-6;
}
#define CLOCK_RATE 1328500 /* Modify from below */
__device__ void sleep(float t) {
clock_t t0 = clock64();
clock_t t1 = t0;
while ((t1 - t0)/(CLOCK_RATE*1000.0f) < t)
t1 = clock64();
}
__global__ void mykernel() {
sleep(1.0);
}
int main(int argc, char* argv[]) {
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
int mp = prop.multiProcessorCount;
//clock_t clock_rate = prop.clockRate;
int num_blocks = atoi(argv[1]);
dim3 block(1);
dim3 grid(num_blocks); /* N blocks */
double start = cpuSecond();
mykernel<<<grid,block>>>();
cudaDeviceSynchronize();
double etime = cpuSecond() - start;
printf("mp %10d\n",mp);
printf("blocks/SM %10.2f\n",num_blocks/((double)mp));
printf("time %10.2f\n",etime);
cudaDeviceReset();
}
Результаты:
% srun -p gpuq sm_short 1792
mp 56
blocks/SM 32.00
time 1.16
% srun -p gpuq sm_short 1793
mp 56
blocks/SM 32.02
time 2.16
% srun -p gpuq sm_short 3584
mp 56
blocks/SM 64.00
time 2.16
% srun -p gpuq sm_short 3585
mp 56
blocks/SM 64.02
time 3.16