Лекция 9. Драйвер компиляции (1265190), страница 2

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

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

Для файлов программы компилируются хостовые объектныефайлы, в которые встроены cubin-ы с relocalable-кодом2. Утилита nvlink вытаскивает из объектных файлов cubin-ыдля заданной архитектуры и линкует, в результатегенерируется новый cubin3. Этот cubin запаковывается в fatbin и встраивается вфиктивный хостовый объектный файл4. Фиктивный объектный файл линкуется с остальнымиобъектными файламиСхема линковкиОпции NVCC –rdc={true,false}Прокинуть в ptxas опцию --compile-only для генерациилинкуемых *.cubin-ов -dcТо же самое, что “–c –rdc=true”.

Остановить цепочку наэтапе генерации объектных файлов. В объектные файлы будутвстроены кубины с relocalable кодом -dlink –arch=…Остановить цепочку на этапе генерации фиктивного объектногофайла с встроенным cubin-результатом линковкиJIT линковка Начиная с cuda 5.5 поддерживается JIT-линковка JIT-компилируется бинарный код для текущейархитектуры, затем линкуется До cuda 5.5 такой возможности не было Нет бинарного кода – ошибка запускаЭффективность кода Ядра с влинкованными зависимостями эффективнее,чем ядра с вызовами функций-зависимостей Нет накладных расходов на вызов Эффективное совместное распределение регистров Поэтому по умолчанию происходит whole programcompilationПлатформонезависимость NVCC Компиляция device-кода происходит независимо откомпиляции хост-кода Можно использовать любой компилятор для хост-кода Скомпилированный device-код встраивается в объектныемодули как текстовая переменная, прозрачно для хосткомпилятора Можно использовать любой формат хранения данных,необходимых для выполнения работы на устройстве Этот формат может быть единым для всех платформПлатформонезависимость NVCC Для компиляции host-кода, создания объектных модулей иитоговой линковки может быть использован любойкомпилятор C/C++ На Windows используем cl Для получения COFF, на *nix gcc для получения ELF В итоге получаем обычный исполняемый файл,специфический для ОС хоста *.exe на Windows, Unix Executable на *nixПромежуточное представлениеАссемблерPTX PTX – “Paralle Thread Execution” Псевдо-ассемблерный язык Привязан к «Виртуальной архитектуре» Генерируется в текстовом формате компилятором cicc(ранее использовался nvopencc) из CUDA C кода Может быть получен по ключу nvcc –ptx, nvcc –keep, cuobjdump -ptx Подробная документация в файлеhttp://docs.nvidia.com/cuda/parallel-threadexecution/index.htmlОсобенности PTX Треахдресные инструкции с указанием типовоперандов shr.u64 %rd14, %rd12, 32;побитовый сдвиг вправо числа из %rd12 на 32 позиции,результат записать в %rd14 cvt.u64.u32 %rd142, %r112;преобразовать unsigned 32-bit integer к 64-bitРегистры Виртуальные регистры (будут отображены на реальныепри компиляции под конкретную архитектуру) reg .u32 %r<335>;выделить 335 регистров %r0, %r1, ..., %r334 типаunsigned 32-bit integer Специальные предопределенные регистры, такие как%tid, %ntid, %ctaid, и %nctaid содержат индексынити, размеры блока, индексы блока, размеры гридаВыделение памяти Выделение статической общей памяти .shared .align 8 .b8 pbatch_cache[15744];выделить 15744 общей памяти Выделение статической константной памяти .const .align 4 .b8 staticKernel[16900]; Выделение статической глобальной памяти .global .align 4 .b8 staticKernel[16900];Сравнения и переходы Инструкция сравнения setp.lt.s32 %p9, %r14, %r9;%p9= %r14 < r%9 Переход по значению предикатного регистра @%p9 bra $label;Перейти на метку label, если %p9 != 0Обращения в памятьop.space.type [куда], [откуда]; op =[ld | st] space= [.reg, .sreg, .local, .global, .param, .shared, .tex, .const] type = [u64 | f32 | … ]ld.global.f32 %f8, [%rd32];st.global.f32 [%rd30], %f9;Movmov.u32 %r2, %ctaid.x;Переслать из специального регистра в обычныйmov.f32 %f9, 0f00000000;%f9 = 0f00000000;mov.u32 %r31, %r22;%r31=%r22mov.u64 %r24, staticKernel;%r24 = &staticKernelПример ptx для суммированияПример ptx для суммированияПример ptx для суммированияПример ptx для суммированияПример ptx для суммированияISA ISA – Instruction Set Architecture Fermi ISA – ассемблер для видеокарт архитектурыферми, бинарный код для выполнения на GPU Компилируется из ptx компилятором ptxas Инструкции трёхадресные Очень кратко описан в http://docs.nvidia.com/cuda/cuda-binary-utilities/index.htmlПолучение ISA $cuobjdump –sass [cubin] $cuobjdump –sass [fatbin] $cuobjdump –sass [elf с встроенным fatbin]Пример ISA для суммированияСмещение/*0000*//*0008*//*0010*//*0018*//*0020*//*0028*//*0030*//*0038*//*0040*//*0048*//*0050*//*0058*//*0060*//*0068*//*0070*//*0078*//*0080*/Бинарный код/*0x00005de428004404*//*0x94001c042c000000*//*0x84009c042c000000*//*0x10015de218000000*//*0x2000dca320044000*//*0x10311ce35000c000*//*0x80329ca3200b8000*//*0x9042dc4348004000*//*0xa0321ca3200b8000*//*0x00a01c8584000000*//*0xb0425c4348004000*//*0xc0319ca3200b8000*//*0x00809c8584000000*//*0xd041dc4348004000*//*0x00201c0348000000*//*0x00601c8594000000*//*0x00001de780000000*/Текстовое представлениеMOV R1, c [0x1] [0x100];S2R R0, SR_CTAid_X;S2R R2, SR_Tid_X;MOV32I R5, 0x4;IMAD R3, R0, c [0x0] [0x8], R2;IMUL.HI R4, R3, 0x4;IMAD R10.CC, R3, R5, c [0x0] [0x20];IADD.X R11, R4, c [0x0] [0x24];IMAD R8.CC, R3, R5, c [0x0] [0x28];LD.E R0, [R10];IADD.X R9, R4, c [0x0] [0x2c];IMAD R6.CC, R3, R5, c [0x0] [0x30];LD.E R2, [R8];IADD.X R7, R4, c [0x0] [0x34];IADD R0, R2, R0;ST.E [R6], R0;EXIT;ISA vs PTX Начинать низкоуровневый анализ программы следует с PTX Код короче и понятнее, привязан к строчкам при помощидиректив .loc НО: в PTX не распределены регистры => неизвестно скольколокальной памяти будет использовано В PTX не применены платформо-зависимые преобразования ISA – именно тот код, который будет исполнятьсяISA vs PTX PTX достаточно подробно документирован PTX может быть скомпилирован компилятором ptxas На PTX можно писать глубоко оптимизированные ядра /редактировать ядра, созданные cicc ISA не документирован. Есть дизассемблер (cuobjdump), но ассемблера от Nvidia нет. Тем не менее, в проекте https://code.google.com/p/asfermi/ISA декодирован и можно транслировать текстовоепредставление ассемблера в бинарноеПрактика Постепенно усложняя ядра и исследуя ассемблер,найти ответы на следующие вопросы:Где расположен стек?Как происходят вызовы функций? А рекурсия?Как передаются параметры в ядра? А грид?Как выглядят операции с общей/константной памятью?Как определить где используется локальная память? Посмотреть на device relocalable code Посмотреть на эффект от добавления__attribute__((noinline)) к __device__ функциям,вызываемым из ядерCUDA Driver API Более низкоуровневое API Позволяет явно управлять стеком контекстов устройств,загрузкой объектным модулей, загрузкой символов изобъектных модулей и т.д. CUDA Runtime API реализовано через использованиеCUDA Driver API Кратко описан в секции H.

