Skip to content

Commit

Permalink
[SYCLomatic] Support migration for 4 surface write APIs with bindless…
Browse files Browse the repository at this point in the history 
… APIs. (#2517)


Signed-off-by: Chen, Sheng S <sheng.s.chen@intel.com>
  • Loading branch information
@ShengchenJ
ShengchenJ authored Dec 12, 2024
1 parent 7a4b708 commit c018d33
Show file tree
Hide file tree
Showing 8 changed files with 212 additions and 56 deletions.
7 changes: 0 additions & 7 deletions clang/lib/DPCT/AnalysisInfo.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -4927,13 +4927,6 @@ void DeviceFunctionDecl::emplaceReplacement() {
if (Obj) {
Obj->merge(FuncInfo->getTextureObject((Obj->getParamIdx())));
if (DpctGlobalInfo::useExtBindlessImages()) {
DpctGlobalInfo::getInstance().addReplacement(
std::make_shared<ExtReplacement>(
Obj->getFilePath(), Obj->getOffset(),
strlen("cudaTextureObject_t"),
MapNames::getClNamespace() +
"ext::oneapi::experimental::sampled_image_handle",
nullptr));
continue;
}
if (!Obj->getType()) {
Expand Down
12 changes: 8 additions & 4 deletions clang/lib/DPCT/RuleInfra/MapNames.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -519,7 +519,7 @@ void MapNames::setExplicitNamespaceMap(
std::make_shared<TypeNameRule>(
DpctGlobalInfo::useExtBindlessImages()
? getClNamespace() +
"ext::oneapi::experimental::sampled_image_handle"
"ext::oneapi::experimental::unsampled_image_handle"
: getDpctNamespace() + "image_wrapper_base_p",
HelperFeatureEnum::device_ext)},
{"textureReference",
Expand Down Expand Up @@ -650,9 +650,13 @@ void MapNames::setExplicitNamespaceMap(
std::make_shared<TypeNameRule>(getClNamespace() + "image_channel_type")},
{"CUarray_format_enum",
std::make_shared<TypeNameRule>(getClNamespace() + "image_channel_type")},
{"CUtexObject", std::make_shared<TypeNameRule>(
getDpctNamespace() + "image_wrapper_base_p",
HelperFeatureEnum::device_ext)},
{"CUtexObject",
std::make_shared<TypeNameRule>(
DpctGlobalInfo::useExtBindlessImages()
? getClNamespace() +
"ext::oneapi::experimental::sampled_image_handle"
: getDpctNamespace() + "image_wrapper_base_p",
HelperFeatureEnum::device_ext)},
{"CUDA_RESOURCE_DESC",
std::make_shared<TypeNameRule>(getDpctNamespace() + "image_data",
HelperFeatureEnum::device_ext)},
Expand Down
68 changes: 65 additions & 3 deletions clang/lib/DPCT/RulesLang/APINamesTexture.inc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -341,13 +341,75 @@ FEATURE_REQUEST_FACTORY(
MEMBER_CALL(ARG(0), true, "set_channel_type", ARG(1)),
MEMBER_CALL(ARG(0), true, "set_channel_num", ARG(2))))

CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
CALL_FACTORY_ENTRY("surf1Dwrite",
CALL(MapNames::getClNamespace() +
"ext::oneapi::experimental::write_image",
ARG(1),
BO(BinaryOperatorKind::BO_Div, CALL("int", ARG(2)),
CALL("sizeof", ARG(0))),
ARG(0))),
UNSUPPORT_FACTORY_ENTRY("surf1Dwrite",
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
ARG("surf1Dwrite"),
ARG("--use-experimental-features=bindless_images")))

CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
CALL_FACTORY_ENTRY("surf2Dwrite",
CALL(MapNames::getClNamespace() +
"ext::oneapi::experimental::write_image",
ARG(1),
CALL(MapNames::getClNamespace() + "int2",
BO(BinaryOperatorKind::BO_Div, ARG(2),
CALL("sizeof", ARG(0))),
ARG(3)),
ARG(0))),
UNSUPPORT_FACTORY_ENTRY("surf2Dwrite",
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
ARG("surf2Dwrite"),
ARG("--use-experimental-features=bindless_images")))

CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
CALL_FACTORY_ENTRY("surf2DLayeredwrite",
CALL(MapNames::getClNamespace() +
"ext::oneapi::experimental::write_image_array",
ARG(1),
CALL(MapNames::getClNamespace() + "int2",
BO(BinaryOperatorKind::BO_Div, ARG(2),
CALL("sizeof", ARG(0))),
ARG(3)),
ARG(4), ARG(0))),
UNSUPPORT_FACTORY_ENTRY("surf2DLayeredwrite",
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
ARG("surf2DLayeredwrite"),
ARG("--use-experimental-features=bindless_images")))

CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
CALL_FACTORY_ENTRY("surf3Dwrite",
CALL(MapNames::getClNamespace() +
"ext::oneapi::experimental::write_image",
ARG(1),
CALL(MapNames::getClNamespace() + "int3",
BO(BinaryOperatorKind::BO_Div, ARG(2),
CALL("sizeof", ARG(0))),
ARG(3), ARG(4)),
ARG(0))),
UNSUPPORT_FACTORY_ENTRY("surf3Dwrite",
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
ARG("surf3Dwrite"),
ARG("--use-experimental-features=bindless_images")))

