-
Notifications
You must be signed in to change notification settings - Fork 154
/
backend.hpp
150 lines (132 loc) · 5.23 KB
/
backend.hpp
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
/*
* This file is part of hipSYCL, a SYCL implementation based on CUDA/HIP
*
* Copyright (c) 2020 Aksel Alpay
* All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are permitted provided that the following conditions are met:
*
* 1. Redistributions of source code must retain the above copyright notice,
* this list of conditions and the following disclaimer.
* 2. Redistributions in binary form must reproduce the above copyright notice,
* this list of conditions and the following disclaimer in the documentation
* and/or other materials provided with the distribution.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
* AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
* ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
* LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
* CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
* SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
* INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
* CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
* ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
* POSSIBILITY OF SUCH DAMAGE.
*/
#ifndef HIPSYCL_LIBKERNEL_BACKEND_HPP
#define HIPSYCL_LIBKERNEL_BACKEND_HPP
#include "cuda/cuda_backend.hpp"
#include "hip/hip_backend.hpp"
#include "spirv/spirv_backend.hpp"
#include "host/host_backend.hpp"
// define (legacy?) platform identification macros
#if HIPSYCL_LIBKERNEL_COMPILER_SUPPORTS_HIP
#define HIPSYCL_PLATFORM_ROCM
#define HIPSYCL_PLATFORM_HIP
#endif
#if HIPSYCL_LIBKERNEL_COMPILER_SUPPORTS_CUDA
#define HIPSYCL_PLATFORM_CUDA
#endif
#if HIPSYCL_LIBKERNEL_COMPILER_SUPPORTS_SPIRV
#define HIPSYCL_PLATFORM_SPIRV
#endif
#ifndef HIPSYCL_LIBKERNEL_IS_UNIFIED_HOST_DEVICE_PASS
#define HIPSYCL_LIBKERNEL_IS_UNIFIED_HOST_DEVICE_PASS 0
#endif
#if HIPSYCL_LIBKERNEL_COMPILER_SUPPORTS_HIP || \
HIPSYCL_LIBKERNEL_COMPILER_SUPPORTS_CUDA || \
HIPSYCL_LIBKERNEL_COMPILER_SUPPORTS_SPIRV
#define HIPSYCL_LIBKERNEL_COMPILER_SUPPORTS_DEVICE 1
#else
#define HIPSYCL_LIBKERNEL_COMPILER_SUPPORTS_DEVICE 0
#endif
#if HIPSYCL_LIBKERNEL_IS_DEVICE_PASS_HOST
#define HIPSYCL_PLATFORM_CPU
#endif
#ifdef HIPSYCL_LIBKERNEL_DEVICE_PASS
#define HIPSYCL_LIBKERNEL_IS_DEVICE_PASS 1
#else
#define HIPSYCL_LIBKERNEL_IS_DEVICE_PASS 0
#endif
#if HIPSYCL_LIBKERNEL_IS_DEVICE_PASS && \
!HIPSYCL_LIBKERNEL_IS_UNIFIED_HOST_DEVICE_PASS
#define SYCL_DEVICE_ONLY
#ifndef __SYCL_DEVICE_ONLY__
#define __SYCL_DEVICE_ONLY__ 1
#endif
#endif
#if !defined(HIPSYCL_LIBKERNEL_IS_UNIFIED_HOST_DEVICE_PASS)
#define HIPSYCL_LIBKERNEL_IS_UNIFIED_HOST_DEVICE_PASS 0
#endif
#if HIPSYCL_LIBKERNEL_IS_UNIFIED_HOST_DEVICE_PASS
#define __SYCL_SINGLE_SOURCE__ 1
#endif
#ifdef __clang__
#define HIPSYCL_FORCE_INLINE \
__attribute__((always_inline)) __attribute__((flatten)) inline
#define HIPSYCL_LOOP_SPLIT_ND_KERNEL [[clang::annotate("hipsycl_nd_kernel")]]
#define HIPSYCL_LOOP_SPLIT_ND_KERNEL_LOCAL_SIZE_ARG [[clang::annotate("hipsycl_nd_kernel_local_size_arg")]]
#define HIPSYCL_LOOP_SPLIT_BARRIER [[clang::annotate("hipsycl_barrier")]]
#else
#define HIPSYCL_FORCE_INLINE inline
#define HIPSYCL_LOOP_SPLIT_ND_KERNEL
#define HIPSYCL_LOOP_SPLIT_BARRIER
#define HIPSYCL_LOOP_SPLIT_ND_KERNEL_LOCAL_SIZE_ARG
#endif
#define HIPSYCL_BUILTIN HIPSYCL_UNIVERSAL_TARGET HIPSYCL_FORCE_INLINE
#if HIPSYCL_LIBKERNEL_COMPILER_SUPPORTS_CUDA || \
HIPSYCL_LIBKERNEL_COMPILER_SUPPORTS_HIP
#define HIPSYCL_HIPLIKE_BUILTIN __device__ HIPSYCL_FORCE_INLINE
#endif
#ifndef __hipsycl_if_target_host
#if !HIPSYCL_LIBKERNEL_IS_DEVICE_PASS
#define __hipsycl_if_target_host(...) __VA_ARGS__
#else
#define __hipsycl_if_target_host(...)
#endif
#endif
#ifndef __hipsycl_if_target_device
#if HIPSYCL_LIBKERNEL_IS_DEVICE_PASS
#define __hipsycl_if_target_device(...) __VA_ARGS__
#else
#define __hipsycl_if_target_device(...)
#endif
#endif
#if HIPSYCL_LIBKERNEL_IS_DEVICE_PASS_CUDA
#define __hipsycl_if_target_cuda(...) __hipsycl_if_target_device(__VA_ARGS__)
#else
#define __hipsycl_if_target_cuda(...)
#endif
#if HIPSYCL_LIBKERNEL_IS_DEVICE_PASS_HIP
#define __hipsycl_if_target_hip(...) __hipsycl_if_target_device(__VA_ARGS__)
#else
#define __hipsycl_if_target_hip(...)
#endif
#if HIPSYCL_LIBKERNEL_IS_DEVICE_PASS_HIP || \
HIPSYCL_LIBKERNEL_IS_DEVICE_PASS_CUDA
#define __hipsycl_if_target_hiplike(...) \
__hipsycl_if_target_device(__VA_ARGS__)
#else
#define __hipsycl_if_target_hiplike(...)
#endif
#if HIPSYCL_LIBKERNEL_IS_DEVICE_PASS_SPIRV
#define __hipsycl_if_target_spirv(...) __hipsycl_if_target_device(__VA_ARGS__)
#else
#define __hipsycl_if_target_spirv(...)
#endif
#define HIPSYCL_LIBKERNEL_IS_EXCLUSIVE_PASS(backend) \
((HIPSYCL_LIBKERNEL_IS_DEVICE_PASS_##backend) && \
!HIPSYCL_LIBKERNEL_IS_UNIFIED_HOST_DEVICE_PASS)
#endif