Лекция 4. Константная память. Регистры и локальная память (1265185), страница 2

Просмтор этого файла доступен только зарегистрированным пользователям. Но у нас супер быстрая регистрация: достаточно только электронной почты!

Текст из файла (страница 2)

в локальной памяти моделируется стек фреймов прирекурсивных вызовахПример__device__ int deviceFunc(int *a) {int x; // скорее всего на регистре (если нет спиллинга)int array[10]; // может в локальной, т.к. далее этот массивиндексируется неизвестным при компиляции индексом...x = sinf(threadIdx.x) // sinf можетиспользовать локальную памятьx = x + array[ threadIdx.x % 10 ];if (x < 100) {x = x + deviceFunc(a); // фрейм вызова будет// сохранен в стеке, расположенном в локальной памяти}return x;}nvcc -Xptxas -v Выводит количество регистров, константной памяти, локальнойпамяти и статической общей памяти, используемые ядром:$:~/programming/testMod$ nvcc -arch=sm_20 -Xptxas -v test.cuptxas info: 0 bytes gmem, 8 bytes cmem[2]ptxas info: Compiling entry function '_Z13matmul_kernelv' for'sm_20'ptxas info: Function properties for _Z13matmul_kernelv8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loadsptxas info: Used 8 registers, 4 bytes smem, 32 bytes cmem[0]Статические ресурсы и occupancy Факторы, влияющие на occupancy: Не более 1536 нитей на sm, не более 8 блоков Не более 48KB общей памяти на sm 32768 регистров на smПусть ядро использует 63 регистра и размер блока 384 нити32768 / 63 = 520 – максимум нитейна sm будет работать всего один блок из 384 нитей(occupancy = 0.25)Статические ресурсы и occupancy Факторы, влияющие на occupancy: Не более 1536 нитей на sm, не более 8 блоков Не более 48KB общей памяти на sm 32768 регистров на smПусть на блок из 512 нитей нужно 32KB общей памятина sm будет работать всего один блок из 512 нитей(occupancy = 0.33)CUDA occupancy calculator http://developer.download.nvidia.com/compute/cuda/CUDA_Occupancy_calculator.xlsCUDA occupancy calculator http://developer.download.nvidia.com/compute/cuda/CUDA_Occupancy_calculator.xlsCUDA occupancy calculatorCUDA occupancy calculatorCUDA occupancy calculatorНехватка ресурсов для запуска Если на мультрипроцессоре не хватает регистров даже для одногоблока, то произойдет ошибка запуска Нужно уменьшить размер блока Пример: Ядро, потребляющее 37 регистров, запустится на блоке из 1024нитей на Kepler и не запустится на Fermi То же самое и в отношении общей памяти – для успешного запуска еёдолжно хватать хотя бы на один блокОграничение числа регистров$nvcc –maxrregcount 20 –c –arch=sm_20 kernel.cu При компиляции принудительно ограничить максимальноечисло используемых регистров Можно контролировать occupancy Приводит к спиллингу регистровОганичение числа регистровПолучаем использование локальной памяти:$ nvcc -arch=sm_20 -Xptxas -v -c kernel.cuptxas info: 0 bytes gmemptxas info: Compiling entry function '_Z6matmul14cudaPitchedPtrS_S_' for 'sm_20‘ptxas info: Function properties for _Z6matmul14cudaPitchedPtrS_S_0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loadsptxas info: Used 32 registers, 128 bytes cmem[0]$ nvcc -arch=sm_20 -maxrregcount 21 -Xptxas -v -c kernel.cuptxas info: 0 bytes gmemptxas info: Compiling entry function '_Z6matmul14cudaPitchedPtrS_S_' for 'sm_20‘ptxas info: Function properties for _Z6matmul14cudaPitchedPtrS_S_96 bytes stack frame, 132 bytes spill stores, 112 bytes spill loadsptxas info: Used 21 registers, 128 bytes cmem[0]Выводы Всегда компилировать с -Xptxas –v и –arch=sm_?? -Xptxas –v выводит использование ресурсов только дляархитектуры, для которой происходит компиляция Учитывать потребление регистров и общей памяти при задании гридадля максимизации occupancy Использовать для этого occupancy calculator Минимизировать использование локальной памяти Надежный способ узнать из-за чего возникло использованиелокальной памяти – посмотреть ptx/ассемблер, об этом в другойчастиThe end.

Характеристики

Список файлов лекций