Skip to content

Fix device-agnostic fp32 GEMM compile on oneAPI 2025.2+ #852

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

Open
ssheorey wants to merge 3 commits into
base: main
Choose a base branch
from
Open

Fix device-agnostic fp32 GEMM compile on oneAPI 2025.2+ #852

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
7636f4c
Fix device-agnostic GEMM SYCL compile on oneAPI 2025.2+
ssheorey Aug 21, 2026
6eb9c78
Merge branch 'main' into fix/device-agnostic-fp32-gemm-oneapi
tdeng5 Sep 22, 2026
2d466f4
Merge branch 'main' into fix/device-agnostic-fp32-gemm-oneapi
tdeng5 Sep 22, 2026
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
2 changes: 1 addition & 1 deletion include/cutlass/arch/cache_operation.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,7 +42,7 @@ namespace arch {

/// Controls PTX cache operations
struct CacheOperation {
enum Kind {
enum Kind : unsigned char {
/// Cache at all levels - accessed again
Always,
/// Cache at global level
Expand Down
32 changes: 32 additions & 0 deletions include/cutlass/arch/memory.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -109,6 +109,8 @@ struct global_load<AccessType,
: "l"(ptr), "r"((int)pred_guard), "r"(data[0].x), "r"(data[0].y),
"r"(data[0].z), "r"(data[0].w), "r"(data[1].x), "r"(data[1].y),
"r"(data[1].z), "r"(data[1].w), "l"(((uint8_t *)ptr) + 16));
#elif defined(__SYCL_DEVICE_ONLY__)
if (pred_guard) D = *(reinterpret_cast<AccessType const *>(ptr));
#else
CUTE_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
Expand Down Expand Up @@ -145,6 +147,8 @@ struct global_load<AccessType,
: "l"(ptr), "r"((int)pred_guard), "r"(data[0].x), "r"(data[0].y),
"r"(data[0].z), "r"(data[0].w), "r"(data[1].x), "r"(data[1].y),
"r"(data[1].z), "r"(data[1].w), "l"(((uint8_t *)ptr) + 16));
#elif defined(__SYCL_DEVICE_ONLY__)
if (pred_guard) D = *(reinterpret_cast<AccessType const *>(ptr));
#else
CUTE_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
Expand Down Expand Up @@ -177,6 +181,8 @@ struct global_load<AccessType,
"}\n"
: "=r"(data.x), "=r"(data.y), "=r"(data.z), "=r"(data.w)
: "l"(ptr), "r"((int)pred_guard), "r"(data.x), "r"(data.y), "r"(data.z), "r"(data.w));
#elif defined(__SYCL_DEVICE_ONLY__)
if (pred_guard) D = *(reinterpret_cast<AccessType const *>(ptr));
#else
CUTE_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
Expand Down Expand Up @@ -205,6 +211,8 @@ struct global_load<AccessType,
"}\n"
: "=r"(data.x), "=r"(data.y), "=r"(data.z), "=r"(data.w)
: "l"(ptr), "r"((int)pred_guard), "r"(data.x), "r"(data.y), "r"(data.z), "r"(data.w));
#elif defined(__SYCL_DEVICE_ONLY__)
if (pred_guard) D = *(reinterpret_cast<AccessType const *>(ptr));
#else
CUTE_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
Expand Down Expand Up @@ -235,6 +243,8 @@ struct global_load<AccessType,
"}\n"
: "=r"(data.x), "=r"(data.y)
: "l"(ptr), "r"((int)pred_guard), "r"(data.x), "r"(data.y));
#elif defined(__SYCL_DEVICE_ONLY__)
if (pred_guard) D = *(reinterpret_cast<AccessType const *>(ptr));
#else
CUTE_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
Expand All @@ -261,6 +271,8 @@ struct global_load<AccessType,
"}\n"
: "=r"(data.x), "=r"(data.y)
: "l"(ptr), "r"((int)pred_guard), "r"(data.x), "r"(data.y));
#elif defined(__SYCL_DEVICE_ONLY__)
if (pred_guard) D = *(reinterpret_cast<AccessType const *>(ptr));
#else
CUTE_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
Expand Down Expand Up @@ -290,6 +302,8 @@ struct global_load<AccessType,
"}\n"
: "=r"(data)
: "l"(ptr), "r"((int)pred_guard), "r"(data));
#elif defined(__SYCL_DEVICE_ONLY__)
if (pred_guard) D = *(reinterpret_cast<AccessType const *>(ptr));
#else
CUTE_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
Expand All @@ -315,6 +329,8 @@ struct global_load<AccessType,
"}\n"
: "=r"(data)
: "l"(ptr), "r"((int)pred_guard), "r"(data));
#elif defined(__SYCL_DEVICE_ONLY__)
if (pred_guard) D = *(reinterpret_cast<AccessType const *>(ptr));
#else
CUTE_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
Expand Down Expand Up @@ -344,6 +360,8 @@ struct global_load<AccessType,
"}\n"
: "=h"(data)
: "l"(ptr), "r"((int)pred_guard), "h"(data));
#elif defined(__SYCL_DEVICE_ONLY__)
if (pred_guard) D = *(reinterpret_cast<AccessType const *>(ptr));
#else
CUTE_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
Expand All @@ -369,6 +387,8 @@ struct global_load<AccessType,
"}\n"
: "=h"(data)
: "l"(ptr), "r"((int)pred_guard), "h"(data));
#elif defined(__SYCL_DEVICE_ONLY__)
if (pred_guard) D = *(reinterpret_cast<AccessType const *>(ptr));
#else
CUTE_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
Expand Down Expand Up @@ -427,6 +447,8 @@ struct global_store<AccessType, 64> {
"r"(data[2].x), "r"(data[2].y), "r"(data[2].z), "r"(data[2].w),
"l"(((uint8_t *)ptr) + 48),
"r"(data[3].x), "r"(data[3].y), "r"(data[3].z), "r"(data[3].w));
#elif defined(__SYCL_DEVICE_ONLY__)
if (pred_guard) *(reinterpret_cast<AccessType *>(ptr)) = D;
#else
CUTE_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
Expand All @@ -452,6 +474,8 @@ struct global_store<AccessType, 32> {
: "l"(ptr), "r"(data[0].x), "r"(data[0].y), "r"(data[0].z),
"r"(data[0].w), "r"((int)pred_guard), "l"(((uint8_t *)ptr) + 16),
"r"(data[1].x), "r"(data[1].y), "r"(data[1].z), "r"(data[1].w));
#elif defined(__SYCL_DEVICE_ONLY__)
if (pred_guard) *(reinterpret_cast<AccessType *>(ptr)) = D;
#else
CUTE_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
Expand All @@ -473,6 +497,8 @@ struct global_store<AccessType, 16> {
"}\n"
:
: "l"(ptr), "r"(data.x), "r"(data.y), "r"(data.z), "r"(data.w), "r"((int)pred_guard));
#elif defined(__SYCL_DEVICE_ONLY__)
if (pred_guard) *(reinterpret_cast<AccessType *>(ptr)) = D;
#else
CUTE_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
Expand All @@ -494,6 +520,8 @@ struct global_store<AccessType, 8> {
"}\n"
:
: "l"(ptr), "r"(data.x), "r"(data.y), "r"((int)pred_guard));
#elif defined(__SYCL_DEVICE_ONLY__)
if (pred_guard) *(reinterpret_cast<AccessType *>(ptr)) = D;
#else
CUTE_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
Expand All @@ -515,6 +543,8 @@ struct global_store<AccessType, 4> {
"}\n"
:
: "l"(ptr), "r"(data), "r"((int)pred_guard));
#elif defined(__SYCL_DEVICE_ONLY__)
if (pred_guard) *(reinterpret_cast<AccessType *>(ptr)) = D;
#else
CUTE_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
Expand All @@ -536,6 +566,8 @@ struct global_store<AccessType, 2> {
"}\n"
:
: "l"(ptr), "h"(data), "r"((int)pred_guard));
#elif defined(__SYCL_DEVICE_ONLY__)
if (pred_guard) *(reinterpret_cast<AccessType *>(ptr)) = D;
#else
CUTE_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
Expand Down
2 changes: 1 addition & 1 deletion include/cutlass/arch/mma.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -206,7 +206,7 @@ struct Mma<gemm::GemmShape<1, 1, 1>, 1, ElementA, LayoutA, ElementB, LayoutB, El

