Skip to content
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.

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

Add support for SYCL on example 35 #142

Conversation

aacostadiaz
Copy link
Collaborator

This PR adds support for SYCL to the original Cutlass 2 example 35. The focus of this PR is only on functionality. Performance optimisations will be addressed in a separate PR to ensure that execution with SYCL can match the performance achieved with cuda.

@@ -159,6 +165,8 @@ struct Options {
/// Returns true if the environment and Toolkit support this
bool supported(bool verbose = true) const {

#if !defined(CUTLASS_ENABLE_SYCL)
Copy link
Collaborator

Choose a reason for hiding this comment

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

This is only temporary needed until we rework how __CUDACC_VER_MAJOR__ works for syclcompat, right? If so a code comment would be good so that we can find those easily in the future.

@@ -513,6 +525,10 @@ struct Testbed {
ElementCompute(0)
);

#if defined(CUTLASS_ENABLE_SYCL)
syclcompat::wait();
Copy link
Collaborator

Choose a reason for hiding this comment

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

why is that extra wait needed? Should probably be a code comment.

auto gemm_event = launch<cutlass::Kernel<GemmKernel>>(launch_policy{
sycl_grid, sycl_block, local_mem_size{static_cast<std::size_t>(gemm_smem_size)}},
params_.gemm);
EventManager::getInstance().addEvent(gemm_event);
Copy link
Collaborator

@rolandschulz rolandschulz Oct 16, 2024

Choose a reason for hiding this comment

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

why does the event need to be recorded? (a general question for each launch)

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

3 participants