@@ -13,111 +13,109 @@ using namespace sycl;
1313
1414int main () {
1515 queue Queue;
16- if (!Queue.is_host ()) {
17- context Context = Queue.get_context ();
18-
19- cl_context ClContext = Context.get ();
20-
21- const size_t CountSources = 3 ;
22- const char *Sources[CountSources] = {
23- " kernel void foo1(global float* Array, global int* Value) { *Array = "
24- " 42; *Value = 1; }\n "
25- " kernel void foo2(global float* Array) { int id = get_global_id(0); "
26- " Array[id] = id; }\n "
27- " kernel void foo3(global float* Array, local float* LocalArray) { "
28- " (void)LocalArray; (void)Array; }\n "
29- };
30-
31- cl_int Err;
32- cl_program ClProgram = clCreateProgramWithSource (ClContext, CountSources,
33- Sources, nullptr , &Err);
34- assert (Err == CL_SUCCESS);
35-
36- Err = clBuildProgram (ClProgram, 0 , nullptr , nullptr , nullptr , nullptr );
37- assert (Err == CL_SUCCESS);
38-
39- cl_kernel FirstCLKernel = clCreateKernel (ClProgram, " foo1"
40- assert (Err == CL_SUCCESS);
41-
42- cl_kernel SecondCLKernel = clCreateKernel (ClProgram, " foo2"
43- assert (Err == CL_SUCCESS);
44-
45- cl_kernel ThirdCLKernel = clCreateKernel (ClProgram, " foo3"
46- assert (Err == CL_SUCCESS);
47-
48- const size_t Count = 100 ;
49- float Array[Count];
50-
51- kernel FirstKernel (FirstCLKernel, Context);
52- kernel SecondKernel (SecondCLKernel, Context);
53- kernel ThirdKernel (ThirdCLKernel, Context);
54- int Value;
55- {
56- buffer<float , 1 > FirstBuffer (Array, range<1 >(1 ));
57- buffer<int , 1 > SecondBuffer (&Value, range<1 >(1 ));
58- Queue.submit ([&](handler &CGH) {
59- CGH.
set_arg (
0 , FirstBuffer.
get_access <access::mode::write>(CGH));
60- CGH.
set_arg (
1 , SecondBuffer.
get_access <access::mode::write>(CGH));
61- CGH.single_task (FirstKernel);
62- });
63- }
64- Queue.wait_and_throw ();
16+ context Context = Queue.get_context ();
17+
18+ cl_context ClContext = Context.get ();
19+
20+ const size_t CountSources = 3 ;
21+ const char *Sources[CountSources] = {
22+ " kernel void foo1(global float* Array, global int* Value) { *Array = "
23+ " 42; *Value = 1; }\n "
24+ " kernel void foo2(global float* Array) { int id = get_global_id(0); "
25+ " Array[id] = id; }\n "
26+ " kernel void foo3(global float* Array, local float* LocalArray) { "
27+ " (void)LocalArray; (void)Array; }\n "
28+ };
29+
30+ cl_int Err;
31+ cl_program ClProgram = clCreateProgramWithSource (ClContext, CountSources,
32+ Sources, nullptr , &Err);
33+ assert (Err == CL_SUCCESS);
34+
35+ Err = clBuildProgram (ClProgram, 0 , nullptr , nullptr , nullptr , nullptr );
36+ assert (Err == CL_SUCCESS);
37+
38+ cl_kernel FirstCLKernel = clCreateKernel (ClProgram, " foo1"
39+ assert (Err == CL_SUCCESS);
40+
41+ cl_kernel SecondCLKernel = clCreateKernel (ClProgram, " foo2"
42+ assert (Err == CL_SUCCESS);
43+
44+ cl_kernel ThirdCLKernel = clCreateKernel (ClProgram, " foo3"
45+ assert (Err == CL_SUCCESS);
46+
47+ const size_t Count = 100 ;
48+ float Array[Count];
49+
50+ kernel FirstKernel (FirstCLKernel, Context);
51+ kernel SecondKernel (SecondCLKernel, Context);
52+ kernel ThirdKernel (ThirdCLKernel, Context);
53+ int Value;
54+ {
55+ buffer<float , 1 > FirstBuffer (Array, range<1 >(1 ));
56+ buffer<int , 1 > SecondBuffer (&Value, range<1 >(1 ));
57+ Queue.submit ([&](handler &CGH) {
58+ CGH.
set_arg (
0 , FirstBuffer.
get_access <access::mode::write>(CGH));
59+ CGH.
set_arg (
1 , SecondBuffer.
get_access <access::mode::write>(CGH));
60+ CGH.single_task (FirstKernel);
61+ });
62+ }
63+ Queue.wait_and_throw ();
64+
65+ assert (Array[0 ] == 42 );
66+ assert (Value == 1 );
67+
68+ {
69+ buffer<float , 1 > FirstBuffer (Array, range<1 >(Count));
70+ Queue.submit ([&](handler &CGH) {
71+ auto Acc = FirstBuffer.
get_access <access::mode::read_write>(CGH);
72+ CGH.
set_arg (
0 , FirstBuffer.
get_access <access::mode::read_write>(CGH));
73+ CGH.parallel_for (range<1 >{Count}, SecondKernel);
74+ });
75+ }
76+ Queue.wait_and_throw ();
77+
78+ for (size_t I = 0 ; I < Count; ++I) {
79+ assert (Array[I] == I);
80+ }
6581
66- assert (Array[0 ] == 42 );
67- assert (Value == 1 );
82+ {
83+ auto dev = Queue.get_device ();
84+ auto ctxt = Queue.get_context ();
85+ if (dev.get_info <info::device::usm_shared_allocations>()) {
86+ float *data =
87+ static_cast <float *>(malloc_shared (Count * sizeof (float ), dev, ctxt));
6888
69- {
70- buffer<float , 1 > FirstBuffer (Array, range<1 >(Count));
7189 Queue.submit ([&](handler &CGH) {
72- auto Acc = FirstBuffer.
get_access <access::mode::read_write>(CGH);
73- CGH.
set_arg (
0 , FirstBuffer.
get_access <access::mode::read_write>(CGH));
90+ CGH.set_arg (0 , data);
7491 CGH.parallel_for (range<1 >{Count}, SecondKernel);
7592 });
76- }
77- Queue.wait_and_throw ();
93+ Queue.wait_and_throw ();
7894
79- for (size_t I = 0 ; I < Count; ++I) {
80- assert (Array[I] == I);
81- }
82-
83- {
84- auto dev = Queue.get_device ();
85- auto ctxt = Queue.get_context ();
86- if (dev.get_info <info::device::usm_shared_allocations>()) {
87- float *data = static_cast <float *>(
88- malloc_shared (Count * sizeof (float ), dev, ctxt));
89-
90- Queue.submit ([&](handler &CGH) {
91- CGH.set_arg (0 , data);
92- CGH.parallel_for (range<1 >{Count}, SecondKernel);
93- });
94- Queue.wait_and_throw ();
95-
96- for (size_t I = 0 ; I < Count; ++I) {
97- assert (data[I] == I);
98- }
99- free (data, ctxt);
95+ for (size_t I = 0 ; I < Count; ++I) {
96+ assert (data[I] == I);
10097 }
98+ free (data, ctxt);
10199 }
100+ }
102101
103- {
104- buffer<float , 1 > FirstBuffer (Array, range<1 >(Count));
105- Queue.submit ([&](handler &CGH) {
106- auto Acc = FirstBuffer.
get_access <access::mode::read_write>(CGH);
107- CGH.
set_arg (
0 , FirstBuffer.
get_access <access::mode::read_write>(CGH));
108- CGH.set_arg (1 , sycl::accessor<float , 1 , sycl::access::mode::read_write,
109- sycl::access::target::local>(
110- sycl::range<1 >(Count), CGH));
111- CGH.parallel_for (range<1 >{Count}, ThirdKernel);
112- });
113- }
114- Queue.wait_and_throw ();
115-
116- clReleaseContext (ClContext);
117- clReleaseKernel (FirstCLKernel);
118- clReleaseKernel (SecondCLKernel);
119- clReleaseKernel (ThirdCLKernel);
120- clReleaseProgram (ClProgram);
102+ {
103+ buffer<float , 1 > FirstBuffer (Array, range<1 >(Count));
104+ Queue.submit ([&](handler &CGH) {
105+ auto Acc = FirstBuffer.
get_access <access::mode::read_write>(CGH);
106+ CGH.
set_arg (
0 , FirstBuffer.
get_access <access::mode::read_write>(CGH));
107+ CGH.set_arg (1 , sycl::accessor<float , 1 , sycl::access::mode::read_write,
108+ sycl::access::target::local>(
109+ sycl::range<1 >(Count), CGH));
110+ CGH.parallel_for (range<1 >{Count}, ThirdKernel);
111+ });
121112 }
113+ Queue.wait_and_throw ();
114+
115+ clReleaseContext (ClContext);
116+ clReleaseKernel (FirstCLKernel);
117+ clReleaseKernel (SecondCLKernel);
118+ clReleaseKernel (ThirdCLKernel);
119+ clReleaseProgram (ClProgram);
122120 return 0 ;
123121}
0 commit comments