/// Specifies internal data type for computation
struct SPFormatType {
enum Kind {
enum Kind : unsigned char {
Thread
};
};
Expand Down
2 changes: 1 addition & 1 deletion include/cutlass/epilogue/thread/scale_type.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,7 +49,7 @@ namespace thread {
/// 1. Scalar means alpha/beta is a single value from host(constant param) or device memory.
/// 2. Vector means alpha/beta is a vector always from device memory.
struct ScaleType {
enum Kind {
enum Kind : unsigned char {
Default, // D = scalar_alpha x Acc + scalar_beta x C
NoBetaScaling, // D = scalar_alpha x Acc + C
OnlyAlphaScaling, // D = scalar_alpha x Acc
Expand Down
7 changes: 5 additions & 2 deletions include/cutlass/gemm/device/gemm_universal_adapter.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -605,10 +605,13 @@ class GemmUniversalAdapter<
sycl_grid, sycl_block, launch_props, kernel_props
};
#if defined(CUTLASS_SYCL_PROFILING_ENABLED)
auto event = compat::experimental::launch<device_kernel<GemmKernel>, GemmKernel>(policy, q, params);
auto event = compat::experimental::launch<device_kernel<GemmKernel>>(

@taozha2 taozha2 Sep 22, 2026

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

I saw compile error here, this API maybe different for different compiler, please check this.

policy, q, params);
EventManager::getInstance().addEvent(event);
#else
compat::experimental::launch<device_kernel<GemmKernel>, GemmKernel, false>(policy, q, params);
compat::experimental::launch<device_kernel<GemmKernel>,
sycl::detail::auto_name, false>(policy,
q, params);
#endif
#endif // !defined(SYCL_EXT_ONEAPI_WORK_GROUP_SCRATCH_MEMORY)
#else
Expand Down
Loading