Driver API cudaprogramming guideИнициализация, контекстcuInit(0); // Явная инициализацияint deviceCount;// Число устройствcuDeviceGetCount(&deviceCount);if (deviceCount == 0) exit(0);CUdevice cuDevice; // получить указатель на устройствоcuDeviceGet(&cuDevice, 0);CUcontext cuContext; // создать контекстcuCtxCreate(&cuContext, 0, cuDevice);Загрузка кода для GPUCUresult cuModuleLoad (CUmodule ∗module, const char∗fname) Грузит модуль с кодом для GPU из файла – текстовыйдокумент с ptx, cubin, fatbinCUresult cuModuleLoadData (CUmodule ∗ module, constvoid ∗ image) Грузит модуль из бинарного представления – загруженныйв память кубин / NULL-terminated строка с PTXJIT-компиляция cuModuleLoad и cuModuleLoadData грузят модули с кодом втекущий контекст.

При отсутствии бинарного кода дляустройства, на котором создан контекст, компилируется PTX Для явного управления компиляцией PTX:CUresult cuModuleLoadDataEx (CUmodule ∗module,const void ∗image,unsigned int numOptions, // Число опцийCUjit_option ∗options,// опцииvoid ∗∗optionValues// значения опцийОпции JIT-компиляцииCU_JIT_MAX_REGISTERSCU_JIT_THREADS_PER_BLOCKCU_JIT_INFO_LOG_BUFFERCU_JIT_INFO_LOG_BUFFER_SIZE_BYTESCU_JIT_ERROR_LOG_BUFFERCU_JIT_ERROR_LOG_BUFFER_SIZE_BYTESCU_JIT_OPTIMIZATION_LEVELCU_JIT_TARGET_FROM_CUCONTEXTCU_JIT_TARGET CU_TARGET_COMPUTE_10, CU_TARGET_COMPUTE_11... CU_JIT_FALLBACK_STRATEGY: CU_PREFER_PTX CU_PREFER_BINARYПолучение символов из модуля СUresult cuModuleGetFunction (CUfunction ∗ hfunc,CUmodule hmod, const char ∗ name) Вытащить из модуля ядро или device-функцию CUresult cuModuleGetGlobal ( CUdeviceptr ∗ dptr,size_t ∗bytes, CUmodule hmod, const char ∗ name) Размер и указатель на память, выделенную под статическуюпеременную с атрибутом __device__ или __constant__Запуск ядерCUresult cuLaunchKernel (CUfunction f,unsigned int gridDimX,unsigned int gridDimY,unsigned int gridDimZ,unsigned int blockDimX,unsigned int blockDimY,unsigned int blockDimZ,unsigned int sharedMemBytes,CUstream hStream, void ∗∗kernelParams,void ∗∗ extra)Пример запуска// Загрузка модуляCUmodule module;cuModuleLoad(&module, "sum_kernel.cubin");// Загрузка ядра из модуляCUfunction kernel;cuModuleGetFunction(&kernel, module, "kernel");// Запуск ядраvoid *args[] = {(void *)&aDev, (void * )&bDev, (void *)&cDev};cuLaunchKernel(kernel, n / BLOCK_SIZE, 1, 1,BLOCK_SIZE, 1, 1, 0, 0, args, NULL);Параметры ядра Во время компиляции: -Xptxas -v$ nvcc -Xptxas -v sub_kernel.cuptxas info: Compiling entry function '_Z10sub_kernelPfS_S_' for 'sm_10'ptxas info: Used 4 registers, 24+16 bytes smem В рантайме:cuFuncGetAttribute ( int ∗ pi, CUfunction_attributeattrib, CUfunction hfunc);cudaFuncGetAttributes ( struct cudaFuncAttributes ∗attr,const char ∗ func)Параметры ядраcudaFuncAttributes fieldCUfunction_attributeint binaryVersionCU_FUNC_ATTRIBUTE_BINARY_VERSIONsize_t constSizeBytesCU_FUNC_ATTRIBUTE_CONST_SIZE_BYTESsize_t localSizeBytesCU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTESint maxThreadsPerBlockCU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCKint numRegsCU_FUNC_ATTRIBUTE_NUM_REGSint ptxVersionCU_FUNC_ATTRIBUTE_PTX_VERSIONsize_t sharedSizeBytesCU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTESПрактика Получить cubin или ptx для некоторого ядра по ключу cubin, -ptx, -keepСоздать контекстЗагрузка модуля cuModuleLoad, cuModuleLoadDataЗагрузка ядра из модуля cuModuleGetFunctionПолучение параметров ядра cuFuncGetAttribute CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES CU_FUNC_ATTRIBUTE_NUM_REGSВывести в консоль параметрыThe end.

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

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