Skip to content

Commit

Permalink
Formatting
Browse files Browse the repository at this point in the history
  • Loading branch information
@joeatodd
joeatodd committed Dec 16, 2024
1 parent 6438826 commit 5cbb157
Show file tree
Hide file tree
Showing 5 changed files with 86 additions and 47 deletions.
12 changes: 10 additions & 2 deletions include/cute/arch/xe_copy_1B.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -37,13 +37,21 @@
#ifdef __SYCL_DEVICE_ONLY__
#define SYCL_DEVICE_BUILTIN(x) SYCL_EXTERNAL extern "C" x
#else
#define SYCL_DEVICE_BUILTIN(x) inline x { CUTE_INVALID_CONTROL_PATH("Attempting to use a device built-in in host code.");}
#define SYCL_DEVICE_BUILTIN(x) \
inline x { \
CUTE_INVALID_CONTROL_PATH( \
"Attempting to use a device built-in in host code."); \
}
#endif

#ifdef __SYCL_DEVICE_ONLY__
#define SYCL_DEVICE_OCL(x) SYCL_EXTERNAL x
#else
#define SYCL_DEVICE_OCL(x) inline x { CUTE_INVALID_CONTROL_PATH("Attempting to use a device built-in in host code.");}
#define SYCL_DEVICE_OCL(x) \
inline x { \
CUTE_INVALID_CONTROL_PATH( \
"Attempting to use a device built-in in host code."); \
}
#endif

using namespace cute;
Expand Down
12 changes: 10 additions & 2 deletions include/cute/arch/xe_copy_2B.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -38,13 +38,21 @@
#ifdef __SYCL_DEVICE_ONLY__
#define SYCL_DEVICE_BUILTIN(x) SYCL_EXTERNAL extern "C" x
#else
#define SYCL_DEVICE_BUILTIN(x) inline x { CUTE_INVALID_CONTROL_PATH("Attempting to use a device built-in in host code.");}
#define SYCL_DEVICE_BUILTIN(x) \
inline x { \
CUTE_INVALID_CONTROL_PATH( \
"Attempting to use a device built-in in host code."); \
}
#endif

#ifdef __SYCL_DEVICE_ONLY__
#define SYCL_DEVICE_OCL(x) SYCL_EXTERNAL x
#else
#define SYCL_DEVICE_OCL(x) inline x { CUTE_INVALID_CONTROL_PATH("Attempting to use a device built-in in host code.");}
#define SYCL_DEVICE_OCL(x) \
inline x { \
CUTE_INVALID_CONTROL_PATH( \
"Attempting to use a device built-in in host code."); \
}
#endif

using namespace cute;
Expand Down
12 changes: 10 additions & 2 deletions include/cute/arch/xe_copy_4B.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -37,13 +37,21 @@
#ifdef __SYCL_DEVICE_ONLY__
#define SYCL_DEVICE_BUILTIN(x) SYCL_EXTERNAL extern "C" x
#else
#define SYCL_DEVICE_BUILTIN(x) inline x { CUTE_INVALID_CONTROL_PATH("Attempting to use a device built-in in host code.");}
#define SYCL_DEVICE_BUILTIN(x) \
inline x { \
CUTE_INVALID_CONTROL_PATH( \
"Attempting to use a device built-in in host code."); \
}
#endif

#ifdef __SYCL_DEVICE_ONLY__
#define SYCL_DEVICE_OCL(x) SYCL_EXTERNAL x
#else
#define SYCL_DEVICE_OCL(x) inline x { CUTE_INVALID_CONTROL_PATH("Attempting to use a device built-in in host code.");}
#define SYCL_DEVICE_OCL(x) \
inline x { \
CUTE_INVALID_CONTROL_PATH( \
"Attempting to use a device built-in in host code."); \
}
#endif