CONDITIONAL_FACTORY_ENTRY(
UseExtBindlessImages,
FEATURE_REQUEST_FACTORY(
HelperFeatureEnum::device_ext,
ENTRY_TEXTURE("surf1Dread", 0x01,
MapNames::getLibraryHelperNamespace() +
"experimental::sample_image_by_byte",
"experimental::fetch_image_by_byte",
1)),
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
ENTRY_TEXTURE("surf1Dread", 0x01, "read_byte", 1)))
Expand All @@ -358,7 +420,7 @@ CONDITIONAL_FACTORY_ENTRY(
HelperFeatureEnum::device_ext,
ENTRY_TEXTURE("surf2Dread", 0x02,
MapNames::getLibraryHelperNamespace() +
"experimental::sample_image_by_byte",
"experimental::fetch_image_by_byte",
1, 2)),
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
ENTRY_TEXTURE("surf2Dread", 0x02, "read_byte", 1,
Expand All @@ -369,7 +431,7 @@ CONDITIONAL_FACTORY_ENTRY(
HelperFeatureEnum::device_ext,
ENTRY_TEXTURE("surf3Dread", 0x03,
MapNames::getLibraryHelperNamespace() +
"experimental::sample_image_by_byte",
"experimental::fetch_image_by_byte",
1, 2, 3)),
FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
ENTRY_TEXTURE("surf3Dread", 0x03, "read_byte", 1, 2,
Expand Down
3 changes: 3 additions & 0 deletions clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -139,6 +139,9 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
TAI = TAL[0];
}
std::string VecTypeName = "float";
if(Source.find("surf") != std::string::npos) {
VecTypeName = "int";
}
if (getDim() != 1)
VecTypeName =
MapNames::getClNamespace() + VecTypeName + std::to_string(getDim());
Expand Down
25 changes: 10 additions & 15 deletions clang/lib/DPCT/RulesLang/RulesLangTexture.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -502,10 +502,6 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
"cudaTextureObject_t", "cudaSurfaceObject_t", "CUtexObject"))))))
.bind("texObj"),
this);
MF.addMatcher(typeLoc(loc(qualType(hasDeclaration(typedefDecl(hasAnyName(
"cudaTextureObject_t", "CUtexObject"))))))
.bind("texObj"),
this);
MF.addMatcher(
memberExpr(
hasObjectExpression(hasType(type(hasUnqualifiedDesugaredType(
Expand Down Expand Up @@ -560,6 +556,10 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
"tex1Dfetch",
"tex1DLayered",
"tex2DLayered",
"surf1Dwrite",
"surf2Dwrite",
"surf3Dwrite",
"surf2DLayeredwrite",
"surf1Dread",
"surf2Dread",
"surf3Dread",
Expand Down Expand Up @@ -986,20 +986,15 @@ void TextureRule::runRule(const MatchFinder::MatchResult &Result) {
return;
}
if (auto FD = DpctGlobalInfo::getParentFunction(TL)) {
if (FD->hasAttr<CUDAGlobalAttr>() || FD->hasAttr<CUDADeviceAttr>()) {
return;
}
} else if (auto VD = DpctGlobalInfo::findAncestor<VarDecl>(TL)) {
if (!VD->hasGlobalStorage()) {
if ((FD->hasAttr<CUDAGlobalAttr>() || FD->hasAttr<CUDADeviceAttr>()) &&
!DpctGlobalInfo::useExtBindlessImages()) {
return;
}
}
emplaceTransformation(new ReplaceToken(
TL->getBeginLoc(), TL->getEndLoc(),
DpctGlobalInfo::useExtBindlessImages()
? MapNames::getClNamespace() +
"ext::oneapi::experimental::sampled_image_handle"
: MapNames::getDpctNamespace() + "image_wrapper_base_p"));
ExprAnalysis A;
A.analyze(*TL);
emplaceTransformation(A.getReplacement());
A.applyAllSubExprRepl();
requestFeature(HelperFeatureEnum::device_ext);
}
}
Expand Down
8 changes: 4 additions & 4 deletions clang/lib/DPCT/SrcAPI/APINames.inc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1485,15 +1485,15 @@ ENTRY(texCubemapLayered, texCubemapLayered, false, NO_FLAG, P4, "comment")
ENTRY(texCubemapLayeredLod, texCubemapLayeredLod, false, NO_FLAG, P4, "comment")
ENTRY(tex2Dgather, tex2Dgather, false, NO_FLAG, P0, "comment")
ENTRY(surf1Dread, surf1Dread, true, NO_FLAG, P4, "Successful")
ENTRY(surf1Dwrite, surf1Dwrite, false, NO_FLAG, P0, "comment")
ENTRY(surf1Dwrite, surf1Dwrite, true, NO_FLAG, P0, "Successful")
ENTRY(surf2Dread, surf2Dread, true, NO_FLAG, P4, "Successful")
ENTRY(surf2Dwrite, surf2Dwrite, false, NO_FLAG, P0, "comment")
ENTRY(surf2Dwrite, surf2Dwrite, true, NO_FLAG, P0, "Successful")
ENTRY(surf3Dread, surf3Dread, true, NO_FLAG, P4, "Successful")
ENTRY(surf3Dwrite, surf3Dwrite, false, NO_FLAG, P0, "comment")
ENTRY(surf3Dwrite, surf3Dwrite, true, NO_FLAG, P0, "Successful")
ENTRY(surf1DLayeredread, surf1DLayeredread, false, NO_FLAG, P4, "comment")
ENTRY(surf1DLayeredwrite, surf1DLayeredwrite, false, NO_FLAG, P4, "comment")
ENTRY(surf2DLayeredread, surf2DLayeredread, false, NO_FLAG, P4, "comment")
ENTRY(surf2DLayeredwrite, surf2DLayeredwrite, false, NO_FLAG, P0, "comment")
ENTRY(surf2DLayeredwrite, surf2DLayeredwrite, true, NO_FLAG, P0, "Successful")
ENTRY(surfCubemapread, surfCubemapread, false, NO_FLAG, P4, "comment")
ENTRY(surfCubemapwrite, surfCubemapwrite, false, NO_FLAG, P4, "comment")
ENTRY(surfCubemapLayeredread, surfCubemapLayeredread, false, NO_FLAG, P4, "comment")
Expand Down
100 changes: 86 additions & 14 deletions clang/runtime/dpct-rt/include/dpct/bindless_images.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -528,9 +528,7 @@ class external_mem_wrapper : public external_mem_wrapper_base {

namespace detail {
struct sampled_image_handle_compare {
bool
operator()(sycl::ext::oneapi::experimental::sampled_image_handle L,
sycl::ext::oneapi::experimental::sampled_image_handle R) const {
template <class T> bool operator()(T L, T R) const {
return L.raw_handle < R.raw_handle;
}
};
Expand All @@ -544,6 +542,14 @@ inline std::pair<image_data, sampling_info> &get_img_info_map(
return img_info_map[handle];
}

inline std::pair<image_data, sampling_info> &get_img_info_map(
const sycl::ext::oneapi::experimental::unsampled_image_handle handle) {
static std::map<sycl::ext::oneapi::experimental::unsampled_image_handle,
std::pair<image_data, sampling_info>,
sampled_image_handle_compare>
img_info_map;
return img_info_map[handle];
}
inline image_mem_wrapper *&get_img_mem_map(
const sycl::ext::oneapi::experimental::sampled_image_handle handle) {
static std::map<sycl::ext::oneapi::experimental::sampled_image_handle,
Expand All @@ -552,6 +558,13 @@ inline image_mem_wrapper *&get_img_mem_map(
return img_mem_map[handle];
}

inline image_mem_wrapper *&get_img_mem_map(
const sycl::ext::oneapi::experimental::unsampled_image_handle handle) {
static std::map<sycl::ext::oneapi::experimental::unsampled_image_handle,
image_mem_wrapper *, sampled_image_handle_compare>
img_mem_map;
return img_mem_map[handle];
}
static inline size_t
get_ele_size(const sycl::ext::oneapi::experimental::image_descriptor &decs) {
size_t channel_size;
Expand Down Expand Up @@ -848,7 +861,7 @@ inline void unmap_resources(int count, external_mem_wrapper **handles,
/// \param [in] q The queue where the image creation be executed.
/// \returns The sampled image handle of created bindless image.
static inline sycl::ext::oneapi::experimental::sampled_image_handle
create_bindless_image(image_data data, sampling_info info = {},
create_bindless_image(image_data data, sampling_info info,
sycl::queue q = get_default_queue()) {
auto samp = sycl::ext::oneapi::experimental::bindless_image_sampler(
info.get_addressing_mode(), info.get_coordinate_normalization_mode(),
Expand Down Expand Up @@ -914,12 +927,69 @@ create_bindless_image(image_data data, sampling_info info = {},
return sycl::ext::oneapi::experimental::sampled_image_handle();
}

/// Create bindless image according to image data.
/// \param [in] data The image data used to create bindless image.
/// \param [in] q The queue where the image creation be executed.
/// \returns The sampled image handle of created bindless image.
static inline sycl::ext::oneapi::experimental::unsampled_image_handle
create_bindless_image(image_data data, sycl::queue q = get_default_queue()) {
switch (data.get_data_type()) {
case image_data_type::linear: {
// TODO: Use pointer to create image when bindless image support.
auto mem = new image_mem_wrapper(
data.get_channel(), data.get_x() / data.get_channel().get_total_size());
auto img = sycl::ext::oneapi::experimental::create_image(
mem->get_handle(), mem->get_desc(), q);
detail::get_img_mem_map(img) = mem;
auto ptr = data.get_data_ptr();
#ifdef DPCT_USM_LEVEL_NONE
q.ext_oneapi_copy(get_buffer(ptr).get_host_access().get_pointer(),
mem->get_handle(), mem->get_desc())
.wait();
#else
q.ext_oneapi_copy(ptr, mem->get_handle(), mem->get_desc()).wait();
#endif
return img;
}
case image_data_type::pitch: {
auto mem =
new image_mem_wrapper(data.get_channel(), data.get_x(), data.get_y());
auto img = sycl::ext::oneapi::experimental::create_image(
mem->get_handle(), mem->get_desc(), q);
detail::get_img_mem_map(img) = mem;
#ifdef DPCT_USM_LEVEL_NONE
q.ext_oneapi_copy(
get_buffer(data.get_data_ptr()).get_host_access().get_pointer(),
mem->get_handle(), mem->get_desc())
.wait();
#else
q.ext_oneapi_copy(data.get_data_ptr(), mem->get_handle(), mem->get_desc())
.wait();

#endif
return img;
}
case image_data_type::matrix: {
const auto mem = static_cast<image_mem_wrapper *>(data.get_data_ptr());
auto img = sycl::ext::oneapi::experimental::create_image(
mem->get_handle(), mem->get_desc(), q);
return img;
}
default:
throw std::runtime_error(
"Unsupported image_data_type in create_bindless_image!");
break;
}
// Must not reach here.
return sycl::ext::oneapi::experimental::unsampled_image_handle();
}

/// Destroy bindless image.
/// \param [in] handle The bindless image should be destroyed.
/// \param [in] q The queue where the image destruction be executed.
static inline void destroy_bindless_image(
sycl::ext::oneapi::experimental::sampled_image_handle handle,
sycl::queue q = get_default_queue()) {
template <class T>
static inline void destroy_bindless_image(T handle,
sycl::queue q = get_default_queue()) {
auto &mem = detail::get_img_mem_map(handle);
if (mem) {
delete mem;
Expand All @@ -928,11 +998,14 @@ static inline void destroy_bindless_image(
sycl::ext::oneapi::experimental::destroy_image_handle(handle, q);
}



/// Get the image data according to sampled image handle.
/// \param [in] handle The bindless image handle.
/// \returns The image data of sampled image.
template <class T>
static inline image_data
get_data(const sycl::ext::oneapi::experimental::sampled_image_handle handle) {
get_data(const T handle) {
return detail::get_img_info_map(handle).first;
}

Expand Down Expand Up @@ -1355,18 +1428,17 @@ static inline void dpct_memcpy(image_mem_wrapper *dest, size_t w_offset_dest,
dpct_memcpy(dest, w_offset_dest, h_offset_dest, temp, s, q);
sycl::free(temp, q);
}

// A wrapper for sycl sample_image function for the byte addressing image.
// A wrapper for sycl fetch_image function for the byte addressing image.
template <typename DataT, typename HintT = DataT, typename CoordT>
DataT sample_image_by_byte(
const sycl::ext::oneapi::experimental::sampled_image_handle &imageHandle,
DataT fetch_image_by_byte(
const sycl::ext::oneapi::experimental::unsampled_image_handle &imageHandle,
CoordT &&coords) {
if constexpr (std::is_scalar_v<CoordT>) {
return sycl::ext::oneapi::experimental::sample_image<DataT, HintT, CoordT>(
return sycl::ext::oneapi::experimental::fetch_image<DataT, HintT, CoordT>(
imageHandle, coords / sizeof(DataT));
} else {
coords[0] = coords[0] / sizeof(DataT);
return sycl::ext::oneapi::experimental::sample_image<DataT, HintT, CoordT>(
return sycl::ext::oneapi::experimental::fetch_image<DataT, HintT, CoordT>(
imageHandle, coords);
}
}
Expand Down
Loading

0 comments on commit c018d33

Please sign in to comment.