Рассмотрим следующую ситуацию: есть приложение, которое использует AMD GPU для своих вычислений. Как правило, на GPU выносятся самые ресурсоемкие операции. Поэтому если приложение работает быстрее своих конкурентов, то может возникнуть желание узнать, какой же алгоритм реализован в этой программе. Но что делать, если программа является проприетарной и распространяется по лицензии, запрещающей reverse engineering и дизассемблирование?
Чтобы не нарушать лицензии, можно воспользоваться одной маленькой хитростью, оставленной разработчиками AMD APP SDK. Однако чтобы эта хитрость сработала, необходимо выполнение еще одного условия (помимо использования разработчиками приложения указанного SDK): приложение должно использовать OpenCL для вычислений на GPU.
Если внимательно почитать документацию AMD Accelerated Parallel Processing OpenCL™ Programming Guide (v1.3f), то в разделе «4.2.1 Intermediate Language and GPU Disassembly» можно найти одну замечательную переменную окружения: GPU_DUMP_DEVICE_KERNEL. Она может принимать 3 значения:
- 1 — сохранить dump на языке AMD IL в локальной директории
- 2 — дизассемблировать ISA файл и сохранить результат в локальной директории
- 3 — сделать оба действия
Эксперимент
Установим значение переменной окржения GPU_DUMP_DEVICE_KERNEL=3. В качестве подопытной программы возьмем пример из AMD APP SDK — программу бинарного поиска BinarySearch.exe. Этот пример не самый интересный, так как рядом уже лежит файл с исходником ядра для GPU: BinarySearch_Kernels.cl. Однако в реальной жизни программы не хранят такую ценную информацию в открытом виде, ее либо шифруют, либо хранят внутри программы.
Итак, после запуска BinarySearch.exe рядом с программой появляются файлы дампов ядра.
Вот исходное ядро, написанное на OpenCL (файл BinarySearch_Kernels.cl):
__kernel void binarySearch( __global uint4 * outputArray, __const __global uint * sortedArray, const unsigned int findMe, const unsigned int globalLowerBound, const unsigned int globalUpperBound, const unsigned int subdivSize) { unsigned int tid = get_global_id(0); /* lower bound and upper bound are computed from segment and total search space for this pass * The total search space is known from global lower and upper bounds for this pass. */ unsigned int lowerBound = globalLowerBound + subdivSize * tid; unsigned int upperBound = lowerBound + subdivSize - 1; /* Then we find the elements at the two ends of the search space for this thread */ unsigned int lowerBoundElement = sortedArray[lowerBound]; unsigned int upperBoundElement = sortedArray[upperBound]; /* If the element to be found does not lie between them, then nothing left to do in this thread */ if( (lowerBoundElement > findMe) || (upperBoundElement < findMe)) { return; } else { /* However, if the element does lie between the lower and upper bounds of this thread's searchspace * we need to narrow down the search further in this search space */ /* The search space for this thread is marked in the output as being the total search space for the next pass */ outputArray[0].x = lowerBound; outputArray[0].y = upperBound; outputArray[0].w = 1; } } /* Другие глобальные функции */ Вот сгенерированный дамп этого ядра на языке AMD IL (файл binarySearch_Juniper.il):
mdef(16383)_out(1)_in(2) mov r0, in0 mov r1, in1 div_zeroop(infinity) r0.x___, r0.x, r1.x mov out0, r0 mend il_cs_2_0 dcl_cb cb0[15] ; Constant buffer that holds ABI data dcl_literal l0, 0x00000004, 0x00000001, 0x00000002, 0x00000003 dcl_literal l1, 0x00FFFFFF, 0xFFFFFFFF, 0xFFFFFFFE, 0xFFFFFFFD dcl_literal l2, 0x0000FFFF, 0xFFFFFFFE, 0x000000FF, 0xFFFFFFFC dcl_literal l3, 0x00000018, 0x00000010, 0x00000008, 0xFFFFFFFF dcl_literal l4, 0xFFFFFF00, 0xFFFF0000, 0xFF00FFFF, 0xFFFF00FF dcl_literal l5, 0x00000000, 0x00000004, 0x00000008, 0x0000000C dcl_literal l6, 0x00000020, 0x00000020, 0x00000020, 0x00000020 dcl_literal l7, 0x00000018, 0x0000001F, 0x00000010, 0x0000001F dcl_literal l8, 0x80000000, 0x80000000, 0x80000000, 0x80000000 call 1024;$ endmain func 1024 ; __OpenCL_binarySearch_kernel mov r1013, cb0[8].x mov r1019, l1.0000 dcl_max_thread_per_group 256 dcl_raw_uav_id(11) dcl_arena_uav_id(8) mov r0.__z_, vThreadGrpIdFlat0.x mov r1022.xyz0, vTidInGrp0.xyz mov r1023.xyz0, vThreadGrpId0.xyz imad r1021.xyz0, r1023.xyzz, cb0[1].xyzz, r1022.xyzz iadd r1021.xyz0, r1021.xyz0, cb0[6].xyz0 iadd r1023.xyz0, r1023.xyz0, cb0[7].xyz0 mov r1023.___w, r0.z ishl r1023.___w, r1023.w, l0.z mov r1018.x___, l0.0000 udiv r1024.xyz_, r1021.xyzz, cb0[10].xyzz imad r1025.xyz0, r1023.xyzz, cb0[10].xyzz, r1022.xyzz dcl_literal l13, 0x00000001, 0x00000001, 0x00000001, 0x00000001; f32:i32 1 dcl_literal l11, 0x00000002, 0x00000002, 0x00000002, 0x00000002; f32:i32 2 dcl_literal l12, 0xFFFFFFFF, 0xFFFFFFFF, 0xFFFFFFFF, 0xFFFFFFFF; f32:i32 4294967295 dcl_cb cb1[6] ; Kernel arg setup: outputArray mov r1.x, cb1[0].x ; Kernel arg setup: sortedArray mov r1.y, cb1[1].x ; Kernel arg setup: findMe mov r1.z, cb1[2].x ; Kernel arg setup: globalLowerBound mov r1.w, cb1[3].x ; Kernel arg setup: globalUpperBound ; Kernel arg setup: subdivSize mov r2.y, cb1[5].x call 1029 ; binarySearch ret endfunc ; __OpenCL_binarySearch_kernel ;ARGSTART:__OpenCL_binarySearch_kernel ;version:2:0:88 ;device:juniper ;uniqueid:1024 ;memory:hwprivate:0 ;memory:hwregion:0 ;memory:hwlocal:0 ;pointer:outputArray:i32:1:1:0:uav:11:16:RW ;pointer:sortedArray:i32:1:1:16:uav:11:4:RO ;value:findMe:i32:1:1:32 ;value:globalLowerBound:i32:1:1:48 ;value:globalUpperBound:i32:1:1:64 ;value:subdivSize:i32:1:1:80 ;function:1:1029 ;uavid:11 ;privateid:1 ;ARGEND:__OpenCL_binarySearch_kernel func 1029 ; binarySearch ; @__OpenCL_binarySearch_kernel ; BB#0: ; %entry mov r65.x___, r2.y mov r65.__z_, r1.z mov r65.___w, r1.y mov r66, r1021.xyz0 mov r66.x___, r66.x000 imul r66.x___, r66.x, r65.x iadd r65._y__, r66.x, r1.w mov r66.x___, l11 ishl r66._y__, r65.y, r66.x iadd r66._y__, r65.w, r66.y mov r1010.x___, r66.y uav_raw_load_id(11)_cached r1011.x___, r1010.x mov r66._y__, r1011.x uge r66._y__, r65.z, r66.y if_logicalnz r66.y iadd r65.x___, r65.x, r65.y mov r66._y__, l12 iadd r65.x___, r65.x, r66.y ishl r66.x___, r65.x, r66.x iadd r65.___w, r65.w, r66.x mov r1010.x___, r65.w uav_raw_load_id(11)_cached r1011.x___, r1010.x mov r65.___w, r1011.x ult r65.__z_, r65.w, r65.z if_logicalnz r65.z else mov r1010.x___, r1.x uav_raw_load_id(11)_cached r1011, r1010 mov r66, r1011 iadd r66, r66.0yzw, r65.y000 iadd r66, r66.x0zw, r65.0x00 mov r65.x___, l13 iadd r66, r66.xyz0, r65.000x mov r1011, r66 mov r1010.x___, r1.x uav_raw_store_id(11) mem0, r1010.x, r1011 endif else endif ret endfunc ; binarySearch ;ARGSTART:binarySearch ;uniqueid:1029 ;ARGEND:binarySearch end Вот сгенерированный дизассемблированный ISA файл (файл binarySearch_Juniper.isa):
ShaderType = IL_SHADER_COMPUTE TargetChip = c ; ------------- SC_SRCSHADER Dump ------------------ SC_SHADERSTATE: u32NumIntVSConst = 0 SC_SHADERSTATE: u32NumIntPSConst = 0 SC_SHADERSTATE: u32NumIntGSConst = 0 SC_SHADERSTATE: u32NumBoolVSConst = 0 SC_SHADERSTATE: u32NumBoolPSConst = 0 SC_SHADERSTATE: u32NumBoolGSConst = 0 SC_SHADERSTATE: u32NumFloatVSConst = 0 SC_SHADERSTATE: u32NumFloatPSConst = 0 SC_SHADERSTATE: u32NumFloatGSConst = 0 fConstantsAvailable = 1237488 iConstantsAvailable = 1237456 bConstantsAvailable = 1237520 u32SCOptions[0] = 0x01A00000 SCOption_IGNORE_SAMPLE_L_BUG SCOption_FLOAT_DO_NOT_DIST SCOption_FLOAT_DO_NOT_REASSOC u32SCOptions[1] = 0x00202000 SCOption_R600_ERROR_ON_DOUBLE_MEMEXP SCOption_SET_VPM_FOR_SCATTER u32SCOptions[2] = 0x00000041 SCOption_R800_UAV_NONARRAY_FIXUP SCOption_R800_UAV_NONUAV_SYNC_WORKAROUND_BUG216513_1 ; -------- Disassembly -------------------- 00 ALU: ADDR(32) CNT(12) KCACHE0(CB0:0-15) KCACHE1(CB1:0-15) 0 x: LSHR R1.x, KC1[0].x, 2 t: MULLO_INT ____, R1.x, KC0[1].x 1 y: ADD_INT ____, R0.x, PS0 2 w: ADD_INT ____, PV1.y, KC0[6].x 3 t: MULLO_INT ____, PV2.w, KC1[5].x 4 y: ADD_INT R1.y, KC1[3].x, PS3 5 x: LSHL ____, PV4.y, 2 6 w: ADD_INT ____, KC1[1].x, PV5.x 7 y: LSHR R0.y, PV6.w, 2 01 TEX: ADDR(64) CNT(1) 8 VFETCH R0.x___, R0.y, fc153 MEGA(4) FETCH_TYPE(NO_INDEX_OFFSET) 02 ALU_PUSH_BEFORE: ADDR(44) CNT(2) KCACHE0(CB1:0-15) 9 z: SETGE_UINT R0.z, KC0[2].x, R0.x 10 x: PREDNE_INT ____, R0.z, 0.0f UPDATE_EXEC_MASK UPDATE_PRED 03 JUMP POP_CNT(1) ADDR(13) 04 ALU: ADDR(46) CNT(7) KCACHE0(CB1:0-15) 11 w: ADD_INT ____, KC0[5].x, R1.y 12 z: ADD_INT R1.z, -1, PV11.w 13 x: LSHL ____, PV12.z, 2 14 z: ADD_INT ____, KC0[1].x, PV13.x 15 y: LSHR R0.y, PV14.z, 2 05 TEX: ADDR(66) CNT(1) 16 VFETCH R0.x___, R0.y, fc153 MEGA(4) FETCH_TYPE(NO_INDEX_OFFSET) 06 ALU_PUSH_BEFORE: ADDR(53) CNT(2) KCACHE0(CB1:0-15) 17 w: SETGT_UINT R0.w, KC0[2].x, R0.x 18 x: PREDE_INT ____, R0.w, 0.0f UPDATE_EXEC_MASK UPDATE_PRED 07 JUMP POP_CNT(2) ADDR(13) 08 ALU: ADDR(55) CNT(2) KCACHE0(CB1:0-15) 19 z: LSHR R0.z, KC0[0].x, 4 09 TEX: ADDR(68) CNT(1) 20 VFETCH R0, R0.z, fc175 FORMAT(32_32_32_32_FLOAT) MEGA(16) FETCH_TYPE(NO_INDEX_OFFSET) 10 ALU: ADDR(57) CNT(4) 21 x: MOV R0.x, R1.y y: MOV R0.y, R1.z w: MOV R0.w, (0x00000001, 1.401298464e-45f).x 11 MEM_RAT_CACHELESS_STORE_RAW: RAT(11)[R1], R0, ARRAY_SIZE(4) MARK VPM 12 POP (2) ADDR(13) 13 NOP NO_BARRIER END_OF_PROGRAM ; ----------------- CS Data ------------------------ ; Input Semantic Mappings ; No input mappings GprPoolSize = 0 CodeLen = 560;Bytes PGM_END_CF = 0; words(64 bit) PGM_END_ALU = 0; words(64 bit) PGM_END_FETCH = 0; words(64 bit) MaxScratchRegsNeeded = 0 ;AluPacking = 0.0 ;AluClauses = 0 ;PowerThrottleRate = 0.0 ; texResourceUsage[0] = 0x00000000 ; texResourceUsage[1] = 0x00000000 ; texResourceUsage[2] = 0x00000000 ; texResourceUsage[3] = 0x00000000 ; texResourceUsage[4] = 0x00000000 ; texResourceUsage[5] = 0x00000000 ; texResourceUsage[6] = 0x00000000 ; texResourceUsage[7] = 0x00000000 ; fetch4ResourceUsage[0] = 0x00000000 ; fetch4ResourceUsage[1] = 0x00000000 ; fetch4ResourceUsage[2] = 0x00000000 ; fetch4ResourceUsage[3] = 0x00000000 ; fetch4ResourceUsage[4] = 0x00000000 ; fetch4ResourceUsage[5] = 0x00000000 ; fetch4ResourceUsage[6] = 0x00000000 ; fetch4ResourceUsage[7] = 0x00000000 ; texSamplerUsage = 0x00000000 ; constBufUsage = 0x00000000 ResourcesAffectAlphaOutput[0] = 0x00000000 ResourcesAffectAlphaOutput[1] = 0x00000000 ResourcesAffectAlphaOutput[2] = 0x00000000 ResourcesAffectAlphaOutput[3] = 0x00000000 ResourcesAffectAlphaOutput[4] = 0x00000000 ResourcesAffectAlphaOutput[5] = 0x00000000 ResourcesAffectAlphaOutput[6] = 0x00000000 ResourcesAffectAlphaOutput[7] = 0x00000000 ;SQ_PGM_RESOURCES = 0x30000102 SQ_PGM_RESOURCES:NUM_GPRS = 2 SQ_PGM_RESOURCES:STACK_SIZE = 1 SQ_PGM_RESOURCES:PRIME_CACHE_ENABLE = 1 ;SQ_PGM_RESOURCES_2 = 0x000000C0 SQ_LDS_ALLOC:SIZE = 0x00000000 ; RatOpIsUsed = 0x800 ; NumThreadPerGroupFlattened = 256 ; SetBufferForNumGroup = true Не знаю, как вам, а мне было бы неприятно, если бы мой супер алгоритм для GPU можно было бы так легко выцепить из программы и проанализировать. Особенно если в этом алгоритме заключалась бы вся суть программы (smile).
Анализ ситуации
Такое поведение характерно только для компилятора OpenCL от AMD и только при запуске приложения на AMD GPU. Если в системе установлен компилятор OpenCL от Nvidia, то, естественно, никакие файлы не генерируются на диске.
Как вы понимаете, данная возможность была оставлена для анализа разработчиками своего кода на OpenCL. Ведь полученные файлы можно запихнуть в профайлер и увидеть, какие операции будут бутылочным горлышком в программе. Однако если не знать про данную глобальную переменную, можно лишиться своей интеллектуальной собственности довольно быстро.
Если внимательно посмотреть на сгенерированный файл binarySearch_Juniper.il, то волосы могут встать дыбом от этого кода: исходное ядро для OpenCL можно переписать на языке AMD IL в 20 строчек, но никак не в 100! Это наталкивает на мысли, что приложения, написанные на OpenCL для AMD GPU на данный момент не будут такими же быстрыми, как и приложения, использующие технологию AMD IL для взаимодействия с GPU.
Как разобраться в написанном в файле binarySearch_Juniper.il, рассказано здесь.
Как можно использовать файл binarySearch_Juniper.il в своей программе, рассказано здесь.
Автор: BrainHacker
