From 6445dcac807345c3b17bc0407bc7a9067edfe8af Mon Sep 17 00:00:00 2001 From: "Ramesh Errabolu (xN/A) TX" Date: Fri, 30 Oct 2015 17:53:25 -0500 Subject: [PATCH] Update Binarysearch and BlackScholes Hsa Sample to support FULL and BASE Profiles [git-p4: depot-paths = "//depot/stg/hsa/drivers/hsa/runtime/": change = 1206169] [ROCm/ROCR-Runtime commit: 2f0425d3542898dc43004d2809a7a540ffb9f970] --- .../samples/common/hsa_rsrc_factory.cpp | 127 +++++++++++++++--- .../samples/common/hsa_rsrc_factory.hpp | 25 ++++ 2 files changed, 136 insertions(+), 16 deletions(-) diff --git a/projects/rocr-runtime/samples/common/hsa_rsrc_factory.cpp b/projects/rocr-runtime/samples/common/hsa_rsrc_factory.cpp index 3bbba59e8e..d1bac57c77 100755 --- a/projects/rocr-runtime/samples/common/hsa_rsrc_factory.cpp +++ b/projects/rocr-runtime/samples/common/hsa_rsrc_factory.cpp @@ -21,7 +21,7 @@ uint32_t hsa_cmdline_arg_cnt; char **hsa_cmdline_arg_list; // Callback function to find and bind kernarg region of an agent -static hsa_status_t find_kernarg(hsa_region_t region, void *data) { +static hsa_status_t find_memregions(hsa_region_t region, void *data) { hsa_region_global_flag_t flags; hsa_region_segment_t segment_id; @@ -31,9 +31,13 @@ static hsa_status_t find_kernarg(hsa_region_t region, void *data) { return HSA_STATUS_SUCCESS; } + AgentInfo *agent_info = (AgentInfo *)data; hsa_region_get_info(region, HSA_REGION_INFO_GLOBAL_FLAGS, &flags); + if (flags & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) { + agent_info->coarse_region = region; + } + if (flags & HSA_REGION_GLOBAL_FLAG_KERNARG) { - AgentInfo *agent_info = (AgentInfo *)data; agent_info->kernarg_region = region; } @@ -41,7 +45,7 @@ static hsa_status_t find_kernarg(hsa_region_t region, void *data) { } // Callback function to get the number of agents -static hsa_status_t get_gpu_agents(hsa_agent_t agent, void *data) { +static hsa_status_t get_hsa_agents(hsa_agent_t agent, void *data) { // Copy handle of agent and increment number of agents reported HsaRsrcFactory *rsrcFactory = reinterpret_cast(data); @@ -50,21 +54,36 @@ static hsa_status_t get_gpu_agents(hsa_agent_t agent, void *data) { hsa_status_t status; hsa_device_type_t type; status = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type); - if (type != HSA_DEVICE_TYPE_GPU) { + if (type == HSA_DEVICE_TYPE_DSP) { + return HSA_STATUS_SUCCESS; + } + + if (type == HSA_DEVICE_TYPE_CPU) { + AgentInfo *agent_info = reinterpret_cast(malloc(sizeof(AgentInfo))); + agent_info->dev_id = agent; + agent_info->dev_type = HSA_DEVICE_TYPE_CPU; + rsrcFactory->AddAgentInfo(agent_info); return HSA_STATUS_SUCCESS; } // Device is a Gpu agent, build an instance of AgentInfo AgentInfo *agent_info = reinterpret_cast(malloc(sizeof(AgentInfo))); agent_info->dev_id = agent; + agent_info->dev_type = HSA_DEVICE_TYPE_GPU; hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, agent_info->name); agent_info->max_wave_size = 0; hsa_agent_get_info(agent, HSA_AGENT_INFO_WAVEFRONT_SIZE, &agent_info->max_wave_size); agent_info->max_queue_size = 0; hsa_agent_get_info(agent, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &agent_info->max_queue_size); + agent_info->profile = hsa_profile_t(108); + hsa_agent_get_info(agent, HSA_AGENT_INFO_PROFILE, &agent_info->profile); + + // Initialize memory regions to zero + agent_info->kernarg_region.handle = 0; + agent_info->coarse_region.handle = 0; - // Find and Bind Kernarg regions of the Gpu agent - hsa_agent_iterate_regions(agent, find_kernarg, agent_info); + // Find and Bind Memory regions of the Gpu agent + hsa_agent_iterate_regions(agent, find_memregions, agent_info); // Save the instance of AgentInfo rsrcFactory->AddAgentInfo(agent_info); @@ -94,9 +113,9 @@ HsaRsrcFactory::HsaRsrcFactory( ) { assert(status == HSA_STATUS_SUCCESS); // Discover the set of Gpu devices available on the platform - status = hsa_iterate_agents(get_gpu_agents, this); + status = hsa_iterate_agents(get_hsa_agents, this); check("Error Calling hsa_iterate_agents", status); - + // Process command line arguments ProcessCmdline( ); } @@ -131,7 +150,32 @@ bool HsaRsrcFactory::GetGpuAgentInfo(uint32_t idx, AgentInfo **agent_info) { } // Copy AgentInfo from specified index - *agent_info = gpu_list_[idx]; + AgentInfo *agent = NULL; + for (int indx = 0; indx < size; indx++) { + agent = gpu_list_[indx]; + if (agent->dev_type == HSA_DEVICE_TYPE_GPU) { + *agent_info = agent; + } + } + return true; +} + +bool HsaRsrcFactory::GetCpuAgentInfo(uint32_t idx, AgentInfo **agent_info) { + + // Determine if request is valid + uint32_t size = uint32_t(gpu_list_.size()); + if (idx >= size) { + return false; + } + + // Copy AgentInfo from specified index + AgentInfo *agent = NULL; + for (int indx = 0; indx < size; indx++) { + agent = gpu_list_[indx]; + if (agent->dev_type == HSA_DEVICE_TYPE_CPU) { + *agent_info = agent; + } + } return true; } @@ -181,7 +225,33 @@ bool HsaRsrcFactory::CreateSignal(uint32_t value, hsa_signal_t *signal) { // // @return uint8_t* Pointer to buffer, null if allocation fails. // -uint8_t* HsaRsrcFactory::AllocateMemory(AgentInfo *agent_info, size_t size) { +uint8_t* HsaRsrcFactory::AllocateLocalMemory(AgentInfo *agent_info, size_t size) { + + hsa_status_t status; + uint8_t *buffer = NULL; + + // Allocate in local memory only if it is available + if (agent_info->coarse_region.handle != 0) { + std::cout << "Allocating in local memory" << std::endl; + status = hsa_memory_allocate(agent_info->coarse_region, size, (void **)&buffer); + return (status == HSA_STATUS_SUCCESS) ? buffer : NULL; + } + + // Allocate in system memory if local memory is not available + std::cout << "Allocating in system memory" << std::endl; + status = hsa_memory_allocate(agent_info->kernarg_region, size, (void **)&buffer); + return (status == HSA_STATUS_SUCCESS) ? buffer : NULL; +} + +// Allocate memory tp pass kernel parameters. +// +// @param agent_info Agent from whose memory region to allocate +// +// @param size Size of memory in terms of bytes +// +// @return uint8_t* Pointer to buffer, null if allocation fails. +// +uint8_t* HsaRsrcFactory::AllocateSysMemory(AgentInfo *agent_info, size_t size) { hsa_status_t status; uint8_t *buffer = NULL; @@ -189,7 +259,31 @@ uint8_t* HsaRsrcFactory::AllocateMemory(AgentInfo *agent_info, size_t size) { return (status == HSA_STATUS_SUCCESS) ? buffer : NULL; } +bool HsaRsrcFactory::TransferData(uint8_t *dest_buff, uint8_t *src_buff, + uint32_t length, bool host_to_dev) { + hsa_status_t status; + + AgentInfo *agent_info; + GetGpuAgentInfo(0, &agent_info); + void *buffer = (host_to_dev) ? dest_buff : src_buff; + status = hsa_memory_assign_agent(buffer, agent_info->dev_id, HSA_ACCESS_PERMISSION_RW); + if (status != HSA_STATUS_SUCCESS) { + return false; + } + status = hsa_memory_copy(dest_buff, src_buff, length); + return (status == HSA_STATUS_SUCCESS); + +} + +// Fake method for compilation steps only +uint8_t* HsaRsrcFactory::AllocateMemory(AgentInfo *agent_info, size_t size) { + + hsa_status_t status; + uint8_t *buffer = NULL; + status = hsa_memory_allocate(agent_info->kernarg_region, size, (void **)&buffer); + return (status == HSA_STATUS_SUCCESS) ? buffer : NULL; +} // Loads an Assembled Brig file and Finalizes it into Device Isa // @@ -234,9 +328,9 @@ bool HsaRsrcFactory::LoadAndFinalize(AgentInfo *agent_info, // Create hsail program. hsa_ext_program_t hsailProgram; status = hsa_ext_program_create(HSA_MACHINE_MODEL_LARGE, - HSA_PROFILE_FULL, - HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO, - NULL, &hsailProgram); + agent_info->profile, + HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO, + NULL, &hsailProgram); check("Error in creating program object", status); // Add hsail module. @@ -266,8 +360,9 @@ bool HsaRsrcFactory::LoadAndFinalize(AgentInfo *agent_info, // Create executable. hsa_executable_t hsaExecutable; - status = hsa_executable_create(HSA_PROFILE_FULL, - HSA_EXECUTABLE_STATE_UNFROZEN, "", &hsaExecutable); + status = hsa_executable_create(agent_info->profile, + HSA_EXECUTABLE_STATE_UNFROZEN, + "", &hsaExecutable); check("Error in creating executable object", status); // Load code object. @@ -341,7 +436,7 @@ bool HsaRsrcFactory::PrintGpuAgents( ) { std::cout << "Hsa Gpu Agent Name: " << agent_info->name << std::endl; std::cout << "Hsa Gpu Agent Max Wave Size: " << agent_info->max_wave_size << std::endl; std::cout << "Hsa Gpu Agent Max Queue Size: " << agent_info->max_queue_size << std::endl; - std::cout << "Hsa Gpu Agent Kernarg Region Id: " << agent_info->kernarg_region.handle << std::endl; + std::cout << "Hsa Gpu Agent Kernarg Region Id: " << agent_info->coarse_region.handle << std::endl; std::cout << std::endl; } return true; diff --git a/projects/rocr-runtime/samples/common/hsa_rsrc_factory.hpp b/projects/rocr-runtime/samples/common/hsa_rsrc_factory.hpp index 0f9a41c8cb..82fefe195f 100755 --- a/projects/rocr-runtime/samples/common/hsa_rsrc_factory.hpp +++ b/projects/rocr-runtime/samples/common/hsa_rsrc_factory.hpp @@ -93,6 +93,9 @@ typedef struct { // Handle of Agent hsa_agent_t dev_id; + + // Agent type - Cpu = 0, Gpu = 1 or Dsp = 2 + uint32_t dev_type; // Name of Agent whose length is less than 64 char name[64]; @@ -103,8 +106,15 @@ typedef struct { // Max size of Queue buffer uint32_t max_queue_size; + // Hsail profile supported by agent + hsa_profile_t profile; + + // Memory region supporting kernel parameters + hsa_region_t coarse_region; + // Memory region supporting kernel arguments hsa_region_t kernarg_region; + } AgentInfo; class HsaRsrcFactory { @@ -133,6 +143,7 @@ class HsaRsrcFactory { // @return bool true if successful, false otherwise // bool GetGpuAgentInfo(uint32_t idx, AgentInfo **agent_info); + bool GetCpuAgentInfo(uint32_t idx, AgentInfo **agent_info); // Create a Queue object and return its handle. The queue object is expected // to support user requested number of Aql dispatch packets. @@ -168,8 +179,22 @@ class HsaRsrcFactory { // // @return uint8_t* Pointer to buffer, null if allocation fails. // + uint8_t* AllocateLocalMemory(AgentInfo *agent_info, size_t size); uint8_t* AllocateMemory(AgentInfo *agent_info, size_t size); + bool TransferData(uint8_t *dest_buff, uint8_t *src_buff, + uint32_t length, bool host_to_dev); + + // Allocate memory tp pass kernel parameters. + // + // @param agent_info Agent from whose memory region to allocate + // + // @param size Size of memory in terms of bytes + // + // @return uint8_t* Pointer to buffer, null if allocation fails. + // + uint8_t* AllocateSysMemory(AgentInfo *agent_info, size_t size); + // Loads an Assembled Brig file and Finalizes it into Device Isa // // @param agent_info Gpu device for which to finalize