codekingpro/portable-devtools
114k
1/*
2 * Copyright 2018-2023 NVIDIA Corporation. All rights reserved.
3 *
4 * NOTICE TO LICENSEE:
5 *
6 * This source code and/or documentation ("Licensed Deliverables") are
7 * subject to NVIDIA intellectual property rights under U.S. and
8 * international Copyright laws.
9 *
10 * These Licensed Deliverables contained herein is PROPRIETARY and
11 * CONFIDENTIAL to NVIDIA and is being provided under the terms and
12 * conditions of a form of NVIDIA software license agreement by and
13 * between NVIDIA and Licensee ("License Agreement") or electronically
14 * accepted by Licensee. Notwithstanding any terms or conditions to
15 * the contrary in the License Agreement, reproduction or disclosure
16 * of the Licensed Deliverables to any third party without the express
17 * written consent of NVIDIA is prohibited.
18 *
19 * NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE
20 * LICENSE AGREEMENT, NVIDIA MAKES NO REPRESENTATION ABOUT THE
21 * SUITABILITY OF THESE LICENSED DELIVERABLES FOR ANY PURPOSE. IT IS
22 * PROVIDED "AS IS" WITHOUT EXPRESS OR IMPLIED WARRANTY OF ANY KIND.
23 * NVIDIA DISCLAIMS ALL WARRANTIES WITH REGARD TO THESE LICENSED
24 * DELIVERABLES, INCLUDING ALL IMPLIED WARRANTIES OF MERCHANTABILITY,
25 * NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
26 * NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE
27 * LICENSE AGREEMENT, IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY
28 * SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL DAMAGES, OR ANY
29 * DAMAGES WHATSOEVER RESULTING FROM LOSS OF USE, DATA OR PROFITS,
30 * WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE OR OTHER TORTIOUS
31 * ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE OR PERFORMANCE
32 * OF THESE LICENSED DELIVERABLES.
33 *
34 * U.S. Government End Users. These Licensed Deliverables are a
35 * "commercial item" as that term is defined at 48 C.F.R. 2.101 (OCT
36 * 1995), consisting of "commercial computer software" and "commercial
37 * computer software documentation" as such terms are used in 48
38 * C.F.R. 12.212 (SEPT 1995) and is provided to the U.S. Government
39 * only as a commercial end item. Consistent with 48 C.F.R.12.212 and
40 * 48 C.F.R. 227.7202-1 through 227.7202-4 (JUNE 1995), all
41 * U.S. Government End Users acquire the Licensed Deliverables with
42 * only those rights set forth herein.
43 *
44 * Any use of the Licensed Deliverables in individual and commercial
45 * software must include, in the user documentation and internal
46 * comments to the code, the above Disclaimer and U.S. Government End
47 * Users Notice.
48 */
49
50#if !defined(__SANITIZER_PATCHING_H__)
51#define __SANITIZER_PATCHING_H__
52
53#include <sanitizer_memory.h>
54#include <sanitizer_result.h>
55
56#include <cuda.h>
57
58#include <stdint.h>
59
60#ifndef SANITIZERAPI
61#ifdef _WIN32
62#define SANITIZERAPI __stdcall
63#else
64#define SANITIZERAPI
65#endif
66#endif
67
68#if defined(__cplusplus)
69extern "C" {
70#endif
71
72/**
73 * \defgroup SANITIZER_PATCHING_API Sanitizer Patching API
74 * Functions, types, and enums that implement the Sanitizer Patching API.
75 * @{
76 */
77
78/**
79 * \addtogroup SANITIZER_PATCHING_API
80 * @{
81 */
82
83typedef struct Sanitizer_Launch_st *Sanitizer_LaunchHandle;
84
85/**
86 * \brief Load a module containing patches that can be used by the
87 * patching API.
88 *
89 * \note \b Thread-safety: an API user must serialize access to
90 * sanitizerAddPatchesFromFile, sanitizerAddPatches, sanitizerPatchInstructions,
91 * and sanitizerPatchModule. For example if sanitizerAddPatchesFromFile(filename)
92 * and sanitizerPatchInstruction(*, *, cbName) are called concurrently and
93 * cbName is intended to be found in the loaded module, the results are
94 * undefined.
95 *
96 * \note The patches loaded are only valid for the specified CUDA context.
97 *
98 * \param filename Path to the module file. This API supports the same module
99 * formats as the cuModuleLoad function from the CUDA driver API.
100 * \param ctx CUDA context in which to load the patches. If ctx is NULL, the
101 * current context will be used.
102 *
103 * \retval SANITIZER_SUCCESS on success
104 * \retval SANITIZER_ERROR_NOT_INITIALIZED if unable to initialize the sanitizer
105 * \retval SANITIZER_ERROR_INVALID_PARAMETER if \p filename is not a path to
106 * a valid CUDA module.
107 */
108SanitizerResult SANITIZERAPI sanitizerAddPatchesFromFile(const char* filename,
109 CUcontext ctx);
110
111/**
112 * \brief Load a module containing patches that can be used by the
113 * patching API.
114 *
115 * \note \b Thread-safety: an API user must serialize access to
116 * sanitizerAddPatchesFromFile, sanitizerAddPatches, sanitizerPatchInstructions,
117 * and sanitizerPatchModule. For example if sanitizerAddPatches(image) and
118 * sanitizerPatchInstruction(*, *, cbName) are called concurrently and cbName
119 * is intended to be found in the loaded image, the results are undefined.
120 *
121 * \note The patches loaded are only valid for the specified CUDA context.
122 *
123 * \param image Pointer to module data to load. This API supports the same
124 * module formats as the cuModuleLoadData and cuModuleLoadFatBinary functions
125 * from the CUDA driver API.
126 * \param ctx CUDA context in which to load the patches. If ctx is NULL, the
127 * current context will be used.
128 *
129 * \retval SANITIZER_SUCCESS on success
130 * \retval SANITIZER_ERROR_NOT_INITIALIZED if unable to initialize the sanitizer
131 * \retval SANITIZER_ERROR_INVALID_PARAMETER if \p image does not point to a
132 * valid CUDA module.
133 */
134SanitizerResult SANITIZERAPI sanitizerAddPatches(const void* image,
135 CUcontext ctx);
136
137/**
138 * \brief Sanitizer patch result codes
139 *
140 * Error and result codes returned by Sanitizer patches.
141 * If a patch returns SANITIZER_PATCH_ERROR, the thread
142 * will be exited. On Volta and newer architectures, the
143 * full warp which the thread belongs to will be exited.
144 */
145typedef enum {
146 /**
147 * No error.
148 */
149 SANITIZER_PATCH_SUCCESS = 0,
150
151 /**
152 * An error was detected in the patch.
153 */
154 SANITIZER_PATCH_ERROR = 1,
155
156 SANITIZER_PATCH_FORCE_INT = 0x7fffffff
157} SanitizerPatchResult;
158
159/**
160 * \brief Function type for a CUDA block enter callback.
161 *
162 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
163 * \p pc is the program counter of the entry point of the block
164 */
165typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackBlockEnter)(void* userdata, uint64_t pc);
166
167/**
168 * \brief Function type for a CUDA block exit callback.
169 *
170 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
171 * \p pc is the program counter of the patched instruction
172 */
173typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackBlockExit)(void* userdata, uint64_t pc);
174
175/**
176 * \brief Flags describing a memory access
177 *
178 * Flags describing a memory access. These values are to be or-combined in the
179 * value of \b flags for a SanitizerCallbackMemoryAccess callback.
180 */
181typedef enum {
182 /**
183 * Empty flag.
184 */
185 SANITIZER_MEMORY_DEVICE_FLAG_NONE = 0,
186
187 /**
188 * Specifies that the access is a read.
189 */
190 SANITIZER_MEMORY_DEVICE_FLAG_READ = 0x1,
191
192 /**
193 * Specifies that the access is a write.
194 */
195 SANITIZER_MEMORY_DEVICE_FLAG_WRITE = 0x2,
196
197 /**
198 * Specifies that the access is a system-scoped atomic.
199 */
200 SANITIZER_MEMORY_DEVICE_FLAG_ATOMSYS = 0x4,
201
202 /**
203 * Specifies that the access is a cache prefetch.
204 */
205 SANITIZER_MEMORY_DEVICE_FLAG_PREFETCH = 0x8,
206
207 SANITIZER_MEMORY_DEVICE_FLAG_FORCE_INT = 0x7fffffff
208} Sanitizer_DeviceMemoryFlags;
209
210/**
211 * \brief Function type for a memory access callback.
212 *
213 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
214 * \p pc is the program counter of the patched instruction
215 * \p ptr is the address of the memory being accessed.
216 * For local or shared memory access, this is the offset within the local
217 * or shared memory window.
218 * \p accessSize is the size of the access in bytes. Valid values are 1, 2, 4,
219 * 8, and 16.
220 * \p flags contains information about the type of access. See
221 * Sanitizer_DeviceMemoryFlags to interpret this value.
222 * \p pData is a pointer which value depends on the type of access:
223 * - If the access is a write, \p pData points to the new value being written.
224 * - If the access is a read and \p pData is not \p NULL, then it points to a
225 * 32-bit mask of loaded bytes being used (padding bytes will not appear).
226 * - If the access is an atomic, the pointer will be \p NULL.
227 */
228typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackMemoryAccess)(void* userdata, uint64_t pc, void* ptr, uint32_t accessSize, uint32_t flags, const void* pData);
229
230/**
231 * \brief Flags describing a barrier
232 *
233 * Flags describing a barrier. These values are to be or-combined in the
234 * value of \b flags for a SanitizerCallbackBarrier callback.
235 */
236typedef enum {
237 /**
238 * Empty flag.
239 */
240 SANITIZER_BARRIER_FLAG_NONE = 0,
241
242 /**
243 * Specifies that the barrier can be called unaligned.
244 * This flag is only valid on SM 7.0 and above.
245 */
246 SANITIZER_BARRIER_FLAG_UNALIGNED_ALLOWED = 0x1,
247
248 SANITIZER_BARRIER_FLAG_FORCE_INT = 0x7fffffff
249} Sanitizer_BarrierFlags;
250
251/**
252 * \brief Function type for a barrier callback.
253 *
254 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
255 * \p pc is the program counter of the patched instruction
256 * \p barIndex is the barrier index.
257 * \p threadCount is the number of expected threads (must be a multiple of the warp size).
258 * \p flags contains information about the barrier.
259 * See Sanitizer_BarrierFlags to interpret this value.
260 * 0 means that all threads are participating in the barrier.
261 */
262typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackBarrier)(void* userdata, uint64_t pc, uint32_t barIndex, uint32_t threadCount, uint32_t flags);
263
264/**
265 * \brief Function type for a syncwarp callback.
266 *
267 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
268 * \p pc is the program counter of the patched instruction
269 * \p mask is the thread mask passed to __syncwarp().
270 */
271typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackSyncwarp)(void* userdata, uint64_t pc, uint32_t mask);
272
273/**
274 * \brief Function type for a shfl callback.
275 *
276 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
277 * \p pc is the program counter of the patched instruction
278 *
279 */
280typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackShfl)(void* userdata, uint64_t pc);
281
282/**
283 * \brief Flags describing a function call
284 *
285 * Flags describing a function call. These values are to be or-combined in the
286 * value of \b flags for a SanitizerCallbackCall callback.
287 */
288typedef enum {
289 /**
290 * Empty flag.
291 */
292 SANITIZER_CALL_FLAG_NONE = 0,
293
294 /**
295 * Specifies that barriers within this function call can be called unaligned.
296 * This flag is only valid on SM 7.0 and above.
297 */
298 SANITIZER_CALL_FLAG_UNALIGNED_ALLOWED = 0x1,
299
300 SANITIZER_CALL_FLAG_FORCE_INT = 0x7fffffff
301} Sanitizer_CallFlags;
302
303/**
304 * \brief Function type for a function call callback.
305 *
306 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
307 * \p pc is the program counter of the patched instruction
308 * \p targetPc is the PC where the called function is located.
309 * \p flags contains information about the function call.
310 */
311typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackCall)(void* userdata, uint64_t pc, uint64_t targetPc, uint32_t flags);
312
313/**
314 * \brief Function type for a function return callback.
315 *
316 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
317 * \p pc is the program counter of the patched instruction
318 *
319 */
320typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackRet)(void* userdata, uint64_t pc);
321
322/**
323 * \brief Function type for a device-side malloc call.
324 *
325 * \note This is called after the call has completed.
326 *
327 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
328 * \p pc is the program counter of the patched instruction
329 * \p allocatedPtr is the pointer returned by device-side malloc
330 * \p allocatedSize is the size requested by the user to device-side malloc.
331 */
332typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackDeviceSideMalloc)(void* userdata, uint64_t pc, void* allocatedPtr, uint64_t allocatedSize);
333
334/**
335 * \brief Function type for a device-side free call.
336 *
337 * \note This is called prior to the actual call.
338 *
339 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
340 * \p pc is the program counter of the patched instruction
341 * \p ptr is the pointer passed to device-side free.
342 */
343typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackDeviceSideFree)(void* userdata, uint64_t pc, void* ptr);
344
345/**
346 * \brief CUDA Barrier action kind.
347 *
348 * Refer to the CUDA Barrier interface section of the CUDA toolkit documentation for
349 * a more extensive description of these actions.
350 */
351typedef enum {
352
353 /**
354 * Invalid action ID.
355 */
356 SANITIZER_CUDA_BARRIER_INVALID = 0,
357
358 /**
359 * Barrier initialization.
360 */
361 SANITIZER_CUDA_BARRIER_INIT = 1,
362
363 /**
364 * Barrier arrive operation. On Hopper and newer architectures,
365 * barrier data is the count argument to the arrive-on operation.
366 */
367 SANITIZER_CUDA_BARRIER_ARRIVE = 2,
368
369 /**
370 * Barrier arrive and drop operation. On Hopper and newer architectures,
371 * barrier data is the count argument to the arrive-on operation.
372 */
373 SANITIZER_CUDA_BARRIER_ARRIVE_DROP = 3,
374
375 /**
376 * Barrier arrive operation without phase completion.
377 * Barrier data is the count argument to the arrive-on operation.
378 */
379 SANITIZER_CUDA_BARRIER_ARRIVE_NOCOMPLETE = 4,
380
381 /**
382 * Barrier arrive and drop operation without phase completion.
383 * Barrier data is the count argument to the arrive-on operation.
384 */
385 SANITIZER_CUDA_BARRIER_ARRIVE_DROP_NOCOMPLETE = 5,
386
387 /**
388 * Barrier wait operation.
389 */
390 SANITIZER_CUDA_BARRIER_WAIT = 6,
391
392 /**
393 * Barrier invalidation.
394 */
395 SANITIZER_CUDA_BARRIER_INVALIDATE = 7,
396
397 SANITIZER_CUDA_BARRIER_FORCE_INT = 0x7fffffff
398} Sanitizer_CudaBarrierInstructionKind;
399
400/**
401 * \brief Function type for a CUDA Barrier action callback.
402 *
403 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
404 * \p pc is the program counter of the patched instruction
405 * \p barrier Barrier address which can be used as a unique identifier
406 * \p kind Barrier action type. See \ref Sanitizer_CudaBarrierInstructionKind
407 * \p data Barrier data. This is specific to each action type, refer to \ref Sanitizer_CudaBarrierInstructionKind
408 */
409typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackCudaBarrier)(void* userdata, uint64_t pc, void* barrier, uint32_t kind, uint32_t data);
410
411/**
412 * \brief Function type for a global to shared memory asynchronous copy.
413 *
414 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
415 * \p pc is the program counter of the patched instruction
416 * \p src is the address of the global memory being read.
417 * This can be NULL if src-size is 0.
418 * \p dst is the address of the shared memory being written.
419 * This is an offset within the shared memory window
420 * \p accessSize is the size of the access in bytes. Valid values are 4, 8 and 16.
421 */
422typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackMemcpyAsync)(void* userdata, uint64_t pc, void* src, uint32_t dst, uint32_t accessSize);
423
424/**
425 * \brief Function type for a pipeline commit
426 *
427 * This can be generated by a pipeline::producer_commit (C++ API), a pipeline_commit (C API)
428 * or a cp.async.commit_group (PTX API).
429 *
430 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
431 * \p pc is the program counter of the patched instruction
432 */
433typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackPipelineCommit)(void* userdata, uint64_t pc);
434
435/**
436 * \brief Function type for a pipeline wait
437 *
438 * This can be generated by a pipeline::consumer_wait (C++ API), a pipeline_wait_prior (C API),
439 * cp.async.wait_group or cp.async.wait_all (PTX API).
440 *
441 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
442 * \p pc is the program counter of the patched instruction
443 * \p groups is the number of groups the pipeline will wait for. 0 is used to wait for all groups.
444 */
445typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackPipelineWait)(void* userdata, uint64_t pc, uint32_t groups);
446
447/**
448 * \brief Function type for a matrix shared memory access callback
449 *
450 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
451 * \p pc is the program counter of the patched instruction
452 * \p address is the address of the shared memory being read or written.
453 * This is an offset within the shared memory window
454 * \p accessSize is the size of the access in bytes. Valid value is 16.
455 * \p flags contains information about the type of access. See
456 * Sanitizer_DeviceMemoryFlags to interpret this value.
457 * \p count is the number of matrices accessed.
458 * \p pNewValue is a pointer to the new value being written if the access is a
459 * write. If the access is a read or an atomic, the pointer will be NULL.
460 */
461typedef SanitizerPatchResult(SANITIZERAPI* SanitizerCallbackMatrixMemoryAccess)(void* userdata, uint64_t pc, uint32_t address, uint32_t accessSize, uint32_t flags, uint32_t count, const void* pNewValue);
462
463/**
464 * \brief Cache control action
465 */
466typedef enum {
467
468 /**
469 * Invalid action ID.
470 */
471 SANITIZER_CACHE_CONTROL_INVALID = 0,
472
473 /**
474 * Prefetch to L1.
475 */
476 SANITIZER_CACHE_CONTROL_L1_PREFETCH = 1,
477
478 /**
479 * Prefetch to L2.
480 */
481 SANITIZER_CACHE_CONTROL_L2_PREFETCH = 2,
482
483 SANITIZER_CACHE_CONTROL_FORCE_INT = 0x7fffffff
484} Sanitizer_CacheControlInstructionKind;
485
486/**
487 * \brief Function type for a cache control instruction callback
488 *
489 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
490 * \p pc is the program counter of the patched instruction
491 * \p address is the address of the memory being controlled
492 * \p kind is the type of cache control. See \ref Sanitizer_CacheControlInstructionKind
493 */
494typedef SanitizerPatchResult(SANITIZERAPI* SanitizerCallbackCacheControl)(void* userdata, uint64_t pc, void* address, Sanitizer_CacheControlInstructionKind kind);
495
496/**
497 * \brief Function type for a cluster barrier arrive.
498 *
499 * This can be generated by a cg::this_cluster().sync() (C++ API), or a
500 * barrier.cluster.arrive (PTX API).
501 *
502 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
503 * \p pc is the program counter of the patched instruction
504 */
505typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackClusterBarrierArrive)(void* userdata, uint64_t pc);
506
507/**
508 * \brief Function type for a cluster barrier wait.
509 *
510 * This can be generated by a cg::this_cluster().sync() (C++ API), or a
511 * barrier.cluster.wait (PTX API).
512 *
513 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
514 * \p pc is the program counter of the patched instruction
515 */
516typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackClusterBarrierArrive)(void* userdata, uint64_t pc);
517
518/**
519 * \brief Flags describing a warpgroup aligned MMA async.
520 *
521 * Flags describing a warpgroup aligned MMA async. These values are to be or-combined in the
522 * value of \b flags for a SanitizerCallbackWarpgroupMMAAsync callback.
523 */
524typedef enum {
525 /**
526 * Empty flag.
527 */
528 SANITIZER_WARPGROUP_MMA_ASYNC_FLAG_NONE = 0,
529
530 /**
531 * Specifies that the MMA async delimits a MMA async group of which it is
532 * the last instruction. Please refer to the PTX documentation for wgmma_async.commit_group
533 * for more details. This property is valid even if the warpMask is zero.
534 */
535 SANITIZER_WARPGROUP_MMA_ASYNC_FLAG_COMMIT_GROUP = 0x1,
536
537 SANITIZER_WARPGROUP_MMA_ASYNC_FLAG_FORCE_INT = 0x7fffffff
538} Sanitizer_WarpgroupMMAAsyncFlags;
539
540/**
541 * \brief Function type for a warpgroup aligned async MMA.
542 *
543 * This can be generated by a wgmma.mma_async in PTX.
544 *
545 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule
546 * \p pc is the program counter of the patched instruction
547 * \p addressMatrixA is the address in shared memory of the matrix A being read. This field is only valid if sizeMatrixA is non-zero and warpMask is full.
548 * \p sizeMatrixA is the size of the matrix A in shared memory. A value of 0 means that the matrix A is read from registers instead.
549 * \p addressMatrixB is the address in shared memory of the matrix B being read. This field is only valid if warpMask is full.
550 * \p sizeMatrixB is the size of the matrix B in shared memory. The value will always be non-zero.
551 * \p flags of type Sanitizer_WarpgroupMMAAsyncFlags provide information about the access. These flags are to be taken into account even if the warpMask is zero.
552 * \p warpMask is a mask of threads that will perform the operation and read the operands. Expected values are either 0x0 or 0xffffffff (full). The value is expected to be the same across the warpgroup.
553 * Other values can be reported but signal a programming error in the target application.
554 */
555typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackWarpgroupMMAAsync)(void* userdata, uint64_t pc, uint32_t addressMatrixA, uint32_t sizeMatrixA, uint32_t addressMatrixB, uint32_t sizeMatrixB, uint32_t flags, uint32_t warpMask);
556
557/**
558 * \brief Function type for a warpgroup MMA wait group
559 *
560 * This can be generated by a wgmma.wait_group in PTX.
561 *
562 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule.
563 * \p pc is the program counter of the patched instruction.
564 * \p numGroups is the maximum number of group that will be left pending after the operation. A value of zero means that all MMA async of the warpgroup are guaranteed to have completed after the operation.
565 * \p warpMask is a mask of threads for which the expected values are either 0x0 or 0xffffffff (full). The value is expected to be the same across the warpgroup.
566 * Other values can be reported but signal a programming error in the target application. If the value is valid, the value has no influence on the operation.
567 */
568typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackWarpgroupWaitGroup)(void* userdata, uint64_t pc, uint32_t numGroups, uint32_t warpMask);
569
570/**
571 * \brief Function type for a warpgroup MMA fence
572 *
573 * This can be generated by a wgmma.fence in PTX.
574 *
575 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule.
576 * \p pc is the program counter of the patched instruction.
577 * \p warpMask is a mask of threads that will perform the fence operation. Expected values are either 0x0 or 0xffffffff (full). The value is expected to be the same across the warpgroup.
578 * Other values can be reported but signal a programming error in the target application.
579 */
580typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackWarpgroupFence)(void* userdata, uint64_t pc, uint32_t warpMask);
581
582/**
583 * \brief Function type for an asynchronous store operation on shared memory
584 *
585 * This can be generated by a st.async PTX instruction
586 *
587 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule.
588 * \p pc is the program counter of the patched instruction.
589 * \p address is the destination address in shared memory.
590 * \p mbarAddress is the address of the mbarrier object.
591 * \p pNewValue is a pointer to the new value being written.
592 * \p accessSize is the size of the access in bytes. Valid values are 4 and 8.
593 *
594 */
595typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackAsyncStore)(void* userdata, uint64_t pc, uint32_t address, uint32_t mbarAddress, void* pNewValue, uint32_t accessSize);
596
597/**
598 * \brief Function type for an asynchronous reduction operation on shared memory
599 *
600 * This can be generated by a red.async PTX instruction
601 *
602 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule.
603 * \p pc is the program counter of the patched instruction.
604 * \p address is the destination address in shared memory.
605 * \p mbarAddress is the address of the mbarrier object.
606 * \p accessSize is the size of the access in bytes. Valid values are 4 and 8.
607 *
608 */
609typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackAsyncReduction)(void* userdata, uint64_t pc, uint32_t address, uint32_t mbarAddress, uint32_t accessSize);
610
611/**
612 * \brief Function type for setting the shared memory size allocated to a block
613 *
614 * This can be generated by a setsmemsize.sync instruction
615 *
616 * \p userdata is a pointer to user data. See \ref sanitizerPatchModule.
617 * \p pc is the program counter of the patched instruction.
618 * \p size is the requested size in bytes.
619 *
620 */
621typedef SanitizerPatchResult (SANITIZERAPI *SanitizerCallbackSetSmemSize)(void* userdata, uint64_t pc, uint32_t size);
622
623/**
624 * \brief Instrumentation.
625 *
626 * Instrumentation. Every entry represent an instruction type or a function
627 * call where a callback patch can be inserted.
628 */
629typedef enum {
630 /**
631 * Invalid instruction ID.
632 */
633 SANITIZER_INSTRUCTION_INVALID = 0,
634
635 /**
636 * CUDA block enter. This is called prior to any user code. The type of the
637 * callback must be SanitizerCallbackBlockEnter.
638 */
639 SANITIZER_INSTRUCTION_BLOCK_ENTER = 1,
640
641 /**
642 * CUDA block exit. This is called after all user code has executed. The type of
643 * the callback must be SanitizerCallbackBlockExit.
644 */
645 SANITIZER_INSTRUCTION_BLOCK_EXIT = 2,
646
647 /**
648 * Global Memory Access. This can be a store, load or atomic operation. The type
649 * of the callback must be SanitizerCallbackMemoryAccess.
650 */
651 SANITIZER_INSTRUCTION_GLOBAL_MEMORY_ACCESS = 3,
652
653 /**
654 * Shared Memory Access. This can be a store, load or atomic operation. The type
655 * of the callback must be SanitizerCallbackMemoryAccess.
656 */
657 SANITIZER_INSTRUCTION_SHARED_MEMORY_ACCESS = 4,
658
659 /**
660 * Local Memory Access. This can be a store or load operation. The type
661 * of the callback must be SanitizerCallbackMemoryAccess.
662 */
663 SANITIZER_INSTRUCTION_LOCAL_MEMORY_ACCESS = 5,
664
665 /**
666 * Barrier. The type of the callback must be SanitizerCallbackBarrier.
667 */
668 SANITIZER_INSTRUCTION_BARRIER = 6,
669
670 /**
671 * Syncwarp. The type of the callback must be SanitizerCallbackSyncwarp.
672 */
673 SANITIZER_INSTRUCTION_SYNCWARP = 7,
674
675 /**
676 * Shfl. The type of the callback must be SanitizerCallbackShfl.
677 */
678 SANITIZER_INSTRUCTION_SHFL = 8,
679
680 /**
681 * Function call. The type of the callback must be SanitizerCallbackCall.
682 */
683 SANITIZER_INSTRUCTION_CALL = 9,
684
685 /**
686 * Function return. The type of the callback must be SanitizerCallbackRet.
687 */
688 SANITIZER_INSTRUCTION_RET = 10,
689
690 /**
691 * Device-side malloc. The type of the callback must be
692 * SanitizerCallbackDeviceSideMalloc.
693 */
694 SANITIZER_INSTRUCTION_DEVICE_SIDE_MALLOC = 11,
695
696 /**
697 * Device-side free. The type of the callback must be
698 * SanitizerCallbackDeviceSideFree.
699 */
700 SANITIZER_INSTRUCTION_DEVICE_SIDE_FREE = 12,
701
702 /**
703 * CUDA Barrier operation. The type of the callback must be
704 * SanitizerCallbackCudaBarrier.
705 */
706 SANITIZER_INSTRUCTION_CUDA_BARRIER = 13,
707
708 /**
709 * Global to shared memory asynchronous copy. The type of the
710 * callback must be SanitizerCallbackMemcpyAsync.
711 */
712 SANITIZER_INSTRUCTION_MEMCPY_ASYNC = 14,
713
714 /**
715 * Pipeline commit. The type of the callback must be
716 * SanitizerCallbackPipelineCommit.
717 */
718 SANITIZER_INSTRUCTION_PIPELINE_COMMIT = 15,
719
720 /**
721 * Pipeline wait. The type of the callback must be
722 * SanitizerCallbackPipelineWait.
723 */
724 SANITIZER_INSTRUCTION_PIPELINE_WAIT = 16,
725
726 /**
727 * Remote Shared Memory Access. This can be a store or load operation. The
728 * type of the callback must be SanitizerCallbackMemoryAccess.
729 */
730 SANITIZER_INSTRUCTION_REMOTE_SHARED_MEMORY_ACCESS = 17,
731
732 /**
733 * Device-side aligned malloc. The type of the callback must be
734 * SanitizerCallbackDeviceSideMalloc.
735 */
736 SANITIZER_INSTRUCTION_DEVICE_ALIGNED_MALLOC = 18,
737
738 /**
739 * Matrix shared memory access. The type of the callback must
740 * be SanitizerCallbackMatrixMemoryAccess.
741 */
742 SANITIZER_INSTRUCTION_MATRIX_MEMORY_ACCESS = 19,
743
744 /**
745 * Cache control instruction. The type of the callback must
746 * be SanitizerCallbackCacheControl.
747 */
748 SANITIZER_INSTRUCTION_CACHE_CONTROL = 20,
749
750 /**
751 * Cluster barrier arrive instruction. The type of the callback must
752 * be SanitizerCallbackClusterBarrierArrive.
753 */
754 SANITIZER_INSTRUCTION_CLUSTER_BARRIER_ARRIVE = 21,
755
756 /**
757 * Cluster barrier wait instruction. The type of the callback must
758 * be SanitizerCallbackClusterBarrierWait.
759 */
760 SANITIZER_INSTRUCTION_CLUSTER_BARRIER_WAIT = 22,
761
762 /**
763 * Warpgroup aligned async MMA instruction. The type of the callback must
764 * be SanitizerCallbackWarpgroupMMAAsync.
765 */
766 SANITIZER_INSTRUCTION_WARPGROUP_MMA_ASYNC = 23,
767
768 /**
769 * Warpgroup wait MMA group instruction. The type of the callback must
770 * be SanitizerCallbackWarpgroupWaitGroup.
771 */
772 SANITIZER_INSTRUCTION_WARPGROUP_WAIT_GROUP = 24,
773
774 /**
775 * Warpgroup fence instruction. The type of the callback must
776 * be SanitizerCallbackWarpgroupFence.
777 */
778 SANITIZER_INSTRUCTION_WARPGROUP_FENCE = 25,
779
780 /**
781 * Asynchronous store instruction. The type of the callback must
782 * be SanitizerCallbackAsyncStore.
783 */
784 SANITIZER_INSTRUCTION_ASYNC_STORE = 26,
785
786 /**
787 * Asynchronous reduction instruction. The type of the callback must
788 * be SanitizerCallbackAsyncReduction.
789 */
790 SANITIZER_INSTRUCTION_ASYNC_REDUCTION = 27,
791
792 /**
793 * Set the shared memory size allocated to a block instruction.
794 * The type of the callback must SanitizerCallbackSetSmemSize
795 */
796 SANITIZER_INSTRUCTION_SET_SHARED_MEMORY_SIZE = 28,
797
798 /**
799 * Barrier after it is released. The type of the callback must be
800 * SanitizerCallbackBarrier.
801 */
802 SANITIZER_INSTRUCTION_BARRIER_RELEASE = 29,
803
804 SANITIZER_INSTRUCTION_FORCE_INT = 0x7fffffff
805} Sanitizer_InstructionId;
806
807/**
808 * \brief Set instrumentation points and patches to be applied in a module.
809 *
810 * Mark that all instrumentation points matching instructionId are to be
811 * patched in order to call the device function identified by
812 * deviceCallbackName. It is up to the API client to ensure that this
813 * device callback exists and match the correct callback format for
814 * this instrumentation point.
815 * \note \b Thread-safety: an API user must serialize access to
816 * sanitizerAddPatchesFromFile, sanitizerAddPatches, sanitizerPatchInstructions,
817 * and sanitizerPatchModule. For example if sanitizerAddPatches(fileName) and
818 * sanitizerPatchInstruction(*, *, cbName) are called concurrently and cbName
819 * is intended to be found in the loaded module, the results are undefined.
820 *
821 * \param instructionId Instrumentation point for which to insert patches
822 * \param module CUDA module to instrument
823 * \param deviceCallbackName Name of the device function callback that the
824 * inserted patch will call at the instrumented points. This function is
825 * expected to be found in code previously loaded by sanitizerAddPatchesFromFile
826 * or sanitizerAddPatches.
827 *
828 * \retval SANITIZER_SUCCESS on success
829 * \retval SANITIZER_ERROR_NOT_INITIALIZED if unable to initialize the sanitizer
830 * \retval SANITIZER_ERROR_INVALID_PARAMETER if \p module is not a CUDA module
831 * or if \p deviceCallbackName function cannot be located.
832 */
833SanitizerResult SANITIZERAPI sanitizerPatchInstructions(const Sanitizer_InstructionId instructionId,
834 CUmodule module,
835 const char* deviceCallbackName);
836
837/**
838 *
839 * \brief Perform the actual instrumentation of a module.
840 *
841 * Perform the instrumentation of a CUDA module based on previous calls to
842 * sanitizerPatchInstructions. This function also specifies the device memory
843 * buffer to be passed in as userdata to all callback functions.
844 * \note \b Thread-safety: an API user must serialize access to
845 * sanitizerAddPatchesFromFile, sanitizerAddPatches, sanitizerPatchInstructions,
846 * and sanitizerPatchModule. For example if sanitizerPatchModule(mod, *) and
847 * sanitizerPatchInstruction(*, mod, *) are called concurrently, the results
848 * are undefined.
849 *
850 * \param module CUDA module to instrument
851 *
852 * \retval SANITIZER_SUCCESS on success
853 * \retval SANITIZER_ERROR_INVALID_PARAMETER if \p module is not a CUDA module
854 */
855SanitizerResult SANITIZERAPI sanitizerPatchModule(CUmodule module);
856
857/**
858 * \brief Specifies the user data pointer for callbacks
859 *
860 * Mark all subsequent launches of \p kernel to use \p userdata
861 * pointer as the device memory buffer to pass in to callback functions.
862 *
863 * \param kernel CUDA function to link to user data. Callbacks in subsequent
864 * launches on this kernel will use \p userdata as callback data.
865 * \param userdata Device memory buffer. This data will be passed to callback
866 * functions via the \p userdata parameter.
867 *
868 * \retval SANITIZER_SUCCESS on success
869 */
870SanitizerResult SANITIZERAPI sanitizerSetCallbackData(CUfunction kernel,
871 const void* userdata);
872
873/**
874 * \brief Specifies the user data pointer for callbacks
875 *
876 * Mark \p launch to use \p userdata pointer as the device memory buffer
877 * to pass in to callback functions. This function is only available if
878 * the driver version is 455 or newer.
879 *
880 * \param launch Kernel launch to link to user data. Callbacks in this kernel
881 * launch will use \p userdata as callback data.
882 * \param kernel CUDA function associated with the kernel launch.
883 * \param stream CUDA stream associated with the stream launch.
884 * \param userdata Device memory buffer. This data will be passed to callback
885 * functions via the \p userdata parameter.
886 *
887 * \retval SANITIZER_SUCCESS on success
888 */
889SanitizerResult SANITIZERAPI sanitizerSetLaunchCallbackData(Sanitizer_LaunchHandle launch,
890 CUfunction kernel,
891 Sanitizer_StreamHandle stream,
892 const void* userdata);
893
894/**
895 * \brief Specifies the user data pointer accessible from callbacks in the
896 * device-launched graphs launched by the specified host-launched graphExec.
897 *
898 * Mark all subsequent launch of \p graphExec to make available \p userdata
899 * in device callbacks from device-launched graphs. \p userdata will not
900 * be set in the callback userdata parameter but must be accessed through
901 * another mean instead. Please refer to the Sanitizer API reference manual.
902 * This function is only available if the driver version is 535 or newer.
903 *
904 * \param graphExec CUDA graphExec that will launch CUDA graphs from the device.
905 * \param stream CUDA stream associated with the stream launch.
906 * \param userdata Device memory buffer.
907 *
908 * \retval SANITIZER_SUCCESS on success
909 */
910SanitizerResult SANITIZERAPI sanitizerSetDeviceGraphData(CUgraphExec graphExec,
911 Sanitizer_StreamHandle stream,
912 const void* userdata);
913
914/**
915 *
916 * \brief Remove existing instrumentation of a module
917 *
918 * Remove any instrumentation of a CUDA module performed by previous calls
919 * to sanitizerPatchModule.
920 * \note \b Thread-safety: an API user must serialize access to
921 * sanitizerPatchModule and sanitizerUnpatchModule on the same module.
922 * For example, if sanitizerPatchModule(mod) and sanitizerUnpatchModule(mod)
923 * are called concurrently, the results are undefined.
924 *
925 * \param module CUDA module on which to remove instrumentation
926 *
927 * \retval SANITIZER_SUCCESS on success
928 */
929SanitizerResult SANITIZERAPI sanitizerUnpatchModule(CUmodule module);
930
931/**
932 *
933 * \brief Get PC and size of a CUDA function
934 *
935 * \param[in] module CUDA module containing the function
936 * \param[in] deviceCallbackName CUDA function name
937 * \param[out] pc Function start program counter (PC) returned
938 * \param[out] size Function size in bytes returned
939 *
940 * \retval SANITIZER_SUCCESS on success
941 * \retval SANITIZER_ERROR_INVALID_PARAMETER if \p functionName function
942 * cannot be located, if pc is NULL or if size is NULL.
943 *
944 */
945SanitizerResult SANITIZERAPI sanitizerGetFunctionPcAndSize(CUmodule module,
946 const char* functionName,
947 uint64_t* pc,
948 uint64_t* size);
949
950/**
951 *
952 * \brief Get PC and size of a device callback
953 *
954 * \param[in] ctx CUDA context in which the patches were loaded.
955 * If ctx is NULL, the current context will be used.
956 * \param[in] deviceCallbackName device function callback name
957 * \param[out] pc Callback PC returned
958 * \param[out] size Callback size returned
959 *
960 * \retval SANITIZER_SUCCESS on success
961 * \retval SANITIZER_ERROR_INVALID_PARAMETER if \p deviceCallbackName function
962 * cannot be located, if pc is NULL or if size is NULL.
963 *
964 */
965SanitizerResult SANITIZERAPI sanitizerGetCallbackPcAndSize(CUcontext ctx,
966 const char* deviceCallbackName,
967 uint64_t* pc,
968 uint64_t* size);
969
970typedef enum {
971 /**
972 * The function is not loaded.
973 */
974 SANITIZER_FUNCTION_NOT_LOADED = 0x0,
975
976 /**
977 * The function is being loaded.
978 */
979 SANITIZER_FUNCTION_PARTIALLY_LOADED = 0x1,
980
981 /**
982 * The function is fully loaded.
983 */
984 SANITIZER_FUNCTION_LOADED = 0x2,
985
986 SANITIZER_FUNCTION_LOADED_FORCE_INT = 0x7fffffff
987} Sanitizer_FunctionLoadedStatus;
988
989/**
990 *
991 * \brief Get the loading status of a function. Requires a driver version >=515.
992 *
993 * \param[in] func CUDA function for which the loading status is queried.
994 * \param[out] loadingStatus Loading status returned
995 *
996 * \retval SANITIZER_SUCCESS on success
997 * \retval SANITIZER_ERROR_INVALID_PARAMETER if \p func is NULL or if
998 * loadingStatus is NULL.
999 * \retval SANITIZER_ERROR_NOT_SUPPORTED if the loading status cannot be queried
1000 * with this driver version.
1001 *
1002 */
1003SanitizerResult SANITIZERAPI sanitizerGetFunctionLoadedStatus(CUfunction func,
1004 Sanitizer_FunctionLoadedStatus* loadingStatus);
1005
1006
1007/** @} */ /* END SANITIZER_PATCHING_API */
1008
1009#if defined(__cplusplus)
1010}
1011#endif
1012
1013#endif /* __SANITIZER_PATCHING_H__ */
1014