Skip to content

Commit fc60063

Browse files
SYCL Unbound Teamisaacault
SYCL Unbound Team
authored and
isaacault
committed
[SYCL][Bindless][2/4] Add experimental implementation of SYCL bindless images extension
This commit stands as the second commit of four to make code review easier, covering the changes made to the PI. Co-authored-by: Isaac Ault <isaac.ault@codeplay.com> Co-authored-by: Hugh Bird <hugh.bird@codeplay.com> Co-authored-by: Duncan Brawley <duncan.brawley@codeplay.com> Co-authored-by: Przemek Malon <przemek.malon@codeplay.com> Co-authored-by: Chedy Najjar <chedy.najjar@codeplay.com> Co-authored-by: Sean Stirling <sean.stirling@codeplay.com> Co-authored-by: Peter Zuzek <peter@codeplay.com> Implement revision 4 of the bindless images extension proposal: intel#9842
1 parent 14a9812 commit fc60063

27 files changed

+2855
-28
lines changed

‎sycl/include/sycl/detail/pi.def

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -126,6 +126,7 @@ _PI_API(piEnqueueMemUnmap)
126126
_PI_API(piextUSMHostAlloc)
127127
_PI_API(piextUSMDeviceAlloc)
128128
_PI_API(piextUSMSharedAlloc)
129+
_PI_API(piextUSMPitchedAlloc)
129130
_PI_API(piextUSMFree)
130131
_PI_API(piextUSMEnqueueMemset)
131132
_PI_API(piextUSMEnqueueMemcpy)
@@ -141,6 +142,27 @@ _PI_API(piextKernelSetArgSampler)
141142

142143
_PI_API(piextPluginGetOpaqueData)
143144

145+
// Bindless Images
146+
_PI_API(piextMemUnsampledImageHandleDestroy)
147+
_PI_API(piextMemSampledImageHandleDestroy)
148+
_PI_API(piextMemImageAllocate)
149+
_PI_API(piextMemImageFree)
150+
_PI_API(piextMemUnsampledImageCreate)
151+
_PI_API(piextMemSampledImageCreate)
152+
_PI_API(piextMemImageCopy)
153+
_PI_API(piextMemImageGetInfo)
154+
_PI_API(piextMemMipmapGetLevel)
155+
_PI_API(piextMemMipmapFree)
156+
157+
// Interop
158+
_PI_API(piextMemImportOpaqueFD)
159+
_PI_API(piextMemReleaseInterop)
160+
_PI_API(piextMemMapExternalArray)
161+
_PI_API(piextImportExternalSemaphoreOpaqueFD)
162+
_PI_API(piextDestroyExternalSemaphore)
163+
_PI_API(piextWaitExternalSemaphore)
164+
_PI_API(piextSignalExternalSemaphore)
165+
144166
_PI_API(piPluginGetLastError)
145167

146168
_PI_API(piTearDown)

‎sycl/include/sycl/detail/pi.h

Lines changed: 265 additions & 5 deletions
Large diffs are not rendered by default.

‎sycl/include/sycl/detail/pi.hpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -148,6 +148,13 @@ using PiMemObjectType = ::pi_mem_type;
148148
using PiMemImageChannelOrder = ::pi_image_channel_order;
149149
using PiMemImageChannelType = ::pi_image_channel_type;
150150
using PiKernelCacheConfig = ::pi_kernel_cache_config;
151+
using PiImageHandle = ::pi_image_handle;
152+
using PiImageMemHandle = ::pi_image_mem_handle;
153+
using PiImageCopyFlags = ::pi_image_copy_flags;
154+
using PiInteropMemHandle = ::pi_interop_mem_handle;
155+
using PiInteropSemaphoreHandle = ::pi_interop_semaphore_handle;
156+
using PiImageOffset = ::pi_image_offset_struct;
157+
using PiImageRegion = ::pi_image_region_struct;
151158

