Skip to content

[SYCL][USM] Enable per-context USM behavior. #467

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
wants to merge 3 commits into from
Closed

[SYCL][USM] Enable per-context USM behavior. #467

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
e7e8c1a
Enable per-context USM behavior.
jbrodman Aug 6, 2019
d3fd40d
Update var names and minor fixes.
jbrodman Aug 9, 2019
74952a7
Fix typo
jbrodman Aug 9, 2019
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
65 changes: 32 additions & 33 deletions sycl/include/CL/sycl/detail/clusm.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,35 +24,35 @@ namespace usm {

class CLUSM {
public:
static bool Create(CLUSM *&pCLUSM);
static void Delete(CLUSM *&pCLUSM);
CLUSM() = default;
~CLUSM() = default;

void initExtensions(cl_platform_id platform);
void initExtensions(cl_context Context, cl_platform_id Platform);

void *hostMemAlloc(cl_context context, cl_mem_properties_intel *properties,
size_t size, cl_uint alignment, cl_int *errcode_ret);
void *deviceMemAlloc(cl_context context, cl_device_id device,
cl_mem_properties_intel *properties, size_t size,
cl_uint alignment, cl_int *errcode_ret);
void *sharedMemAlloc(cl_context context, cl_device_id device,
cl_mem_properties_intel *properties, size_t size,
cl_uint alignment, cl_int *errcode_ret);
void *hostMemAlloc(cl_context Context, cl_mem_properties_intel *Properties,
size_t Size, cl_uint Alignment, cl_int *Errcode_ret);
void *deviceMemAlloc(cl_context Context, cl_device_id Device,
cl_mem_properties_intel *Properties, size_t Size,
cl_uint Alignment, cl_int *Errcode_ret);
void *sharedMemAlloc(cl_context Context, cl_device_id Device,
cl_mem_properties_intel *Properties, size_t Size,
cl_uint Alignment, cl_int *Errcode_ret);

cl_int memFree(cl_context context, const void *ptr);
cl_int memFree(cl_context Context, const void *Ptr);

cl_int getMemAllocInfoINTEL(cl_context context, const void *ptr,
cl_mem_info_intel param_name,
size_t param_value_size, void *param_value,
size_t *param_value_size_ret);
cl_int getMemAllocInfoINTEL(cl_context Context, const void *Ptr,
cl_mem_info_intel Param_name,
size_t Param_value_size, void *Param_value,
size_t *Param_value_size_ret);

cl_int setKernelExecInfo(cl_kernel kernel, cl_kernel_exec_info param_name,
size_t param_value_size, const void *param_value);
cl_int setKernelExecInfo(cl_kernel Kernel, cl_kernel_exec_info Param_name,
size_t Param_value_size, const void *Param_value);

cl_int setKernelIndirectUSMExecInfo(cl_command_queue queue, cl_kernel kernel);
cl_int setKernelIndirectUSMExecInfo(cl_command_queue Queue, cl_kernel Kernel);

template <class T>
cl_int writeParamToMemory(size_t param_value_size, T param,
size_t *param_value_size_ret, T *pointer) const;
cl_int writeParamToMemory(size_t Param_value_size, T Param,
size_t *Param_value_size_ret, T *Pointer) const;

bool useCLUSM() { return mEnableCLUSM; }

Expand All @@ -63,9 +63,6 @@ class CLUSM {
bool mInitialized = false;
std::mutex mLock;

CLUSM() = default;
~CLUSM() = default;

struct SUSMAllocInfo {
SUSMAllocInfo() = default;

Expand Down Expand Up @@ -108,23 +105,25 @@ class CLUSM {
} // namespace usm

namespace cliext {
bool initializeExtensions(cl_platform_id platform);
bool initializeExtensions(cl_context context, cl_platform_id platform);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please, follow llvm coding guidelines. name variables starting from big letter, functions from small.

} // namespace cliext

} // namespace detail
} // namespace sycl
} // namespace cl

__SYCL_EXPORTED extern cl::sycl::detail::usm::CLUSM *gCLUSM;
inline cl::sycl::detail::usm::CLUSM *GetCLUSM() {
if (gCLUSM == nullptr) {
cl::sycl::detail::usm::CLUSM::Create(gCLUSM);
__SYCL_EXPORTED extern std::map<cl_context, cl::sycl::detail::usm::CLUSM *>
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can't you store pointer to CLUSM in the context object?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yep. Wasn't sure if you guys preferred it to be that invasive.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actually - not sure I can really do this.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

See below in the opencl_shim part.

gCLUSM;
inline cl::sycl::detail::usm::CLUSM *GetCLUSM(cl_context ctxt) {
if (!cl::sycl::detail::pi::piUseBackend(
cl::sycl::detail::pi::PiBackend::SYCL_BE_PI_OPENCL)) {
// Bail if we're not using a CL backend. CLUSM is not relevant.
return nullptr;
}

cl::sycl::detail::usm::CLUSM *retVal = nullptr;
if (cl::sycl::detail::pi::piUseBackend(
cl::sycl::detail::pi::PiBackend::SYCL_BE_PI_OPENCL)) {
retVal = gCLUSM;
cl::sycl::detail::usm::CLUSM *&retVal = gCLUSM[ctxt];
if (retVal == nullptr) {
retVal = new cl::sycl::detail::usm::CLUSM();
}
return retVal;
}
9 changes: 6 additions & 3 deletions sycl/source/detail/context_impl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -40,9 +40,12 @@ context_impl::context_impl(const vector_class<cl::sycl::device> Devices,
RT::piContextCreate(0, DeviceIds.size(), DeviceIds.data(), 0, 0, &Err),
Err));

if (usm::CLUSM* clusm = GetCLUSM()) {
cl_platform_id id = m_Platform.get();
clusm->initExtensions(id);
if (pi::piUseBackend(pi::PiBackend::SYCL_BE_PI_OPENCL)) {
cl_context ctxt = pi::pi_cast<cl_context>(m_Context);
if (usm::CLUSM *clusm = GetCLUSM(ctxt)) {
cl_platform_id id = m_Platform.get();
clusm->initExtensions(ctxt, id);
}
}
}

Expand Down
65 changes: 35 additions & 30 deletions sycl/source/detail/scheduler/commands.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -676,36 +676,41 @@ cl_int ExecCGCommand::enqueueImp() {
MQueue->get_device())->getHandleRef());

// TODO: Replace CL with PI
auto clusm = GetCLUSM();
if (usesUSM && clusm) {
cl_bool t = CL_TRUE;
auto theKernel = pi::pi_cast<cl_kernel>(Kernel);
// Enable USM Indirect Access for Kernels
if (clusm->useCLUSM()) {
CHECK_OCL_CODE(clusm->setKernelExecInfo(
theKernel, CL_KERNEL_EXEC_INFO_INDIRECT_HOST_ACCESS_INTEL,
sizeof(cl_bool), &t));
CHECK_OCL_CODE(clusm->setKernelExecInfo(
theKernel, CL_KERNEL_EXEC_INFO_INDIRECT_DEVICE_ACCESS_INTEL,
sizeof(cl_bool), &t));
CHECK_OCL_CODE(clusm->setKernelExecInfo(
theKernel, CL_KERNEL_EXEC_INFO_INDIRECT_SHARED_ACCESS_INTEL,
sizeof(cl_bool), &t));

// This passes all the allocations we've tracked as SVM Pointers
CHECK_OCL_CODE(clusm->setKernelIndirectUSMExecInfo(
pi::pi_cast<cl_command_queue>(MQueue->getHandleRef()), theKernel));
} else if (clusm->isInitialized()) {
// Sanity check that nothing went wrong setting up clusm
CHECK_OCL_CODE(clSetKernelExecInfo(
theKernel, CL_KERNEL_EXEC_INFO_INDIRECT_HOST_ACCESS_INTEL,
sizeof(cl_bool), &t));
CHECK_OCL_CODE(clSetKernelExecInfo(
theKernel, CL_KERNEL_EXEC_INFO_INDIRECT_DEVICE_ACCESS_INTEL,
sizeof(cl_bool), &t));
CHECK_OCL_CODE(clSetKernelExecInfo(
theKernel, CL_KERNEL_EXEC_INFO_INDIRECT_SHARED_ACCESS_INTEL,
sizeof(cl_bool), &t));
if (pi::piUseBackend(pi::PiBackend::SYCL_BE_PI_OPENCL)) {
cl_context CLContext = pi::pi_cast<cl_context>(
detail::getSyclObjImpl(Context)->getHandleRef());
auto Clusm = GetCLUSM(CLContext);
if (usesUSM && Clusm) {
cl_bool TrueVal = CL_TRUE;
auto TheKernel = pi::pi_cast<cl_kernel>(Kernel);
// Enable USM Indirect Access for Kernels
if (Clusm->useCLUSM()) {
CHECK_OCL_CODE(Clusm->setKernelExecInfo(
TheKernel, CL_KERNEL_EXEC_INFO_INDIRECT_HOST_ACCESS_INTEL,
sizeof(cl_bool), &TrueVal));
CHECK_OCL_CODE(Clusm->setKernelExecInfo(
TheKernel, CL_KERNEL_EXEC_INFO_INDIRECT_DEVICE_ACCESS_INTEL,
sizeof(cl_bool), &TrueVal));
CHECK_OCL_CODE(Clusm->setKernelExecInfo(
TheKernel, CL_KERNEL_EXEC_INFO_INDIRECT_SHARED_ACCESS_INTEL,
sizeof(cl_bool), &TrueVal));

// This passes all the allocations we've tracked as SVM Pointers
CHECK_OCL_CODE(Clusm->setKernelIndirectUSMExecInfo(
pi::pi_cast<cl_command_queue>(MQueue->getHandleRef()),
TheKernel));
} else if (Clusm->isInitialized()) {
// Sanity check that nothing went wrong setting up clusm
CHECK_OCL_CODE(clSetKernelExecInfo(
TheKernel, CL_KERNEL_EXEC_INFO_INDIRECT_HOST_ACCESS_INTEL,
sizeof(cl_bool), &TrueVal));
CHECK_OCL_CODE(clSetKernelExecInfo(
TheKernel, CL_KERNEL_EXEC_INFO_INDIRECT_DEVICE_ACCESS_INTEL,
sizeof(cl_bool), &TrueVal));
CHECK_OCL_CODE(clSetKernelExecInfo(
TheKernel, CL_KERNEL_EXEC_INFO_INDIRECT_SHARED_ACCESS_INTEL,
sizeof(cl_bool), &TrueVal));
}
}
}

Expand Down
20 changes: 3 additions & 17 deletions sycl/source/detail/usm/clusm.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,33 +17,19 @@
#include <stdarg.h>
#include <time.h> // strdate

cl::sycl::detail::usm::CLUSM *gCLUSM = nullptr;
std::map<cl_context,cl::sycl::detail::usm::CLUSM *> gCLUSM;

namespace cl {
namespace sycl {
namespace detail {
namespace usm {

bool CLUSM::Create(CLUSM *&pCLUSM) {
pCLUSM = new CLUSM();
if (pCLUSM) {
return true;
}

return false;
}

void CLUSM::Delete(CLUSM *&pCLUSM) {
delete pCLUSM;
pCLUSM = nullptr;
}

void CLUSM::initExtensions(cl_platform_id platform) {
void CLUSM::initExtensions(cl_context context, cl_platform_id platform) {
// If OpenCL supports the USM Extension, don't enable CLUSM.
std::lock_guard<std::mutex> guard(mLock);

if (!mInitialized) {
mEnableCLUSM = !cliext::initializeExtensions(platform);
mEnableCLUSM = !cliext::initializeExtensions(context, platform);
mInitialized = true;
}
}
Expand Down
Loading