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));