Skip to content

[SYCL] Add new aspect ext_oneapi_virtual_functions #15577

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 3 commits into from
Oct 10, 2024
Merged

[SYCL] Add new aspect ext_oneapi_virtual_functions #15577

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
40ee3b9
[SYCL] Add new aspect ext_oneapi_virtual_functions
KornevNikita Oct 2, 2024
b457119
apply suggestions
KornevNikita Oct 7, 2024
f03ddbe
apply suggestions
KornevNikita Oct 8, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 3 additions & 1 deletion llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -85,6 +85,7 @@ def AspectExt_oneapi_virtual_mem : Aspect<"ext_oneapi_virtual_mem">;
def AspectExt_oneapi_cuda_cluster_group : Aspect<"ext_oneapi_cuda_cluster_group">;
def AspectExt_intel_fpga_task_sequence : Aspect<"ext_intel_fpga_task_sequence">;
def AspectExt_oneapi_atomic16 : Aspect<"ext_oneapi_atomic16">;
def AspectExt_oneapi_virtual_functions : Aspect<"ext_oneapi_virtual_functions">;
// Deprecated aspects
def AspectInt64_base_atomics : Aspect<"int64_base_atomics">;
def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics">;
Expand Down Expand Up @@ -148,7 +149,8 @@ def : TargetInfo<"__TestAspectList",
AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph, AspectExt_oneapi_private_alloca,
AspectExt_oneapi_queue_profiling_tag, AspectExt_oneapi_virtual_mem, AspectExt_oneapi_cuda_cluster_group,
AspectExt_intel_fpga_task_sequence,
AspectExt_oneapi_atomic16],
AspectExt_oneapi_atomic16,
AspectExt_oneapi_virtual_functions],
[]>;
// This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
// match.
Expand Down
10 changes: 10 additions & 0 deletions sycl/include/sycl/device_aspect_macros.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -395,6 +395,11 @@
#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_atomic16__ 0
#endif

#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_virtual_functions__
//__SYCL_ASPECT(ext_oneapi_virtual_functions, 81)
#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_virtual_functions__ 0
#endif

#ifndef __SYCL_ANY_DEVICE_HAS_host__
// __SYCL_ASPECT(host, 0)
#define __SYCL_ANY_DEVICE_HAS_host__ 0
Expand Down Expand Up @@ -779,3 +784,8 @@
//__SYCL_ASPECT(ext_oneapi_oneapi_atomic16, 80)
#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_atomic16__ 0
#endif

