bash> amdclang -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx90a -O0 -g tmp.c bash> ./a.out ========================== BEGIN ompt_start_tool_result_t *ompt_start_tool(unsigned int, const char *) END ompt_start_tool_result_t *ompt_start_tool(unsigned int, const char *) ========================== ========================== BEGIN int ompt_initialize(ompt_function_lookup_t, int, ompt_data_t *) device_initialize : always device_finalize : always device_load : always device_unload : never target_data_op : always target : always target_map : never target_submit : always END int ompt_initialize(ompt_function_lookup_t, int, ompt_data_t *) ========================== ========================== BEGIN void target() #pragma omp target [ print_ompt_device_initialize] device_num = 0 | type = AMD gfx90a | device = 0x103ddb0 | lookup = 0x14ea97eb6b90 [ print_ompt_device_load] device_num = 0 | filename = (null) | offset_in_line = 0 | vma_in_file = (nil) | bytes = 38304 | host_addr = 0x2011e0 | device_addr = (nil) | module_id = 0 [ print_ompt_device_target] kind = ompt_target | endpoint = ompt_scope_begin | device_num = 0 | task_data->value = 0 | task_data->ptr = (nil) | target_id = 1 | codeptr_ra = 0x20c733 [print_ompt_device_target_submit] target_id = 1 | host_op_id = 2 | requested_num_teams = 0 [ print_ompt_device_target] kind = ompt_target | endpoint = ompt_scope_end | device_num = 0 | task_data->value = 0 | task_data->ptr = (nil) | target_id = 1 | codeptr_ra = 0x20c733 #pragma omp target map(A[:N]) [ print_ompt_device_target] kind = ompt_target | endpoint = ompt_scope_begin | device_num = 0 | task_data->value = 0 | task_data->ptr = (nil) | target_id = 3 | codeptr_ra = 0x20c804 [print_ompt_device_target_data_op] target_id = 3 | host_op_id = 4 | optype = ompt_target_data_alloc (1) | src_addr = 0x20ec58 | src_device_num = 8 | dest_addr = (nil) | dest_device_num = 0 | bytes = 8 | codeptr_ra = 0x14ea993e8897 [print_ompt_device_target_data_op] target_id = 3 | host_op_id = 5 | optype = ompt_target_data_alloc (1) | src_addr = 0x1063040 | src_device_num = 8 | dest_addr = (nil) | dest_device_num = 0 | bytes = 40 | codeptr_ra = 0x14ea993e8897 [print_ompt_device_target_data_op] target_id = 3 | host_op_id = 6 | optype = ompt_target_data_transfer_to_device (2) | src_addr = 0x1063040 | src_device_num = 8 | dest_addr = 0x14ea8cc01000 | dest_device_num = 0 | bytes = 40 | codeptr_ra = 0x14ea993e9a70 [print_ompt_device_target_data_op] target_id = 3 | host_op_id = 7 | optype = ompt_target_data_transfer_to_device (2) | src_addr = 0x1086360 | src_device_num = 8 | dest_addr = 0x14ea8cc00000 | dest_device_num = 0 | bytes = 8 | codeptr_ra = 0x14ea993e9a70 [print_ompt_device_target_submit] target_id = 3 | host_op_id = 8 | requested_num_teams = 0 [print_ompt_device_target_data_op] target_id = 3 | host_op_id = 9 | optype = ompt_target_data_transfer_from_device (3) | src_addr = 0x14ea8cc01000 | src_device_num = 0 | dest_addr = 0x1063040 | dest_device_num = 8 | bytes = 40 | codeptr_ra = 0x14ea993ead70 [print_ompt_device_target_data_op] target_id = 3 | host_op_id = 10 | optype = ompt_target_data_delete (4) | src_addr = 0x14ea8cc01000 | src_device_num = 0 | dest_addr = (nil) | dest_device_num = -1 | bytes = 0 | codeptr_ra = 0x14ea993ea2fd [ print_ompt_device_target] kind = ompt_target | endpoint = ompt_scope_end | device_num = 0 | task_data->value = 0 | task_data->ptr = (nil) | target_id = 3 | codeptr_ra = 0x20c804 END void target() ========================== ========================== BEGIN void target() #pragma omp target [ print_ompt_device_target] kind = ompt_target | endpoint = ompt_scope_begin | device_num = 0 | task_data->value = 0 | task_data->ptr = (nil) | target_id = 11 | codeptr_ra = 0x20c733 [print_ompt_device_target_submit] target_id = 11 | host_op_id = 12 | requested_num_teams = 0 [ print_ompt_device_target] kind = ompt_target | endpoint = ompt_scope_end | device_num = 0 | task_data->value = 0 | task_data->ptr = (nil) | target_id = 11 | codeptr_ra = 0x20c733 #pragma omp target map(A[:N]) [ print_ompt_device_target] kind = ompt_target | endpoint = ompt_scope_begin | device_num = 0 | task_data->value = 0 | task_data->ptr = (nil) | target_id = 13 | codeptr_ra = 0x20c804 [print_ompt_device_target_data_op] target_id = 13 | host_op_id = 14 | optype = ompt_target_data_alloc (1) | src_addr = 0x1086b30 | src_device_num = 8 | dest_addr = (nil) | dest_device_num = 0 | bytes = 40 | codeptr_ra = 0x14ea993e8897 [print_ompt_device_target_data_op] target_id = 13 | host_op_id = 15 | optype = ompt_target_data_transfer_to_device (2) | src_addr = 0x1086b30 | src_device_num = 8 | dest_addr = 0x14ea8cc01000 | dest_device_num = 0 | bytes = 40 | codeptr_ra = 0x14ea993e9a70 [print_ompt_device_target_submit] target_id = 13 | host_op_id = 16 | requested_num_teams = 0 [print_ompt_device_target_data_op] target_id = 13 | host_op_id = 17 | optype = ompt_target_data_transfer_from_device (3) | src_addr = 0x14ea8cc01000 | src_device_num = 0 | dest_addr = 0x1086b30 | dest_device_num = 8 | bytes = 40 | codeptr_ra = 0x14ea993ead70 [print_ompt_device_target_data_op] target_id = 13 | host_op_id = 18 | optype = ompt_target_data_delete (4) | src_addr = 0x14ea8cc01000 | src_device_num = 0 | dest_addr = (nil) | dest_device_num = -1 | bytes = 0 | codeptr_ra = 0x14ea993ea2fd [ print_ompt_device_target] kind = ompt_target | endpoint = ompt_scope_end | device_num = 0 | task_data->value = 0 | task_data->ptr = (nil) | target_id = 13 | codeptr_ra = 0x20c804 END void target() ========================== [ print_ompt_device_finalize] device_num = 0 ========================== BEGIN void ompt_finalize(ompt_data_t *) END void ompt_finalize(ompt_data_t *) ==========================