AMD APP SDKを使用して記述されたOpenCL上の独自アプリケーションを分析します

次の状況を考慮してください。計算にAMD GPUを使用するアプリケーションがあります。 原則として、最もリソースを消費する操作はGPUで実行されます。 したがって、アプリケーションが競合他社よりも高速に実行される場合、このプログラムに実装されているアルゴリズムを知りたい場合があります。 しかし、プログラムが独自仕様であり、リバースエンジニアリングと逆アセンブリを禁止するライセンスの下で配布されている場合はどうでしょうか。



ライセンスに違反しないように、AMD APP SDKの開発者が残した1つの小さなトリックを使用できます。 ただし、このトリックを機能させるには、アプリケーション開発者が指定されたSDKを使用することに加えて、もう1つの条件を満たしている必要があります。アプリケーションはGPUコンピューティングにOpenCLを使用する必要があります。



AMD Accelerated Parallel Processing OpenCLプログラミングガイド（v1.3f）のドキュメントを注意深く読んだ場合、 「4.2.1中間言語とGPUの逆アセンブリ」セクションで、GPU_DUMP_DEVICE_KERNELというすばらしい環境変数を見つけることができます。 次の3つの値を取ることができます。

• 1-ローカルディレクトリの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.yy: MOV R0.y, R1.zw: 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のスーパーアルゴリズムをプログラムから簡単に抽出して分析できるとしたら、私にとっては不愉快です。 特に、プログラムの本質（笑顔）がこのアルゴリズムにある場合。

状況分析

この動作は、AMD OpenCLコンパイラでのみ、AMD GPUでアプリケーションが起動された場合にのみ典型的です。 システムにNvidiaのOpenCLコンパイラがインストールされている場合、もちろん、ディスク上にファイルは生成されません。



ご存じのとおり、この機能はOpenCLコードの開発者による分析のために残されました。 結局、受信したファイルをプロファイラーにプッシュして、プログラムのボトルネックになる操作を確認できます。 ただし、このグローバル変数について知らない場合は、知的財産をすぐに失う可能性があります。



生成されたbinarySearch_Juniper.ilファイルを注意深く見ると、髪はこのコードの端に立つことができます：OpenCLの元のカーネルは、AMD ILで100行ではなく20行で書き直すことができます！ これは、現時点でAMD GPU用のOpenCLで書かれたアプリケーションは、AMD ILテクノロジーを使用してGPUと対話するアプリケーションほど高速ではないことを示唆しています。



ファイルに記述されたbinarySearch_Juniper.ilを理解する方法は、 ここで説明されています 。

ここで、プログラムでbinarySearch_Juniper.ilファイルを使用する方法について説明します 。