#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_virtual_functions__
//__SYCL_ASPECT(ext_oneapi_virtual_functions, 81)
#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_virtual_functions__ 0
#endif
1 change: 1 addition & 0 deletions sycl/include/sycl/info/aspects.def
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -71,3 +71,4 @@ __SYCL_ASPECT(ext_oneapi_unique_addressing_per_dim, 77)
__SYCL_ASPECT(ext_oneapi_bindless_images_sample_1d_usm, 78)
__SYCL_ASPECT(ext_oneapi_bindless_images_sample_2d_usm, 79)
__SYCL_ASPECT(ext_oneapi_atomic16, 80)
__SYCL_ASPECT(ext_oneapi_virtual_functions, 81)
7 changes: 7 additions & 0 deletions sycl/source/detail/device_impl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -768,6 +768,13 @@ bool device_impl::has(aspect Aspect) const {
// Likely L0 doesn't check it properly. Need to double-check.
return has_extension("cl_ext_float_atomics");
}
case aspect::ext_oneapi_virtual_functions: {
// TODO: move to UR like e.g. aspect::ext_oneapi_virtual_mem
backend BE = getBackend();
bool isCompatibleBE = BE == sycl::backend::ext_oneapi_level_zero ||
BE == sycl::backend::opencl;
return (is_cpu() || is_gpu()) && isCompatibleBE;
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is this something we'd want to move to UR eventually?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What is the criteria of moving things to UR?

The only HW which we don't expect to support function pointers is FPGA. For now we also exclude CUDA & HIP, because we haven't tested those configurations. However, I can imagine that 3rd-party backends may not support necessary SPIR-V extensions we use for virtual functions and therefore the ultimate query should be about SPIR-V extension support here.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

My rule of thumb is: If you have to check for the backend, it should have been done in UR. On a more general note though, my concern is another implementation of L0 or OpenCL could come in and it wouldn't support these, then this would be invalid. In a case like that, it would have to be addressed in the UR adapters anyway.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

My rule of thumb is: If you have to check for the backend, it should have been done in UR.

Good point. Even though I consider exclusion of CUDA & HIP backends temporary, I don't know if there are any other queries that should be performed for those backends (like checking CUDA SM version or something).

Do we have some generic API which we can extend (device info, I assume), i.e. what is the path of extending UR to support such query?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It would be a query like for aspect::ext_oneapi_virtual_mem a couple of lines above. I don't remember if UR has some tooling to help generate more of these enums. @kbenzie is there?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It would be a query like for aspect::ext_oneapi_virtual_mem a couple of lines above.

Thanks. Long-term, I think that we should probably move this into UR, but I'm not entirely sure on the name yet. We may add support for generic function pointers later and I believe that aspect for them will be checking for the very same things as we check for virtual functions. Simply because virtual functions are function pointers under the hood.

Therefore, it probably makes sense to have a single generic UR query for both (something like "indirect calls"). However, since we are not yet there with function pointers, I'm hesitant to suggest such generic name just yet.

If there are no objections, I would suggest to have this aspect handled at SYCL level for now and leave a TODO comment to move it into UR later.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It would be a query like for aspect::ext_oneapi_virtual_mem a couple of lines above. I don't remember if UR has some tooling to help generate more of these enums. @kbenzie is there?

This would entail adding an entry to the ur_device_info_t enum here and runnign the generate target to produce the header and validation layer source changes from that. Then it would be a case of implementing the query in each adapter.

If there are no objections, I would suggest to have this aspect handled at SYCL level for now and leave a TODO comment to move it into UR later.

👍

}
}

return false; // This device aspect has not been implemented yet.
Expand Down
3 changes: 3 additions & 0 deletions sycl/test-e2e/Basic/aspects.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -90,6 +90,9 @@ int main() {
if (plt.has(aspect::ext_oneapi_atomic16)) {
std::cout << " ext_oneapi_atomic16" << std::endl;
}
if (plt.has(aspect::ext_oneapi_virtual_functions)) {
std::cout << " ext_oneapi_virtual_functions" << std::endl;
}
}
std::cout << "Passed." << std::endl;
return 0;
Expand Down
3 changes: 0 additions & 3 deletions sycl/test-e2e/VirtualFunctions/2/1/1/missing-overrides.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,3 @@
// UNSUPPORTED: cuda, hip, acc
// FIXME: replace unsupported with an aspect check once we have it
//
// RUN: %{build} -o %t.out %helper-includes
// RUN: %{run} %t.out

Expand Down
3 changes: 0 additions & 3 deletions sycl/test-e2e/VirtualFunctions/2/1/1/more-complex-hierarchy.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,3 @@
// UNSUPPORTED: cuda, hip, acc
// FIXME: replace unsupported with an aspect check once we have it
//
// RUN: %{build} -o %t.out %helper-includes
// RUN: %{run} %t.out

Expand Down
3 changes: 0 additions & 3 deletions sycl/test-e2e/VirtualFunctions/2/1/1/simple-hierarchy.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,3 @@
// UNSUPPORTED: cuda, hip, acc
// FIXME: replace unsupported with an aspect check once we have it
//
// RUN: %{build} -o %t.out %helper-includes
// RUN: %{run} %t.out

Expand Down
3 changes: 0 additions & 3 deletions sycl/test-e2e/VirtualFunctions/2/2/single-construct-single-use.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,3 @@
// UNSUPPORTED: cuda, hip, acc
// FIXME: replace unsupported with an aspect check once we have it
//
// RUN: %{build} -o %t.out %helper-includes
// RUN: %{run} %t.out

Expand Down
1 change: 1 addition & 0 deletions sycl/test-e2e/VirtualFunctions/lit.local.cfg
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,3 +4,4 @@ import os
# paths like "../../../helper.hpp" in them, so let's just register a
# substitution to add directory with helper headers into include search path
config.substitutions.append(("%helper-includes", "-I {}".format(os.path.dirname(os.path.abspath(__file__)))))
config.required_features += ['aspect-ext_oneapi_virtual_functions']
Loading