1 /*
2 //@HEADER
3 // ************************************************************************
4 //
5 // Kokkos v. 3.0
6 // Copyright (2020) National Technology & Engineering
7 // Solutions of Sandia, LLC (NTESS).
8 //
9 // Under the terms of Contract DE-NA0003525 with NTESS,
10 // the U.S. Government retains certain rights in this software.
11 //
12 // Redistribution and use in source and binary forms, with or without
13 // modification, are permitted provided that the following conditions are
14 // met:
15 //
16 // 1. Redistributions of source code must retain the above copyright
17 // notice, this list of conditions and the following disclaimer.
18 //
19 // 2. Redistributions in binary form must reproduce the above copyright
20 // notice, this list of conditions and the following disclaimer in the
21 // documentation and/or other materials provided with the distribution.
22 //
23 // 3. Neither the name of the Corporation nor the names of the
24 // contributors may be used to endorse or promote products derived from
25 // this software without specific prior written permission.
26 //
27 // THIS SOFTWARE IS PROVIDED BY NTESS "AS IS" AND ANY
28 // EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
29 // IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
30 // PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NTESS OR THE
31 // CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
32 // EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
33 // PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
34 // PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
35 // LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
36 // NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
37 // SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
38 //
39 // Questions? Contact Christian R. Trott (crtrott@sandia.gov)
40 //
41 // ************************************************************************
42 //@HEADER
43 */
44
45 #include <Kokkos_Core.hpp>
46 #include <TestCuda_Category.hpp>
47
48 namespace Test {
49
test_abort()50 __global__ void test_abort() { Kokkos::abort("test_abort"); }
51
test_cuda_spaces_int_value(int * ptr)52 __global__ void test_cuda_spaces_int_value(int *ptr) {
53 if (*ptr == 42) {
54 *ptr = 2 * 42;
55 }
56 }
57
TEST(cuda,space_access)58 TEST(cuda, space_access) {
59 static_assert(Kokkos::Impl::MemorySpaceAccess<Kokkos::HostSpace,
60 Kokkos::HostSpace>::assignable,
61 "");
62
63 static_assert(
64 Kokkos::Impl::MemorySpaceAccess<Kokkos::HostSpace,
65 Kokkos::CudaHostPinnedSpace>::assignable,
66 "");
67
68 static_assert(!Kokkos::Impl::MemorySpaceAccess<Kokkos::HostSpace,
69 Kokkos::CudaSpace>::assignable,
70 "");
71
72 static_assert(!Kokkos::Impl::MemorySpaceAccess<Kokkos::HostSpace,
73 Kokkos::CudaSpace>::accessible,
74 "");
75
76 static_assert(
77 !Kokkos::Impl::MemorySpaceAccess<Kokkos::HostSpace,
78 Kokkos::CudaUVMSpace>::assignable,
79 "");
80
81 static_assert(
82 Kokkos::Impl::MemorySpaceAccess<Kokkos::HostSpace,
83 Kokkos::CudaUVMSpace>::accessible,
84 "");
85
86 //--------------------------------------
87
88 static_assert(Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaSpace,
89 Kokkos::CudaSpace>::assignable,
90 "");
91
92 static_assert(
93 Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaSpace,
94 Kokkos::CudaUVMSpace>::assignable,
95 "");
96
97 static_assert(
98 !Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaSpace,
99 Kokkos::CudaHostPinnedSpace>::assignable,
100 "");
101
102 static_assert(
103 Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaSpace,
104 Kokkos::CudaHostPinnedSpace>::accessible,
105 "");
106
107 static_assert(!Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaSpace,
108 Kokkos::HostSpace>::assignable,
109 "");
110
111 static_assert(!Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaSpace,
112 Kokkos::HostSpace>::accessible,
113 "");
114
115 //--------------------------------------
116
117 static_assert(
118 Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaUVMSpace,
119 Kokkos::CudaUVMSpace>::assignable,
120 "");
121
122 static_assert(!Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaUVMSpace,
123 Kokkos::CudaSpace>::assignable,
124 "");
125
126 static_assert(Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaUVMSpace,
127 Kokkos::CudaSpace>::accessible,
128 "");
129
130 static_assert(!Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaUVMSpace,
131 Kokkos::HostSpace>::assignable,
132 "");
133
134 static_assert(!Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaUVMSpace,
135 Kokkos::HostSpace>::accessible,
136 "");
137
138 static_assert(
139 !Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaUVMSpace,
140 Kokkos::CudaHostPinnedSpace>::assignable,
141 "");
142
143 static_assert(
144 Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaUVMSpace,
145 Kokkos::CudaHostPinnedSpace>::accessible,
146 "");
147
148 //--------------------------------------
149
150 static_assert(
151 Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaHostPinnedSpace,
152 Kokkos::CudaHostPinnedSpace>::assignable,
153 "");
154
155 static_assert(!Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaHostPinnedSpace,
156 Kokkos::HostSpace>::assignable,
157 "");
158
159 static_assert(Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaHostPinnedSpace,
160 Kokkos::HostSpace>::accessible,
161 "");
162
163 static_assert(!Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaHostPinnedSpace,
164 Kokkos::CudaSpace>::assignable,
165 "");
166
167 static_assert(!Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaHostPinnedSpace,
168 Kokkos::CudaSpace>::accessible,
169 "");
170
171 static_assert(
172 !Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaHostPinnedSpace,
173 Kokkos::CudaUVMSpace>::assignable,
174 "");
175
176 static_assert(
177 Kokkos::Impl::MemorySpaceAccess<Kokkos::CudaHostPinnedSpace,
178 Kokkos::CudaUVMSpace>::accessible,
179 "");
180
181 //--------------------------------------
182
183 static_assert(
184 !Kokkos::Impl::SpaceAccessibility<Kokkos::Cuda,
185 Kokkos::HostSpace>::accessible,
186 "");
187
188 static_assert(Kokkos::Impl::SpaceAccessibility<Kokkos::Cuda,
189 Kokkos::CudaSpace>::accessible,
190 "");
191
192 static_assert(
193 Kokkos::Impl::SpaceAccessibility<Kokkos::Cuda,
194 Kokkos::CudaUVMSpace>::accessible,
195 "");
196
197 static_assert(
198 Kokkos::Impl::SpaceAccessibility<Kokkos::Cuda,
199 Kokkos::CudaHostPinnedSpace>::accessible,
200 "");
201
202 static_assert(
203 !Kokkos::Impl::SpaceAccessibility<Kokkos::HostSpace,
204 Kokkos::CudaSpace>::accessible,
205 "");
206
207 static_assert(
208 Kokkos::Impl::SpaceAccessibility<Kokkos::HostSpace,
209 Kokkos::CudaUVMSpace>::accessible,
210 "");
211
212 static_assert(
213 Kokkos::Impl::SpaceAccessibility<Kokkos::HostSpace,
214 Kokkos::CudaHostPinnedSpace>::accessible,
215 "");
216
217 static_assert(std::is_same<Kokkos::Impl::HostMirror<Kokkos::CudaSpace>::Space,
218 Kokkos::HostSpace>::value,
219 "");
220
221 static_assert(
222 std::is_same<Kokkos::Impl::HostMirror<Kokkos::CudaUVMSpace>::Space,
223 Kokkos::Device<Kokkos::HostSpace::execution_space,
224 Kokkos::CudaUVMSpace>>::value,
225 "");
226
227 static_assert(
228 std::is_same<Kokkos::Impl::HostMirror<Kokkos::CudaHostPinnedSpace>::Space,
229 Kokkos::CudaHostPinnedSpace>::value,
230 "");
231
232 static_assert(std::is_same<Kokkos::Device<Kokkos::HostSpace::execution_space,
233 Kokkos::CudaUVMSpace>,
234 Kokkos::Device<Kokkos::HostSpace::execution_space,
235 Kokkos::CudaUVMSpace>>::value,
236 "");
237
238 static_assert(Kokkos::Impl::SpaceAccessibility<
239 Kokkos::Impl::HostMirror<Kokkos::Cuda>::Space,
240 Kokkos::HostSpace>::accessible,
241 "");
242
243 static_assert(Kokkos::Impl::SpaceAccessibility<
244 Kokkos::Impl::HostMirror<Kokkos::CudaSpace>::Space,
245 Kokkos::HostSpace>::accessible,
246 "");
247
248 static_assert(Kokkos::Impl::SpaceAccessibility<
249 Kokkos::Impl::HostMirror<Kokkos::CudaUVMSpace>::Space,
250 Kokkos::HostSpace>::accessible,
251 "");
252
253 static_assert(
254 Kokkos::Impl::SpaceAccessibility<
255 Kokkos::Impl::HostMirror<Kokkos::CudaHostPinnedSpace>::Space,
256 Kokkos::HostSpace>::accessible,
257 "");
258 #ifdef KOKKOS_ENABLE_CUDA_UVM
259 using uvm_view = Kokkos::View<double *, Kokkos::CudaUVMSpace>;
260 static_assert(std::is_same<uvm_view::HostMirror::execution_space,
261 Kokkos::DefaultHostExecutionSpace>::value,
262 "Verify HostMirror execution space is really a host space");
263 #endif
264 }
265
TEST(cuda,uvm)266 TEST(cuda, uvm) {
267 if (Kokkos::CudaUVMSpace::available()) {
268 int *uvm_ptr = (int *)Kokkos::kokkos_malloc<Kokkos::CudaUVMSpace>(
269 "uvm_ptr", sizeof(int));
270
271 *uvm_ptr = 42;
272
273 Kokkos::Cuda().fence();
274 test_cuda_spaces_int_value<<<1, 1>>>(uvm_ptr);
275 Kokkos::Cuda().fence();
276
277 EXPECT_EQ(*uvm_ptr, int(2 * 42));
278
279 Kokkos::kokkos_free<Kokkos::CudaUVMSpace>(uvm_ptr);
280 }
281 }
282
283 template <class MemSpace, class ExecSpace>
284 struct TestViewCudaAccessible {
285 enum { N = 1000 };
286
287 using V = Kokkos::View<double *, MemSpace>;
288
289 V m_base;
290
291 struct TagInit {};
292 struct TagTest {};
293
294 KOKKOS_INLINE_FUNCTION
operator ()Test::TestViewCudaAccessible295 void operator()(const TagInit &, const int i) const { m_base[i] = i + 1; }
296
297 KOKKOS_INLINE_FUNCTION
operator ()Test::TestViewCudaAccessible298 void operator()(const TagTest &, const int i, long &error_count) const {
299 if (m_base[i] != i + 1) ++error_count;
300 }
301
TestViewCudaAccessibleTest::TestViewCudaAccessible302 TestViewCudaAccessible() : m_base("base", N) {}
303
runTest::TestViewCudaAccessible304 static void run() {
305 TestViewCudaAccessible self;
306 Kokkos::parallel_for(
307 Kokkos::RangePolicy<typename MemSpace::execution_space, TagInit>(0, N),
308 self);
309 typename MemSpace::execution_space().fence();
310
311 // Next access is a different execution space, must complete prior kernel.
312 long error_count = -1;
313 Kokkos::parallel_reduce(Kokkos::RangePolicy<ExecSpace, TagTest>(0, N), self,
314 error_count);
315 EXPECT_EQ(error_count, 0);
316 }
317 };
318
TEST(cuda,impl_view_accessible)319 TEST(cuda, impl_view_accessible) {
320 TestViewCudaAccessible<Kokkos::CudaSpace, Kokkos::Cuda>::run();
321
322 TestViewCudaAccessible<Kokkos::CudaUVMSpace, Kokkos::Cuda>::run();
323 TestViewCudaAccessible<Kokkos::CudaUVMSpace,
324 Kokkos::HostSpace::execution_space>::run();
325
326 TestViewCudaAccessible<Kokkos::CudaHostPinnedSpace, Kokkos::Cuda>::run();
327 TestViewCudaAccessible<Kokkos::CudaHostPinnedSpace,
328 Kokkos::HostSpace::execution_space>::run();
329 }
330
331 template <class MemSpace>
332 struct TestViewCudaTexture {
333 enum { N = 1000 };
334
335 using V = Kokkos::View<double *, MemSpace>;
336 using T = Kokkos::View<const double *, MemSpace, Kokkos::MemoryRandomAccess>;
337
338 V m_base;
339 T m_tex;
340
341 struct TagInit {};
342 struct TagTest {};
343
344 KOKKOS_INLINE_FUNCTION
operator ()Test::TestViewCudaTexture345 void operator()(const TagInit &, const int i) const { m_base[i] = i + 1; }
346
347 KOKKOS_INLINE_FUNCTION
operator ()Test::TestViewCudaTexture348 void operator()(const TagTest &, const int i, long &error_count) const {
349 if (m_tex[i] != i + 1) ++error_count;
350 }
351
TestViewCudaTextureTest::TestViewCudaTexture352 TestViewCudaTexture() : m_base("base", N), m_tex(m_base) {}
353
runTest::TestViewCudaTexture354 static void run() {
355 EXPECT_TRUE((std::is_same<typename V::reference_type, double &>::value));
356 EXPECT_TRUE(
357 (std::is_same<typename T::reference_type, const double>::value));
358
359 EXPECT_TRUE(V::reference_type_is_lvalue_reference); // An ordinary view.
360 EXPECT_FALSE(T::reference_type_is_lvalue_reference); // Texture fetch
361 // returns by value.
362
363 TestViewCudaTexture self;
364 Kokkos::parallel_for(Kokkos::RangePolicy<Kokkos::Cuda, TagInit>(0, N),
365 self);
366
367 long error_count = -1;
368 Kokkos::parallel_reduce(Kokkos::RangePolicy<Kokkos::Cuda, TagTest>(0, N),
369 self, error_count);
370 EXPECT_EQ(error_count, 0);
371 }
372 };
373
TEST(cuda,impl_view_texture)374 TEST(cuda, impl_view_texture) {
375 TestViewCudaTexture<Kokkos::CudaSpace>::run();
376 TestViewCudaTexture<Kokkos::CudaUVMSpace>::run();
377 }
378
379 } // namespace Test
380