Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Add support to select compatible kernel profile based on agents capabilities #7

Open
wants to merge 1 commit into
base: master
Choose a base branch
from
Open

Add support to select compatible kernel profile based on agents capabilities #7

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 4 additions & 1 deletion src/core/api/test_hsa_code_object_get_info.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -100,7 +100,10 @@ int test_hsa_code_object_get_info() {

// Load the BRIG module
hsa_ext_module_t module;
ASSERT(0 == load_module_from_file("no_op.brig", &module));
ASSERT(0 == load_base_or_full_module_from_file(agent,
"no_op_base_large.brig",
"no_op.brig",
&module));

hsa_machine_model_t agent_machine_model;
status = hsa_agent_get_info(agent, HSA_AGENT_INFO_MACHINE_MODEL, &agent_machine_model);
Expand Down
11 changes: 7 additions & 4 deletions src/core/api/test_hsa_code_symbol_get_info.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -164,17 +164,20 @@ int test_hsa_code_symbol_get_info() {
status = hsa_init();
ASSERT(HSA_STATUS_SUCCESS == status);

// Load a valid brig module
hsa_ext_module_t module;
ASSERT(0 == load_module_from_file("vector_copy.brig", &module));

// Get the kernel dispatch agent
hsa_agent_t agent;
agent.handle = (uint64_t)-1;
status = hsa_iterate_agents(
callback_get_kernel_dispatch_agent, &agent);
ASSERT((uint64_t)-1 != agent.handle);

// Load a valid brig module
hsa_ext_module_t module;
ASSERT(0 == load_base_or_full_module_from_file(agent,
"vector_copy_base_large.brig",
"vector_copy.brig",
&module));

// Get machine model and profile to create a program
hsa_machine_model_t machine_model;
status = hsa_agent_get_info(agent, HSA_AGENT_INFO_MACHINE_MODEL, &machine_model);
Expand Down
5 changes: 4 additions & 1 deletion src/core/api/test_hsa_executable_load_code_object.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -94,7 +94,10 @@ void load_module_finalize_program(
ASSERT((uint64_t)-1 != agent_isa.handle);

// Load the BRIG module
ASSERT(0 == load_module_from_file("no_op.brig", module_ptr));
ASSERT(0 == load_base_or_full_module_from_file(agent,
"no_op_base_large.brig",
"no_op.brig",
module_ptr));

hsa_machine_model_t machine_model;
status = hsa_agent_get_info(agent, HSA_AGENT_INFO_MACHINE_MODEL, &machine_model);
Expand Down
15 changes: 12 additions & 3 deletions src/core/api/test_hsa_isa_compatible.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -103,7 +103,10 @@ int test_hsa_isa_compatible() {

// Load the BRIG module
hsa_ext_module_t module;
ASSERT(0 == load_module_from_file("no_op.brig", &module));
ASSERT(0 == load_base_or_full_module_from_file(agent,
"no_op_base_large.brig",
"no_op.brig",
&module));

hsa_machine_model_t machine_model;
status = hsa_agent_get_info(agent, HSA_AGENT_INFO_MACHINE_MODEL, &machine_model);
Expand Down Expand Up @@ -193,7 +196,10 @@ int test_hsa_isa_compatible_invalid_isa() {

// Load the BRIG module
hsa_ext_module_t module;
ASSERT(0 == load_module_from_file("no_op.brig", &module));
ASSERT(0 == load_base_or_full_module_from_file(agent,
"no_op_base_large.brig",
"no_op.brig",
&module));

hsa_machine_model_t machine_model;
status = hsa_agent_get_info(agent, HSA_AGENT_INFO_MACHINE_MODEL, &machine_model);
Expand Down Expand Up @@ -291,7 +297,10 @@ int test_hsa_isa_compatible_null_result() {

// Load the BRIG module
hsa_ext_module_t module;
ASSERT(0 == load_module_from_file("no_op.brig", &module));
ASSERT(0 == load_base_or_full_module_from_file(agent,
"no_op_base_large.brig",
"no_op.brig",
&module));

hsa_machine_model_t machine_model;
status = hsa_agent_get_info(agent, HSA_AGENT_INFO_MACHINE_MODEL, &machine_model);
Expand Down
17 changes: 10 additions & 7 deletions src/core/code/test_code_define_global_agent.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -91,10 +91,6 @@ int test_code_define_global_agent() {
return 0;
}

// Load the BRIG module
hsa_ext_module_t module;
ASSERT(0 == load_module_from_file("global_agent_vector_copy.brig", &module));

// Get a list of agents, and iterate throught the list
struct agent_list_s agent_list;
get_agent_list(&agent_list);
Expand All @@ -109,6 +105,13 @@ int test_code_define_global_agent() {
continue;
}

// Load the BRIG module
hsa_ext_module_t module;
ASSERT(0 == load_base_or_full_module_from_file(agent_list.agents[ii],
"global_agent_vector_copy_base_large.brig",
"global_agent_vector_copy.brig",
&module));

// Get the ISA from this agent
hsa_isa_t agent_isa;
agent_isa.handle = (uint64_t)-1;
Expand Down Expand Up @@ -232,13 +235,13 @@ int test_code_define_global_agent() {
// Destroy the code object
status = hsa_code_object_destroy(code_object);
ASSERT(HSA_STATUS_SUCCESS == status);

// Destroy the loaded module
destroy_module(module);
}

free_agent_list(&agent_list);

// Destroy the loaded module
destroy_module(module);

// Shutdown runtime
status = hsa_shut_down();
ASSERT(HSA_STATUS_SUCCESS == status);
Expand Down
17 changes: 10 additions & 7 deletions src/core/code/test_code_define_global_program.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -91,10 +91,6 @@ int test_code_define_global_program() {
return 0;
}

// Load the BRIG module
hsa_ext_module_t module;
ASSERT(0 == load_module_from_file("global_vector_copy.brig", &module));

// Get a list of agents, and iterate throught the list
struct agent_list_s agent_list;
get_agent_list(&agent_list);
Expand All @@ -109,6 +105,13 @@ int test_code_define_global_program() {
continue;
}

// Load the BRIG module
hsa_ext_module_t module;
ASSERT(0 == load_base_or_full_module_from_file(agent_list.agents[ii],
"global_vector_copy_base_large.brig",
"global_vector_copy.brig",
&module));

// Get the ISA from this agent
hsa_isa_t agent_isa;
agent_isa.handle = (uint64_t)-1;
Expand Down Expand Up @@ -231,13 +234,13 @@ int test_code_define_global_program() {
// Destroy the code object
status = hsa_code_object_destroy(code_object);
ASSERT(HSA_STATUS_SUCCESS == status);

// Destroy the loaded module
destroy_module(module);
}

free_agent_list(&agent_list);

// Destroy the loaded module
destroy_module(module);

// Shutdown runtime
status = hsa_shut_down();
ASSERT(HSA_STATUS_SUCCESS == status);
Expand Down
17 changes: 10 additions & 7 deletions src/core/code/test_code_define_readonly_agent.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -92,10 +92,6 @@ int test_code_define_readonly_agent() {
return 0;
}

// Load the BRIG module
hsa_ext_module_t module;
ASSERT(0 == load_module_from_file("readonly_vector_copy.brig", &module));

// Get a list of agents, and iterate throught the list
struct agent_list_s agent_list;
get_agent_list(&agent_list);
Expand All @@ -110,6 +106,13 @@ int test_code_define_readonly_agent() {
continue;
}

// Load the BRIG module
hsa_ext_module_t module;
ASSERT(0 == load_base_or_full_module_from_file(agent_list.agents[ii],
"readonly_vector_copy_base_large.brig",
"readonly_vector_copy.brig",
&module));

// Get the ISA from this agent
hsa_isa_t agent_isa;
agent_isa.handle = (uint64_t)-1;
Expand Down Expand Up @@ -232,13 +235,13 @@ int test_code_define_readonly_agent() {
// Destroy the code object
status = hsa_code_object_destroy(code_object);
ASSERT(HSA_STATUS_SUCCESS == status);

// Destroy the loaded module
destroy_module(module);
}

free_agent_list(&agent_list);

// Destroy the loaded module
destroy_module(module);

// Shutdown runtime
status = hsa_shut_down();
ASSERT(HSA_STATUS_SUCCESS == status);
Expand Down
14 changes: 8 additions & 6 deletions src/core/code/test_code_iterate_symbols.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -388,10 +388,6 @@ int test_code_iterate_symbols() {
status = hsa_init();
ASSERT(HSA_STATUS_SUCCESS == status);

// Load the BRIG module
hsa_ext_module_t module;
ASSERT(0 == load_module_from_file("vector_copy.brig", &module));

// Get a list of agents, and iterate throught the list
struct agent_list_s agent_list;
get_agent_list(&agent_list);
Expand All @@ -405,6 +401,12 @@ int test_code_iterate_symbols() {
continue;
}

// Load the BRIG module
hsa_ext_module_t module;
ASSERT(0 == load_base_or_full_module_from_file(agent_list.agents[ii],
"vector_copy_base_large.brig",
"vector_copy.brig",
&module));
// Get the ISA from this agent
hsa_isa_t agent_isa;
agent_isa.handle = (uint64_t)-1;
Expand Down Expand Up @@ -472,12 +474,12 @@ int test_code_iterate_symbols() {
ASSERT(HSA_STATUS_SUCCESS == status);
status = hsa_code_object_destroy(code_object);
ASSERT(HSA_STATUS_SUCCESS == status);

destroy_module(module);
}

free_agent_list(&agent_list);

destroy_module(module);

status = hsa_shut_down();
ASSERT(HSA_STATUS_SUCCESS == status);

Expand Down
15 changes: 11 additions & 4 deletions src/core/code/test_code_mixed_scope.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -119,7 +119,14 @@ int test_code_mixed_scope() {
if (HSA_AGENT_FEATURE_KERNEL_DISPATCH != feature) {
continue;
}


// Load the BRIG module
hsa_ext_module_t module;
ASSERT(0 == load_base_or_full_module_from_file(agent_list.agents[ii],
"mixed_scope_base_large.brig",
"mixed_scope.brig",
&module));

// Create a queue for dispatching
hsa_queue_t* queue;
status = hsa_queue_create(agent_list.agents[ii], 1024, HSA_QUEUE_TYPE_SINGLE, NULL, NULL,UINT32_MAX,UINT32_MAX, &queue);
Expand Down Expand Up @@ -275,12 +282,12 @@ int test_code_mixed_scope() {
// Destroy the queue
status = hsa_queue_destroy(queue);
ASSERT(HSA_STATUS_SUCCESS == status);

// Destroy the loaded module
destroy_module(module);
}

free_agent_list(&agent_list);

// Destroy the loaded module
destroy_module(module);

// Shutdown runtime
status = hsa_shut_down();
Expand Down
19 changes: 11 additions & 8 deletions src/core/code/test_code_module_scope_symbol.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -90,11 +90,7 @@ int test_code_module_scope_symbol() {
if(HSA_STATUS_SUCCESS != status) {
return 0;
}

// Load the BRIG module
hsa_ext_module_t module;
ASSERT(0 == load_module_from_file("module_scope.brig", &module));


// Module name and symbol names
const char* module_name = "&module_scope";
const char* symbol_name = "&__vector_copy_kernel";
Expand All @@ -113,6 +109,13 @@ int test_code_module_scope_symbol() {
continue;
}

// Load the BRIG module
hsa_ext_module_t module;
ASSERT(0 == load_base_or_full_module_from_file(agent_list.agents[ii],
"module_scope_base_large.brig",
"module_scope.brig",
&module));

// Create a queue for dispatching
hsa_queue_t* queue;
status = hsa_queue_create(agent_list.agents[ii], 1024, HSA_QUEUE_TYPE_SINGLE, NULL, NULL,UINT32_MAX,UINT32_MAX, &queue);
Expand Down Expand Up @@ -262,12 +265,12 @@ int test_code_module_scope_symbol() {
// Destroy the queue
status = hsa_queue_destroy(queue);
ASSERT(HSA_STATUS_SUCCESS == status);

// Destroy the loaded module
destroy_module(module);
}

free_agent_list(&agent_list);

// Destroy the loaded module
destroy_module(module);

// Shutdown runtime
status = hsa_shut_down();
Expand Down
17 changes: 10 additions & 7 deletions src/core/code/test_code_multiple_executables.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -89,10 +89,6 @@ int test_code_multiple_executables() {
return 0;
}

// Load the BRIG module
hsa_ext_module_t module;
ASSERT(0 == load_module_from_file("global_vector_copy.brig", &module));

// Get a list of agents, and iterate throught the list
struct agent_list_s agent_list;
get_agent_list(&agent_list);
Expand All @@ -107,6 +103,14 @@ int test_code_multiple_executables() {
continue;
}

// Load the BRIG module
hsa_ext_module_t module;
ASSERT(0 == load_base_or_full_module_from_file(agent_list.agents[ii],
"global_vector_copy_base_large.brig",
"global_vector_copy.brig",
&module));


// Get the ISA from this agent
hsa_isa_t agent_isa;
agent_isa.handle = (uint64_t)-1;
Expand Down Expand Up @@ -237,13 +241,12 @@ int test_code_multiple_executables() {
// Destroy the code object
status = hsa_code_object_destroy(code_object);
ASSERT(HSA_STATUS_SUCCESS == status);
// Destroy the loaded module
destroy_module(module);
}

free_agent_list(&agent_list);

// Destroy the loaded module
destroy_module(module);

// Shutdown runtime
status = hsa_shut_down();
ASSERT(HSA_STATUS_SUCCESS == status);
Expand Down
7 changes: 7 additions & 0 deletions src/core/code/test_code_program_scope_symbol.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -113,6 +113,13 @@ int test_code_program_scope_symbol() {
continue;
}

// Load the BRIG module
hsa_ext_module_t module;
ASSERT(0 == load_base_or_full_module_from_file(agent_list.agents[ii],
"program_scope_base_large.brig",
"program_scope.brig",
&module));

// Create a queue for dispatching
hsa_queue_t* queue;
status = hsa_queue_create(agent_list.agents[ii], 1024, HSA_QUEUE_TYPE_SINGLE, NULL, NULL,UINT32_MAX,UINT32_MAX, &queue);
Expand Down
Loading