1 /*
2 * Copyright (C) 2021 Alyssa Rosenzweig <alyssa@rosenzweig.io>
3 * Copyright 2019 Collabora, Ltd.
4 *
5 * Permission is hereby granted, free of charge, to any person obtaining a
6 * copy of this software and associated documentation files (the "Software"),
7 * to deal in the Software without restriction, including without limitation
8 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
9 * and/or sell copies of the Software, and to permit persons to whom the
10 * Software is furnished to do so, subject to the following conditions:
11 *
12 * The above copyright notice and this permission notice (including the next
13 * paragraph) shall be included in all copies or substantial portions of the
14 * Software.
15 *
16 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
17 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
19 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
20 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
21 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
22 * SOFTWARE.
23 */
24
25 #include <inttypes.h>
26 #include "agx_device.h"
27 #include "agx_bo.h"
28 #include "decode.h"
29
30 unsigned AGX_FAKE_HANDLE = 0;
31 uint64_t AGX_FAKE_LO = 0;
32 uint64_t AGX_FAKE_HI = (1ull << 32);
33
34 static void
agx_bo_free(struct agx_device * dev,struct agx_bo * bo)35 agx_bo_free(struct agx_device *dev, struct agx_bo *bo)
36 {
37 #if __APPLE__
38 const uint64_t handle = bo->handle;
39
40 kern_return_t ret = IOConnectCallScalarMethod(dev->fd,
41 AGX_SELECTOR_FREE_MEM,
42 &handle, 1, NULL, NULL);
43
44 if (ret)
45 fprintf(stderr, "error freeing BO mem: %u\n", ret);
46 #else
47 free(bo->ptr.cpu);
48 #endif
49
50 /* Reset the handle */
51 memset(bo, 0, sizeof(*bo));
52 }
53
54 void
agx_shmem_free(struct agx_device * dev,unsigned handle)55 agx_shmem_free(struct agx_device *dev, unsigned handle)
56 {
57 #if __APPLE__
58 const uint64_t input = handle;
59 kern_return_t ret = IOConnectCallScalarMethod(dev->fd,
60 AGX_SELECTOR_FREE_SHMEM,
61 &input, 1, NULL, NULL);
62
63 if (ret)
64 fprintf(stderr, "error freeing shmem: %u\n", ret);
65 #else
66 #endif
67 }
68
69 struct agx_bo
agx_shmem_alloc(struct agx_device * dev,size_t size,bool cmdbuf)70 agx_shmem_alloc(struct agx_device *dev, size_t size, bool cmdbuf)
71 {
72 struct agx_bo bo;
73
74 #if __APPLE__
75 struct agx_create_shmem_resp out = {};
76 size_t out_sz = sizeof(out);
77
78 uint64_t inputs[2] = {
79 size,
80 cmdbuf ? 1 : 0 // 2 - error reporting, 1 - no error reporting
81 };
82
83 kern_return_t ret = IOConnectCallMethod(dev->fd,
84 AGX_SELECTOR_CREATE_SHMEM, inputs, 2, NULL, 0, NULL,
85 NULL, &out, &out_sz);
86
87 assert(ret == 0);
88 assert(out_sz == sizeof(out));
89 assert(out.size == size);
90 assert(out.map != 0);
91
92 bo = (struct agx_bo) {
93 .type = cmdbuf ? AGX_ALLOC_CMDBUF : AGX_ALLOC_MEMMAP,
94 .handle = out.id,
95 .ptr.cpu = out.map,
96 .size = out.size,
97 .guid = 0, /* TODO? */
98 };
99 #else
100 bo = (struct agx_bo) {
101 .type = cmdbuf ? AGX_ALLOC_CMDBUF : AGX_ALLOC_MEMMAP,
102 .handle = AGX_FAKE_HANDLE++,
103 .ptr.cpu = calloc(1, size),
104 .size = size,
105 .guid = 0, /* TODO? */
106 };
107 #endif
108
109 if (dev->debug & AGX_DBG_TRACE)
110 agxdecode_track_alloc(&bo);
111
112 return bo;
113 }
114
115 static struct agx_bo *
agx_bo_alloc(struct agx_device * dev,size_t size,uint32_t flags)116 agx_bo_alloc(struct agx_device *dev, size_t size,
117 uint32_t flags)
118 {
119 struct agx_bo *bo;
120 unsigned handle = 0;
121
122 #if __APPLE__
123 uint32_t mode = 0x430; // shared, ?
124
125 uint32_t args_in[24] = { 0 };
126 args_in[4] = 0x4000101; //0x1000101; // unk
127 args_in[5] = mode;
128 args_in[16] = size;
129 args_in[20] = flags;
130
131 uint64_t out[10] = { 0 };
132 size_t out_sz = sizeof(out);
133
134 kern_return_t ret = IOConnectCallMethod(dev->fd,
135 AGX_SELECTOR_ALLOCATE_MEM, NULL, 0, args_in,
136 sizeof(args_in), NULL, 0, out, &out_sz);
137
138 assert(ret == 0);
139 assert(out_sz == sizeof(out));
140 handle = (out[3] >> 32ull);
141 #else
142 /* Faked software path until we have a DRM driver */
143 handle = (++AGX_FAKE_HANDLE);
144 #endif
145
146 pthread_mutex_lock(&dev->bo_map_lock);
147 bo = agx_lookup_bo(dev, handle);
148 pthread_mutex_unlock(&dev->bo_map_lock);
149
150 /* Fresh handle */
151 assert(!memcmp(bo, &((struct agx_bo) {}), sizeof(*bo)));
152
153 bo->type = AGX_ALLOC_REGULAR;
154 bo->size = size;
155 bo->flags = flags;
156 bo->dev = dev;
157 bo->handle = handle;
158
159 ASSERTED bool lo = (flags & 0x08000000);
160
161 #if __APPLE__
162 bo->ptr.gpu = out[0];
163 bo->ptr.cpu = (void *) out[1];
164 bo->guid = out[5];
165 #else
166 if (lo) {
167 bo->ptr.gpu = AGX_FAKE_LO;
168 AGX_FAKE_LO += bo->size;
169 } else {
170 bo->ptr.gpu = AGX_FAKE_HI;
171 AGX_FAKE_HI += bo->size;
172 }
173
174 bo->ptr.gpu = (((uint64_t) bo->handle) << (lo ? 16 : 24));
175 bo->ptr.cpu = calloc(1, bo->size);
176 #endif
177
178 assert(bo->ptr.gpu < (1ull << (lo ? 32 : 40)));
179
180 return bo;
181 }
182
183 void
agx_bo_reference(struct agx_bo * bo)184 agx_bo_reference(struct agx_bo *bo)
185 {
186 if (bo) {
187 ASSERTED int count = p_atomic_inc_return(&bo->refcnt);
188 assert(count != 1);
189 }
190 }
191
192 void
agx_bo_unreference(struct agx_bo * bo)193 agx_bo_unreference(struct agx_bo *bo)
194 {
195 if (!bo)
196 return;
197
198 /* Don't return to cache if there are still references */
199 if (p_atomic_dec_return(&bo->refcnt))
200 return;
201
202 struct agx_device *dev = bo->dev;
203
204 pthread_mutex_lock(&dev->bo_map_lock);
205
206 /* Someone might have imported this BO while we were waiting for the
207 * lock, let's make sure it's still not referenced before freeing it.
208 */
209 if (p_atomic_read(&bo->refcnt) == 0) {
210 if (dev->debug & AGX_DBG_TRACE)
211 agxdecode_track_free(bo);
212
213 /* TODO: cache */
214 agx_bo_free(dev, bo);
215
216 }
217 pthread_mutex_unlock(&dev->bo_map_lock);
218 }
219
220 struct agx_bo *
agx_bo_create(struct agx_device * dev,unsigned size,unsigned flags)221 agx_bo_create(struct agx_device *dev, unsigned size, unsigned flags)
222 {
223 struct agx_bo *bo;
224 assert(size > 0);
225
226 /* To maximize BO cache usage, don't allocate tiny BOs */
227 size = ALIGN_POT(size, 4096);
228
229 /* TODO: Cache fetch */
230 bo = agx_bo_alloc(dev, size, flags);
231
232 if (!bo) {
233 fprintf(stderr, "BO creation failed\n");
234 return NULL;
235 }
236
237 p_atomic_set(&bo->refcnt, 1);
238
239 if (dev->debug & AGX_DBG_TRACE)
240 agxdecode_track_alloc(bo);
241
242 return bo;
243 }
244
245 static void
agx_get_global_ids(struct agx_device * dev)246 agx_get_global_ids(struct agx_device *dev)
247 {
248 #if __APPLE__
249 uint64_t out[2] = {};
250 size_t out_sz = sizeof(out);
251
252 ASSERTED kern_return_t ret = IOConnectCallStructMethod(dev->fd,
253 AGX_SELECTOR_GET_GLOBAL_IDS,
254 NULL, 0, &out, &out_sz);
255
256 assert(ret == 0);
257 assert(out_sz == sizeof(out));
258 assert(out[1] > out[0]);
259
260 dev->next_global_id = out[0];
261 dev->last_global_id = out[1];
262 #else
263 dev->next_global_id = 0;
264 dev->last_global_id = 0x1000000;
265 #endif
266 }
267
268 uint64_t
agx_get_global_id(struct agx_device * dev)269 agx_get_global_id(struct agx_device *dev)
270 {
271 if (unlikely(dev->next_global_id >= dev->last_global_id)) {
272 agx_get_global_ids(dev);
273 }
274
275 return dev->next_global_id++;
276 }
277
278 /* Tries to open an AGX device, returns true if successful */
279
280 bool
agx_open_device(void * memctx,struct agx_device * dev)281 agx_open_device(void *memctx, struct agx_device *dev)
282 {
283 #if __APPLE__
284 kern_return_t ret;
285
286 /* TODO: Support other models */
287 CFDictionaryRef matching = IOServiceNameMatching("AGXAcceleratorG13G_B0");
288
289 io_service_t service =
290 IOServiceGetMatchingService(kIOMasterPortDefault, matching);
291
292 if (!service)
293 return false;
294
295 ret = IOServiceOpen(service, mach_task_self(), AGX_SERVICE_TYPE, &dev->fd);
296
297 if (ret)
298 return false;
299
300 const char *api = "Equestria";
301 char in[16] = { 0 };
302 assert(strlen(api) < sizeof(in));
303 memcpy(in, api, strlen(api));
304
305 ret = IOConnectCallStructMethod(dev->fd, AGX_SELECTOR_SET_API, in,
306 sizeof(in), NULL, NULL);
307
308 /* Oddly, the return codes are flipped for SET_API */
309 if (ret != 1)
310 return false;
311 #else
312 /* Only open a fake AGX device on other operating systems if forced */
313 if (!getenv("AGX_FAKE_DEVICE"))
314 return false;
315 #endif
316
317 dev->memctx = memctx;
318 util_sparse_array_init(&dev->bo_map, sizeof(struct agx_bo), 512);
319
320 dev->queue = agx_create_command_queue(dev);
321 dev->cmdbuf = agx_shmem_alloc(dev, 0x4000, true); // length becomes kernelCommandDataSize
322 dev->memmap = agx_shmem_alloc(dev, 0x10000, false);
323 agx_get_global_ids(dev);
324
325 return true;
326 }
327
328 void
agx_close_device(struct agx_device * dev)329 agx_close_device(struct agx_device *dev)
330 {
331 util_sparse_array_finish(&dev->bo_map);
332
333 #if __APPLE__
334 kern_return_t ret = IOServiceClose(dev->fd);
335
336 if (ret)
337 fprintf(stderr, "Error from IOServiceClose: %u\n", ret);
338 #endif
339 }
340
341 #if __APPLE__
342 static struct agx_notification_queue
agx_create_notification_queue(mach_port_t connection)343 agx_create_notification_queue(mach_port_t connection)
344 {
345 struct agx_create_notification_queue_resp resp;
346 size_t resp_size = sizeof(resp);
347 assert(resp_size == 0x10);
348
349 ASSERTED kern_return_t ret = IOConnectCallStructMethod(connection,
350 AGX_SELECTOR_CREATE_NOTIFICATION_QUEUE,
351 NULL, 0, &resp, &resp_size);
352
353 assert(resp_size == sizeof(resp));
354 assert(ret == 0);
355
356 mach_port_t notif_port = IODataQueueAllocateNotificationPort();
357 IOConnectSetNotificationPort(connection, 0, notif_port, resp.unk2);
358
359 return (struct agx_notification_queue) {
360 .port = notif_port,
361 .queue = resp.queue,
362 .id = resp.unk2
363 };
364 }
365 #endif
366
367 struct agx_command_queue
agx_create_command_queue(struct agx_device * dev)368 agx_create_command_queue(struct agx_device *dev)
369 {
370 #if __APPLE__
371 struct agx_command_queue queue = {};
372
373 {
374 uint8_t buffer[1024 + 8] = { 0 };
375 const char *path = "/tmp/a.out";
376 assert(strlen(path) < 1022);
377 memcpy(buffer + 0, path, strlen(path));
378
379 /* Copy to the end */
380 unsigned END_LEN = MIN2(strlen(path), 1024 - strlen(path));
381 unsigned SKIP = strlen(path) - END_LEN;
382 unsigned OFFS = 1024 - END_LEN;
383 memcpy(buffer + OFFS, path + SKIP, END_LEN);
384
385 buffer[1024] = 0x2;
386
387 struct agx_create_command_queue_resp out = {};
388 size_t out_sz = sizeof(out);
389
390 ASSERTED kern_return_t ret = IOConnectCallStructMethod(dev->fd,
391 AGX_SELECTOR_CREATE_COMMAND_QUEUE,
392 buffer, sizeof(buffer),
393 &out, &out_sz);
394
395 assert(ret == 0);
396 assert(out_sz == sizeof(out));
397
398 queue.id = out.id;
399 assert(queue.id);
400 }
401
402 queue.notif = agx_create_notification_queue(dev->fd);
403
404 {
405 uint64_t scalars[2] = {
406 queue.id,
407 queue.notif.id
408 };
409
410 ASSERTED kern_return_t ret = IOConnectCallScalarMethod(dev->fd,
411 0x1D,
412 scalars, 2, NULL, NULL);
413
414 assert(ret == 0);
415 }
416
417 {
418 uint64_t scalars[2] = {
419 queue.id,
420 0x1ffffffffull
421 };
422
423 ASSERTED kern_return_t ret = IOConnectCallScalarMethod(dev->fd,
424 0x31,
425 scalars, 2, NULL, NULL);
426
427 assert(ret == 0);
428 }
429
430 return queue;
431 #else
432 return (struct agx_command_queue) {
433 0
434 };
435 #endif
436 }
437
438 void
agx_submit_cmdbuf(struct agx_device * dev,unsigned cmdbuf,unsigned mappings,uint64_t scalar)439 agx_submit_cmdbuf(struct agx_device *dev, unsigned cmdbuf, unsigned mappings, uint64_t scalar)
440 {
441 #if __APPLE__
442 struct agx_submit_cmdbuf_req req = {
443 .count = 1,
444 .command_buffer_shmem_id = cmdbuf,
445 .segment_list_shmem_id = mappings,
446 .notify_1 = 0xABCD,
447 .notify_2 = 0x1234,
448 };
449
450 ASSERTED kern_return_t ret = IOConnectCallMethod(dev->fd,
451 AGX_SELECTOR_SUBMIT_COMMAND_BUFFERS,
452 &scalar, 1,
453 &req, sizeof(req),
454 NULL, 0, NULL, 0);
455 assert(ret == 0);
456 return;
457 #endif
458 }
459
460 /*
461 * Wait for a frame to finish rendering.
462 *
463 * The macOS kernel indicates that rendering has finished using a notification
464 * queue. The kernel will send two messages on the notification queue. The
465 * second message indicates that rendering has completed. This simple routine
466 * waits for both messages. It's important that IODataQueueDequeue is used in a
467 * loop to flush the entire queue before calling
468 * IODataQueueWaitForAvailableData. Otherwise, we can race and get stuck in
469 * WaitForAvailabaleData.
470 */
471 void
agx_wait_queue(struct agx_command_queue queue)472 agx_wait_queue(struct agx_command_queue queue)
473 {
474 #if __APPLE__
475 uint64_t data[4];
476 unsigned sz = sizeof(data);
477 unsigned message_id = 0;
478 uint64_t magic_numbers[2] = { 0xABCD, 0x1234 };
479
480 while (message_id < 2) {
481 IOReturn ret = IODataQueueWaitForAvailableData(queue.notif.queue, queue.notif.port);
482
483 if (ret) {
484 fprintf(stderr, "Error waiting for available data\n");
485 return;
486 }
487
488 while (IODataQueueDequeue(queue.notif.queue, data, &sz) == kIOReturnSuccess) {
489 assert(sz == sizeof(data));
490 assert(data[0] == magic_numbers[message_id]);
491 message_id++;
492 }
493 }
494 #endif
495 }
496