AMD ISA 코드 읽는 데 문제가 있습니다.

Question

here 문서를 사용하여 AMD Southern Island GPU의 어셈블리 언어를 읽는 데 문제가 있습니다.

1 shader main 
2  asic(SI_ASIC) 
3  type(CS) 
4  
5  s_buffer_load_dword s0, s[8:11], 0x00     // what is the purpose of s[8:11] ? 
6  s_waitcnt  lgkmcnt(0) 
7  v_add_i32  v1, vcc, s0, v0       // I guess v0 initially contains the local IDs ? 
8  v_add_i32  v0, vcc, s0, v0       // wouldn't a v_mov v0, v1 performs better ? 
9  buffer_load_ubyte v2, v1, s[4:7], 0 offen    // s[4:7] ? 
10  buffer_load_ubyte v0, v0, s[4:7], 0 offen offset:256 
11  s_waitcnt  vmcnt(0) 
12  v_max_u32  v0, v2, v0 
13  buffer_store_byte v0, v1, s[4:7], 0 offen glc 
14  s_endpgm 
15 end 

• 내가 이해할 수없는 것은 내가 가지고 수있는 방법입니다 :

  1 __attribute__((reqd_work_group_size(256, 1, 1))) 
  2 void kernel foo(global uchar* data) { 
  3  const uint block_size = get_local_size(0); 
  4  const uint lid = get_local_id(0); 
  5   
  6  data[lid] = max(data[lid], data[lid + block_size]); 
  7 } 
  

  그리고 (DIS) 어셈블리 AMD의 컴파일러에 의해 생성 : 여기

  샘플 인 OpenCL 코드 예를 들어, ISA 설명서 (12.6 벡터 메모리 버퍼 명령어 참조) 만있는 * buffer_load_ubyte * 명령어를 작성했습니다.
• 마이크로 코드 정보를 어떻게 읽어야합니까? (예 : * v_add_i32 * 명령어의 경우 161 페이지)
• 프로세서 아키텍처를 명시하는 표준 방법이 있습니까?

감사합니다.

시는 : 어셈블리 코드는 세계 여기이 게시

Answer 1

논평에서 약간의 보너스 질문이 있습니다. 이것은 제가 수강하는 수업의 일환으로 주어진 SI 어셈블리입니다. 이 코드가 Multi2sim 위에 실행되면 ABI에서 어떤 일이 벌어지고 있는지 이해할 수 있습니다.

.global vector_add 
.args 
    i32* src1 0 uav10 RO 
    i32* src2 16 uav11 RO 
    i32* dst 32 uav12 RW 
.metadata 
    uavprivate = 0 
    hwregion = 0 
    hwlocal = 0 
    FloatMode = 192 
    IeeeMode = 0 
    # Loads UAV table in s2...s3 
    userElements[0] = PTR_UAV_TABLE, 0, s[2:3] 
    # Loads constant buffer 0 descriptor in s4...s7 
    userElements[1] = IMM_CONST_BUFFER, 0, s[4:7] 
    # Loads constant buffer 1 descriptor in s8...s11 
    userElements[2] = IMM_CONST_BUFFER, 1, s[8:11] 
    # Forces wg_id[0] (work-group ID in dimension 0) to be available in s12 
    COMPUTE_PGM_RSRC2:USER_SGPR = 12 
    COMPUTE_PGM_RSRC2:TGID_X_EN = 1 
.text 
    # Load lsize[0] into s0 
    s_buffer_load_dword s0, s[4:7], 0x04 
    # Load src1, src2, and dst base addresses (arguments) from CB1 
    s_buffer_load_dword s4, s[8:11], 0x00 
    s_buffer_load_dword s5, s[8:11], 0x04 
    s_buffer_load_dword s6, s[8:11], 0x08 
    # Load UAVs from UAV table 
    s_load_dwordx4 s[20:23], s[2:3], 0x50 
    s_load_dwordx4 s[24:27], s[2:3], 0x58 
    s_load_dwordx4 s[28:31], s[2:3], 0x60 
    # Waits for memory operations to complete 
    s_waitcnt lgkmcnt(0) 
    # v1 <= lsize[0] 
    v_mov_b32 v1, s0 
    # v1 <= lsize[0] * wg_id[0] 
    v_mul_i32_i24 v1, s12, v1 
    # v2 <= lsize[0] * wg_id[0] + lid[0] = gid[0] 
    v_add_i32 v2, vcc, v0, v1 
    # v3 <= gid[0] * 4 
    v_lshlrev_b32 v3, 2, v2 
    # Calcaulte effective addresses 
    v_add_i32 v10, vcc, s4, v3 
    v_add_i32 v11, vcc, s5, v3 
    v_add_i32 v12, vcc, s6, v3 
    # Load src1[id] and src2[id] 
    tbuffer_load_format_x v20, v10, s[20:23], 0 offen format:[BUF_DATA_FORMAT_32, BUF_NUM_FORMAT_FLOAT] 
    tbuffer_load_format_x v21, v11, s[24:27], 0 offen format:[BUF_DATA_FORMAT_32,BUF_NUM_FORMAT_FLOAT] 
    # Waits for memory operations to complete 
    s_waitcnt vmcnt(0) 
    # Add source elements 
    v_add_i32 v22, vcc, v20, v21 
    # Store result in dst[id] 
    tbuffer_store_format_x v22, v12, s[28:31], 0 offen format:[BUF_DATA_FORMAT_32,BUF_NUM_FORMAT_FLOAT] 
    # End program 
    s_endpgm 

