Анализируем проприетарное приложение на OpenCL, написанное с использованием AMD APP SDK

    Рассмотрим следующую ситуацию: есть приложение, которое использует 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 в своей программе, рассказано здесь.
    Ads
    AdBlock has stolen the banner, but banners are not teeth — they will be back

    More

    Comments 17

      0
      Осталось понять, баг это все-таки или фича:)

      Я правильно поднял что для любого приложния достаточно воткнуть видеокарту от MAD и посмотреть как оно работает, т.е. зависит это от системы пользователя, а не от самой программы, и это на совести OpenCL для AMD, а не самих приложений? Есть способ для разработчика обезопасить себя от этого?
        0
        Для любого приложения, использующего OpenCl и для AMD, а не MAD, конечно же.
          +1
          Это фича компилятора от AMD, так как это есть в документации и даже описаны цели, зачем это было сделано.
          Да, берем любое приложение с OpenCL, ставим компилятор от AMD, устанавливаем переменную окружения и смотрим на результат.

          > зависит это от системы пользователя, а не от самой программы, и это на совести OpenCL для AMD, а не самих приложений?
          Да.

          > Есть способ для разработчика обезопасить себя от этого?
          1) Проверять значение переменной в начале работы, возможно менять это значение в ходе работы.
          2) Проверять наличие этих файлов на диске и удалять.
          3) Не использовать OpenCL, как крайняя мера. :)
        0
        Точно так же можно перехватить вызов CUDA dll и сдампить себе ядро, которое потом дизассемблировать. Да, конечно, это немного сложнее.
          0
          Так же можно сделать и для AMD dll. Но я написал, что есть лицензии, запрещающие дизассемблирование и reverse engineering. А здесь достаточно установить одну переменную окружения и вуаля!
          Я просто к тому, что против дизассемблирования и reverse engineering'a можно бороться и довольно успешно. А тут, если не знать о такой фиче компилятора, можно наступить на грабли, аккуратно разложенные создателями компилятора.
            +1
            Если лицензия запрещает дизассеблирование — то и подход, описанный в статье, тоже противоречит этой лицензии. То, что код получен штатными средствами компилятора никого не волнует.

            Что же на счёт защиты, то в большинстве случаев пытаются хотя бы заставить работать ядра на GPU на порядок быстрее, чем на обычном CPU. Если туда вставить всякие защиты от отладки — то производительность ядер упадёт и толку от использования GPU не будет никакого.
            Впрочем, сильно оптимизированный код читается и реверсится тоже тяжко. Помню, была одна замечательная статья про умножение разряженных матриц на ATI, результирующее ядро у них получилось вот такое, тут в исходниках то сложно сходу разобраться, как оно работает.
              0
              То, что генерирует компилятор, не является частью исходного кода приложения. Это, скажем так, временные файлы компилятора. Он их просто сохранил на диск. Разве это попадает под нарушение лицензии?

              Естественно, сам код, выполняющийся на GPU, обфусцировать не стоит. Скорость, конечно, упадет на порядок.
          0
          >>против дизассемблирования и reverse engineering'a можно бороться и довольно успешно

          как?
            0
            >>против дизассемблирования
            Шифрование участков программы, применение обфусцирующих средств (виртуальных машин, например).

            >>против reverse engineering'a
            Проверка целостности памяти кода программы (чтоб break нельзя было поставить), проверка времени выполнения участков кода, применение обфусцирующих средств.

            Если задаться целью узнать алгоритм программы, то это, конечно, всегда можно сделать. Вопрос только времени и ресурсов. Но против непрофессионалов указанные методы работают. А вот приведенный в статье метод может выполнить и непрофессионал.
              0
              >>>>>против дизассемблирования
              Шифрование участков программы, применение обфусцирующих средств (виртуальных машин, например).

              >>против reverse engineering'a
              Проверка целостности памяти кода программы (чтоб break нельзя было поставить), проверка времени выполнения участков кода, применение обфусцирующих средств.

              это ненадёжно. Поэтому гпу интересен именно как аппаратная ВМ.
              Вам извесны дизассемблеры шейдеров (для ати нвидиа)?

                0
                >> это ненадёжно
                Достаточно, чтобы усложнить жизнь взломщику. :)

                >> Вам извесны дизассемблеры шейдеров (для ати нвидиа)?
                Nvidia CUDA не оперирует шейдерами, как таковыми.
                Дизассемблер кода шейдера в OpenCL? Не, не слышал. Просто шейдеры написаны на своем языке, который не нужно дизассемблировать, чтобы понять его суть.
            0
            Пообщавшись некоторое время с профессионалами, которые в моём BarsWF до его open-source релиза добавляли свой функционал и интересовались некоторыми деталями имплементации (вроде зачем я тащу тригонометрию в ядро криптоанализа) я точно знаю, что бинарник — открытая книга для специалиста.

            С GPU кодом это просто иногда чуть проще.
              0
              Я и не спорю, что специалист может если не все, то очень многое.
              Просто хотел предупредить разработчиков об этой лазейке.
                +1
                Какими, к чёрту, профессионалами?)
                откройте любымым hex вювером ваш винарник и поищите: il_ps_2_0
                всё найдётся всё: pastebin.com/XNdxZHZ4

                  0
                  Там ковыряли намного глубже, в SSE2 коде и добавлением нового функционала который не ограничивается il кодом )
                    0
                    А что добавляли?
                    Или удаляли проверку на принадлежность процесору к семейства Intel?
                      0
                      Ограниченную поддержку соли.

              Only users with full accounts can post comments. Log in, please.