Skip to content

Commit d922264

Browse files
authored
Update 03-sycl-advanced.md
1 parent a845339 commit d922264

File tree

1 file changed

+38
-9
lines changed

1 file changed

+38
-9
lines changed

sycl/docs/03-sycl-advanced.md

Lines changed: 38 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -17,6 +17,14 @@ lang: en
1717

1818
# Unified Shared Memory II
1919

20+
| Function | Location | Device Accessible
21+
------------------+--------------------+--------------------
22+
| malloc_device | Device | Yes
23+
| malloc_shared | Dynamic migration | Yes
24+
| malloc_host | Host | Device can read
25+
26+
# `malloc_device`
27+
2028
<small>
2129
```cpp
2230
std::vector<int> y(N, 1);
@@ -33,24 +41,45 @@ lang: en
3341
}).wait();
3442
// Copy results back to host
3543
q.memcpy(y.data(), d_y, N * sizeof(int)).wait();
36-
37-
// Free the device memory
38-
sycl::free(d_y, q);
3944

4045
// Verify the results
4146
for (int i = 0; i < N; i++) {
4247
assert(y[i] == 2);
4348
}
49+
50+
// Free the device memory
51+
sycl::free(d_y, q);
4452
```
4553
</small>
4654
47-
# Unified Shared Memory III
4855
49-
| Function | Location | Device Accessible
50-
------------------+--------------------+--------------------
51-
| malloc_device | Device | Yes
52-
| malloc_shared | Dynamic migration | Yes
53-
| malloc_host | Host | Device can read
56+
# `malloc_shared`
57+
58+
<small>
59+
```cpp
60+
int* y = malloc_shared<int>(N, q);
61+
// Initialize from host
62+
for (int i = 0; i < N; i++) {
63+
y[i] = 1;
64+
}
65+
66+
q.submit([&](handler& cgh) {
67+
cgh.parallel_for(range<1>(N), [=](id<1> idx) {
68+
y[idx] += 1;
69+
});
70+
}).wait();
71+
72+
// No memcpy needed — host can read directly
73+
for (int i = 0; i < N; i++) {
74+
assert(y[i] == 2);
75+
std::cout << "y[" << i << "] = " << y[i] << "\n";
76+
}
77+
78+
// Free the shared memory
79+
sycl::free(y, q);
80+
```
81+
</small>
82+
5483

5584
# Parallel launch with **nd-range** I
5685

0 commit comments

Comments
 (0)