이것은 간단한 벡터 추가 커널이며 3 개의 인수를 취합니다. src1, src2 및 dst.

Answer 2

나는 약 1 년 늦었지만 어쩌면 다른 사람을 도울 것입니다. AMD의 GCN을 이해

모든이 내려 온다 : 커널에

Address = BASE + offset + lane 

는 다음 레지스터가 사전로드

• 들입니다 출시 [4 : 7] 기본 주소
• 입니다 s [8:11]은 매개 변수에 대한 포인터입니다.
• 오프셋 v0에는 레인 번호가 미리로드되어 있습니다 (0-63)

간단히 말하면, 우리는 (1) 주소 * 데이터를 얻으려고합니다 (2) * 데이터에서 값을 얻으십시오 (3) 오프셋에서 값을 얻으십시오 : 256 (4) 두 숫자를 더하십시오 (5) 결과가 * 데이터에 좋지 않습니다. 는 "v_mov의 V0, V1"대신 사용

shader main 
    asic(SI_ASIC) 
    type(CS) 

    s_buffer_load_dword s0, s[8:11], 0x00 // s[8:11] is the pointer to the params 
    s_waitcnt  lgkmcnt(0)   // wait for s0 to be filled 
    v_add_i32  v1, vcc, s0, v0  // s0=offset v0=lane We just need the base now. 
    v_add_i32  v0, vcc, s0, v0  // wouldn't a v_mov v0, v1 performs better 
    buffer_load_ubyte v2, v1, s[4:7], 0 offen //Get value at Base(s[4:7]) + v1(offset & lane) 
    buffer_load_ubyte v0, v0, s[4:7], 0 offen offset:256 // like above but address+256 
    s_waitcnt  vmcnt(0)      //wait for the memory transfer to complete 
    v_max_u32  v0, v2, v0      // do the MAX operation 
    buffer_store_byte v0, v1, s[4:7], 0 offen glc //save v0 using the base+v1(offset+lane) 
    s_endpgm          //stop kernel 
end 

"v_add_i32의 V0, VCC를, S0, V0"라인 10은 빠르게 v_mov 때문에 가서 모두 동일한 시간을 v_add하지 것이다. 그러나 이것이 CPU에 있다면 v_mov는 이전의 명령에 의존하고 한 번에 여러 개의 명령을 수행 할 수 없으므로 속도가 느려집니다. GPP는 한 번에 코어에서 여러 명령어를 수행 할 수 없으므로 두 방법 모두 같은 속도가됩니다.

내가 이해할 수없는 것은, 예를 들어, ISA 설명서 (12.6 벡터 메모리 버퍼 명령어 참조)가있는 buffer_load_ubyte 명령어를 작성한 것일 수 있습니다. 따라하기가 어렵습니다. 당신은 거의 예를 통해 배울 필요가 있습니다.

어떻게 마이크로 코드 정보를 읽어야합니까? (예 : v_add_i32 명령어의 경우 161 페이지) 어렵습니다. 당신은 거의 흔적과 오류가 거의 필요합니다. 매뉴얼은 실제로 몇 군데 정확하지 않고 포럼에서 AMD에보고했습니다. 나는 컴파일러 (asm4gcn)를 만들었고 나는 이것에 대해 고심해야만했다. 나는 또한 참고로 다른 민족 프로젝트를 사용했다.

프로세서 아키텍처를 설명하는 표준 방법이 있습니까? 나는 그들이 어떤면에서 모두 다른 것 같아요. 그러나 대부분은 캐시 섹션, 부동 소수점 유닛, 제어 흐름 유닛 및 레지스터를 가지고 있습니다.