diff --git a/SYCL/Plugin/interop-level-zero-image-get-native-mem.cpp b/SYCL/Plugin/interop-level-zero-image-get-native-mem.cpp new file mode 100644 index 0000000000..aedadcf77d --- /dev/null +++ b/SYCL/Plugin/interop-level-zero-image-get-native-mem.cpp @@ -0,0 +1,106 @@ +// REQUIRES: level_zero, level_zero_dev_kit +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out 2>&1 | FileCheck %s + +// we use the interop to get the native image handle and then use that to make a +// new image and enumerate the pixels. + +// CHECK: (0 0) -- { 0 0 0 0 } +// CHECK-NEXT: (1 0) -- { 1 1 1 1 } +// CHECK-NEXT: (2 0) -- { 2 2 2 2 } +// CHECK-NEXT: (3 0) -- { 3 3 3 3 } +// CHECK-NEXT: (0 1) -- { 4 4 4 4 } +// CHECK-NEXT: (1 1) -- { 5 5 5 5 } +// CHECK-NEXT: (2 1) -- { 6 6 6 6 } +// CHECK-NEXT: (3 1) -- { 7 7 7 7 } + +// clang++ -fsycl -o las.bin -I$SYCL_HOME/build/install/include/sycl -lze_loader +// interop-level-zero-image-get-native-mem.cpp + +#include +#include +#include +using namespace sycl; + +int main() { +#ifdef SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO + constexpr auto BE = sycl::backend::ext_oneapi_level_zero; + sycl::device D = + sycl::ext::oneapi::filter_selector("level_zero:gpu").select_device(); + + sycl::context Ctx{D}; + sycl::queue Q(Ctx, D); + auto ZeContext = sycl::get_native(Ctx); + auto ZeDevice = sycl::get_native(D); + + // ----------- IMAGE STUFF + using pixelT = sycl::uint4; // accessor + using ChannelDataT = std::uint8_t; // allocator + constexpr long width = 4; + constexpr long height = 2; + constexpr long numPixels = width * height; + ChannelDataT *sourceData = + (ChannelDataT *)std::calloc(numPixels * 4, sizeof(ChannelDataT)); + // initialize data: [ (0 0 0 0) (1 1 1 1) ...] + for (size_t i = 0; i < numPixels; i++) { + for (size_t chan = 0; chan < 4; chan++) { + size_t idx = (i * 4) + chan; + sourceData[idx] = (ChannelDataT)i; + } + } + // 8 bits per channel, four per pixel. + sycl::image_channel_order ChanOrder = sycl::image_channel_order::rgba; + sycl::image_channel_type ChanType = sycl::image_channel_type::unsigned_int8; + + const sycl::range<2> ImgRange_2D(width, height); + { // closure + // 1 - Create simple image. + sycl::image<2> image_2D(sourceData, ChanOrder, ChanType, ImgRange_2D); + + // 2 - Grab it's image handle via the get_native_mem interop. + using nativeH = sycl::backend_return_t>; + sycl::buffer passBack(range<1>{1}); + + Q.submit([&](handler &cgh) { + auto image_acc = + image_2D.get_access(cgh); + auto passBackAcc = passBack.get_host_access(sycl::write_only); + cgh.host_task([=](const interop_handle &IH) { + // There is nothing with image handles in the L0 API except + // create and destroy. So let's do that. + auto ZeImageH = IH.get_native_mem(image_acc); + passBackAcc[0] = ZeImageH; + }); + }).wait(); + + // Now we have the ZeImageH, so let's make a new SYCL image from it. + auto passBackAcc = passBack.get_host_access(sycl::read_only); + nativeH ZeImageH = passBackAcc[0]; + sycl::backend_input_t> imageData{ + ZeImageH, ChanOrder, ChanType, ImgRange_2D, + sycl::ext::oneapi::level_zero::ownership::keep}; + sycl::image<2> NewImg = sycl::make_image(imageData, Ctx); + + // Then use that image to read and stream out the data. + Q.submit([&](handler &cgh) { + auto read_acc = NewImg.get_access(cgh); + sycl::stream out(2024, 400, cgh); + cgh.single_task([=]() { + for (unsigned y = 0; y < height; y++) { + for (unsigned x = 0; x < width; x++) { + auto location = sycl::int2{x, y}; + pixelT somePixel = read_acc.read(location); + out << "(" << x << " " << y << ") -- { " << somePixel[0] << " " + << somePixel[1] << " " << somePixel[2] << " " << somePixel[3] + << " }" << sycl::endl; + } + } + }); + }).wait(); + } // ~image + +#else + std::cout << "Missing Level-Zero backend. Test skipped." << std::endl; +#endif + return 0; +} \ No newline at end of file diff --git a/SYCL/Plugin/interop-level-zero-image-ownership.cpp b/SYCL/Plugin/interop-level-zero-image-ownership.cpp new file mode 100644 index 0000000000..329ac64a47 --- /dev/null +++ b/SYCL/Plugin/interop-level-zero-image-ownership.cpp @@ -0,0 +1,137 @@ +// REQUIRES: level_zero, level_zero_dev_kit +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out +// RUN: env ZE_DEBUG=1 %GPU_RUN_PLACEHOLDER %t.out 2>&1 | FileCheck %s + +// This test verifies that ownership is working correctly. +// If ownership is ::transfer then the ~image destructor will end up calling +// zeImageDestroy +// CHECK: test ownership::transfer +// CHECK: ZE ---> zeImageDestroy + +// With ownership ::keep it is must be called manually. +// CHECK: test ownership::keep +// CHECK: zeImageDestroy MANUAL + +// No other calls should appear. +// CHECK-NOT: zeImageDestroy + +// clang++ -fsycl -o wfd.bin -I$SYCL_HOME/build/install/include/sycl -lze_loader +// interop-level-zero-image-ownership.cpp + +#include +#include +#include + +using namespace sycl; + +void test(sycl::ext::oneapi::level_zero::ownership Ownership) { + + constexpr auto BE = sycl::backend::ext_oneapi_level_zero; + + platform Plt{gpu_selector_v}; + + auto Devices = Plt.get_devices(); + + if (Devices.size() < 1) { + std::cout << "Devices not found" << std::endl; + return; + } + + device Device = Devices[0]; + context Context{Device}; + queue Queue{Context, Device}; + + // Get native Level Zero handles + auto ZeContext = get_native(Context); + auto ZeDevice = get_native(Device); + + // ----------- Image Fundamentals + using pixelT = sycl::uint4; // accessor + using ChannelDataT = std::uint8_t; // allocator + sycl::image_channel_order ChanOrder = sycl::image_channel_order::rgba; + sycl::image_channel_type ChanType = sycl::image_channel_type::unsigned_int8; + constexpr uint32_t numChannels = 4; // L0 only supports RGBA at this time. + + constexpr uint32_t width = 8; + constexpr uint32_t height = 4; + constexpr uint32_t depth = 1; + + const sycl::range<2> ImgRange_2D(width, height); + + // ----------- Basic LevelZero Description + ze_image_format_type_t ZeImageFormatType = ZE_IMAGE_FORMAT_TYPE_UINT; + size_t ZeImageFormatTypeSize = 8; + ze_image_format_layout_t ZeImageFormatLayout = ZE_IMAGE_FORMAT_LAYOUT_8_8_8_8; + ze_image_format_t ZeFormatDesc = { + ZeImageFormatLayout, ZeImageFormatType, + ZE_IMAGE_FORMAT_SWIZZLE_R, ZE_IMAGE_FORMAT_SWIZZLE_G, + ZE_IMAGE_FORMAT_SWIZZLE_B, ZE_IMAGE_FORMAT_SWIZZLE_A}; + + ze_image_desc_t ZeImageDesc_base; + ZeImageDesc_base.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; + ZeImageDesc_base.pNext = nullptr; + ZeImageDesc_base.flags = ZE_IMAGE_FLAG_KERNEL_WRITE; + // ZeImageDesc_base.flags = 0; + ZeImageDesc_base.arraylevels = 0; + ZeImageDesc_base.miplevels = 0; + ZeImageDesc_base.format = ZeFormatDesc; + + { + // ------ 2D ------ + ze_image_desc_t ZeImageDesc_2D = ZeImageDesc_base; + ZeImageDesc_2D.type = ZE_IMAGE_TYPE_2D; + ZeImageDesc_2D.width = width; + ZeImageDesc_2D.height = height; + ZeImageDesc_2D.depth = 1; + + ze_image_handle_t ZeHImage_2D; + ze_result_t res = + zeImageCreate(ZeContext, ZeDevice, &ZeImageDesc_2D, &ZeHImage_2D); + if (res != ZE_RESULT_SUCCESS) { + std::cout << "unable to create image " << res << std::endl; + return; + } + + { // closure + sycl::backend_input_t> ImageInteropInput_2D{ + ZeHImage_2D, ChanOrder, ChanType, ImgRange_2D, Ownership}; + auto Image_2D = sycl::make_image(ImageInteropInput_2D, Context); + + Queue.submit([&](sycl::handler &cgh) { + auto write_acc = + Image_2D.get_access(cgh); + + cgh.parallel_for(ImgRange_2D, [=](sycl::item<2> Item) { + auto location = sycl::int2{Item[0], Item[1]}; + auto sum = Item[0] + Item[1]; + const pixelT somepixel = {sum, sum, sum, sum}; + write_acc.write(location, somepixel); + }); + }); + Queue.wait_and_throw(); + + } // ~image + // if ownership was transfer, then the ZeHImage_2D was destroyed as part of + // the ~image destruction (or deferred) + + if (Ownership == sycl::ext::oneapi::level_zero::ownership::keep) { + zeImageDestroy(ZeHImage_2D); + std::cout << "zeImageDestroy MANUAL" << std::endl; + } + + } // closure +} + +int main() { +#ifdef SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO + std::cout << "test ownership::transfer" << std::endl; + test(sycl::ext::oneapi::level_zero::ownership::transfer); + + std::cout << "test ownership::keep" << std::endl; + test(sycl::ext::oneapi::level_zero::ownership::keep); +#else + std::cout << "Missing Level-Zero backend. Test skipped." << std::endl; +#endif + std::cout << "chau" << std::endl; + return 0; +} \ No newline at end of file diff --git a/SYCL/Plugin/interop-level-zero-image.cpp b/SYCL/Plugin/interop-level-zero-image.cpp new file mode 100644 index 0000000000..4c392b1e03 --- /dev/null +++ b/SYCL/Plugin/interop-level-zero-image.cpp @@ -0,0 +1,220 @@ +// REQUIRES: level_zero, level_zero_dev_kit +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +// This test verifies that make_image is working for 1D, 2D and 3D images. +// We instantiate an image with L0, set its body, then use a host accessor to +// verify that the pixels are set correctly. + +// clang++ -fsycl -o ilzi.bin -I$SYCL_HOME/build/install/include/sycl +// -lze_loader interop-level-zero-image.cpp + +#include +#include +#include + +using namespace sycl; + +int main() { +#ifdef SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO + constexpr auto BE = sycl::backend::ext_oneapi_level_zero; + + platform Plt{gpu_selector_v}; + + auto Devices = Plt.get_devices(); + + if (Devices.size() < 1) { + std::cout << "Devices not found" << std::endl; + return 0; + } + + device Device = Devices[0]; + context Context{Device}; + queue Queue{Context, Device}; + + // Get native Level Zero handles + auto ZeContext = get_native(Context); + auto ZeDevice = get_native(Device); + + // ----------- Image Fundamentals + using pixelT = sycl::uint4; // accessor + using ChannelDataT = std::uint8_t; // allocator + sycl::image_channel_order ChanOrder = sycl::image_channel_order::rgba; + sycl::image_channel_type ChanType = sycl::image_channel_type::unsigned_int8; + constexpr uint32_t numChannels = 4; // L0 only supports RGBA at this time. + + constexpr uint32_t width = 8; + constexpr uint32_t height = 4; + constexpr uint32_t depth = 2; + + const sycl::range<1> ImgRange_1D(width); + const sycl::range<2> ImgRange_2D(width, height); + const sycl::range<3> ImgRange_3D(width, height, depth); + + // ----------- Basic LevelZero Description + ze_image_format_type_t ZeImageFormatType = ZE_IMAGE_FORMAT_TYPE_UINT; + size_t ZeImageFormatTypeSize = 8; + ze_image_format_layout_t ZeImageFormatLayout = ZE_IMAGE_FORMAT_LAYOUT_8_8_8_8; + ze_image_format_t ZeFormatDesc = { + ZeImageFormatLayout, ZeImageFormatType, + ZE_IMAGE_FORMAT_SWIZZLE_R, ZE_IMAGE_FORMAT_SWIZZLE_G, + ZE_IMAGE_FORMAT_SWIZZLE_B, ZE_IMAGE_FORMAT_SWIZZLE_A}; + + ze_image_desc_t ZeImageDesc_base; + ZeImageDesc_base.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; + ZeImageDesc_base.pNext = nullptr; + ZeImageDesc_base.flags = ZE_IMAGE_FLAG_KERNEL_WRITE; + // ZeImageDesc_base.flags = 0; // <-- for read only + ZeImageDesc_base.arraylevels = 0; + ZeImageDesc_base.miplevels = 0; + ZeImageDesc_base.format = ZeFormatDesc; + + // ------ 1D ------ + { + std::cout << "glorious 1D" << std::endl; + // 1D image + ze_image_desc_t ZeImageDesc_1D = ZeImageDesc_base; + ZeImageDesc_1D.type = ZE_IMAGE_TYPE_1D; + ZeImageDesc_1D.width = width; + ZeImageDesc_1D.height = 1; + ZeImageDesc_1D.depth = 1; + + ze_image_handle_t ZeHImage_1D; + zeImageCreate(ZeContext, ZeDevice, &ZeImageDesc_1D, &ZeHImage_1D); + + { // closure + sycl::backend_input_t> ImageInteropInput_1D{ + ZeHImage_1D, ChanOrder, ChanType, ImgRange_1D, + sycl::ext::oneapi::level_zero::ownership::keep}; + auto Image_1D = sycl::make_image(ImageInteropInput_1D, Context); + + Queue.submit([&](sycl::handler &cgh) { + auto write_acc = + Image_1D.get_access(cgh); + + cgh.parallel_for(ImgRange_1D, [=](sycl::item<1> Item) { + int x = Item[0]; + const pixelT somePixel = {x, x, x, x}; + write_acc.write(x, somePixel); + }); + }); + Queue.wait_and_throw(); + + // now check with host accessor. + auto read_acc = Image_1D.get_access(); + for (int col = 0; col < width; col++) { + const pixelT somePixel = read_acc.read(col); + // const pixelT expectedPixel = {col,col,col,col}; + // assert(somePixel == expectedPixel); + assert(somePixel[0] == col && somePixel[1] == col && + somePixel[2] == col && somePixel[3] == col); + } + + } // ~image + } // closure + + { + // ------ 2D ------ + std::cout << "glorious 2D" << std::endl; + // 2D image + ze_image_desc_t ZeImageDesc_2D = ZeImageDesc_base; + ZeImageDesc_2D.type = ZE_IMAGE_TYPE_2D; + ZeImageDesc_2D.width = width; + ZeImageDesc_2D.height = height; + ZeImageDesc_2D.depth = 1; + + ze_image_handle_t ZeHImage_2D; + zeImageCreate(ZeContext, ZeDevice, &ZeImageDesc_2D, &ZeHImage_2D); + + { // closure + sycl::backend_input_t> ImageInteropInput_2D{ + ZeHImage_2D, ChanOrder, ChanType, ImgRange_2D, + sycl::ext::oneapi::level_zero::ownership::keep}; + auto Image_2D = sycl::make_image(ImageInteropInput_2D, Context); + + Queue.submit([&](sycl::handler &cgh) { + auto write_acc = + Image_2D.get_access(cgh); + + cgh.parallel_for(ImgRange_2D, [=](sycl::item<2> Item) { + auto location = sycl::int2{Item[0], Item[1]}; + auto sum = Item[0] + Item[1]; + const pixelT somepixel = {sum, sum, sum, sum}; + write_acc.write(location, somepixel); + }); + }); + Queue.wait_and_throw(); + + // now check with host accessor. + auto read_acc = Image_2D.get_access(); + for (int row = 0; row < height; row++) { + for (int col = 0; col < width; col++) { + auto location = sycl::int2{col, row}; + const pixelT somePixel = read_acc.read(location); + auto sum = col + row; + // const pixelT expectedPixel = {sum,sum,sum,sum}; + // assert(somePixel == expectedPixel); + assert(somePixel[0] == sum && somePixel[1] == sum && + somePixel[2] == sum && somePixel[3] == sum); + } + } + + } // ~image + } // closure + + { + // ------ 3D ------ + std::cout << "glorious 3D" << std::endl; + // 3D image + ze_image_desc_t ZeImageDesc_3D = ZeImageDesc_base; + ZeImageDesc_3D.type = ZE_IMAGE_TYPE_3D; + ZeImageDesc_3D.width = width; + ZeImageDesc_3D.height = height; + ZeImageDesc_3D.depth = depth; + + ze_image_handle_t ZeHImage_3D; + zeImageCreate(ZeContext, ZeDevice, &ZeImageDesc_3D, &ZeHImage_3D); + + { // closure + sycl::backend_input_t> ImageInteropInput_3D{ + ZeHImage_3D, ChanOrder, ChanType, ImgRange_3D, + sycl::ext::oneapi::level_zero::ownership::keep}; + auto Image_3D = sycl::make_image(ImageInteropInput_3D, Context); + + Queue.submit([&](sycl::handler &cgh) { + auto write_acc = + Image_3D.get_access(cgh); + + cgh.parallel_for(ImgRange_3D, [=](sycl::item<3> Item) { + auto location = sycl::int4{Item[0], Item[1], Item[2], 0}; + auto sum = Item[0] + Item[1] + Item[2]; + const pixelT somepixel = {sum, sum, sum, sum}; + write_acc.write(location, somepixel); + }); + }); + Queue.wait_and_throw(); + + // now check with host accessor. + auto read_acc = Image_3D.get_access(); + for (int row = 0; row < height; row++) { + for (int col = 0; col < width; col++) { + for (int z = 0; z < depth; z++) { + auto location = sycl::int4{col, row, z, 0}; + const pixelT somePixel = read_acc.read(location); + auto sum = col + row + z; + // const pixelT expectedPixel = {sum,sum,sum,sum}; + // assert(somePixel == expectedPixel); + assert(somePixel[0] == sum && somePixel[1] == sum && + somePixel[2] == sum && somePixel[3] == sum); + } + } + } + + } // ~image + } // closure + +#else + std::cout << "Missing Level-Zero backend. Test skipped." << std::endl; +#endif + return 0; +} \ No newline at end of file