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

@@ -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<cl_ulong>(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);

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<size_t>(arg_index));