50 lines
39 KiB
Plaintext
50 lines
39 KiB
Plaintext
|
ELF > p[ @ ˜’ @ 8 @ $ " @ @ @ h h ¨ ¨ ¨ pK pK pK p[ p[ ° ° ] } } 8_ 8� 8� H Ø @] @} @} Ð Ð Råtd ] } } à Påtd H H H l l Qåtd Ä Ä Ä /lib64/ld-linux-x86-64.so.2 GNU " " 8 " G " _ y “ Ë Ð × " e l s … ’ ™ ¥ ¯ ¹ ¿ Ô ï ÷ ü g u „ ´ Á Ø î ‰ ‘ œ Àl m
|
|||
|
|
P 5 P x € Ž € ¤ € r' - t)’ D ‰)
|
|||
|
|
P qø— ^ Ó¯k m ui ‚ P&y œ ui ‚ ”‘– ® ´‘– ¹ % ! % ˜ C(EÕLyIk¶ __libc_start_main __gmon_start__ __register_frame_info __cxa_finalize __deregister_frame_info __hipPopCallConfiguration hipLaunchKernel _ZSt4cout _ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc exit strlen _ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE9_M_createERmm memcpy _ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE7compareEPKc strtol _ZdlPv hipGetDeviceCount hipSetDevice printf malloc puts hipMalloc hipMemcpy ceilf hipDeviceSynchronize __hipPushCallConfiguration hipFree free _ZSt19__throw_logic_errorPKc _ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l __assert_fail _Unwind_Resume _ZNSt8ios_base4InitC1Ev _ZNSt8ios_base4InitD1Ev __cxa_atexit __hipRegisterFatBinary __hipRegisterFunction __hipUnregisterFatBinary __gxx_personality_v0 libamdhip64.so.6 hip_4.2 libstdc++.so.6 GLIBCXX_3.4 GLIBCXX_3.4.9 GLIBCXX_3.4.21 CXXABI_1.3 libm.so.6 GLIBC_2.2.5 libgcc_s.so.1 GCC_3.0 libc.so.6 GLIBC_2.14 GLIBC_2.34 /opt/rocm-6.0.2/lib/llvm/bin/../../../lib ( 0 ‘ % `� h� p� x� €� ˆ� �� ˜�
|
|||
|
|
� ¨� °�
|
|||
|
|
¸� À� È� Ð� Ø� à� è� ð� ø� � � � � � (� 0� 8� @� H� P� & X� ! `� " h� # p� $ x� ' ÿÿÿÿÿÿÿÿÿÿÿÿÿÿ ñh㈵øä>@\ à\ Finished allocating vectors on the GPU -b (hipMalloc(&d_c, bytes))==hipSuccess basic_string::_M_construct null not valid Finished copying vectors to the GPU -d/--dev <value> Set the device ID [Default: 0]
|
|||
|
|
-b/--blockSize <value> Set the block size
|
|||
|
|
--device (hipFree(d_a))==hipSuccess Optional:
|
|||
|
|
(hipMemcpy(d_b, h_b, bytes, hipMemcpyHostToDevice))==hipSuccess (hipMemcpy( h_c, d_c, bytes, hipMemcpyDeviceToHost))==hipSuccess (hipSetDevice(devId))==hipSuccess (hipFree(d_c))==hipSuccess --vec -d sw thinks it moved %f KB per wave
|
|||
|
|
Error at position i %d, Expected: %f, Found: %f
|
|||
|
|
-n/--numThreads <value> Set the num of threads
|
|||
|
|
-h _Z7vecCopyPdS_S_ii Finished copying the output vector from the GPU to the CPU --multikernel Releasing GPU memory --help vcopy testing on GCD %d
|
|||
|
|
Usage: vcopy [OPTIONS]
|
|||
|
|
-h/--help Display this help message
|
|||
|
|
Launching the kernel on the GPU n > 0 (hipMalloc(&d_b, bytes))==hipSuccess Finished allocating vectors on the CPU Required:
|
|||
|
|
--blockSize int main(int, char **) blockSize > 0 ./sample/vcopy.cpp --iter (hipMalloc(&d_a, bytes))==hipSuccess (hipFree(d_b))==hipSuccess -i/--iter <value> Set the num of iterations [Default: 1]
|
|||
|
|
-n Total threads: %d, Grid Size: %d block Size:%d, Wavefronts:%d:
|
|||
|
|
(hipMemcpy(d_a, h_a, bytes, hipMemcpyHostToDevice))==hipSuccess (hipGetDeviceCount(&numGpuDevices))==hipSuccess Finished executing kernel _Z9vecCopy_2PdS_S_ii Releasing CPU memory -i ÿÿ œ œ›Š ·í __CLANG_OFFLOAD_BUNDLE__ host-x86_64-unknown-linux-- ð hipv4-amdgcn-amd-amdhsa--gfx908
|
|||
|
|
„
|
|||
|
|
m
|
|||
|
|
AMDGPU ƒ®amdhsa.kernels’Þ «.agpr_count ¥.argsÜ „®.address_space¦global§.offset ¥.size«.value_kindglobal_buffer„®.address_space¦global§.offset¥.size«.value_kindglobal_buffer„®.address_space¦global§.offset¥.size«.value_kindglobal_bufferƒ§.offset¥.size«.value_kind¨by_valueƒ§.offset ¥.size«.value_kind¨by_valueƒ§.offset ¥.size«.value_kind´hidden_block_count_xƒ§.offset$¥.size«.value_kind´hidden_block_count_yƒ§.offset(¥.size«.value_kind´hidden_block_count_zƒ§.offset,¥.size«.value_kind³hidden_group_size_xƒ§.offset.¥.size«.value_kind³hidden_group_size_yƒ§.offset0¥.size«.value_kind³hidden_group_size_zƒ§.offset2¥.size«.value_kind²hidden_remainder_xƒ§.offset4¥.size«.value_kind²hidden_remainder_yƒ§.offset6¥.size«.value_kind²hidden_remainder_zƒ§.offsetH¥.size«.value_kind¶hidden_global_offset_xƒ§.offsetP¥.size«.value_kind¶hidden_global_offset_yƒ§.offsetX¥.size«.value_kind¶hidden_global_offset_zƒ§.offset`¥.size«.value_kind°hidden_grid_dims¹.group_segment_fixed_size ¶.kernarg_segment_alignµ.kernarg_segment_sizeÍ ©.language¨OpenCL C±.language_version’ ¸.max_flat_workgroup_sizeÍ ¥.name²_Z7vecCopyPdS_S_ii».private_segment_fixed_size «.sgpr_count ±.sgpr_spill_count §.symbolµ_Z7vecCopyPdS_S_ii.kd¸.uniform_work_group_size³.uses_dynamic_stack«.vgpr_count±.vgpr_spill_count ¯.wavefront_size@Þ «.agpr_count ¥.argsÜ „®.address_space¦global§.offset ¥.size«.value_kindglobal_buffer„®.address_space¦global§.offset¥.size«.value_kindglobal_buffer„®.address_space¦global§.offset¥.size«.value_kindglobal_bufferƒ§.offset¥.size«.value_kind¨by_valueƒ§.offset ¥.size«.value_kind¨by_valueƒ§.offset ¥.size«.value_kind´hidden_block_count_xƒ§.offset$¥.size«.value_kind´hidden_block_count_yƒ§.offset(¥.size«.value_kind´hidden_block_count_zƒ§.offset,¥.size«.value_kind³hidden_group_size_xƒ§.offset.¥.size«.value_kind³hidden_group_size_yƒ§.offset0¥.size«.value_kind³hidden_group_size_zƒ§.offset2¥.size«.value_kind²hidden_remainder_xƒ§.offset4¥.size«.value_kind²hidden_remainder_yƒ§.offset6¥.size«.value_kind²hidden_remainder_zƒ§.offsetH¥.size«.value_kind¶hidden_global_offset_xƒ§.offsetP¥.size«.value_kind¶hidden_global_offset_yƒ§.offsetX¥.size«.value_kind¶hidden_global_offset_zƒ§.offset`¥.size«.value_kind°hidden_grid_dims¹.group_segment_fixed_size ¶.kernarg_segment_alignµ.kernarg_segment_sizeÍ ©.language¨OpenCL C±.language_version’ ¸.max_flat_workgroup_sizeÍ ¥.name´_Z9vecCopy_2PdS_S_ii».private_segment_fixed_size «.sgpr_count ±.sgpr_spill_count §.symbol·_Z9vecCopy_2PdS_S_ii.kd¸.uniform_work_group_size³.uses_dynamic_stack«.vgpr_count±.vgpr_spill_count ¯.wavefront_size@amdhsa.target¹amdgcn-amd-amdhsa--gfx908®amdhsa.version’ € À @ * ! € ? @ W ! @ ƒ ! A ¯ ! B
|
|||
|
|
|
|||
|
|
X1s4¤ tŠ pŒˆh0‡|[äPéã‘Ë—QN _Z7vecCopyPdS_S_ii _Z7vecCopyPdS_S_ii.kd _Z9vecCopy_2PdS_S_ii _Z9vecCopy_2PdS_S_ii.kd _ZN17__HIP_CoordinatesI14__HIP_BlockIdxE1xE _ZN17__HIP_CoordinatesI14__HIP_BlockDimE1xE _ZN17__HIP_CoordinatesI15__HIP_ThreadIdxE1xE @ A ¯ Œ A ¯ Œ À, B À ÀŒ¿ ÿ †ÿÿ ’ h ˆ}j €¾ ˆ¿ À ‚ À Ÿ " �Òƒ ÀŒ¿~ 28 €TÜ ~ 28pŒ¿ €tÜ �¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ €¿ À, B À ÀŒ¿ ÿ †ÿÿ ’ h ˆ}j €¾ ˆ¿ À ‚ À Ÿ " �Òƒ ÀŒ¿~ 28 €TÜ ~ 28pŒ¿ €tÜ �¿ ˆ Ð
|
|||
|
|
|
|||
|
|
Ü õþÿo H
|
|||
|
|
ˆ
|
|||
|
|
Linker: AMD LLD 17.0.0 AMD clang version 17.0.0 (https://github.com/RadeonOpenCompute/llvm-project roc-6.0.2 24012 af27734ed982b52a9f1be0f035ac91726fc697e4) Ü €1 € À @ * ! € ? @ W ! @ ƒ ! A ¯ ! B .note .dynsym .gnu.hash .hash .dynstr .rodata .text .dynamic .comment .symtab .shstrtab .strtab _Z7vecCopyPdS_S_ii _Z7vecCopyPdS_S_ii.kd _Z9vecCopy_2PdS_S_ii _Z9vecCopy_2PdS_S_ii.kd _ZN17__HIP_CoordinatesI14__HIP_BlockIdxE1xE _ZN17__HIP_CoordinatesI14__HIP_BlockDimE1xE _ZN17__HIP_CoordinatesI15__HIP_ThreadIdxE1xE _DYNAMIC „
|
|||
|
|
ˆ ˆ À öÿÿo H
|
|||
|
|
H
|
|||
|
|
@ ˆ
|
|||
|
|
ˆ
|
|||
|
|
H Ð
|
|||
|
|
Ð
|
|||
|
|
Ü ' À À ƒ @ / € 5 €1 € p > 0 ð ž G � Ø O h a Y É å FPIH ;l h ˆ ˜ � ´ á Ô 8 ô Ø x < Ø ¸! h H" ˆ h" P ˆ" zR x � Ø & D 4 ð H á D A†C
|
|||
|
|
h U A†C
|
|||
|
|
P ˆ < – D€�ˆD�L ¬ ¸ – D€�ˆD�L Ð 4 ƒ A ä ! AZ ü H Œ AD0ƒ�A ¸ A] 4 à zPLR xm � ˆ $ ( ¤ ABB B(B0A8G ƒŒ�Ž�†’. Y. D. a. Q. a. m. T. a. k. Q. a. n. Q. a. n. Q. U. h. . f. O. o. K. K. K. K. J. U. U. U. «. j. ]. f. F. L. \. t. E. \. ™. {.Q¨G°I . E. X. ™. {.Q¨G°I J. b. J. ^. `. O. R. R. M. J. H. H. Q8A0B(B BBAA . J. T. T. T. T. T. T. T. O. J. B. O. O. O. O. O. O. O. O. J. Y. Y. Y. Y. Y. Y. Y. Y. Y. Y. Y. Y. f. E. ó ú1íI‰Ñ^H‰âHƒäðPTE1À1ÉH�=� ÿ{# ôÌÌÌÌÌÌÌÌÌÌó úÃUH‰å¶Ð4 ¶ÀH…Àu.ÆÁ4 H‹Z# H…Àt H�Î4 H‰ÆH�œìÿÿH‰Çèü ë�]ÃUH‰å¶ì4 ¶ÀH…Àu?ÆÝ4 H‹ # H…ÀtH‹*3 H‰ÇèÒ H‹ # H…ÀtH�GìÿÿH‰ÇèÇ ë�]ÃÌÌHƒìxH‰|$HH‰t$@H‰T$8‰L$D‰$H�D$HH‰D$PH�D$@H‰D$XH�D$8H‰D$`H�D$H‰D$hH‰àH‰D$pH�|$(H�t$H�T$H�L$èl H‹t$(‹T$0H‹L$D‹D$ L�L$P¿H ÿt$ÿt$èR H�Ĉ Ãf. „ HƒìxH‰|$HH‰t$@H‰T$8‰L$D‰$H�D$HH‰D$PH�D$@H‰D$XH�D$8H‰D$`H�D$H‰D$hH‰àH‰D$pH�|$(H�t$H�T$H�L$èÌ
|
|||
|
|
H‹t$(‹T$0H‹L$D‹D$ L�L$P¿P ÿt$ÿt$è²
|
|||
|
|
H�Ĉ Ãf. „ P¿ ‘ ¾u è
|
|||
|
|
¿ ‘ ¾7 è‘
|
|||
|
|
¿ ‘ ¾ è‚
|
|||
|
|
¿ ‘ ¾1 ès
|
|||
|
|
¿ ‘ ¾… èd
|
|||
|
|
¿ ‘ ¾õ èU
|
|||
|
|
¿ ‘ ¾Í èF
|
|||
|
|
¿ ‘ ¾� è7
|
|||
|
|
¿ è=
|
|||
|
|
ffff. „ UAWAVAUATSH�ìè H‰t$…ÿŽ£ H�l$p‰øH‰D$ H÷ØH‰„$˜ ¸ H‰D$H1ÀH‰„$à A½ L�d$`H‰D$H‰„$� 1ÀH‰D$@ë!f D H‹„$˜ LèHÿÀIÿÅHƒø„† H‹D$N‹tèøH‰l$`M…ö„o L‰÷è• H‰ÃH‰D$(I‰ïHƒør!L‰çH�t$(1Òè… H‰ÅH‰D$`H‹L$(H‰L$pH…Ût,HƒûuA¶ˆM ë ffffff. „ H‰ïL‰öH‰ÚèR H‹D$(H‰D$hH‹L$`Æ ¾B L‰çèB …Àt¾ L‰çè1 …Àu=L;l$ s6H‹D$J‹<è1öº
|
|||
|
|
è! H‰„$� L‰ýH‹|$`H9ï…v éóþÿÿ ¾N L‰çèã …ÀL‰ýt¾ L‰çèÏ …Àu;L;l$ s4H‹D$J‹<è1öº
|
|||
|
|
è¿ H‰D$H‹|$`H9ï… é—þÿÿ € ¾a L‰çèƒ …Àt¾T L‰çèr …Àu>L;l$ s7H‹D$J‹<è1öº
|
|||
|
|
èb H‰D$@H‹|$`H9ï…½ é:þÿÿf. „ ¾† L‰çè# …Àt¾ L‰çè …Àu>L;l$ s7H‹D$J‹<è1öº
|
|||
|
|
è H‰D$HH‹|$`H9ïuaéÞýÿÿfffff. „ ¾2 L‰çèÃ
|