diff --git a/samples/Makefile b/samples/Makefile index 5340ae3dd7..c7a9143e8f 100644 --- a/samples/Makefile +++ b/samples/Makefile @@ -6,12 +6,12 @@ SUBDIRS = NBody SUBDIRS += DwtHarr1D SUBDIRS += BitionicSort SUBDIRS += BinarySearch -SUBDIRS += BlackScholes -SUBDIRS += FloydWarshall -SUBDIRS += FastWalshTransform +#SUBDIRS += BlackScholes +#SUBDIRS += FloydWarshall +#SUBDIRS += FastWalshTransform SUBDIRS += MatrixTranspose SUBDIRS += MatrixMultiplication -SUBDIRS += MonteCarloAsian -SUBDIRS += SimpleConvolution +#SUBDIRS += MonteCarloAsian +#SUBDIRS += SimpleConvolution include $(OPENCL_DEPTH)/runtimenew/runtimerules diff --git a/samples/common/hsa_rsrc_factory.cpp b/samples/common/hsa_rsrc_factory.cpp index ea8ca3f2da..2937100690 100755 --- a/samples/common/hsa_rsrc_factory.cpp +++ b/samples/common/hsa_rsrc_factory.cpp @@ -12,6 +12,7 @@ #include "hsa.h" #include "elf_utils.h" #include "hsa_rsrc_factory.hpp" +#include "hsa_ext_alt_finalize.h" using namespace std; @@ -232,7 +233,7 @@ uint8_t* HsaRsrcFactory::AllocateMemory(AgentInfo *agent_info, size_t size) { // bool HsaRsrcFactory::LoadAndFinalize(AgentInfo *agent_info, const char *brig_path, char *kernel_name, - hsa_ext_code_descriptor_t **code_desc) { + hsa_executable_symbol_t *code_desc) { // Load BRIG, encapsulated in an ELF container, into a BRIG module. status_t build_err; @@ -246,6 +247,67 @@ bool HsaRsrcFactory::LoadAndFinalize(AgentInfo *agent_info, status = hsa_find_symbol_offset(brig_obj, kernel_name, &kernel_symbol); check("Error in Finding the Symbol Offset for the Kernel", status); + // Copy handle of Brig object + hsa_ext_alt_module_t brig_module_v3; + brig_module_v3.handle = uint64_t(brig_obj); + + // Create hsail program. + hsa_ext_alt_program_t hsailProgram; + status = hsa_ext_alt_program_create(HSA_MACHINE_MODEL_LARGE, + HSA_PROFILE_FULL, + HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO, + NULL, &hsailProgram); + check("Error in creating program object", status); + + // Add hsail module. + status = hsa_ext_alt_program_add_module(hsailProgram, brig_module_v3); + check("Error in adding module to program object", status); + + // Finalize hsail program. + hsa_isa_t isa; + memset(&isa, 0, sizeof(hsa_isa_t)); + + hsa_ext_alt_control_directives_t control_directives; + memset(&control_directives, 0, sizeof(hsa_ext_alt_control_directives_t)); + + hsa_code_object_t code_object; + status = hsa_ext_alt_program_finalize(hsailProgram, + isa, + 0, + control_directives, + NULL, //"-g -O0 -dump-isa", + HSA_CODE_OBJECT_TYPE_PROGRAM, + &code_object); + check("Error in finalizing program object", status); + + //status = hsa_ext_alt_program_destroy(hsailProgram); + //check("Error in destroying program object", status); + + // Create executable. + hsa_executable_t hsaExecutable; + status = hsa_executable_create(HSA_PROFILE_FULL, + HSA_EXECUTABLE_STATE_UNFROZEN, "", &hsaExecutable); + check("Error in creating executable object", status); + + // Load code object. + status = hsa_executable_load_code_object(hsaExecutable, agent_info->dev_id, code_object, ""); + check("Error in loading executable object", status); + + // Freeze executable. + status = hsa_executable_freeze(hsaExecutable, ""); + check("Error in freezing executable object", status); + + // Get symbol handle. + hsa_executable_symbol_t kernelSymbol; + status = hsa_executable_get_symbol(hsaExecutable, "", + kernel_name, agent_info->dev_id, 0, &kernelSymbol); + + // Update output parameter + *code_desc = kernelSymbol; + return true; + + /** + // Create Hsa Program hsa_ext_program_handle_t program; status = hsa_ext_program_create(&agent_info->dev_id, 1, @@ -278,6 +340,7 @@ bool HsaRsrcFactory::LoadAndFinalize(AgentInfo *agent_info, check("Error Querying the Kernel Descriptor Address", status); return true; + **/ } // Add an instance of AgentInfo representing a Hsa Gpu agent diff --git a/samples/common/hsa_rsrc_factory.hpp b/samples/common/hsa_rsrc_factory.hpp index 72ac8f9ca5..2c0dce83eb 100755 --- a/samples/common/hsa_rsrc_factory.hpp +++ b/samples/common/hsa_rsrc_factory.hpp @@ -181,7 +181,7 @@ class HsaRsrcFactory { // bool LoadAndFinalize(AgentInfo *agent_info, const char *brig_path, char *kernel_name, - hsa_ext_code_descriptor_t **code_desc); + hsa_executable_symbol_t *code_desc); // Add an instance of AgentInfo representing a Hsa Gpu agent void AddAgentInfo(AgentInfo *agent_info);