152159
__SYCL_EXPORT void contextSetExtendedDeleter(const sycl::context &constext,
153160
pi_context_extended_deleter func,

‎sycl/include/sycl/sampler.hpp

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -30,6 +30,11 @@ enum class filtering_mode : unsigned int {
3030
linear = PI_SAMPLER_FILTER_MODE_LINEAR
3131
};
3232

33+
enum class mipmap_filtering_mode : unsigned int {
34+
nearest = PI_SAMPLER_MIP_FILTER_MODE_NEAREST,
35+
linear = PI_SAMPLER_MIP_FILTER_MODE_LINEAR
36+
};
37+
3338
enum class coordinate_normalization_mode : unsigned int {
3439
normalized = 1,
3540
unnormalized = 0
@@ -69,6 +74,12 @@ class __SYCL_EXPORT __SYCL_SPECIAL_CLASS __SYCL_TYPE(sampler) sampler {
6974
addressing_mode addressingMode, filtering_mode filteringMode,
7075
const property_list &propList = {});
7176

77+
sampler(coordinate_normalization_mode normalizationMode,
78+
addressing_mode addressingMode, filtering_mode filteringMode,
79+
mipmap_filtering_mode mipmapFilteringMode, float minMipmapLevelClamp,
80+
float maxMipmapLevelClamp, float maxAnisotropy,
81+
const property_list &propList = {});
82+
7283
#ifdef __SYCL_INTERNAL_API
7384
sampler(cl_sampler clSampler, const context &syclContext);
7485
#endif
@@ -102,8 +113,16 @@ class __SYCL_EXPORT __SYCL_SPECIAL_CLASS __SYCL_TYPE(sampler) sampler {
102113

103114
filtering_mode get_filtering_mode() const;
104115

116+
mipmap_filtering_mode get_mipmap_filtering_mode() const;
117+
105118
coordinate_normalization_mode get_coordinate_normalization_mode() const;
106119

120+
float get_min_mipmap_level_clamp() const;
121+
122+
float get_max_mipmap_level_clamp() const;
123+
124+
float get_max_anisotropy() const;
125+
107126
private:
108127
#ifdef __SYCL_DEVICE_ONLY__
109128
detail::sampler_impl impl;

‎sycl/plugins/cuda/CMakeLists.txt

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -64,6 +64,8 @@ add_sycl_plugin(cuda
6464
"../unified_runtime/ur/adapters/cuda/enqueue.cpp"
6565
"../unified_runtime/ur/adapters/cuda/event.cpp"
6666
"../unified_runtime/ur/adapters/cuda/event.hpp"
67+
"../unified_runtime/ur/adapters/cuda/image.cpp"
68+
"../unified_runtime/ur/adapters/cuda/image.hpp"
6769
"../unified_runtime/ur/adapters/cuda/kernel.cpp"
6870
"../unified_runtime/ur/adapters/cuda/kernel.hpp"
6971
"../unified_runtime/ur/adapters/cuda/memory.cpp"

‎sycl/plugins/cuda/pi_cuda.cpp

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -160,6 +160,7 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
160160
_PI_CL(piextUSMHostAlloc, pi2ur::piextUSMHostAlloc)
161161
_PI_CL(piextUSMDeviceAlloc, pi2ur::piextUSMDeviceAlloc)
162162
_PI_CL(piextUSMSharedAlloc, pi2ur::piextUSMSharedAlloc)
163+
_PI_CL(piextUSMPitchedAlloc, pi2ur::piextUSMPitchedAlloc)
163164
_PI_CL(piextUSMFree, pi2ur::piextUSMFree)
164165
_PI_CL(piextUSMEnqueueMemset, pi2ur::piextUSMEnqueueMemset)
165166
_PI_CL(piextUSMEnqueueMemcpy, pi2ur::piextUSMEnqueueMemcpy)
@@ -186,6 +187,30 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
186187
_PI_CL(piGetDeviceAndHostTimer, pi2ur::piGetDeviceAndHostTimer)
187188
_PI_CL(piPluginGetBackendOption, pi2ur::piPluginGetBackendOption)
188189

190+
// Bindless Images
191+
_PI_CL(piextMemUnsampledImageHandleDestroy,
192+
pi2ur::piextMemUnsampledImageHandleDestroy)
193+
_PI_CL(piextMemSampledImageHandleDestroy,
194+
pi2ur::piextMemSampledImageHandleDestroy)
195+
_PI_CL(piextMemImageAllocate, pi2ur::piextMemImageAllocate)
196+
_PI_CL(piextMemImageFree, pi2ur::piextMemImageFree)
197+
_PI_CL(piextMemUnsampledImageCreate, pi2ur::piextMemUnsampledImageCreate)
198+
_PI_CL(piextMemSampledImageCreate, pi2ur::piextMemSampledImageCreate)
199+
_PI_CL(piextMemImageCopy, pi2ur::piextMemImageCopy)
200+
_PI_CL(piextMemImageGetInfo, pi2ur::piextMemImageGetInfo)
201+
202+
_PI_CL(piextMemMipmapGetLevel, pi2ur::piextMemMipmapGetLevel)
203+
_PI_CL(piextMemMipmapFree, pi2ur::piextMemMipmapFree)
204+
205+
_PI_CL(piextMemImportOpaqueFD, pi2ur::piextMemImportOpaqueFD)
206+
_PI_CL(piextMemReleaseInterop, pi2ur::piextMemReleaseInterop)
207+
_PI_CL(piextMemMapExternalArray, pi2ur::piextMemMapExternalArray)
208+
_PI_CL(piextImportExternalSemaphoreOpaqueFD,
209+
pi2ur::piextImportExternalSemaphoreOpaqueFD)
210+
_PI_CL(piextDestroyExternalSemaphore, pi2ur::piextDestroyExternalSemaphore)
211+
_PI_CL(piextWaitExternalSemaphore, pi2ur::piextWaitExternalSemaphore)
212+
_PI_CL(piextSignalExternalSemaphore, pi2ur::piextSignalExternalSemaphore)
213+
189214
#undef _PI_CL
190215

191216
return PI_SUCCESS;

‎sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1553,7 +1553,7 @@ pi_result piextEventCreateWithNativeHandle(pi_native_handle, pi_context, bool,
15531553
DIE_NO_IMPLEMENTATION;
15541554
}
15551555
pi_result piSamplerCreate(pi_context, const pi_sampler_properties *,
1556-
pi_sampler *) {
1556+
const float, const float, const float, pi_sampler *) {
15571557
DIE_NO_IMPLEMENTATION;
15581558
}
15591559

‎sycl/plugins/hip/pi_hip.hpp

Lines changed: 12 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -937,15 +937,24 @@ struct _pi_kernel {
937937
/// Implementation of samplers for HIP
938938
///
939939
/// Sampler property layout:
940-
/// | 31 30 ... 6 5 | 4 3 2 | 1 | 0 |
941-
/// | N/A | addressing mode | fiter mode | normalize coords |
940+
/// | <bits> | <usage>
941+
/// -----------------------------------
942+
/// | 31 30 ... 6 | N/A
943+
/// | 5 | mip filter mode
944+
/// | 4 3 2 | addressing mode
945+
/// | 1 | filter mode
946+
/// | 0 | normalize coords
942947
struct _pi_sampler {
943948
std::atomic_uint32_t refCount_;
944949
pi_uint32 props_;
950+
float minMipmapLevelClamp_;
951+
float maxMipmapLevelClamp_;
952+
float maxAnisotropy_;
945953
pi_context context_;
946954

947955
_pi_sampler(pi_context context)
948-
: refCount_(1), props_(0), context_(context) {}
956+
: refCount_(1), props_(0), minMipmapLevelClamp_(0.0f),
957+
maxMipmapLevelClamp_(0.0f), maxAnisotropy_(0.0f), context_(context) {}
949958

950959
pi_uint32 increment_reference_count() noexcept { return ++refCount_; }
951960

0 commit comments

Comments
 (0)