Skip to content

Commit

Permalink
[SYCL][Bindless][E2E] Test normalized usm bindless images (#15299)
Browse files Browse the repository at this point in the history 
Modify pre-existing read_norm_types test to add testing for 2D USM
bindless images.

Associated UR PR which fixes a bug in normalized USM image creation:
oneapi-src/unified-runtime#2056
  • Loading branch information
@Seanst98
Seanst98 authored Oct 3, 2024
1 parent 4c223d1 commit 11353c6
Show file tree
Hide file tree
Showing 2 changed files with 57 additions and 19 deletions.
14 changes: 7 additions & 7 deletions sycl/cmake/modules/FetchUnifiedRuntime.cmake
View file
Original file line number Diff line number Diff line change
Expand Up @@ -117,14 +117,14 @@ if(SYCL_UR_USE_FETCH_CONTENT)
endfunction()

set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
# commit 00f958f375205fd86309f95b925141cf664ff955
# Merge: cc2d5909 98a67a2e
# commit 1f13d2ceb0494d84ce7b32f6b453dbb256fb702a
# Merge: 5276c534 bcf2244d
# Author: aarongreig <[email protected]>
# Date: Wed Oct 2 09:51:21 2024 +0100
# Merge pull request #2139 from nrspruit/zeHandle_copy_dependencies
# [L0] Pass and track event dependencies required before executing Memory
# Copy buffer inits
set(UNIFIED_RUNTIME_TAG 00f958f375205fd86309f95b925141cf664ff955)
# Date: Wed Oct 2 15:04:33 2024 +0100
# Merge pull request #2056 from Seanst98/sean/usm-normalized-fix
#
# [CUDA][Bindless] Address USM normalized type image creation failure and functionality
set(UNIFIED_RUNTIME_TAG 1f13d2ceb0494d84ce7b32f6b453dbb256fb702a)

set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES")
# Due to the use of dependentloadflag and no installer for UMF and hwloc we need
Expand Down
62 changes: 50 additions & 12 deletions sycl/test-e2e/bindless_images/read_norm_types.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,7 @@
#include <iostream>
#include <limits>
#include <sycl/detail/core.hpp>
#include <sycl/usm.hpp>

#include "helpers/common.hpp"
#include <sycl/ext/oneapi/bindless_images.hpp>
Expand All @@ -29,7 +30,7 @@ bool run_test(sycl::range<NDims> globalSize, sycl::range<NDims> localSize) {

std::vector<InputType> dataIn(numElems, InputType((DType)dtypeMaxVal));
std::vector<OutputType> dataOut(numElems);
std::vector<OutputType> expected(numElems, OutputType(1.f));
std::vector<OutputType> expected(numElems, OutputType(2.f));

try {

Expand All @@ -47,9 +48,30 @@ bool run_test(sycl::range<NDims> globalSize, sycl::range<NDims> localSize) {
sycl::coordinate_normalization_mode::normalized,
sycl::filtering_mode::nearest};

auto imgIn = syclexp::create_image(imgMemIn, sampler, descIn, q);
auto imgIn1 = syclexp::create_image(imgMemIn, sampler, descIn, q);
auto imgOut = syclexp::create_image(imgMemOut, descOut, q);

void *allocUSM = nullptr;
syclexp::image_mem_handle allocMem;
syclexp::sampled_image_handle imgIn2;

if constexpr (NDims == 2) {
size_t pitch = 0;
allocUSM = syclexp::pitched_alloc_device(&pitch, descIn, q);

if (allocUSM == nullptr) {
std::cerr << "Error allocating 2D USM memory!" << std::endl;
return false;
}
imgIn2 = syclexp::create_image(allocUSM, pitch, sampler, descIn, q);
q.ext_oneapi_copy(dataIn.data(), allocUSM, descIn, pitch);

} else {
allocMem = syclexp::alloc_image_mem(descIn, q);
imgIn2 = syclexp::create_image(allocMem, sampler, descIn, q);
q.ext_oneapi_copy(dataIn.data(), allocMem, descIn);
}

q.ext_oneapi_copy(dataIn.data(), imgMemIn, descIn);
q.wait_and_throw();

Expand All @@ -60,27 +82,35 @@ bool run_test(sycl::range<NDims> globalSize, sycl::range<NDims> localSize) {
if constexpr (NDims == 1) {
size_t dim0 = it.get_global_id(0);
float fdim0 = dim0 / globalSize[0];
OutputType pixel =
syclexp::sample_image<OutputType>(imgIn, fdim0);
syclexp::write_image(imgOut, int(dim0), pixel);
OutputType pixel1 =
syclexp::sample_image<OutputType>(imgIn1, fdim0);
OutputType pixel2 =
syclexp::sample_image<OutputType>(imgIn2, fdim0);
syclexp::write_image(imgOut, int(dim0), pixel1 + pixel2);
} else if constexpr (NDims == 2) {
size_t dim0 = it.get_global_id(0);
size_t dim1 = it.get_global_id(1);
float fdim0 = dim0 / globalSize[0];
float fdim1 = dim1 / globalSize[1];
OutputType pixel = syclexp::sample_image<OutputType>(
imgIn, sycl::float2(fdim0, fdim1));
syclexp::write_image(imgOut, sycl::int2(dim0, dim1), pixel);
OutputType pixel1 = syclexp::sample_image<OutputType>(
imgIn1, sycl::float2(fdim0, fdim1));
OutputType pixel2 = syclexp::sample_image<OutputType>(
imgIn2, sycl::float2(fdim0, fdim1));
syclexp::write_image(imgOut, sycl::int2(dim0, dim1),
pixel1 + pixel2);
} else if constexpr (NDims == 3) {
size_t dim0 = it.get_global_id(0);
size_t dim1 = it.get_global_id(1);
size_t dim2 = it.get_global_id(2);
float fdim0 = dim0 / globalSize[0];
float fdim1 = dim1 / globalSize[1];
float fdim2 = dim2 / globalSize[2];
OutputType pixel = syclexp::sample_image<OutputType>(
imgIn, sycl::float3(fdim0, fdim1, fdim2));
syclexp::write_image(imgOut, sycl::int3(dim0, dim1, dim2), pixel);
OutputType pixel1 = syclexp::sample_image<OutputType>(
imgIn1, sycl::float3(fdim0, fdim1, fdim2));
OutputType pixel2 = syclexp::sample_image<OutputType>(
imgIn2, sycl::float3(fdim0, fdim1, fdim2));
syclexp::write_image(imgOut, sycl::int3(dim0, dim1, dim2),
pixel1 + pixel2);
}
});
});
Expand All @@ -89,12 +119,20 @@ bool run_test(sycl::range<NDims> globalSize, sycl::range<NDims> localSize) {
q.ext_oneapi_copy(imgMemOut, dataOut.data(), descOut);
q.wait_and_throw();

syclexp::destroy_image_handle(imgIn, q);
syclexp::destroy_image_handle(imgIn1, q);
syclexp::destroy_image_handle(imgIn2, q);
syclexp::destroy_image_handle(imgOut, q);

syclexp::free_image_mem(imgMemIn, syclexp::image_type::standard, dev, ctxt);
syclexp::free_image_mem(imgMemOut, syclexp::image_type::standard, dev,
ctxt);

if constexpr (NDims == 2) {
sycl::free(allocUSM, ctxt);
} else {
syclexp::free_image_mem(allocMem, syclexp::image_type::standard, dev,
ctxt);
}
} catch (sycl::exception e) {
std::cerr << "SYCL exception caught! : " << e.what() << "\n";
return false;
Expand Down

0 comments on commit 11353c6

Please sign in to comment.