Skip to content
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

Jump to bottom

[SYCL][Bindless] Add experimental implementation of bindless images extension #9665

Closed
wants to merge 1 commit into from
Closed

[SYCL][Bindless] Add experimental implementation of bindless images extension #9665

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
1,510 changes: 1,510 additions & 0 deletions libclc/ptx-nvidiacl/libspirv/images/image.cl
View file
Open in desktop

Large diffs are not rendered by default.

291 changes: 291 additions & 0 deletions libclc/ptx-nvidiacl/libspirv/images/image_helpers.ll
View file
Open in desktop

Large diffs are not rendered by default.

4 changes: 4 additions & 0 deletions sycl/include/CL/__spirv/spirv_ops.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -258,6 +258,10 @@ template <typename SampledType, typename TempRetT, typename TempArgT>
extern __DPCPP_SYCL_EXTERNAL TempRetT
__spirv_ImageSampleExplicitLod(SampledType, TempArgT, int, float);

template <typename SampledType, typename TempRetT, typename TempArgT>
extern __DPCPP_SYCL_EXTERNAL TempRetT
__spirv_ImageSampleExplicitLod(SampledType, TempArgT, int, TempArgT, TempArgT);

#define __SYCL_OpGroupAsyncCopyGlobalToLocal __spirv_GroupAsyncCopy
#define __SYCL_OpGroupAsyncCopyLocalToGlobal __spirv_GroupAsyncCopy

Expand Down
57 changes: 57 additions & 0 deletions sycl/include/sycl/detail/cg.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -75,6 +75,9 @@ class CG {
CopyToDeviceGlobal = 19,
CopyFromDeviceGlobal = 20,
ReadWriteHostPipe = 21,
CopyImage = 22,
SemaphoreWait = 23,
SemaphoreSignal = 24,
};

struct StorageInitHelper {
Expand Down Expand Up @@ -493,6 +496,60 @@ class CGCopyFromDeviceGlobal : public CG {
size_t getOffset() { return MOffset; }
detail::OSModuleHandle getOSModuleHandle() { return MOSModuleHandle; }
};
/// "Copy Image" command group class.
class CGCopyImage : public CG {
void *MSrc;
void *MDst;
RT::PiMemImageDesc MImageDesc;
RT::PiMemImageFormat MImageFormat;
RT::PiImageCopyFlags MImageCopyFlags;

public:
CGCopyImage(void *Src, void *Dst, RT::PiMemImageDesc ImageDesc,
RT::PiMemImageFormat ImageFormat,
RT::PiImageCopyFlags ImageCopyFlags, CG::StorageInitHelper CGData,
detail::code_location loc = {})
: CG(CopyImage, std::move(CGData), std::move(loc)), MSrc(Src), MDst(Dst),
MImageDesc(ImageDesc), MImageFormat(ImageFormat),
MImageCopyFlags(ImageCopyFlags) {}

void *getSrc() const { return MSrc; }
void *getDst() const { return MDst; }
RT::PiMemImageDesc getDesc() const { return MImageDesc; }
RT::PiMemImageFormat getFormat() const { return MImageFormat; }
RT::PiImageCopyFlags getCopyFlags() const { return MImageCopyFlags; }
};

/// "Semaphore Wait" command group class.
class CGSemaphoreWait : public CG {
RT::PiInteropSemaphoreHandle MInteropSemaphoreHandle;

public:
CGSemaphoreWait(RT::PiInteropSemaphoreHandle InteropSemaphoreHandle,
CG::StorageInitHelper CGData, detail::code_location loc = {})
: CG(SemaphoreWait, std::move(CGData), std::move(loc)),
MInteropSemaphoreHandle(InteropSemaphoreHandle) {}

RT::PiInteropSemaphoreHandle getInteropSemaphoreHandle() const {
return MInteropSemaphoreHandle;
}
};

/// "Semaphore Signal" command group class.
class CGSemaphoreSignal : public CG {
RT::PiInteropSemaphoreHandle MInteropSemaphoreHandle;

public:
CGSemaphoreSignal(RT::PiInteropSemaphoreHandle InteropSemaphoreHandle,
CG::StorageInitHelper CGData,
detail::code_location loc = {})
: CG(SemaphoreSignal, std::move(CGData), std::move(loc)),
MInteropSemaphoreHandle(InteropSemaphoreHandle) {}

RT::PiInteropSemaphoreHandle getInteropSemaphoreHandle() const {
return MInteropSemaphoreHandle;
}
};

} // namespace detail
} // __SYCL_INLINE_VER_NAMESPACE(_V1)
Expand Down
50 changes: 50 additions & 0 deletions sycl/include/sycl/detail/image_ocl_types.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -83,6 +83,56 @@ static RetType __invoke__ImageRead(ImageT Img, CoordT Coords) {
return sycl::detail::convertDataToType<TempRetT, RetType>(Ret);
}

