From d2b85cc11857ad8b184b5896aec80a01b330dca2 Mon Sep 17 00:00:00 2001 From: foreman Date: Thu, 14 Jun 2018 17:51:35 -0400 Subject: [PATCH] P4 to Git Change 1568521 by gandryey@gera-w8 on 2018/06/14 17:43:52 SWDEV-79445 - OCL generic changes and code clean-up - Change LDS setup to account the size, since LC forces 4 bytes for LDS offsets always http://ocltc.amd.com/reviews/r/15197/ Affected files ... ... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_execute.cpp#28 edit ... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_program.cpp#49 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpublit.cpp#130 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpukernel.cpp#327 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palblit.cpp#25 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palkernel.cpp#56 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.cpp#109 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocblit.hpp#11 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rockernel.cpp#38 edit ... //depot/stg/opencl/drivers/opencl/runtime/platform/command.cpp#88 edit ... //depot/stg/opencl/drivers/opencl/runtime/platform/kernel.cpp#34 edit ... //depot/stg/opencl/drivers/opencl/runtime/platform/kernel.hpp#25 edit --- opencl/api/opencl/amdocl/cl_execute.cpp | 10 ---------- opencl/api/opencl/amdocl/cl_program.cpp | 2 +- 2 files changed, 1 insertion(+), 11 deletions(-) diff --git a/opencl/api/opencl/amdocl/cl_execute.cpp b/opencl/api/opencl/amdocl/cl_execute.cpp index c0972f5faf..15415ac3ad 100644 --- a/opencl/api/opencl/amdocl/cl_execute.cpp +++ b/opencl/api/opencl/amdocl/cl_execute.cpp @@ -239,16 +239,6 @@ RUNTIME_ENTRY(cl_int, clEnqueueNDRangeKernel, return CL_INVALID_KERNEL_ARGS; } - // Check that we do not exceed the amount of available local memory. - const size_t align = device.info().minDataTypeAlignSize_; - cl_ulong requiredLocalMemSize = - static_cast(amdKernel->parameters().localMemSize(align)) + - amd::alignUp(devKernel->workGroupInfo()->localMemSize_, align); - - if (requiredLocalMemSize > device.info().localMemSize_) { - return CL_OUT_OF_RESOURCES; - } - amd::Command::EventWaitList eventWaitList; cl_int err = amd::clSetEventWaitList(eventWaitList, hostQueue, num_events_in_wait_list, event_wait_list); diff --git a/opencl/api/opencl/amdocl/cl_program.cpp b/opencl/api/opencl/amdocl/cl_program.cpp index f0b3831df4..836e7b9b69 100644 --- a/opencl/api/opencl/amdocl/cl_program.cpp +++ b/opencl/api/opencl/amdocl/cl_program.cpp @@ -1535,7 +1535,7 @@ RUNTIME_ENTRY(cl_int, clSetKernelArg, } const amd::KernelParameterDescriptor& desc = signature.at(arg_index); - const bool is_local = desc.size_ == 0; + const bool is_local = (desc.addressQualifier_ == CL_KERNEL_ARG_ADDRESS_LOCAL); if (((arg_value == NULL) && !is_local && (desc.type_ != T_POINTER)) || ((arg_value != NULL) && is_local)) { as_amd(kernel)->parameters().reset(static_cast(arg_index));