AI/rocm-systems
2
0
フォーク 0
コード イシュー プルリクエスト Actions パッケージ プロジェクト リリース Wiki アクティビティ

P4 to Git Change 1536925 by vsytchen@vsytchen-ocl-win10 on 2018/04/04 17:20:38

ソースを参照 
SWDEV-79445 - OCL generic changes and code clean-up

	1. This change replaces the use of std::map with std::unordered_map to improve lookup/insert time.
	2. Replace the use of std::make_pair and std::pair constructor with uniform initialization for cleaner code.
	3. Replace the use of std::Container::iterator type with the auto keyword for cleaner code.
	4. Use range based for loops where needed.

	ReviewBoardURL = http://ocltc.amd.com/reviews/r/14517/diff/

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#4 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_context.cpp#58 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_d3d10.cpp#16 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_d3d10_amd.hpp#9 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_d3d11.cpp#24 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_d3d11_amd.hpp#13 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_d3d9.cpp#34 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_d3d9_amd.hpp#17 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_gl.cpp#57 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_pipe.cpp#7 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_program.cpp#46 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_svm.cpp#23 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/appprofile.hpp#14 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/cpu/cpuprogram.cpp#72 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/cpu/cpuvirtual.cpp#27 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/device.cpp#216 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/device.hpp#297 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpuappprofile.cpp#13 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpubinary.cpp#59 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpucompiler.cpp#158 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpudevice.cpp#587 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpukernel.cpp#322 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpuprintf.cpp#46 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpuprogram.cpp#237 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpuprogram.hpp#70 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpuresource.cpp#242 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpuvirtual.cpp#415 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpuvirtual.hpp#143 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palappprofile.cpp#3 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palcompiler.cpp#22 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.cpp#79 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palprintf.cpp#9 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palprogram.cpp#59 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palresource.cpp#60 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.cpp#84 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.hpp#46 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/CMakeLists.txt#11 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/pro/prodevice.cpp#4 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/pro/prodevice.hpp#5 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocbinary.hpp#6 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/roccompiler.cpp#42 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/roccounters.cpp#3 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocprintf.cpp#10 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocprogram.cpp#81 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/command.cpp#81 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/command.hpp#89 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/commandqueue.cpp#24 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/context.cpp#49 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/context.hpp#29 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/memory.cpp#129 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/memory.hpp#102 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/perfctr.hpp#7 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/program.cpp#91 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/program.hpp#43 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/sampler.hpp#9 edit
... //depot/stg/opencl/drivers/opencl/runtime/utils/flags.cpp#17 edit


