GV

GeezerVison

The World Wide Web, restored to 800 pixels of glory.

Home|About|Gaming cabinet|Archives|How it works|Starter sites

On this site

  • Sign in
  • ActionsCI/CD and automation
  • CodespacesCloud development environments
  • All Features
  • Documentation
  • Training
  • Certifications
  • Build with Copilot
  • Customize Copilot
  • Integrations
  • Trending repos
  • GitHub Sponsors
  • Developer articles
  • Enterprises

This view

  • Original page
  • New address

GeezerVison is showing github.com — GPU Metal: optionally fail when a kernel ran on an unmapped memory pool · AliceO2Group/AliceO2@5dffc90

File tree

Expand file treeCollapse file tree
Open diff view settings
Filter options
  • GPU/GPUTracking/Base/metal
    • GPUReconstructionMETAL.metal
    • GPUReconstructionMetal.mm
    • GPUReconstructionMetalIncludesHost.h
    • GPUReconstructionMetalKernels.mm
Expand file treeCollapse file tree
Open diff view settings
Collapse file

‎GPU/GPUTracking/Base/metal/GPUReconstructionMETAL.metal‎

Copy file name to clipboardExpand all lines: GPU/GPUTracking/Base/metal/GPUReconstructionMETAL.metal
+9Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -82,5 +82,14 @@ using namespace metal;
8282

8383
#include "GPUReconstructionKernelList.h"
8484

85+
// Reads a word of the memory pool through its address, as the kernels do, and
86+
// counts a read that does not return it; the host fails the run on a count.
87+
kernel void krnl_metalCanary(device atomic_uint* bad [[buffer(0)]], constant uint64_t& word [[buffer(1)]])
88+
{
89+
if (*(device uint32_t*)word != 0xA11CE0A5u) {
90+
atomic_fetch_add_explicit(bad, 1u, memory_order_relaxed);
91+
}
92+
}
93+
8594
// clang-format on
8695
#pragma clang diagnostic pop
Collapse file

‎GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm‎

Copy file name to clipboardExpand all lines: GPU/GPUTracking/Base/metal/GPUReconstructionMetal.mm
+21-1Lines changed: 21 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -185,7 +185,12 @@ uint32_t metalGPUCoreCount()
185185
// The kernels reconstruct device pointers from these bases, so they must be
186186
// the addresses the GPU sees, not the CPU-visible ones of a shared buffer.
187187
mDeviceMemoryBase = (void*)(uintptr_t)mInternals->mem_gpu.gpuAddress;
188-
mDeviceMemorySize = mInternals->mem_gpu.allocatedSize;
188+
mDeviceMemorySize = mInternals->mem_gpu.allocatedSize - sizeof(uint32_t);
189+
*(uint32_t*)((char*)mInternals->mem_gpu.contents + mDeviceMemorySize) = mInternals->canaryMagic;
190+
if (getenv("GPUCA_METAL_CANARY")) {
191+
mInternals->canaryBad = [mInternals->device newBufferWithLength:sizeof(uint32_t) options:MTLResourceStorageModeShared];
192+
*(uint32_t*)mInternals->canaryBad.contents = 0;
193+
}
189194
mDeviceConstantMem = (GPUConstantMem*)(uintptr_t)mInternals->mem_constant.gpuAddress;
190195

191196
if (GetProcessingSettings().debugLevel >= 1) {
@@ -233,6 +238,7 @@ uint32_t metalGPUCoreCount()
233238
if (mDeviceMemoryBase) {
234239
[mInternals->mem_gpu release];
235240
[mInternals->mem_constant release];
241+
[mInternals->canaryPipeline release];
236242
for (uint32_t i = 0; i < mInternals->functions.size(); i++) {
237243
[mInternals->functions[i] release];
238244
}
@@ -252,6 +258,7 @@ uint32_t metalGPUCoreCount()
252258
[mInternals->streamFence[i] release];
253259
}
254260
[mInternals->mem_host release];
261+
[mInternals->canaryBad release];
255262
}
256263
// The completion handlers of the limit still use the internals
257264
std::unique_lock<std::mutex> lock(mInternals->inFlightMutex);
@@ -476,6 +483,10 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const
476483
[cb release];
477484
}
478485
mInternals->submittedBuffers[stream].clear();
486+
if (const uint32_t bad = mInternals->canaryBad ? *(volatile uint32_t*)mInternals->canaryBad.contents : 0) {
487+
GPUError("Metal: %u kernels ran while the memory pool was not mapped for the GPU", bad);
488+
error = error ? error : 1;
489+
}
479490
if (drainStart) {
480491
fprintf(mInternals->timeline, "DR s%d %.6f %.6f\n", stream, drainStart, metalHostTime());
481492
}
@@ -612,5 +623,14 @@ MetalMemRef metalResolve(const GPUReconstructionMetalInternals* internals, const
612623
return 1;
613624
}
614625

