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

Enabed Complex data type for Gemm #462

Merged

Conversation

OuadiElfarouki
Copy link
Collaborator

This PR extends support to sycl::complex<float/double> data types to GEMM operators (gemm & batched extensions).
Main changes :

  • Added a wrapper around sycl::complex data type to be used within fully & partially vectorized gemm kernels to mimic sycl::vec (as we don't have an implementation for it in DPCPP yet).
  • Changes so far rely on the experimental support of sycl::complex in sycl/llvm branch, which is used through the extension header : ext/oneapi/experimental/sycl_complex.hpp
  • Gemm configurations for complex type are put separately from float/double types as they might expect custom optimal values (to be tuned later once benchmarking is ready)
  • Added tests for GEMM & batched derivatives, tested passing on intel, amd and Nvidia GPUs as well as intel CPU.

common/include/common/common_utils.hpp Outdated Show resolved Hide resolved
include/blas_meta.h Outdated Show resolved Hide resolved
Copy link
Collaborator

@s-Nick s-Nick left a comment

Choose a reason for hiding this comment

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

LGTM!

@@ -186,5 +186,149 @@ struct Packetize {
}
};

#ifdef BLAS_ENABLE_COMPLEX
Copy link
Collaborator

Choose a reason for hiding this comment

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

I think it's best if we move this to a separate file to avoid confusion between the normal load/store.

Copy link
Collaborator Author

Choose a reason for hiding this comment

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

Moved the complex section to gemm_load_store_complex.hpp

}

if (batch > 1 && batch_type == gemm_batch_type_t::interleaved) {
Copy link
Collaborator

Choose a reason for hiding this comment

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

We could move this if condition to before the call to reference_blas to avoid running the cgemm unnecessarily.

Copy link
Collaborator Author

Choose a reason for hiding this comment

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

Yes done!

@OuadiElfarouki OuadiElfarouki force-pushed the complex_gemm_pr branch 3 times, most recently from c8b7fbe to 622318e Compare October 17, 2023 09:39
@muhammad-tanvir-1211 muhammad-tanvir-1211 merged commit 19b0fed into codeplaysoftware:master Oct 24, 2023
3 checks passed
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