[ROCm/clr commit: 1aade00b09]
このコミットが含まれているのは:
foreman
2018-04-04 18:00:17 -04:00
親 4cb43aeb3e
コミット 344b507e4f
11個のファイルの変更、66行の追加、92行の削除
Patchファイルをダウンロード Diffファイルをダウンロード すべてのファイルを展開 すべてのファイルを折り畳む
projects/clr/opencl/api/opencl/amdocl/cl_context.cpp
+2 -3
ファイルの表示
@@ -296,9 +296,8 @@ RUNTIME_ENTRY(cl_int, clGetContextInfo,
*not_null(param_value_size_ret) = valueSize;
if (param_value != NULL) {
cl_device_id* device_list = (cl_device_id*)param_value;
std::vector<amd::Device*>::const_iterator it;
for (it = devices.begin(); it != devices.end(); ++it) {
*device_list++ = const_cast<cl_device_id>(as_cl(*it));
for (const auto& it : devices) {
*device_list++ = const_cast<cl_device_id>(as_cl(it));
}
}
return CL_SUCCESS;
projects/clr/opencl/api/opencl/amdocl/cl_d3d10.cpp
+6 -8
ファイルの表示
@@ -125,7 +125,7 @@ RUNTIME_ENTRY(cl_int, clGetDeviceIDsFromD3D10KHR,
break;
}
std::vector<amd::Device*>::iterator it = compatible_devices.begin();
auto it = compatible_devices.cbegin();
cl_uint compatible_count = std::min(num_entries, (cl_uint)compatible_devices.size());
while (compatible_count--) {
@@ -298,9 +298,8 @@ RUNTIME_ENTRY_RET(cl_mem, clCreateImageFromD3D10Resource,
const std::vector<amd::Device*>& devices = as_amd(context)->devices();
bool supportPass = false;
bool sizePass = false;
std::vector<amd::Device*>::const_iterator it;
for (it = devices.begin(); it != devices.end(); ++it) {
if ((*it)->info().imageSupport_) {
for (const auto& it : devices) {
if (it->info().imageSupport_) {
supportPass = true;
}
}
@@ -1104,9 +1103,8 @@ int D3D10Object::initD3D10Object(const Context& amdContext, ID3D10Resource* pRes
ScopedLock sl(resLock_);
// Check if this ressource has already been used for interop
std::vector<std::pair<void*, UINT>>::iterator it;
for (it = resources_.begin(); it != resources_.end(); ++it) {
if ((*it).first == (void*)pRes && (*it).second == subres) {
for (const auto& it : resources_) {
if (it.first == (void*)pRes && it.second == subres) {
return CL_INVALID_D3D10_RESOURCE_KHR;
}
}
@@ -1349,7 +1347,7 @@ int D3D10Object::initD3D10Object(const Context& amdContext, ID3D10Resource* pRes
return CL_INVALID_IMAGE_FORMAT_DESCRIPTOR;
}
}
resources_.push_back(std::make_pair(pRes, subres));
resources_.push_back({pRes, subres});
return CL_SUCCESS;
}
projects/clr/opencl/api/opencl/amdocl/cl_d3d10_amd.hpp
+1 -2
ファイルの表示
@@ -163,9 +163,8 @@ public:
if(pD3D10ResOrig_) pD3D10ResOrig_->Release();
if(pQuery_) pQuery_->Release();
// Check if this resource has already been used for interop
std::vector<std::pair<void*, UINT>>::iterator it;
if(resources_.size()) {
for(it = resources_.begin(); it != resources_.end(); ++it) {
for(auto& it = resources_.cbegin(); it != resources_.cend(); it++) {
if(((pD3D10ResOrig_ && (*it).first == (void*) pD3D10ResOrig_)
|| ((*it).first == (void*) pD3D10Res_))
&& (*it).second == subRes_) {
projects/clr/opencl/api/opencl/amdocl/cl_d3d11.cpp
+7 -9
ファイルの表示
@@ -127,7 +127,7 @@ RUNTIME_ENTRY(cl_int, clGetDeviceIDsFromD3D11KHR,
break;
}
std::vector<amd::Device*>::iterator it = compatible_devices.begin();
auto it = compatible_devices.cbegin();
cl_uint compatible_count = std::min(num_entries, (cl_uint)compatible_devices.size());
while (compatible_count--) {
@@ -300,9 +300,8 @@ RUNTIME_ENTRY_RET(cl_mem, clCreateImageFromD3D11Resource,
const std::vector<amd::Device*>& devices = as_amd(context)->devices();
bool supportPass = false;
bool sizePass = false;
std::vector<amd::Device*>::const_iterator it;
for (it = devices.begin(); it != devices.end(); ++it) {
if ((*it)->info().imageSupport_) {
for (const auto& it : devices) {
if (it->info().imageSupport_) {
supportPass = true;
}
}
@@ -613,10 +612,9 @@ int D3D11Object::initD3D11Object(const Context& amdContext, ID3D11Resource* pRes
ScopedLock sl(resLock_);
// Check if this ressource has already been used for interop
std::vector<std::pair<void*, std::pair<UINT, UINT>>>::iterator it;
for (it = resources_.begin(); it != resources_.end(); ++it) {
if ((*it).first == (void*)pRes && (*it).second.first == subres &&
(*it).second.second == plane) {
for (const auto& it : resources_) {
if (it.first == (void*)pRes && it.second.first == subres &&
it.second.second == plane) {
return CL_INVALID_D3D11_RESOURCE_KHR;
}
}
@@ -889,7 +887,7 @@ int D3D11Object::initD3D11Object(const Context& amdContext, ID3D11Resource* pRes
return CL_INVALID_IMAGE_FORMAT_DESCRIPTOR;
}
}
resources_.push_back(std::make_pair(pRes, std::make_pair(subres, plane)));
resources_.push_back({pRes, {subres, plane}});
return CL_SUCCESS;
}
projects/clr/opencl/api/opencl/amdocl/cl_d3d11_amd.hpp
+1 -2
ファイルの表示
@@ -171,9 +171,8 @@ public:
if(pD3D11ResOrig_) pD3D11ResOrig_->Release();
if(pQuery_) pQuery_->Release();
// Check if this resource has already been used for interop
std::vector<std::pair<void*, std::pair<UINT,UINT>>>::iterator it;
if(resources_.size()) {
for(it = resources_.begin(); it != resources_.end(); ++it) {
for(auto& it = resources_.cbegin(); it != resources_.cend(); it++) {
if(((pD3D11ResOrig_ && (*it).first == (void*) pD3D11ResOrig_)
|| ((*it).first == (void*) pD3D11Res_))
&& (*it).second.first == subRes_
projects/clr/opencl/api/opencl/amdocl/cl_d3d9.cpp
+11 -14
ファイルの表示
@@ -101,7 +101,7 @@ RUNTIME_ENTRY(cl_int, clGetDeviceIDsFromDX9MediaAdapterKHR,
break;
}
std::vector<amd::Device*>::iterator it = compatible_devices.begin();
auto it = compatible_devices.cbegin();
cl_uint compatible_count = std::min(num_entries, (cl_uint)compatible_devices.size());
while (compatible_count--) {
@@ -183,9 +183,8 @@ RUNTIME_ENTRY_RET(cl_mem, clCreateFromDX9MediaSurfaceKHR,
const std::vector<amd::Device*>& devices = as_amd(context)->devices();
bool supportPass = false;
bool sizePass = false;
std::vector<amd::Device*>::const_iterator it;
for (it = devices.begin(); it != devices.end(); ++it) {
if ((*it)->info().imageSupport_) {
for (const auto& it : devices) {
if (it->info().imageSupport_) {
supportPass = true;
}
}
@@ -484,9 +483,8 @@ int D3D9Object::initD3D9Object(const Context& amdContext,
return CL_INVALID_DX9_MEDIA_ADAPTER_KHR; // Not supported yet
}
std::vector<std::pair<TD3D9RESINFO, TD3D9RESINFO>>::iterator it;
for (it = resources_.begin(); it != resources_.end(); ++it) {
if ((*it).first.surfInfo.resource == cl_surf_info->resource && (*it).first.surfPlane == plane) {
for (const auto& it : resources_) {
if (it.first.surfInfo.resource == cl_surf_info->resource && it.first.surfPlane == plane) {
return CL_INVALID_D3D9_RESOURCE_KHR;
}
}
@@ -519,12 +517,11 @@ int D3D9Object::initD3D9Object(const Context& amdContext,
// first check if the format is NV12 or YV12, which we need special handling
if (NULL == shared_handle) {
bool found = false;
std::vector<std::pair<TD3D9RESINFO, TD3D9RESINFO>>::iterator it;
for (it = resources_.begin(); it != resources_.end(); ++it) {
if ((*it).first.surfInfo.resource == cl_surf_info->resource &&
(*it).first.surfPlane != plane) {
obj.handleShared_ = (*it).second.surfInfo.shared_handle;
obj.pD3D9Res_ = (*it).second.surfInfo.resource;
for (const auto& it : resources_) {
if (it.first.surfInfo.resource == cl_surf_info->resource &&
it.first.surfPlane != plane) {
obj.handleShared_ = it.second.surfInfo.shared_handle;
obj.pD3D9Res_ = it.second.surfInfo.resource;
obj.pD3D9Res_->AddRef();
obj.objDesc_ = obj.objDescOrig_;
found = true;
@@ -567,7 +564,7 @@ int D3D9Object::initD3D9Object(const Context& amdContext,
TD3D9RESINFO d3d9ObjShared = {{obj.pD3D9Res_, obj.handleShared_}, plane};
if (errcode == CL_SUCCESS) {
resources_.push_back(std::make_pair(d3d9ObjOri, d3d9ObjShared));
resources_.push_back({d3d9ObjOri, d3d9ObjShared});
}
return errcode;
projects/clr/opencl/api/opencl/amdocl/cl_d3d9_amd.hpp
+1 -2
ファイルの表示
@@ -239,9 +239,8 @@ public:
if(pD3D9Aux_) pD3D9Aux_->Release();
if(pQuery_) pQuery_->Release();
//if the resouce is being used
std::vector<std::pair<TD3D9RESINFO, TD3D9RESINFO>>::iterator it;
if(resources_.size()) {
for(it = resources_.begin(); it != resources_.end(); ++it) {
for(auto& it = resources_.cbegin(); it != resources_.cend(); it++) {
if( surfInfo_.resource &&
((*it).first.surfInfo.resource == surfInfo_.resource) &&
((*it).first.surfPlane == surfPlane_)) {
projects/clr/opencl/api/opencl/amdocl/cl_gl.cpp
+10 -14
ファイルの表示
@@ -199,9 +199,8 @@ RUNTIME_ENTRY_RET(cl_mem, clCreateFromGLTexture,
const std::vector<amd::Device*>& devices = as_amd(context)->devices();
bool supportPass = false;
bool sizePass = false;
std::vector<amd::Device*>::const_iterator it;
for (it = devices.begin(); it != devices.end(); ++it) {
if ((*it)->info().imageSupport_) {
for (const auto& it : devices) {
if (it->info().imageSupport_) {
supportPass = true;
}
}
@@ -285,9 +284,8 @@ RUNTIME_ENTRY_RET(cl_mem, clCreateFromGLTexture2D,
const std::vector<amd::Device*>& devices = as_amd(context)->devices();
bool supportPass = false;
bool sizePass = false;
std::vector<amd::Device*>::const_iterator it;
for (it = devices.begin(); it != devices.end(); ++it) {
if ((*it)->info().imageSupport_) {
for (const auto& it : devices) {
if (it->info().imageSupport_) {
supportPass = true;
}
}
@@ -366,9 +364,8 @@ RUNTIME_ENTRY_RET(cl_mem, clCreateFromGLTexture3D,
const std::vector<amd::Device*>& devices = as_amd(context)->devices();
bool supportPass = false;
bool sizePass = false;
std::vector<amd::Device*>::const_iterator it;
for (it = devices.begin(); it != devices.end(); ++it) {
if ((*it)->info().imageSupport_) {
for (const auto& it : devices) {
if (it->info().imageSupport_) {
supportPass = true;
}
}
@@ -914,9 +911,8 @@ RUNTIME_ENTRY(cl_int, clGetGLContextInfoKHR,
if (param_value != NULL) {
cl_device_id* deviceList = (cl_device_id*)param_value;
std::vector<amd::Device*>::const_iterator it;
for (it = compatible_devices.begin(); it != compatible_devices.end(); ++it) {
*deviceList++ = as_cl(*it);
for (const auto& it : compatible_devices) {
*deviceList++ = as_cl(it);
}
}
@@ -1616,8 +1612,8 @@ cl_mem clCreateFromGLBufferAMD(Context& amdContext, cl_mem_flags flags, GLuint b
// We should come up with a more elegant solution to handle this.
assert(amdContext.devices().size() == 1);
std::vector<amd::Device*>::const_iterator itr = amdContext.devices().begin();
amd::Device& dev = *(*itr);
const auto it = amdContext.devices().cbegin();
const amd::Device& dev = *(*it);
if (dev.type() != CL_DEVICE_TYPE_CPU) {
device::Memory* mem = pBufferGL->getDeviceMemory(dev);
projects/clr/opencl/api/opencl/amdocl/cl_pipe.cpp
+2 -3
ファイルの表示
@@ -81,10 +81,9 @@ RUNTIME_ENTRY_RET(cl_mem, clCreatePipe,
size_t size = sizeof(struct clk_pipe_t) + pipe_packet_size * pipe_max_packets;
const std::vector<amd::Device*>& devices = as_amd(context)->devices();
std::vector<amd::Device*>::const_iterator it;
bool sizePass = false;
for (it = devices.begin(); it != devices.end(); ++it) {
if (((*it)->info().maxMemAllocSize_ >= size)) {
for (const auto& it : devices) {
if (it->info().maxMemAllocSize_ >= size) {
sizePass = true;
break;
}
projects/clr/opencl/api/opencl/amdocl/cl_program.cpp
+16 -24
ファイルの表示
@@ -24,9 +24,8 @@ static amd::Program* createProgram(cl_context context, cl_uint num_devices,
// Add programs for all devices in the context.
if (device_list == NULL) {
const std::vector<amd::Device*>& devices = as_amd(context)->devices();
std::vector<amd::Device*>::const_iterator it;
for (it = devices.begin(); it != devices.end(); ++it) {
if (program->addDeviceProgram(**it) == CL_OUT_OF_HOST_MEMORY) {
for (const auto& it : devices) {
if (program->addDeviceProgram(*it) == CL_OUT_OF_HOST_MEMORY) {
*not_null(errcode_ret) = CL_OUT_OF_HOST_MEMORY;
program->release();
return NULL;
@@ -142,9 +141,8 @@ RUNTIME_ENTRY_RET(cl_program, clCreateProgramWithSource,
// Add programs for all devices in the context.
const std::vector<amd::Device*>& devices = as_amd(context)->devices();
std::vector<amd::Device*>::const_iterator it;
for (it = devices.begin(); it != devices.end(); ++it) {
if (program->addDeviceProgram(**it) == CL_OUT_OF_HOST_MEMORY) {
for (const auto& it : devices) {
if (program->addDeviceProgram(*it) == CL_OUT_OF_HOST_MEMORY) {
*not_null(errcode_ret) = CL_OUT_OF_HOST_MEMORY;
program->release();
return (cl_program)0;
@@ -204,9 +202,8 @@ RUNTIME_ENTRY_RET(cl_program, clCreateProgramWithIL,
// Add programs for all devices in the context.
const std::vector<amd::Device*>& devices = as_amd(context)->devices();
std::vector<amd::Device*>::const_iterator it;
for (it = devices.begin(); it != devices.end(); ++it) {
if (program->addDeviceProgram(**it, il, length) == CL_OUT_OF_HOST_MEMORY) {
for (const auto& it : devices) {
if (program->addDeviceProgram(*it, il, length) == CL_OUT_OF_HOST_MEMORY) {
*not_null(errcode_ret) = CL_OUT_OF_HOST_MEMORY;
program->release();
return (cl_program)0;
@@ -368,9 +365,8 @@ RUNTIME_ENTRY_RET(cl_program, clCreateProgramWithAssemblyAMD,
// Add programs for all devices in the context.
const std::vector<amd::Device*>& devices = as_amd(context)->devices();
std::vector<amd::Device*>::const_iterator it;
for (it = devices.begin(); it != devices.end(); ++it) {
if (program->addDeviceProgram(**it) == CL_OUT_OF_HOST_MEMORY) {
for (const auto& it : devices) {
if (program->addDeviceProgram(*it) == CL_OUT_OF_HOST_MEMORY) {
*not_null(errcode_ret) = CL_OUT_OF_HOST_MEMORY;
program->release();
return (cl_program)0;
@@ -964,9 +960,8 @@ RUNTIME_ENTRY(cl_int, clGetProgramInfo,
*not_null(param_value_size_ret) = valueSize;
if (param_value != NULL) {
cl_device_id* device_list = (cl_device_id*)param_value;
amd::Program::devicelist_t::const_iterator it;
for (it = devices.begin(); it != devices.end(); ++it) {
*device_list++ = const_cast<cl_device_id>(as_cl(*it));
for (const auto& it : devices) {
*device_list++ = const_cast<cl_device_id>(as_cl(it));
}
if (param_value_size > valueSize) {
::memset(static_cast<address>(param_value) + valueSize, '\0',
@@ -991,9 +986,8 @@ RUNTIME_ENTRY(cl_int, clGetProgramInfo,
*not_null(param_value_size_ret) = valueSize;
if (param_value != NULL) {
size_t* binary_sizes = (size_t*)param_value;
amd::Program::devicelist_t::const_iterator it;
for (it = devices.begin(); it != devices.end(); ++it) {
*binary_sizes++ = amdProgram->getDeviceProgram(**it)->binary().second;
for (const auto& it : devices) {
*binary_sizes++ = amdProgram->getDeviceProgram(*it)->binary().second;
}
if (param_value_size > valueSize) {
::memset(static_cast<address>(param_value) + valueSize, '\0',
@@ -1014,9 +1008,8 @@ RUNTIME_ENTRY(cl_int, clGetProgramInfo,
*not_null(param_value_size_ret) = valueSize;
if (param_value != NULL) {
char** binaries = (char**)param_value;
amd::Program::devicelist_t::const_iterator it;
for (it = devices.begin(); it != devices.end(); ++it) {
const device::Program::binary_t& binary = amdProgram->getDeviceProgram(**it)->binary();
for (const auto& it : devices) {
const device::Program::binary_t& binary = amdProgram->getDeviceProgram(*it)->binary();
// If an entry value in the array is NULL,
// then runtime should skip copying the program binary
if (*binaries != NULL) {
@@ -1376,9 +1369,8 @@ RUNTIME_ENTRY(cl_int, clCreateKernelsInProgram, (cl_program program, cl_uint num
const amd::Program::symbols_t& symbols = as_amd(program)->symbols();
cl_kernel* result = kernels;
amd::Program::symbols_t::const_iterator it;
for (it = symbols.begin(); it != symbols.end(); ++it) {
amd::Kernel* kernel = new amd::Kernel(*as_amd(program), it->second, it->first);
for (const auto& it : symbols) {
amd::Kernel* kernel = new amd::Kernel(*as_amd(program), it.second, it.first);
if (kernel == NULL) {
while (--result >= kernels) {
as_amd(*result)->release();
projects/clr/opencl/api/opencl/amdocl/cl_svm.cpp
+9 -11
ファイルの表示
@@ -74,7 +74,7 @@ static bool validateMapFlags(cl_map_flags flags) {
* \param size is the size in bytes of the SVM buffer to be allocated.
*
* \param alignment is the minimum alignment in bytes that is required for the
* newly created buffer’s memory region. It must be a power of two up to the
* newly created buffer’s memory region. It must be a power of two up to the
* largest data type supported by the OpenCL device. For the full profile, the
* largest data type is long16. For the embedded profile, it is long16 if the
* device supports 64-bit integers; otherwise it is int16. If alignment is 0, a
@@ -127,25 +127,24 @@ RUNTIME_ENTRY_RET_NOERRCODE(void*, clSVMAlloc, (cl_context context, cl_svm_mem_f
cl_device_svm_capabilities combinedSvmCapabilities = 0;
const cl_uint hostAddressBits = LP64_SWITCH(32, 64);
cl_uint minContextAlignment = std::numeric_limits<uint>::max();
std::vector<amd::Device*>::const_iterator it;
for (it = devices.begin(); it != devices.end(); ++it) {
cl_device_svm_capabilities svmCapabilities = (*it)->info().svmCapabilities_;
for (const auto& it : devices) {
cl_device_svm_capabilities svmCapabilities = it->info().svmCapabilities_;
if (svmCapabilities == 0) {
continue;
}
combinedSvmCapabilities |= svmCapabilities;
if ((*it)->info().maxMemAllocSize_ >= size) {
if (it->info().maxMemAllocSize_ >= size) {
sizePass = true;
}
if ((*it)->info().addressBits_ < hostAddressBits) {
if (it->info().addressBits_ < hostAddressBits) {
LogWarning("address mode mismatch between host and device");
return NULL;
}
// maximum alignment for a device is given in bits.
cl_uint baseAlignment = (*it)->info().memBaseAddrAlign_ >> 3;
cl_uint baseAlignment = it->info().memBaseAddrAlign_ >> 3;
if (alignment > baseAlignment) {
LogWarning("invalid parameter \"alignment\"");
return NULL;
@@ -877,7 +876,7 @@ RUNTIME_EXIT
* (clEnqueueNDRangeKernel) until the argument value is changed by a call to
* clSetKernelArgSVMPointer for \a kernel. The SVM pointer can only be used for
* arguments that are declared to be a pointer to global or constant memory.
* The SVM pointer value must be aligned according to the argument’s type. For
* The SVM pointer value must be aligned according to the argument�s type. For
* example, if the argument is declared to be global float4 *p, the SVM pointer
* value passed for p must be at a minimum aligned to a float4. The SVM pointer
* value specified as the argument value can be the pointer returned by
@@ -985,9 +984,8 @@ RUNTIME_ENTRY(cl_int, clSetKernelExecInfo, (cl_kernel kernel, cl_kernel_exec_inf
const amd::Context* amdContext = &amdKernel->program().context();
bool foundFineGrainedSystemDevice = false;
const std::vector<amd::Device*>& devices = amdContext->devices();
std::vector<amd::Device*>::const_iterator it;
for (it = devices.begin(); it != devices.end(); ++it) {
if ((*it)->info().svmCapabilities_ & CL_DEVICE_SVM_FINE_GRAIN_SYSTEM) {
for (const auto it : devices) {
if (it->info().svmCapabilities_ & CL_DEVICE_SVM_FINE_GRAIN_SYSTEM) {
foundFineGrainedSystemDevice = true;
break;
}