File size: 7,434 Bytes
6be3106 | 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 108 109 110 111 112 113 114 115 116 117 118 119 120 121 122 123 124 125 126 127 128 129 130 131 132 133 134 135 136 137 138 139 140 141 142 143 144 145 146 147 148 149 150 151 152 153 154 155 156 157 158 159 160 161 162 163 164 165 166 167 168 169 170 171 172 173 174 175 176 177 178 179 180 181 182 183 184 185 186 187 | // Copyright 2021 NVIDIA Corporation. All rights reserved
//
// This sample demonstrates a very simple use case of the Checkpoint API -
// An array is saved to device, a checkpoint is saved capturing these initial
// values, the device memory is update with a new value, then restored to
// initial value using the previously saved checkpoint. By validating that
// the device values return the initial values, this demonstrates that the
// checkpoint API worked as expected.
#include <cuda.h>
#include <iostream>
#include <stdlib.h>
using namespace std;
#include <cupti_checkpoint.h>
using namespace NV::Cupti::Checkpoint;
#define CHECKPOINT_API_CALL(apiFuncCall) \
do { \
CUptiResult _status = apiFuncCall; \
if (_status != CUPTI_SUCCESS) { \
const char *errstr; \
cuptiGetResultString(_status, &errstr); \
fprintf(stderr, "%s:%d: error: function %s failed with error %s.\n", \
__FILE__, __LINE__, #apiFuncCall, errstr); \
exit(EXIT_FAILURE); \
} \
} while (0)
#define RUNTIME_API_CALL(apiFuncCall) \
do { \
cudaError_t _status = apiFuncCall; \
if (_status != cudaSuccess) { \
fprintf(stderr, "%s:%d: error: function %s failed with error %s.\n", \
__FILE__, __LINE__, #apiFuncCall, cudaGetErrorString(_status));\
exit(EXIT_FAILURE); \
} \
} while (0)
#define MEMORY_ALLOCATION_CALL(var) \
do { \
if (var == NULL) { \
fprintf(stderr, "%s:%d: Error: Memory Allocation Failed \n", \
__FILE__, __LINE__); \
exit(EXIT_FAILURE); \
} \
} while (0)
#define DRIVER_API_CALL(apiFuncCall) \
do { \
CUresult _status = apiFuncCall; \
if (_status != CUDA_SUCCESS) { \
fprintf(stderr, "%s:%d: error: function %s failed with error %d.\n", \
__FILE__, __LINE__, #apiFuncCall, _status); \
exit(EXIT_FAILURE); \
} \
} while (0)
// Basic example of a kernel which may overwrite its own input data
// This is not intended to show how to write a well-designed reduction,
// but to demonstrate that a kernel which modifies its input data can be
// replayed using the checkpoint API and get reproducible results.
//
// Sums n values, returning the total sum in data[0]
__global__ void reduce(float * data, size_t n)
{
float thd_sum = 0.0;
// Each thread sums its elements locally
for (int i = threadIdx.x; i < n; i+= blockDim.x)
{
thd_sum += data[i];
}
// And saves the per-thread sum back to the thread's first element
data[threadIdx.x] = thd_sum;
__syncthreads();
// Then, thread 0 reduces those per-thread sums to a single value in data[0]
if (threadIdx.x == 0)
{
float total_sum = 0.0;
size_t set_elems = (blockDim.x < n ? blockDim.x : n);
for (int i = 0; i < set_elems; i++)
{
total_sum += data[i];
}
data[0] = total_sum;
}
}
int main()
{
CUcontext ctx;
// Set up a context for device 0
RUNTIME_API_CALL(cudaSetDevice(0));
DRIVER_API_CALL(cuCtxCreate(&ctx, 0, 0));
// Allocate host and device arrays and initialize to known values
float * d_A;
size_t el_A = 1024 * 1024;
size_t sz_A = el_A * sizeof(float);
RUNTIME_API_CALL(cudaMalloc(&d_A, sz_A));
MEMORY_ALLOCATION_CALL(d_A);
float * h_A = (float *)malloc(sz_A);
MEMORY_ALLOCATION_CALL(h_A);
for (size_t i = 0; i < el_A; i++)
{
h_A[i] = 1.0;
}
RUNTIME_API_CALL(cudaMemcpy(d_A, h_A, sz_A, cudaMemcpyHostToDevice));
cout << "Initially, d_A[0] = " << h_A[0] << endl;
// Demonstrate a case where calling a kernel repeatedly may cause incorrect
// behavior due to internally modifying its input data
cout << "Without checkpoint:" << endl;
for (int repeat = 0; repeat < 3; repeat++)
{
reduce<<<1, 64>>>(d_A, el_A);
// Test return value - should change each iteration due to not resetting input array
float ret;
RUNTIME_API_CALL(cudaMemcpy(&ret, d_A, sizeof(float), cudaMemcpyDeviceToHost));
cout << "After " << (repeat + 1) << " iteration" << (repeat > 0 ? "s" : "") << ", d_A[0] = " << ret << endl;
}
// Re-initialize input array
RUNTIME_API_CALL(cudaMemcpy(d_A, h_A, sz_A, cudaMemcpyHostToDevice));
cout << "Reset device array - d_A[0] = " << h_A[0] << endl;
// Configure a checkpoint object
CUpti_Checkpoint cp = { CUpti_Checkpoint_STRUCT_SIZE };
cp.ctx = ctx;
cp.optimizations = 1;
float expected;
cout << "With checkpoint:" << endl;
for (int repeat = 0; repeat < 3; repeat++)
{
// Save or restore the checkpoint as needed
if (repeat == 0)
{
CHECKPOINT_API_CALL(cuptiCheckpointSave(&cp));
}
else
{
CHECKPOINT_API_CALL(cuptiCheckpointRestore(&cp));
}
// Call reduction kernel that modifies its own input
reduce<<<1, 64>>>(d_A, el_A);
// Check the output value (d_A[0])
float ret;
RUNTIME_API_CALL(cudaMemcpy(&ret, d_A, sizeof(float), cudaMemcpyDeviceToHost));
// The first call to the kernel produces the expected result - with checkpoint, every subsequent call should also return this
if (repeat == 0)
{
expected = ret;
}
cout << "After " << (repeat + 1) << " iteration" << (repeat > 0 ? "s" : "") << ", d_A[0] = " << ret << endl;
// Verify that this iteration's output value matches the expected value from the first iteration
if (ret != expected)
{
cerr << "Error - repeat " << repeat << " did not match expected value (" << ret << " != " << expected << "), did checkpoint not restore input data correctly?" << endl;
CHECKPOINT_API_CALL(cuptiCheckpointFree(&cp));
exit(EXIT_FAILURE);
}
}
// Clean up
CHECKPOINT_API_CALL(cuptiCheckpointFree(&cp));
exit(EXIT_SUCCESS);
}
|