enum class CacheControl {
Expand Down
12 changes: 10 additions & 2 deletions include/cute/arch/xe_copy_8B.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -37,13 +37,21 @@
#ifdef __SYCL_DEVICE_ONLY__
#define SYCL_DEVICE_BUILTIN(x) SYCL_EXTERNAL extern "C" x
#else
#define SYCL_DEVICE_BUILTIN(x) inline x { CUTE_INVALID_CONTROL_PATH("Attempting to use a device built-in in host code.");}
#define SYCL_DEVICE_BUILTIN(x) \
inline x { \
CUTE_INVALID_CONTROL_PATH( \
"Attempting to use a device built-in in host code."); \
}
#endif

#ifdef __SYCL_DEVICE_ONLY__
#define SYCL_DEVICE_OCL(x) SYCL_EXTERNAL x
#else
#define SYCL_DEVICE_OCL(x) inline x { CUTE_INVALID_CONTROL_PATH("Attempting to use a device built-in in host code.");}
#define SYCL_DEVICE_OCL(x) \
inline x { \
CUTE_INVALID_CONTROL_PATH( \
"Attempting to use a device built-in in host code."); \
}
#endif

using namespace cute;
Expand Down
85 changes: 46 additions & 39 deletions include/cutlass/arch/memory.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -109,7 +109,8 @@ struct global_load<AccessType,
"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));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}
};
Expand Down Expand Up @@ -144,7 +145,8 @@ struct global_load<AccessType,
"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));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}
};
Expand Down Expand Up @@ -175,7 +177,8 @@ struct global_load<AccessType,
: "=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));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}
};
Expand All @@ -202,7 +205,8 @@ struct global_load<AccessType,
: "=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));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}
};
Expand Down Expand Up @@ -231,7 +235,8 @@ struct global_load<AccessType,
: "=r"(data.x), "=r"(data.y)
: "l"(ptr), "r"((int)pred_guard), "r"(data.x), "r"(data.y));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}
};
Expand All @@ -256,7 +261,8 @@ struct global_load<AccessType,
: "=r"(data.x), "=r"(data.y)
: "l"(ptr), "r"((int)pred_guard), "r"(data.x), "r"(data.y));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}
};
Expand Down Expand Up @@ -284,7 +290,8 @@ struct global_load<AccessType,
: "=r"(data)
: "l"(ptr), "r"((int)pred_guard), "r"(data));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}
};
Expand All @@ -308,7 +315,8 @@ struct global_load<AccessType,
: "=r"(data)
: "l"(ptr), "r"((int)pred_guard), "r"(data));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}
};
Expand Down Expand Up @@ -336,7 +344,8 @@ struct global_load<AccessType,
: "=h"(data)
: "l"(ptr), "r"((int)pred_guard), "h"(data));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}
};
Expand All @@ -360,7 +369,8 @@ struct global_load<AccessType,
: "=h"(data)
: "l"(ptr), "r"((int)pred_guard), "h"(data));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}
};
Expand Down Expand Up @@ -417,7 +427,8 @@ struct global_store<AccessType, 64> {
"l"(((uint8_t *)ptr) + 48),
"r"(data[3].x), "r"(data[3].y), "r"(data[3].z), "r"(data[3].w));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}
};
Expand All @@ -441,7 +452,8 @@ struct global_store<AccessType, 32> {
"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));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}
};
Expand All @@ -461,7 +473,8 @@ struct global_store<AccessType, 16> {
:
: "l"(ptr), "r"(data.x), "r"(data.y), "r"(data.z), "r"(data.w), "r"((int)pred_guard));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}
};
Expand All @@ -481,7 +494,8 @@ struct global_store<AccessType, 8> {
:
: "l"(ptr), "r"(data.x), "r"(data.y), "r"((int)pred_guard));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}
};
Expand All @@ -501,7 +515,8 @@ struct global_store<AccessType, 4> {
:
: "l"(ptr), "r"(data), "r"((int)pred_guard));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}
};
Expand All @@ -521,7 +536,8 @@ struct global_store<AccessType, 2> {
:
: "l"(ptr), "h"(data), "r"((int)pred_guard));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}
};
Expand Down Expand Up @@ -551,7 +567,8 @@ void shared_load<2>(void *dst, uint32_t ptr) {
: "=h"(*reinterpret_cast<uint16_t *>(dst))
: "r"(ptr));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}

Expand All @@ -564,23 +581,8 @@ void shared_load<4>(void *dst, uint32_t ptr) {
: "=r"(*reinterpret_cast<uint32_t *>(dst))
: "r"(ptr));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}

/// ld.shared - 64b
template <>
CUTLASS_DEVICE
void shared_load<8>(void *dst, uint32_t ptr) {
uint2 *dst_u64 = reinterpret_cast<uint2 *>(dst);
#if defined(__CUDA_ARCH__) || defined(__SYCL_CUDA_ARCH__)
asm volatile("ld.shared.v2.u32 {%0, %1}, [%2];\n"
:
"=r"(dst_u64->x),
"=r"(dst_u64->y)
: "r"(ptr));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}

Expand All @@ -598,7 +600,8 @@ void shared_load<16>(void *dst, uint32_t ptr) {
"=r"(dst_u128->w)
: "r"(ptr));
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}

Expand All @@ -621,7 +624,8 @@ void shared_store<2>(uint32_t ptr, void const *src) {
"h"(*reinterpret_cast<uint16_t const *>(src))
);
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}

Expand All @@ -636,7 +640,8 @@ void shared_store<4>(uint32_t ptr, void const *src) {
"r"(*reinterpret_cast<uint32_t const *>(src))
);
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}

Expand All @@ -653,7 +658,8 @@ void shared_store<8>(uint32_t ptr, void const *src) {
"r"(dst_u64->y)
);
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}

Expand All @@ -672,7 +678,8 @@ void shared_store<16>(uint32_t ptr, void const *src) {
"r"(dst_u128->w)
);
#else
CUTLASS_INVALID_CONTROL_PATH("Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
CUTLASS_INVALID_CONTROL_PATH(
"Attempting to use Nvidia-specific code path on non-Nvidia hardware.");
#endif
}

Expand Down

0 comments on commit 5cbb157

Please sign in to comment.