hcc icon indicating copy to clipboard operation
hcc copied to clipboard

llvm-mc generated hsaco , can't load by hsa_executable_load_code_object.

Open JiniusDnn opened this issue 8 years ago • 1 comments

thanks for everyone who read this , I am new to hcc.

I try to compile asm into hsaco and load it to hsa.

the asm was the example hello world from hcc project:

.hsa_code_object_version 2,0
.hsa_code_object_isa 8, 0, 3, "AMD", "AMDGPU"

.text
.p2align 8
.amdgpu_hsa_kernel hello_world

hello_world:

   .amd_kernel_code_t
      enable_sgpr_kernarg_segment_ptr = 1
      is_ptr64 = 1
      compute_pgm_rsrc1_vgprs = 8
      compute_pgm_rsrc1_sgprs = 8
      compute_pgm_rsrc2_user_sgpr = 2
      kernarg_segment_byte_size = 8
      wavefront_sgpr_count = 2
      workitem_vgpr_count = 3
  .end_amd_kernel_code_t

  s_load_dwordx2 s[0:1], s[0:1] 0x0
  v_mov_b32 v0, 3.14159
  s_waitcnt lgkmcnt(0)
  v_mov_b32 v1, s0
  v_mov_b32 v2, s1
  flat_store_dword v[1:2], v0
  s_waitcnt 0
  s_endpgm

I can compile it by llvm-mc

$(hcc)/bin/llvm-mc -arch=amdgcn -mcpu=gfx803 -filetype=obj -o $(S_FILE).hsaco $(S_FILE).s

but when I load it from hsa api, it failed at:

hsa_status = hsa_executable_load_code_object(hsaExecutable, device, code_object, NULL);

thanks again if you can point out what's wrong here.

JiniusDnn avatar Oct 16 '17 22:10 JiniusDnn

solved by ld.lld .

gonuco avatar Oct 16 '17 22:10 gonuco