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.

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

Fix icpx build and warnings #104

Merged
merged 2 commits into from
Aug 5, 2024
Merged

Fix icpx build and warnings #104

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
bd46540
Fix icpx build and warnings
rolandschulz Jul 25, 2024
002144a
Use enable-global-offset only for nvidia backend
aacostadiaz Aug 5, 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
20 changes: 13 additions & 7 deletions cmake/FindDPCPP.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,30 +38,36 @@ find_library(DPCPP_LIB_DIR NAMES sycl sycl6 PATHS "${DPCPP_BIN_DIR}/../lib")

add_library(DPCPP::DPCPP INTERFACE IMPORTED)

set(DPCPP_FLAGS "-fsycl;-mllvm;-enable-global-offset=false;")
aacostadiaz marked this conversation as resolved.
Show resolved Hide resolved
set(DPCPP_FLAGS "-fsycl;")
set(DPCPP_COMPILE_ONLY_FLAGS "")

if(NOT "${DPCPP_SYCL_TARGET}" STREQUAL "")
list(APPEND DPCPP_FLAGS "-fsycl-targets=${DPCPP_SYCL_TARGET};")
endif()
list(APPEND DPCPP_FLAGS "${DPCPP_USER_FLAGS};")

if(NOT "${DPCPP_USER_FLAGS}" STREQUAL "")
list(APPEND DPCPP_FLAGS "${DPCPP_USER_FLAGS};")
endif()

if(NOT "${DPCPP_SYCL_ARCH}" STREQUAL "")
if("${DPCPP_SYCL_TARGET}" STREQUAL "nvptx64-nvidia-cuda")
list(APPEND DPCPP_FLAGS "-Xsycl-target-backend")
list(APPEND DPCPP_FLAGS "--cuda-gpu-arch=${DPCPP_SYCL_ARCH}")
list(APPEND DPCPP_COMPILE_ONLY_FLAGS; "-mllvm;-enable-global-offset=false;")
AD2605 marked this conversation as resolved.
Show resolved Hide resolved
endif()
endif()

if(UNIX)
set_target_properties(DPCPP::DPCPP PROPERTIES
INTERFACE_COMPILE_OPTIONS "${DPCPP_FLAGS}"
INTERFACE_COMPILE_OPTIONS "${DPCPP_FLAGS};${DPCPP_COMPILE_ONLY_FLAGS}"
INTERFACE_LINK_OPTIONS "${DPCPP_FLAGS}"
INTERFACE_LINK_LIBRARIES ${DPCPP_LIB_DIR}
INTERFACE_INCLUDE_DIRECTORIES "${DPCPP_BIN_DIR}/../include/sycl;${DPCPP_BIN_DIR}/../include")
message(STATUS "DPCPP INCLUDE DIR: ${DPCPP_BIN_DIR}/../include/sycl;${DPCPP_BIN_DIR}/../include")
message(STATUS "Using DPCPP flags: ${DPCPP_FLAGS}")
message(STATUS "Using DPCPP flags: ${DPCPP_FLAGS};${DPCPP_COMPILE_ONLY_FLAGS}")
else()
set_target_properties(DPCPP::DPCPP PROPERTIES
INTERFACE_COMPILE_OPTIONS "${DPCPP_FLAGS}"
INTERFACE_COMPILE_OPTIONS "${DPCPP_FLAGS};${DPCPP_COMPILE_ONLY_FLAGS}"
INTERFACE_LINK_LIBRARIES ${DPCPP_LIB_DIR}
INTERFACE_INCLUDE_DIRECTORIES "${DPCPP_BIN_DIR}/../include/sycl")
endif()
Expand All @@ -88,8 +94,8 @@ function(add_sycl_to_target)
endfunction()

function(add_sycl_include_directories_to_target NAME)
target_include_directories(${NAME}
target_include_directories(${NAME} SYSTEM
PUBLIC ${DPCPP_BIN_DIR}/../include/sycl
PUBLIC ${DPCPP_BIN_DIR}/../include>
PUBLIC ${DPCPP_BIN_DIR}/../include
)
endfunction()
2 changes: 1 addition & 1 deletion include/cute/arch/copy_xe.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -41,7 +41,7 @@ namespace cute
#define SYCL_DEVICE_BUILTIN(x) SYCL_EXTERNAL extern "C" x
#else
#define SYCL_DEVICE_BUILTIN(x) \
inline x { assert(false); }
inline x { CUTE_INVALID_CONTROL_PATH("Trying to use XE built-in on non-XE hardware"); }
#endif

enum class CacheControl {
Expand Down
2 changes: 1 addition & 1 deletion include/cute/arch/mma_xe.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,7 +37,7 @@
#ifdef __SYCL_DEVICE_ONLY__
#define SYCL_DEVICE_OCL(x) SYCL_EXTERNAL x
#else
#define SYCL_DEVICE_OCL(x) inline x { assert(false); }
#define SYCL_DEVICE_OCL(x) inline x { CUTE_INVALID_CONTROL_PATH("Trying to use XE built-in on non-XE hardware"); }
#endif

SYCL_DEVICE_OCL(cute::intel::float8 intel_sub_group_bf16_bf16_matrix_mad_k16(cute::intel::short8 a, cute::intel::int8 b, cute::intel::float8 acc));
Expand Down
2 changes: 2 additions & 0 deletions include/cute/config.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -132,6 +132,8 @@
// Fail and print a message. Typically used for notification of a compiler misconfiguration.
#if defined(__CUDA_ARCH__)
# define CUTE_INVALID_CONTROL_PATH(x) assert(0 && x); printf(x); __brkpt()
#elif defined(__has_builtin) && __has_builtin(__builtin_unreachable)
# define CUTE_INVALID_CONTROL_PATH(x) assert(0 && x); printf(x); __builtin_unreachable()
#else
# define CUTE_INVALID_CONTROL_PATH(x) assert(0 && x); printf(x)
#endif
Expand Down
14 changes: 7 additions & 7 deletions include/cute/util/sycl_vec.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,9 +38,9 @@ namespace cute
namespace intel
{
#ifdef __SYCL_DEVICE_ONLY__
template <class T, int N> using vector_t = typename sycl::vec<T, N>::vector_t;
template <class T, int N> using vector_t = T __attribute__((ext_vector_type(N)));
#else
template <class T, int N> using vector_t = sycl::vec<T, N>;
template <class T, int N> using vector_t = sycl::marray<T, N>;
#endif

using float8 = vector_t<float, 8>;
Expand All @@ -50,11 +50,11 @@ using int16 = vector_t<int, 16>;
using uint8 = vector_t<uint, 8>;
using uint16 = vector_t<uint, 16>;

typedef ushort __attribute__((ext_vector_type(8))) ushort8;
typedef ushort __attribute__((ext_vector_type(16))) ushort16;
typedef ushort __attribute__((ext_vector_type(32))) ushort32;
typedef ushort __attribute__((ext_vector_type(64))) ushort64;
typedef uint __attribute__((ext_vector_type(32))) uint32;
using ushort8 = vector_t<ushort, 8>;
using ushort16 = vector_t<ushort, 16>;
using ushort32 = vector_t<ushort, 32>;
using ushort64 = vector_t<ushort, 64>;
using uint32 = vector_t<uint, 32>;

using coord_t = vector_t<int, 2>;
} // namespace intel end
Expand Down