template <typename RetType, typename SmpImageT, typename CoordT>
static RetType __invoke__ImageReadLod(SmpImageT SmpImg, CoordT Coords,
float Level) {

// Convert from sycl types to builtin types to get correct function mangling.
using TempRetT = sycl::detail::ConvertToOpenCLType_t<RetType>;
using TempArgT = sycl::detail::ConvertToOpenCLType_t<CoordT>;

TempArgT TmpCoords =
sycl::detail::convertDataToType<CoordT, TempArgT>(Coords);

enum ImageOperands { Lod = 0x2 };

// OpImageSampleExplicitLod
// Its components must be the same as Sampled Type of the underlying
// OpTypeImage
// Sampled Image must be an object whose type is OpTypeSampledImage
// Image Operands encodes what operands follow. Either Lod
// or Grad image operands must be present
TempRetT Ret = __spirv_ImageSampleExplicitLod<SmpImageT, TempRetT, TempArgT>(
SmpImg, TmpCoords, ImageOperands::Lod, Level);
return sycl::detail::convertDataToType<TempRetT, RetType>(Ret);
}

template <typename RetType, typename SmpImageT, typename CoordT>
static RetType __invoke__ImageReadGrad(SmpImageT SmpImg, CoordT Coords,
CoordT Dx, CoordT Dy) {

// Convert from sycl types to builtin types to get correct function mangling.
using TempRetT = sycl::detail::ConvertToOpenCLType_t<RetType>;
using TempArgT = sycl::detail::ConvertToOpenCLType_t<CoordT>;

TempArgT TmpCoords =
sycl::detail::convertDataToType<CoordT, TempArgT>(Coords);
TempArgT TmpGraddX = sycl::detail::convertDataToType<CoordT, TempArgT>(Dx);
TempArgT TmpGraddY = sycl::detail::convertDataToType<CoordT, TempArgT>(Dy);

enum ImageOperands { Grad = 0x3 };

// OpImageSampleExplicitLod
// Its components must be the same as Sampled Type of the underlying
// OpTypeImage
// Sampled Image must be an object whose type is OpTypeSampledImage
// Image Operands encodes what operands follow. Either Lod
// or Grad image operands must be present
TempRetT Ret = __spirv_ImageSampleExplicitLod<SmpImageT, TempRetT, TempArgT>(
SmpImg, TmpCoords, ImageOperands::Grad, TmpGraddX, TmpGraddY);
return sycl::detail::convertDataToType<TempRetT, RetType>(Ret);
}

template <typename RetType, typename ImageT, typename CoordT>
static RetType __invoke__ImageReadSampler(ImageT Img, CoordT Coords,
const __ocl_sampler_t &Smpl) {
Expand Down
22 changes: 22 additions & 0 deletions sycl/include/sycl/detail/pi.def
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -126,6 +126,7 @@ _PI_API(piEnqueueMemUnmap)
_PI_API(piextUSMHostAlloc)
_PI_API(piextUSMDeviceAlloc)
_PI_API(piextUSMSharedAlloc)
_PI_API(piextUSMPitchedAlloc)
Copy link
Contributor

@alycm alycm May 31, 2023 •
edited
Loading

Choose a reason for hiding this comment

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

Discussed this internally with the team, but to make the point visible to all reviewers: As this PR changes PI then it will require Unified Runtime changes to match this.

Documentation on how to add experimental features is being written here oneapi-src/unified-runtime#546

You can ask us (@alycm, @jandres742, @kbenzie) any questions and we'll help you out.

You can also see in-progress work of another experimental change (not even a draft PR yet) here: Bensuo/unified-runtime#1. The YAML changes are where to look first, a lot of the rest is autogenerated from that.

_PI_API(piextUSMFree)
_PI_API(piextUSMEnqueueMemset)
_PI_API(piextUSMEnqueueMemcpy)
Expand All @@ -141,6 +142,27 @@ _PI_API(piextKernelSetArgSampler)

_PI_API(piextPluginGetOpaqueData)

// Bindless Images
_PI_API(piextMemUnsampledImageHandleDestroy)
_PI_API(piextMemSampledImageHandleDestroy)
_PI_API(piextMemImageAllocate)
_PI_API(piextMemImageFree)
_PI_API(piextMemUnsampledImageCreate)
_PI_API(piextMemSampledImageCreate)
_PI_API(piextMemImageCopy)
_PI_API(piextMemImageGetInfo)
_PI_API(piextMemMipmapGetLevel)
_PI_API(piextMemMipmapFree)

// Interop
_PI_API(piextMemImportOpaqueFD)
_PI_API(piextMemReleaseInterop)
_PI_API(piextMemMapExternalArray)
_PI_API(piextImportExternalSemaphoreOpaqueFD)
_PI_API(piextDestroyExternalSemaphore)
_PI_API(piextWaitExternalSemaphore)
_PI_API(piextSignalExternalSemaphore)

_PI_API(piPluginGetLastError)

_PI_API(piTearDown)
Expand Down
Loading