code/day94-gpu-sharing/gpu_sharing.cuThis is the source used by the lesson and its recorded evidence. Compile commands and expected output live in the directory README.
// SPDX-License-Identifier: MIT
//
// Day 94: sharing one GPU between two workloads.
//
// What it does: asks the driver how this card's SMs may be split, then runs
// two workloads five ways. The batch job is one long dependent-FMA kernel
// over 4 Mi elements. The short job is 64 launches of the same kernel over
// 64 Ki elements. The five schedules are: each job alone on the whole
// device, both back to back on one stream, both on two ordinary streams of
// the whole device, and both on two green contexts that own disjoint halves
// of the SMs. For each schedule the program reports when the short job
// finished and when everything finished.
//
// What it does not measure: MPS and MIG. MPS needs a control daemon and MIG
// needs root plus a device reconfiguration, so neither can be driven from
// inside one unprivileged process. Green contexts can, which is why they are
// the mechanism this file runs and the other two stay on the page.
//
// Green contexts are a CUDA 12.4 driver API; cuGreenCtxStreamCreate arrived
// in 12.5. The runtime-API spelling, cudaGreenCtxCreate and
// cudaExecutionCtxStreamCreate, is CUDA 13.1 and newer, so this file uses
// the driver API and links -lcuda.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -o gpu_sharing gpu_sharing.cu -lcuda
// Run: ./gpu_sharing
//
// VERIFIED: Tesla T4, driver 580.173.02, CUDA 12.6, 2026-09-02. The separate
// node captures record Default compute mode and successful MPS start/quit.
#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <vector>
#include <cuda.h>
#include <cuda_runtime.h>
// The one error macro. This file is standalone, so it carries its own
// verbatim copy. `err_` has a trailing underscore so it cannot collide with a
// variable at the call site.
#define CUDA_CHECK(call) \
do { \
cudaError_t err_ = (call); \
if (err_ != cudaSuccess) { \
std::fprintf(stderr, "CUDA error %s:%d: %s: %s\n", __FILE__, \
__LINE__, #call, cudaGetErrorString(err_)); \
std::exit(EXIT_FAILURE); \
} \
} while (0)
// The driver API returns CUresult, not cudaError_t, so CUDA_CHECK cannot
// wrap it and CUDA-CODE-STYLE.md allows no second error macro. Every driver
// call in this file goes through this helper and a real `if`, which is the
// shape day 44 uses for cuBLAS.
static bool driverOk(CUresult res, const char* what) {
if (res == CUDA_SUCCESS) {
return true;
}
const char* name = nullptr;
const char* text = nullptr;
cuGetErrorName(res, &name);
cuGetErrorString(res, &text);
std::fprintf(stderr, "%s failed: %s: %s\n", what,
(name != nullptr) ? name : "unknown error",
(text != nullptr) ? text : "no description");
return false;
}
// The batch job: 4 Mi elements, 8192 chained FMAs each, which is about
// 3.4e10 fused multiply-adds and lands in the milliseconds on any card in
// this course's matrix. The short job is the same kernel over 64 Ki elements
// with 1024 steps, launched 64 times, so it is the thing whose finish time
// a tenant would care about.
constexpr size_t kBulkElems = size_t{1} << 22;
constexpr int kBulkIters = 8192;
constexpr size_t kShortElems = size_t{1} << 16;
constexpr int kShortIters = 1024;
constexpr int kShortLaunches = 64;
constexpr int kThreadsPerBlock = 256; // 8 warps
constexpr int kWarmupRuns = 3;
// v = fmaf(v, kMul, c) with kMul exactly 0.5 and c = in[i]. Halving is exact
// in binary floating point, so the iterate is (2 - 2^-k) * c and the chain
// converges to 2 * c. That fixed point is what makes the correctness check
// a closed form instead of a second copy of the loop.
constexpr float kMul = 0.5f;
// About eight ulp of a float in the [2, 4) range the answers land in. The
// recurrence contracts by one half per step, so a rounding error introduced
// at step k is halved at every step after it and nothing accumulates: what
// is left is the last rounding at the fixed point, where 2c - 2^-23 is an
// exact tie and can round down one ulp. Eight ulp covers that tie with room
// to spare and is far too tight to hide a wrong element, because inputValue
// hashes the index so neighbouring answers differ in the first digit, not
// the last.
constexpr float kRelTolerance = 1e-6f;
constexpr int kBulkBlocks =
static_cast<int>((kBulkElems + kThreadsPerBlock - 1) / kThreadsPerBlock);
constexpr int kShortBlocks =
static_cast<int>((kShortElems + kThreadsPerBlock - 1) / kThreadsPerBlock);
static_assert(kThreadsPerBlock % 32 == 0,
"block size must be a whole number of warps");
static_assert(kMul == 0.5f,
"the closed-form reference below is derived for kMul = 0.5");
static_assert(kBulkIters > 32 && kShortIters > 32,
"both chains must be long enough to reach the fixed point, "
"which takes about 25 steps in float");
// out[i] is in[i] driven through `iters` steps of v = v * kMul + in[i].
//
// Memory: a grid-stride pass, so for each step the 32 lanes of a warp read 32
// consecutive floats and the load is one 128-byte request. The loop body
// touches no memory at all, which is the point: the kernel is compute bound
// and its time scales with `iters`, so one launch is a batch job and another
// with a smaller `iters` is a short one.
//
// Launch assumption: none beyond a positive block count. The grid-stride
// loop is its own bounds check.
// snippet: kernel
__global__ void fmaChain(const float* __restrict__ in, float* __restrict__ out,
size_t n, int iters) {
const size_t step = gridDim.x * static_cast<size_t>(blockDim.x);
for (size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
i < n; i += step) {
const float c = in[i];
float v = c;
for (int k = 0; k < iters; ++k) {
v = fmaf(v, kMul, c);
}
out[i] = v;
}
}
// end snippet
// A hash of the index, so no two nearby elements share a value and an index
// slip of any size lands far from the right answer. The result sits in
// [1, 2), so the chain's fixed point sits in [2, 4).
static float inputValue(size_t i) {
const unsigned int h = static_cast<unsigned int>(i) * 2654435761u;
return 1.0f + static_cast<float>(h >> 8) * (1.0f / 16777216.0f);
}
// The reference, and it is not a re-run of the kernel's loop. Solving
// v_{k+1} = v_k / 2 + c with v_0 = c gives v_k = (2 - 2^-k) * c, so past
// about 25 steps the answer is 2 * c and the reference never walks the
// chain. A reference that re-executed the same loop would agree with
// whatever the loop did, including a wrong loop.
static void chainLimitCpu(const float* in, float* out, size_t n) {
for (size_t i = 0; i < n; ++i) {
out[i] = static_cast<float>(2.0 * static_cast<double>(in[i]));
}
}
static size_t firstMismatch(const float* got, const float* want, size_t n,
float relTolerance) {
for (size_t i = 0; i < n; ++i) {
const float scale = (want[i] == 0.0f) ? 1.0f : std::fabs(want[i]);
if (std::fabs(got[i] - want[i]) > relTolerance * scale) {
return i;
}
}
return n;
}
// One tenant: a green context, the CUcontext it converts to, and the stream
// its work goes on. A client with a null context is the whole device, which
// is how the same code runs the unpartitioned schedules.
//
// cudaStream_t and CUstream are the same typedef, `struct CUstream_st *`, so
// a stream the driver API creates goes straight into a <<< >>> launch.
struct Client {
CUgreenCtx green = nullptr;
CUcontext ctx = nullptr;
cudaStream_t stream = nullptr;
unsigned int smCount = 0;
};
struct Buffers {
const float* d_x = nullptr;
float* d_bulkOut = nullptr;
float* d_shortOut = nullptr;
};
static bool pushClient(const Client* c) {
if (c->ctx == nullptr) {
return true;
}
return driverOk(cuCtxPushCurrent(c->ctx), "cuCtxPushCurrent");
}
static bool popClient(const Client* c) {
if (c->ctx == nullptr) {
return true;
}
CUcontext prev = nullptr;
return driverOk(cuCtxPopCurrent(&prev), "cuCtxPopCurrent");
}
// Turn one SM partition into a client. Wrap the resource in a descriptor,
// provision a green context from it, convert that to a CUcontext so the
// runtime can be told which context a launch belongs to, and open the stream
// the work goes on. CU_STREAM_NON_BLOCKING is not a choice: the 12.6
// reference says of the stream flags that it "must be specified".
// snippet: green-ctx
static bool makeClient(CUdevice dev, CUdevResource* part, Client* c) {
CUdevResourceDesc desc = nullptr;
if (!driverOk(cuDevResourceGenerateDesc(&desc, part, 1),
"cuDevResourceGenerateDesc")) {
return false;
}
if (!driverOk(
cuGreenCtxCreate(&c->green, desc, dev, CU_GREEN_CTX_DEFAULT_STREAM),
"cuGreenCtxCreate")) {
return false;
}
if (!driverOk(cuCtxFromGreenCtx(&c->ctx, c->green), "cuCtxFromGreenCtx")) {
return false;
}
return driverOk(
cuGreenCtxStreamCreate(&c->stream, c->green, CU_STREAM_NON_BLOCKING, 0),
"cuGreenCtxStreamCreate");
}
// end snippet
// What the green context actually got, which is never assumed from what was
// requested.
static bool readClientSms(Client* c) {
CUdevResource got{};
if (!driverOk(
cuGreenCtxGetDevResource(c->green, &got, CU_DEV_RESOURCE_TYPE_SM),
"cuGreenCtxGetDevResource")) {
return false;
}
c->smCount = got.sm.smCount;
return true;
}
static bool destroyClient(Client* c) {
bool ok = true;
if (c->stream != nullptr) {
// Runtime stream operations must run with the stream's green context
// current, just like launches in submit(). Destroying it while the
// primary context is current is not a valid cross-context operation.
if (!pushClient(c)) {
ok = false;
} else {
const cudaError_t streamResult = cudaStreamDestroy(c->stream);
if (streamResult != cudaSuccess) {
std::fprintf(stderr, "cudaStreamDestroy failed: %s\n",
cudaGetErrorString(streamResult));
ok = false;
}
if (!popClient(c)) {
ok = false;
}
}
c->stream = nullptr;
}
if (c->green != nullptr) {
ok = driverOk(cuGreenCtxDestroy(c->green), "cuGreenCtxDestroy");
c->green = nullptr;
}
c->ctx = nullptr;
return ok;
}
// One row of the split table: what was asked for, and what the driver did
// with it.
struct SplitRow {
unsigned int minCount = 0;
unsigned int groups = 0;
unsigned int perGroup = 0;
unsigned int remaining = 0;
bool uneven = false;
};
// Two calls, and the first one is the interesting one. Passing a null result
// array asks the driver to simulate the split and report how many groups it
// would make, which is how a program learns this card's partition
// granularity without carrying a table keyed on compute capability.
static bool probeSplit(const CUdevResource* input, unsigned int useFlags,
unsigned int minCount, SplitRow* row) {
row->minCount = minCount;
// snippet: split
unsigned int groups = 0;
if (!driverOk(cuDevSmResourceSplitByCount(nullptr, &groups, input, nullptr,
useFlags, minCount),
"cuDevSmResourceSplitByCount (simulate)")) {
return false;
}
row->groups = groups;
if (groups == 0) {
return true;
}
std::vector<CUdevResource> parts(groups);
CUdevResource rest{};
unsigned int made = groups;
if (!driverOk(cuDevSmResourceSplitByCount(parts.data(), &made, input, &rest,
useFlags, minCount),
"cuDevSmResourceSplitByCount")) {
return false;
}
// end snippet
row->groups = made;
row->perGroup = (made > 0) ? parts[0].sm.smCount : 0;
for (unsigned int g = 1; g < made; ++g) {
if (parts[g].sm.smCount != row->perGroup) {
row->uneven = true;
}
}
row->remaining =
(rest.type == CU_DEV_RESOURCE_TYPE_SM) ? rest.sm.smCount : 0;
return true;
}
static void printSplitTable(const char* label, const CUdevResource* devRes,
unsigned int useFlags, const unsigned int* probes,
int probeCount) {
std::printf("\n%s\n", label);
std::printf(" minCount groups SMs/group remaining equal?\n");
std::printf(" -------- ------ --------- --------- ------\n");
for (int p = 0; p < probeCount; ++p) {
SplitRow row;
if (!probeSplit(devRes, useFlags, probes[p], &row)) {
std::printf(" %8u (the driver refused this request)\n",
probes[p]);
continue;
}
std::printf(" %8u %6u %9u %9u %6s\n", row.minCount, row.groups,
row.perGroup, row.remaining, row.uneven ? "no" : "yes");
}
}
static void launchBulk(cudaStream_t stream, const Buffers& b) {
fmaChain<<<kBulkBlocks, kThreadsPerBlock, 0, stream>>>(
b.d_x, b.d_bulkOut, kBulkElems, kBulkIters);
CUDA_CHECK(cudaGetLastError());
}
static void launchShort(cudaStream_t stream, const Buffers& b) {
for (int k = 0; k < kShortLaunches; ++k) {
fmaChain<<<kShortBlocks, kThreadsPerBlock, 0, stream>>>(
b.d_x, b.d_shortOut, kShortElems, kShortIters);
}
CUDA_CHECK(cudaGetLastError());
}
// Enqueue one job on one client's stream, with the client's context pushed
// current around the launch so the runtime and the stream agree on which
// context the work belongs to. A null client means the job is not part of
// this schedule.
static bool submit(Client* c, const Buffers& b, bool batch) {
if (c == nullptr) {
return true;
}
if (!pushClient(c)) {
return false;
}
if (batch) {
launchBulk(c->stream, b);
} else {
launchShort(c->stream, b);
}
return popClient(c);
}
static bool waitOn(Client* c) {
if (c == nullptr) {
return true;
}
if (!pushClient(c)) {
return false;
}
CUDA_CHECK(cudaStreamSynchronize(c->stream));
return popClient(c);
}
// Times one schedule and returns two numbers: the host moment the short job
// finished, and the host moment everything finished, both relative to the
// moment the schedule started.
//
// Every stream here is non-blocking, so the NULL stream never waits for any
// of them and an event recorded on it completes as soon as the GPU reaches
// it. The elapsed time between two such events is the host wall time between
// the two record calls, which is the clock day 59 used to time a leg that
// ran on the CPU. It keeps every number in this program on the GPU clock.
//
// Passing the same client twice puts both jobs on one stream, which is the
// serial row: there the short job cannot finish before the batch job does,
// and the two columns come out equal.
static bool runSchedule(Client* bulkOn, Client* shortOn, const Buffers& b,
float* shortMs, float* totalMs) {
CUDA_CHECK(cudaDeviceSynchronize());
cudaEvent_t evStart = nullptr;
cudaEvent_t evShort = nullptr;
cudaEvent_t evAll = nullptr;
CUDA_CHECK(cudaEventCreate(&evStart));
CUDA_CHECK(cudaEventCreate(&evShort));
CUDA_CHECK(cudaEventCreate(&evAll));
// snippet: schedule
CUDA_CHECK(cudaEventRecord(evStart, 0));
bool ok = submit(bulkOn, b, true) && submit(shortOn, b, false);
ok = ok && waitOn(shortOn);
CUDA_CHECK(cudaEventRecord(evShort, 0));
ok = ok && waitOn(bulkOn);
CUDA_CHECK(cudaEventRecord(evAll, 0));
CUDA_CHECK(cudaEventSynchronize(evAll));
CUDA_CHECK(cudaGetLastError());
// end snippet
CUDA_CHECK(cudaEventElapsedTime(shortMs, evStart, evShort));
CUDA_CHECK(cudaEventElapsedTime(totalMs, evStart, evAll));
CUDA_CHECK(cudaEventDestroy(evStart));
CUDA_CHECK(cudaEventDestroy(evShort));
CUDA_CHECK(cudaEventDestroy(evAll));
return ok;
}
struct Row {
const char* name = "";
unsigned int bulkSms = 0;
unsigned int shortSms = 0;
float shortMs = 0.0f;
float totalMs = 0.0f;
bool ranBulk = false;
bool ranShort = false;
};
int main() {
CUDA_CHECK(cudaSetDevice(0));
cudaDeviceProp prop;
CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
std::printf("GPU: %s (compute capability %d.%d), %d SMs\n", prop.name,
prop.major, prop.minor, prop.multiProcessorCount);
std::printf("batch job: %zu elements, %d chained FMAs each\n", kBulkElems,
kBulkIters);
std::printf("short job: %d launches over %zu elements, %d FMAs each\n",
kShortLaunches, kShortElems, kShortIters);
// What the 12.6 Green Contexts reference predicts for this card, printed
// beside what the driver does below. The reference calls it "a guideline
// for each architecture and may be subject to change", so it is reported
// here rather than gated on.
const char* rule = "1 SM minimum";
if (prop.major == 7) {
rule = "2 SM minimum, multiple of 2";
} else if (prop.major == 8) {
rule = "4 SM minimum, multiple of 2";
} else if (prop.major >= 9) {
rule = "8 SM minimum, multiple of 8";
}
std::printf("documented split granularity at %d.%d: %s\n", prop.major,
prop.minor, rule);
if (!driverOk(cuInit(0), "cuInit")) {
return EXIT_FAILURE;
}
CUdevice dev = 0;
if (!driverOk(cuDeviceGet(&dev, 0), "cuDeviceGet")) {
return EXIT_FAILURE;
}
CUdevResource devRes{};
if (!driverOk(cuDeviceGetDevResource(dev, &devRes, CU_DEV_RESOURCE_TYPE_SM),
"cuDeviceGetDevResource")) {
return EXIT_FAILURE;
}
const unsigned int smTotal = devRes.sm.smCount;
std::printf("cuDeviceGetDevResource reports %u SMs\n", smTotal);
const unsigned int probes[] = {1, 2, 3, 4, 5, 8, 10, 16, 20, 32};
const int probeCount = static_cast<int>(sizeof(probes) / sizeof(probes[0]));
printSplitTable("Part 1: cuDevSmResourceSplitByCount, useFlags 0", &devRes,
0, probes, probeCount);
printSplitTable(
"Part 1b: the same, with CU_DEV_SM_RESOURCE_SPLIT_IGNORE_SM_"
"COSCHEDULING",
&devRes, CU_DEV_SM_RESOURCE_SPLIT_IGNORE_SM_COSCHEDULING, probes,
probeCount);
// The largest request that still yields two partitions. A card whose SM
// count does not halve cleanly, a 46 SM part asked for 23, rounds up and
// returns one group, so walk down rather than assume the arithmetic.
unsigned int want = (smTotal >= 2) ? smTotal / 2 : 1;
unsigned int groups = 0;
while (want >= 1) {
if (!driverOk(cuDevSmResourceSplitByCount(nullptr, &groups, &devRes,
nullptr, 0, want),
"cuDevSmResourceSplitByCount (simulate)")) {
return EXIT_FAILURE;
}
if (groups >= 2) {
break;
}
--want;
}
if (groups < 2) {
std::fprintf(stderr,
"this device cannot be split into two SM partitions\n");
return EXIT_FAILURE;
}
CUdevResource parts[2] = {};
CUdevResource rest{};
unsigned int made = 2;
if (!driverOk(
cuDevSmResourceSplitByCount(parts, &made, &devRes, &rest, 0, want),
"cuDevSmResourceSplitByCount")) {
return EXIT_FAILURE;
}
if (made != 2) {
std::fprintf(stderr, "asked for 2 partitions and got %u\n", made);
return EXIT_FAILURE;
}
const unsigned int restSms =
(rest.type == CU_DEV_RESOURCE_TYPE_SM) ? rest.sm.smCount : 0;
std::printf(
"\nsplit for the schedules below: minCount %u gives %u and %u SMs, "
"%u left over\n",
want, parts[0].sm.smCount, parts[1].sm.smCount, restSms);
const size_t bulkBytes = kBulkElems * sizeof(float);
const size_t shortBytes = kShortElems * sizeof(float);
std::vector<float> h_x(kBulkElems);
std::vector<float> h_want(kBulkElems);
std::vector<float> h_got(kBulkElems);
std::vector<float> h_first(kBulkElems);
std::vector<float> h_shortWant(kShortElems);
std::vector<float> h_shortGot(kShortElems);
std::vector<float> h_shortFirst(kShortElems);
for (size_t i = 0; i < kBulkElems; ++i) {
h_x[i] = inputValue(i);
}
chainLimitCpu(h_x.data(), h_want.data(), kBulkElems);
chainLimitCpu(h_x.data(), h_shortWant.data(), kShortElems);
float* d_x = nullptr;
float* d_bulkOut = nullptr;
float* d_shortOut = nullptr;
CUDA_CHECK(cudaMalloc(&d_x, bulkBytes));
CUDA_CHECK(cudaMalloc(&d_bulkOut, bulkBytes));
CUDA_CHECK(cudaMalloc(&d_shortOut, shortBytes));
CUDA_CHECK(cudaMemcpy(d_x, h_x.data(), bulkBytes, cudaMemcpyHostToDevice));
Buffers bufs;
bufs.d_x = d_x;
bufs.d_bulkOut = d_bulkOut;
bufs.d_shortOut = d_shortOut;
Client full1;
Client full2;
full1.smCount = static_cast<unsigned int>(prop.multiProcessorCount);
full2.smCount = full1.smCount;
CUDA_CHECK(cudaStreamCreateWithFlags(&full1.stream, cudaStreamNonBlocking));
CUDA_CHECK(cudaStreamCreateWithFlags(&full2.stream, cudaStreamNonBlocking));
int status = EXIT_SUCCESS;
Client tenantA;
Client tenantB;
if (!makeClient(dev, &parts[0], &tenantA) || !readClientSms(&tenantA)) {
status = EXIT_FAILURE;
}
if (status == EXIT_SUCCESS &&
(!makeClient(dev, &parts[1], &tenantB) || !readClientSms(&tenantB))) {
status = EXIT_FAILURE;
}
Row rows[5];
if (status == EXIT_SUCCESS) {
std::printf(
"green contexts provisioned: %u SMs and %u SMs, out of %u\n",
tenantA.smCount, tenantB.smCount, smTotal);
// Warm up every path that will be timed, not just the first one.
// Lazy module loading makes the first launch of a kernel pay its own
// load, and a green context that has never run anything pays its
// own first-launch cost too.
for (int w = 0; w < kWarmupRuns && status == EXIT_SUCCESS; ++w) {
if (!submit(&full1, bufs, false) || !submit(&full2, bufs, false) ||
!submit(&tenantA, bufs, false) ||
!submit(&tenantB, bufs, false) || !submit(&full1, bufs, true)) {
status = EXIT_FAILURE;
}
}
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaGetLastError());
rows[0].name = "alone-batch";
rows[0].ranBulk = true;
rows[1].name = "alone-short";
rows[1].ranShort = true;
rows[2].name = "serial";
rows[2].ranBulk = true;
rows[2].ranShort = true;
rows[3].name = "shared";
rows[3].ranBulk = true;
rows[3].ranShort = true;
rows[4].name = "partitioned";
rows[4].ranBulk = true;
rows[4].ranShort = true;
Client* bulkOn[5] = {&full1, nullptr, &full1, &full1, &tenantA};
Client* shortOn[5] = {nullptr, &full1, &full1, &full2, &tenantB};
for (int r = 0; r < 5 && status == EXIT_SUCCESS; ++r) {
if (!runSchedule(bulkOn[r], shortOn[r], bufs, &rows[r].shortMs,
&rows[r].totalMs)) {
status = EXIT_FAILURE;
break;
}
rows[r].bulkSms = (bulkOn[r] != nullptr) ? bulkOn[r]->smCount : 0;
rows[r].shortSms =
(shortOn[r] != nullptr) ? shortOn[r]->smCount : 0;
// Copy the answers back after every schedule. Sharing a GPU must
// not change a bit of arithmetic, and the only way to know is to
// look at all of it every time.
if (rows[r].ranBulk) {
CUDA_CHECK(cudaMemcpy(h_got.data(), d_bulkOut, bulkBytes,
cudaMemcpyDeviceToHost));
const size_t bad = firstMismatch(h_got.data(), h_want.data(),
kBulkElems, kRelTolerance);
if (bad != kBulkElems) {
std::fprintf(stderr,
"%s: batch job wrong at %zu: got %.9g, "
"want %.9g\n",
rows[r].name, bad,
static_cast<double>(h_got[bad]),
static_cast<double>(h_want[bad]));
status = EXIT_FAILURE;
} else if (r == 0) {
h_first = h_got;
} else {
for (size_t i = 0; i < kBulkElems; ++i) {
if (h_got[i] != h_first[i]) {
std::fprintf(
stderr,
"%s: batch job differs from alone-batch at "
"%zu: %.9g vs %.9g\n",
rows[r].name, i, static_cast<double>(h_got[i]),
static_cast<double>(h_first[i]));
status = EXIT_FAILURE;
break;
}
}
}
}
if (status == EXIT_SUCCESS && rows[r].ranShort) {
CUDA_CHECK(cudaMemcpy(h_shortGot.data(), d_shortOut, shortBytes,
cudaMemcpyDeviceToHost));
const size_t bad =
firstMismatch(h_shortGot.data(), h_shortWant.data(),
kShortElems, kRelTolerance);
if (bad != kShortElems) {
std::fprintf(stderr,
"%s: short job wrong at %zu: got %.9g, "
"want %.9g\n",
rows[r].name, bad,
static_cast<double>(h_shortGot[bad]),
static_cast<double>(h_shortWant[bad]));
status = EXIT_FAILURE;
} else if (r == 1) {
h_shortFirst = h_shortGot;
} else if (r > 1) {
for (size_t i = 0; i < kShortElems; ++i) {
if (h_shortGot[i] != h_shortFirst[i]) {
std::fprintf(
stderr,
"%s: short job differs from alone-short at "
"%zu: %.9g vs %.9g\n",
rows[r].name, i,
static_cast<double>(h_shortGot[i]),
static_cast<double>(h_shortFirst[i]));
status = EXIT_FAILURE;
break;
}
}
}
}
}
}
if (status == EXIT_SUCCESS) {
std::printf(
"\nPart 2: two workloads, five schedules. Times are wall clock "
"from\nthe start of the schedule, measured with events on an "
"idle NULL stream.\n\n");
std::printf(
" schedule batch SMs short SMs short done (ms) "
"all done (ms)\n");
std::printf(
" ------------ --------- --------- --------------- "
"-------------\n");
for (int r = 0; r < 5; ++r) {
char batchCell[16];
char shortCell[16];
std::snprintf(batchCell, sizeof(batchCell), "%u", rows[r].bulkSms);
std::snprintf(shortCell, sizeof(shortCell), "%u", rows[r].shortSms);
std::printf(" %-12s %9s %9s %15.3f %13.3f\n", rows[r].name,
rows[r].ranBulk ? batchCell : "-",
rows[r].ranShort ? shortCell : "-", rows[r].shortMs,
rows[r].totalMs);
}
std::printf(
"\nper short launch: %.3f us alone, %.3f us shared, "
"%.3f us partitioned\n",
static_cast<double>(rows[1].shortMs) * 1000.0 / kShortLaunches,
static_cast<double>(rows[3].shortMs) * 1000.0 / kShortLaunches,
static_cast<double>(rows[4].shortMs) * 1000.0 / kShortLaunches);
}
// The accounting gate. The reference promises "a disjoint set of
// symmetrical partitions", so two partitions of one device must be equal,
// both non-empty, and together no larger than the device. This one is a
// guarantee in the text, not the granularity guideline above, so it is a
// gate rather than a printed row.
if (status == EXIT_SUCCESS) {
if (tenantA.smCount == 0 || tenantB.smCount == 0 ||
tenantA.smCount != tenantB.smCount ||
tenantA.smCount + tenantB.smCount > smTotal) {
std::fprintf(stderr,
"partition accounting is wrong: %u + %u SMs out of "
"%u, and the two must be equal and non-empty\n",
tenantA.smCount, tenantB.smCount, smTotal);
status = EXIT_FAILURE;
}
}
// Both, unconditionally: a short circuit here would leak the second
// green context whenever the first one failed to go away.
const bool freedA = destroyClient(&tenantA);
const bool freedB = destroyClient(&tenantB);
if (!freedA || !freedB) {
status = EXIT_FAILURE;
}
CUDA_CHECK(cudaStreamDestroy(full1.stream));
CUDA_CHECK(cudaStreamDestroy(full2.stream));
CUDA_CHECK(cudaFree(d_x));
CUDA_CHECK(cudaFree(d_bulkOut));
CUDA_CHECK(cudaFree(d_shortOut));
if (status == EXIT_SUCCESS) {
std::printf(
"\nevery schedule matched the closed-form reference and the "
"alone runs, bit for bit\n");
}
return status;
}