357 lines
12 KiB
C++
357 lines
12 KiB
C++
static int load_symbol(void *lib, const char *name, void **out) {
|
|
*out = dlsym(lib, name);
|
|
return *out != NULL;
|
|
}
|
|
|
|
static int load_cuda(struct cuda_api *api) {
|
|
memset(api, 0, sizeof(*api));
|
|
api->lib = dlopen("libcuda.so.1", RTLD_NOW | RTLD_LOCAL);
|
|
if (!api->lib) {
|
|
return 0;
|
|
}
|
|
if (!(
|
|
load_symbol(api->lib, "cuInit", (void **)&api->cuInit) &&
|
|
load_symbol(api->lib, "cuDeviceGetCount", (void **)&api->cuDeviceGetCount) &&
|
|
load_symbol(api->lib, "cuDeviceGet", (void **)&api->cuDeviceGet) &&
|
|
load_symbol(api->lib, "cuDeviceGetName", (void **)&api->cuDeviceGetName) &&
|
|
load_symbol(api->lib, "cuDeviceGetAttribute", (void **)&api->cuDeviceGetAttribute) &&
|
|
load_symbol(api->lib, "cuCtxCreate_v2", (void **)&api->cuCtxCreate) &&
|
|
load_symbol(api->lib, "cuCtxDestroy_v2", (void **)&api->cuCtxDestroy) &&
|
|
load_symbol(api->lib, "cuCtxSynchronize", (void **)&api->cuCtxSynchronize) &&
|
|
load_symbol(api->lib, "cuMemAlloc_v2", (void **)&api->cuMemAlloc) &&
|
|
load_symbol(api->lib, "cuMemFree_v2", (void **)&api->cuMemFree) &&
|
|
load_symbol(api->lib, "cuMemsetD8_v2", (void **)&api->cuMemsetD8) &&
|
|
load_symbol(api->lib, "cuMemcpyHtoD_v2", (void **)&api->cuMemcpyHtoD) &&
|
|
load_symbol(api->lib, "cuMemcpyDtoH_v2", (void **)&api->cuMemcpyDtoH) &&
|
|
load_symbol(api->lib, "cuModuleLoadDataEx", (void **)&api->cuModuleLoadDataEx) &&
|
|
load_symbol(api->lib, "cuModuleGetFunction", (void **)&api->cuModuleGetFunction) &&
|
|
load_symbol(api->lib, "cuLaunchKernel", (void **)&api->cuLaunchKernel))) {
|
|
dlclose(api->lib);
|
|
memset(api, 0, sizeof(*api));
|
|
return 0;
|
|
}
|
|
load_symbol(api->lib, "cuMemGetInfo_v2", (void **)&api->cuMemGetInfo);
|
|
load_symbol(api->lib, "cuStreamCreate", (void **)&api->cuStreamCreate);
|
|
if (!load_symbol(api->lib, "cuStreamDestroy_v2", (void **)&api->cuStreamDestroy)) {
|
|
load_symbol(api->lib, "cuStreamDestroy", (void **)&api->cuStreamDestroy);
|
|
}
|
|
return 1;
|
|
}
|
|
|
|
static const char *cu_error_name(struct cuda_api *api, CUresult rc) {
|
|
const char *value = NULL;
|
|
if (api->cuGetErrorName && api->cuGetErrorName(rc, &value) == CU_SUCCESS && value) {
|
|
return value;
|
|
}
|
|
return "CUDA_ERROR";
|
|
}
|
|
|
|
static const char *cu_error_string(struct cuda_api *api, CUresult rc) {
|
|
const char *value = NULL;
|
|
if (api->cuGetErrorString && api->cuGetErrorString(rc, &value) == CU_SUCCESS && value) {
|
|
return value;
|
|
}
|
|
return "unknown";
|
|
}
|
|
|
|
static int check_rc(struct cuda_api *api, const char *step, CUresult rc) {
|
|
if (rc == CU_SUCCESS) {
|
|
return 1;
|
|
}
|
|
fprintf(stderr, "%s failed: %s (%s)\n", step, cu_error_name(api, rc), cu_error_string(api, rc));
|
|
return 0;
|
|
}
|
|
|
|
static double now_seconds(void) {
|
|
struct timespec ts;
|
|
clock_gettime(CLOCK_MONOTONIC, &ts);
|
|
return (double)ts.tv_sec + ((double)ts.tv_nsec / 1000000000.0);
|
|
}
|
|
|
|
static size_t round_down_size(size_t value, size_t multiple) {
|
|
if (multiple == 0 || value < multiple) {
|
|
return value;
|
|
}
|
|
return value - (value % multiple);
|
|
}
|
|
|
|
static int query_compute_capability(struct cuda_api *api, CUdevice dev, int *major, int *minor) {
|
|
int cc_major = 0;
|
|
int cc_minor = 0;
|
|
if (!check_rc(api,
|
|
"cuDeviceGetAttribute(major)",
|
|
api->cuDeviceGetAttribute(&cc_major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, dev))) {
|
|
return 0;
|
|
}
|
|
if (!check_rc(api,
|
|
"cuDeviceGetAttribute(minor)",
|
|
api->cuDeviceGetAttribute(&cc_minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, dev))) {
|
|
return 0;
|
|
}
|
|
*major = cc_major;
|
|
*minor = cc_minor;
|
|
return 1;
|
|
}
|
|
|
|
static int query_multiprocessor_count(struct cuda_api *api, CUdevice dev, int *count) {
|
|
int mp_count = 0;
|
|
if (!check_rc(api,
|
|
"cuDeviceGetAttribute(multiprocessors)",
|
|
api->cuDeviceGetAttribute(&mp_count, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, dev))) {
|
|
return 0;
|
|
}
|
|
*count = mp_count;
|
|
return 1;
|
|
}
|
|
|
|
static size_t clamp_budget_to_free_memory(struct cuda_api *api, size_t requested_bytes) {
|
|
size_t free_bytes = 0;
|
|
size_t total_bytes = 0;
|
|
size_t max_bytes = requested_bytes;
|
|
|
|
if (!api->cuMemGetInfo) {
|
|
return requested_bytes;
|
|
}
|
|
if (api->cuMemGetInfo(&free_bytes, &total_bytes) != CU_SUCCESS || free_bytes == 0) {
|
|
return requested_bytes;
|
|
}
|
|
|
|
max_bytes = (free_bytes * 9u) / 10u;
|
|
if (max_bytes < (size_t)4u * 1024u * 1024u) {
|
|
max_bytes = (size_t)4u * 1024u * 1024u;
|
|
}
|
|
if (requested_bytes > max_bytes) {
|
|
return max_bytes;
|
|
}
|
|
return requested_bytes;
|
|
}
|
|
|
|
static int choose_stream_count(int mp_count, int planned_profiles, size_t total_budget, int have_streams) {
|
|
int stream_count = 1;
|
|
if (!have_streams || mp_count <= 0 || planned_profiles <= 0) {
|
|
return 1;
|
|
}
|
|
|
|
stream_count = mp_count / 8;
|
|
if (stream_count < 2) {
|
|
stream_count = 2;
|
|
}
|
|
if (stream_count > MAX_STRESS_STREAMS) {
|
|
stream_count = MAX_STRESS_STREAMS;
|
|
}
|
|
|
|
while (stream_count > 1) {
|
|
size_t per_stream_budget = total_budget / ((size_t)planned_profiles * (size_t)stream_count);
|
|
if (per_stream_budget >= MIN_STREAM_BUDGET_BYTES) {
|
|
break;
|
|
}
|
|
stream_count--;
|
|
}
|
|
return stream_count;
|
|
}
|
|
|
|
#if HAVE_CUBLASLT_HEADERS
|
|
static size_t clamp_single_precision_profile_budget(size_t profile_budget_bytes) {
|
|
if (profile_budget_bytes > MAX_SINGLE_PRECISION_PROFILE_BUDGET_BYTES) {
|
|
return MAX_SINGLE_PRECISION_PROFILE_BUDGET_BYTES;
|
|
}
|
|
return profile_budget_bytes;
|
|
}
|
|
#endif
|
|
|
|
static void destroy_streams(struct cuda_api *api, CUstream *streams, int count) {
|
|
if (!api->cuStreamDestroy) {
|
|
return;
|
|
}
|
|
for (int i = 0; i < count; i++) {
|
|
if (streams[i]) {
|
|
api->cuStreamDestroy(streams[i]);
|
|
streams[i] = NULL;
|
|
}
|
|
}
|
|
}
|
|
|
|
#if HAVE_CUBLASLT_HEADERS
|
|
static void append_detail(char *buf, size_t cap, const char *fmt, ...) {
|
|
size_t len = strlen(buf);
|
|
if (len >= cap) {
|
|
return;
|
|
}
|
|
va_list ap;
|
|
va_start(ap, fmt);
|
|
vsnprintf(buf + len, cap - len, fmt, ap);
|
|
va_end(ap);
|
|
}
|
|
#endif
|
|
|
|
static int run_ptx_fallback(struct cuda_api *api,
|
|
CUdevice dev,
|
|
const char *device_name,
|
|
int cc_major,
|
|
int cc_minor,
|
|
int seconds,
|
|
int size_mb,
|
|
struct stress_report *report) {
|
|
CUcontext ctx = NULL;
|
|
CUmodule module = NULL;
|
|
CUfunction kernel = NULL;
|
|
uint32_t sample[256];
|
|
CUdeviceptr device_mem[MAX_STRESS_STREAMS] = {0};
|
|
CUstream streams[MAX_STRESS_STREAMS] = {0};
|
|
uint32_t words[MAX_STRESS_STREAMS] = {0};
|
|
uint32_t rounds[MAX_STRESS_STREAMS] = {0};
|
|
void *params[MAX_STRESS_STREAMS][3];
|
|
size_t bytes_per_stream[MAX_STRESS_STREAMS] = {0};
|
|
unsigned long iterations = 0;
|
|
int mp_count = 0;
|
|
int stream_count = 1;
|
|
|
|
memset(report, 0, sizeof(*report));
|
|
snprintf(report->backend, sizeof(report->backend), "driver-ptx");
|
|
snprintf(report->device, sizeof(report->device), "%s", device_name);
|
|
report->cc_major = cc_major;
|
|
report->cc_minor = cc_minor;
|
|
report->buffer_mb = size_mb;
|
|
|
|
if (!check_rc(api, "cuCtxCreate", api->cuCtxCreate(&ctx, 0, dev))) {
|
|
return 0;
|
|
}
|
|
|
|
size_t requested_bytes = (size_t)size_mb * 1024u * 1024u;
|
|
if (requested_bytes < MIN_PROFILE_BUDGET_BYTES) {
|
|
requested_bytes = MIN_PROFILE_BUDGET_BYTES;
|
|
}
|
|
size_t total_bytes = clamp_budget_to_free_memory(api, requested_bytes);
|
|
if (total_bytes < MIN_PROFILE_BUDGET_BYTES) {
|
|
total_bytes = MIN_PROFILE_BUDGET_BYTES;
|
|
}
|
|
report->buffer_mb = (int)(total_bytes / (1024u * 1024u));
|
|
|
|
if (query_multiprocessor_count(api, dev, &mp_count) &&
|
|
api->cuStreamCreate &&
|
|
api->cuStreamDestroy) {
|
|
stream_count = choose_stream_count(mp_count, 1, total_bytes, 1);
|
|
}
|
|
if (stream_count > 1) {
|
|
int created = 0;
|
|
for (; created < stream_count; created++) {
|
|
if (!check_rc(api, "cuStreamCreate", api->cuStreamCreate(&streams[created], 0))) {
|
|
destroy_streams(api, streams, created);
|
|
stream_count = 1;
|
|
break;
|
|
}
|
|
}
|
|
}
|
|
report->stream_count = stream_count;
|
|
|
|
for (int lane = 0; lane < stream_count; lane++) {
|
|
size_t slice = total_bytes / (size_t)stream_count;
|
|
if (lane == stream_count - 1) {
|
|
slice = total_bytes - ((size_t)lane * (total_bytes / (size_t)stream_count));
|
|
}
|
|
slice = round_down_size(slice, sizeof(uint32_t));
|
|
if (slice < MIN_PROFILE_BUDGET_BYTES) {
|
|
slice = MIN_PROFILE_BUDGET_BYTES;
|
|
}
|
|
bytes_per_stream[lane] = slice;
|
|
words[lane] = (uint32_t)(slice / sizeof(uint32_t));
|
|
|
|
if (!check_rc(api, "cuMemAlloc", api->cuMemAlloc(&device_mem[lane], slice))) {
|
|
goto fail;
|
|
}
|
|
if (!check_rc(api, "cuMemsetD8", api->cuMemsetD8(device_mem[lane], 0, slice))) {
|
|
goto fail;
|
|
}
|
|
rounds[lane] = 2048;
|
|
params[lane][0] = &device_mem[lane];
|
|
params[lane][1] = &words[lane];
|
|
params[lane][2] = &rounds[lane];
|
|
}
|
|
|
|
if (!check_rc(api,
|
|
"cuModuleLoadDataEx",
|
|
api->cuModuleLoadDataEx(&module, ptx_source, 0, NULL, NULL))) {
|
|
goto fail;
|
|
}
|
|
if (!check_rc(api, "cuModuleGetFunction", api->cuModuleGetFunction(&kernel, module, "burn"))) {
|
|
goto fail;
|
|
}
|
|
|
|
unsigned int threads = 256;
|
|
|
|
double deadline = now_seconds() + (double)seconds;
|
|
double next_sync = now_seconds() + 1.0;
|
|
while (now_seconds() < deadline) {
|
|
int launched = 0;
|
|
for (int lane = 0; lane < stream_count; lane++) {
|
|
unsigned int blocks = (unsigned int)((words[lane] + threads - 1) / threads);
|
|
if (!check_rc(api,
|
|
"cuLaunchKernel",
|
|
api->cuLaunchKernel(kernel,
|
|
blocks,
|
|
1,
|
|
1,
|
|
threads,
|
|
1,
|
|
1,
|
|
0,
|
|
streams[lane],
|
|
params[lane],
|
|
NULL))) {
|
|
goto fail;
|
|
}
|
|
launched++;
|
|
iterations++;
|
|
}
|
|
if (launched <= 0) {
|
|
goto fail;
|
|
}
|
|
double now = now_seconds();
|
|
if (now >= next_sync || now >= deadline) {
|
|
if (!check_rc(api, "cuCtxSynchronize", api->cuCtxSynchronize())) {
|
|
goto fail;
|
|
}
|
|
next_sync = now + 1.0;
|
|
}
|
|
}
|
|
api->cuCtxSynchronize();
|
|
|
|
if (!check_rc(api, "cuMemcpyDtoH", api->cuMemcpyDtoH(sample, device_mem[0], sizeof(sample)))) {
|
|
goto fail;
|
|
}
|
|
|
|
for (size_t i = 0; i < sizeof(sample) / sizeof(sample[0]); i++) {
|
|
report->checksum += sample[i];
|
|
}
|
|
report->iterations = iterations;
|
|
snprintf(report->details,
|
|
sizeof(report->details),
|
|
"fallback_int32=OK requested_mb=%d actual_mb=%d streams=%d per_stream_mb=%zu iterations=%lu\n",
|
|
size_mb,
|
|
report->buffer_mb,
|
|
report->stream_count,
|
|
bytes_per_stream[0] / (1024u * 1024u),
|
|
iterations);
|
|
|
|
for (int lane = 0; lane < stream_count; lane++) {
|
|
if (device_mem[lane]) {
|
|
api->cuMemFree(device_mem[lane]);
|
|
}
|
|
}
|
|
destroy_streams(api, streams, stream_count);
|
|
api->cuCtxDestroy(ctx);
|
|
return 1;
|
|
|
|
fail:
|
|
for (int lane = 0; lane < MAX_STRESS_STREAMS; lane++) {
|
|
if (device_mem[lane]) {
|
|
api->cuMemFree(device_mem[lane]);
|
|
}
|
|
}
|
|
destroy_streams(api, streams, MAX_STRESS_STREAMS);
|
|
if (ctx) {
|
|
api->cuCtxDestroy(ctx);
|
|
}
|
|
return 0;
|
|
}
|