AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

P4 to Git Change 1981324 by kjayapra@3_HIPWS_TXT_ROCM on 2019/08/11 18:44:40

Browse Source 
SWDEV-188177 - Texture API implementation and support for extern variables.

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.def.in#18 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.map.in#20 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_internal.hpp#35 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#32 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#37 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_texture.cpp#14 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/device.hpp#340 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/devprogram.cpp#57 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/devprogram.hpp#31 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpudevice.cpp#608 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpudevice.hpp#172 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpuprogram.cpp#250 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpuprogram.hpp#79 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.cpp#152 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.hpp#41 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palprogram.cpp#96 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palprogram.hpp#39 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.cpp#133 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.hpp#39 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocprogram.cpp#105 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocprogram.hpp#48 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/program.cpp#102 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/program.hpp#46 edit
This commit is contained in:
foreman
2019-08-11 18:53:11 -04:00
parent 0036c61a8c
commit 0c83114de0
17 changed files with 246 additions and 77 deletions
Download Patch File Download Diff File Expand all files Collapse all files
rocclr/runtime/device/rocm/rocprogram.cpp
+23 -4
View File
@@ -74,7 +74,7 @@ Program::~Program() {
releaseClBinary();
}
Program::Program(roc::NullDevice& device) : device::Program(device) {
Program::Program(roc::NullDevice& device, amd::Program& owner) : device::Program(device, owner) {
hsaExecutable_.handle = 0;
hsaCodeObjectReader_.handle = 0;
}
@@ -113,6 +113,21 @@ bool Program::initClBinary(char* binaryIn, size_t size) {
return clBinary()->setBinary(bin, sz, (decryptedBin != nullptr));
}
bool Program::defineGlobalVar(const char* name, void* dptr) {
hsa_status_t status = HSA_STATUS_SUCCESS;
hsa_agent_t hsa_device = dev().getBackendDevice();
status = hsa_executable_agent_global_variable_define(hsaExecutable_, hsa_device, name, dptr);
if (status != HSA_STATUS_SUCCESS) {
buildLog_ += "Error: Could not define global variable : ";
buildLog_ += hsa_strerror(status);
buildLog_ += "\n";
}
return (status == HSA_STATUS_SUCCESS);
}
bool Program::createGlobalVarObj(amd::Memory** amd_mem_obj, void** device_pptr,
size_t* bytes, const char* global_name) const {
hsa_status_t status = HSA_STATUS_SUCCESS;
@@ -197,7 +212,7 @@ bool Program::createGlobalVarObj(amd::Memory** amd_mem_obj, void** device_pptr,
return true;
}
HSAILProgram::HSAILProgram(roc::NullDevice& device) : roc::Program(device) {
HSAILProgram::HSAILProgram(roc::NullDevice& device, amd::Program& owner) : roc::Program(device, owner) {
xnackEnabled_ = dev().settings().enableXNACK_;
sramEccEnabled_ = dev().info().sramEccEnabled_;
machineTarget_ = dev().deviceInfo().complibTarget_;
@@ -258,6 +273,8 @@ bool HSAILProgram::setKernels(amd::option::Options* options, void* binary, size_
return false;
}
defineUndefinedVars();
// Load the code object.
hsa_code_object_reader_t codeObjectReader;
status = hsa_code_object_reader_create_from_memory(data, secSize, &codeObjectReader);
@@ -402,8 +419,8 @@ bool HSAILProgram::setKernels(amd::option::Options* options, void* binary, size_
}
LightningProgram::LightningProgram(roc::NullDevice& device)
: roc::Program(device) {
LightningProgram::LightningProgram(roc::NullDevice& device, amd::Program& owner)
: roc::Program(device, owner) {
isLC_ = true;
xnackEnabled_ = dev().settings().enableXNACK_;
sramEccEnabled_ = dev().info().sramEccEnabled_;
@@ -461,6 +478,8 @@ bool LightningProgram::setKernels(amd::option::Options* options, void* binary, s
return false;
}
defineUndefinedVars();
// Load the code object.
status = hsa_code_object_reader_create_from_memory(binary, binSize, &hsaCodeObjectReader_);
if (status != HSA_STATUS_SUCCESS) {