626+
id<MTLFunction> canary = [mInternals->library newFunctionWithName:@"krnl_metalCanary"];
627+
mInternals->canaryPipeline = canary ? [mInternals->device newComputePipelineStateWithFunction:canary error:&error] : nil;
628+
[canary release];
629+
if (mInternals->canaryPipeline == nil) {
630+
NSLog(@"%@", error);
631+
GPUError("Error creating the Metal canary pipeline state");
632+
return 1;
633+
}
634+
615635
return AddKernels();
616636
}
Collapse file

‎GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h‎

Copy file name to clipboardExpand all lines: GPU/GPUTracking/Base/metal/GPUReconstructionMetalIncludesHost.h
+16Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -96,6 +96,12 @@ struct GPUReconstructionMetalInternals {
9696
id<MTLBuffer> mem_gpu;
9797
id<MTLBuffer> mem_constant;
9898
id<MTLBuffer> mem_host;
99+
// Opt-in, GPUCA_METAL_CANARY=1: a kernel that reads a word past the end of the
100+
// device pool, where the host put canaryMagic, and counts the reads that return
101+
// something else. canaryBad is nil when it is off.
102+
static constexpr uint32_t canaryMagic = 0xA11CE0A5;
103+
id<MTLComputePipelineState> canaryPipeline;
104+
id<MTLBuffer> canaryBad = nil;
99105

100106
id<MTLLibrary> library;
101107

@@ -156,6 +162,16 @@ inline void metalTimelineAdd(GPUReconstructionMetalInternals* internals, int32_t
156162
internals->timelineOps[stream] += buf;
157163
}
158164
}
165+
// Runs after each kernel in its encoder, so a kernel that ran while the pool
166+
// was not mapped for the GPU is counted and the run fails instead of going on
167+
// with the results it did not produce.
168+
inline void metalEncodeCanary(GPUReconstructionMetalInternals* internals, id<MTLComputeCommandEncoder> encoder, uint64_t word)
169+
{
170+
[encoder setComputePipelineState:internals->canaryPipeline];
171+
[encoder setBuffer:internals->canaryBad offset:0 atIndex:0];
172+
[encoder setBytes:&word length:sizeof(word) atIndex:1];
173+
[encoder dispatchThreadgroups:MTLSizeMake(1, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
174+
}
159175
template <class E>
160176
inline void metalEndEncoder(GPUReconstructionMetalInternals* internals, E encoder, int32_t stream)
161177
{
Collapse file

‎GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm‎

Copy file name to clipboardExpand all lines: GPU/GPUTracking/Base/metal/GPUReconstructionMetalKernels.mm
+3Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -65,6 +65,9 @@
6565
}
6666
[computeEncoder dispatchThreadgroups:MTLSizeMake(kExec.nBlocks, 1, 1)
6767
threadsPerThreadgroup:MTLSizeMake(threadGroupSize, 1, 1)];
68+
if (mInternals->canaryBad) {
69+
metalEncodeCanary(mInternals, computeEncoder, (uint64_t)mDeviceMemoryBase + mDeviceMemorySize);
70+
}
6871
metalEndEncoder(mInternals, computeEncoder, kExec.stream);
6972
id<MTLCommandBuffer> cb = [mInternals->commandBuffers[kExec.stream] retain];
7073
if (kEvents.ev) {

0 commit comments

Comments
0 (0)
You are visitor 01402520 — GeezerVison does not store the pages you look up.
Best viewed with Netscape Navigator 4.7 or Internet Explorer 6 at 800×600. No cookies required.