Skip to content

Commit f169009

Browse files
authored
[NFC][SYCL][Bindless] apply convert_color for non-nvptx backend as well (#12214)
Same as NVPTX, Intel GPU backend also needs converting a user type to a standard type, otherwise mangled builtin name is unrecognizable.
1 parent cff54c8 commit f169009

File tree

1 file changed

+3
-9
lines changed

1 file changed

+3
-9
lines changed

sycl/include/sycl/ext/oneapi/bindless_images.hpp

Lines changed: 3 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -696,9 +696,8 @@ template <typename CoordT> constexpr size_t coord_size() {
696696
}
697697
}
698698

699-
#if defined(__NVPTX__)
700-
// bit_cast Color to a type the NVPTX backend is known to accept
701-
template <typename DataT> constexpr auto convert_color_nvptx(DataT Color) {
699+
// bit_cast Color to a type the backend is known to accept
700+
template <typename DataT> constexpr auto convert_color(DataT Color) {
702701
constexpr size_t dataSize = sizeof(DataT);
703702
static_assert(
704703
dataSize == 1 || dataSize == 2 || dataSize == 4 || dataSize == 8 ||
@@ -717,7 +716,6 @@ template <typename DataT> constexpr auto convert_color_nvptx(DataT Color) {
717716
return sycl::bit_cast<sycl::vec<uint32_t, 4>>(Color);
718717
}
719718
}
720-
#endif
721719

722720
// assert coords or elements of coords is of an integer type
723721
template <typename CoordT> constexpr void assert_unsampled_coords() {
@@ -1048,18 +1046,14 @@ void write_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
10481046
"for 1D, 2D and 3D images, respectively.");
10491047

10501048
#ifdef __SYCL_DEVICE_ONLY__
1051-
#if defined(__NVPTX__)
10521049
if constexpr (detail::is_recognized_standard_type<DataT>()) {
10531050
__invoke__ImageWrite((uint64_t)imageHandle.raw_handle, coords, color);
10541051
} else {
10551052
// Convert DataT to a supported backend write type when user-defined type is
10561053
// passed
10571054
__invoke__ImageWrite((uint64_t)imageHandle.raw_handle, coords,
1058-
detail::convert_color_nvptx(color));
1055+
detail::convert_color(color));
10591056
}
1060-
#else
1061-
__invoke__ImageWrite((uint64_t)imageHandle.raw_handle, coords, color);
1062-
#endif
10631057
#else
10641058
assert(false); // Bindless images not yet implemented on host
10651059
#endif

0 commit comments

Comments
 (0)