refactor: modularize audit and harden build validation
This commit is contained in:
Mikhail Chusavitin
@@ -0,0 +1,944 @@
|
||||
#if HAVE_CUBLASLT_HEADERS
|
||||
typedef cublasStatus_t (*cublasLtCreate_fn)(cublasLtHandle_t *);
|
||||
typedef cublasStatus_t (*cublasLtDestroy_fn)(cublasLtHandle_t);
|
||||
typedef cublasStatus_t (*cublasLtMatmulDescCreate_fn)(cublasLtMatmulDesc_t *,
|
||||
cublasComputeType_t,
|
||||
cudaDataType_t);
|
||||
typedef cublasStatus_t (*cublasLtMatmulDescDestroy_fn)(cublasLtMatmulDesc_t);
|
||||
typedef cublasStatus_t (*cublasLtMatmulDescSetAttribute_fn)(cublasLtMatmulDesc_t,
|
||||
cublasLtMatmulDescAttributes_t,
|
||||
const void *,
|
||||
size_t);
|
||||
typedef cublasStatus_t (*cublasLtMatrixLayoutCreate_fn)(cublasLtMatrixLayout_t *,
|
||||
cudaDataType_t,
|
||||
uint64_t,
|
||||
uint64_t,
|
||||
int64_t);
|
||||
typedef cublasStatus_t (*cublasLtMatrixLayoutDestroy_fn)(cublasLtMatrixLayout_t);
|
||||
typedef cublasStatus_t (*cublasLtMatmulPreferenceCreate_fn)(cublasLtMatmulPreference_t *);
|
||||
typedef cublasStatus_t (*cublasLtMatmulPreferenceDestroy_fn)(cublasLtMatmulPreference_t);
|
||||
typedef cublasStatus_t (*cublasLtMatmulPreferenceSetAttribute_fn)(cublasLtMatmulPreference_t,
|
||||
cublasLtMatmulPreferenceAttributes_t,
|
||||
const void *,
|
||||
size_t);
|
||||
typedef cublasStatus_t (*cublasLtMatmulAlgoGetHeuristic_fn)(
|
||||
cublasLtHandle_t,
|
||||
cublasLtMatmulDesc_t,
|
||||
cublasLtMatrixLayout_t,
|
||||
cublasLtMatrixLayout_t,
|
||||
cublasLtMatrixLayout_t,
|
||||
cublasLtMatrixLayout_t,
|
||||
cublasLtMatmulPreference_t,
|
||||
int,
|
||||
cublasLtMatmulHeuristicResult_t *,
|
||||
int *);
|
||||
typedef cublasStatus_t (*cublasLtMatmul_fn)(cublasLtHandle_t,
|
||||
cublasLtMatmulDesc_t,
|
||||
const void *,
|
||||
const void *,
|
||||
cublasLtMatrixLayout_t,
|
||||
const void *,
|
||||
cublasLtMatrixLayout_t,
|
||||
const void *,
|
||||
const void *,
|
||||
cublasLtMatrixLayout_t,
|
||||
void *,
|
||||
cublasLtMatrixLayout_t,
|
||||
const cublasLtMatmulAlgo_t *,
|
||||
void *,
|
||||
size_t,
|
||||
cudaStream_t);
|
||||
|
||||
struct cublaslt_api {
|
||||
void *lib;
|
||||
cublasLtCreate_fn cublasLtCreate;
|
||||
cublasLtDestroy_fn cublasLtDestroy;
|
||||
cublasLtMatmulDescCreate_fn cublasLtMatmulDescCreate;
|
||||
cublasLtMatmulDescDestroy_fn cublasLtMatmulDescDestroy;
|
||||
cublasLtMatmulDescSetAttribute_fn cublasLtMatmulDescSetAttribute;
|
||||
cublasLtMatrixLayoutCreate_fn cublasLtMatrixLayoutCreate;
|
||||
cublasLtMatrixLayoutDestroy_fn cublasLtMatrixLayoutDestroy;
|
||||
cublasLtMatmulPreferenceCreate_fn cublasLtMatmulPreferenceCreate;
|
||||
cublasLtMatmulPreferenceDestroy_fn cublasLtMatmulPreferenceDestroy;
|
||||
cublasLtMatmulPreferenceSetAttribute_fn cublasLtMatmulPreferenceSetAttribute;
|
||||
cublasLtMatmulAlgoGetHeuristic_fn cublasLtMatmulAlgoGetHeuristic;
|
||||
cublasLtMatmul_fn cublasLtMatmul;
|
||||
};
|
||||
|
||||
struct profile_desc {
|
||||
const char *name;
|
||||
const char *block_label;
|
||||
int min_cc;
|
||||
int enabled;
|
||||
int needs_scalar_scale;
|
||||
int needs_block_scale;
|
||||
int min_multiple;
|
||||
cudaDataType_t a_type;
|
||||
cudaDataType_t b_type;
|
||||
cudaDataType_t c_type;
|
||||
cudaDataType_t d_type;
|
||||
cublasComputeType_t compute_type;
|
||||
};
|
||||
|
||||
struct prepared_profile {
|
||||
struct profile_desc desc;
|
||||
CUstream stream;
|
||||
cublasLtMatmulDesc_t op_desc;
|
||||
cublasLtMatrixLayout_t a_layout;
|
||||
cublasLtMatrixLayout_t b_layout;
|
||||
cublasLtMatrixLayout_t c_layout;
|
||||
cublasLtMatrixLayout_t d_layout;
|
||||
cublasLtMatmulPreference_t preference;
|
||||
cublasLtMatmulHeuristicResult_t heuristic;
|
||||
CUdeviceptr a_dev;
|
||||
CUdeviceptr b_dev;
|
||||
CUdeviceptr c_dev;
|
||||
CUdeviceptr d_dev;
|
||||
CUdeviceptr a_scale_dev;
|
||||
CUdeviceptr b_scale_dev;
|
||||
CUdeviceptr workspace_dev;
|
||||
size_t workspace_size;
|
||||
uint64_t m;
|
||||
uint64_t n;
|
||||
uint64_t k;
|
||||
unsigned long iterations;
|
||||
int ready;
|
||||
};
|
||||
|
||||
static const struct profile_desc k_profiles[] = {
|
||||
{
|
||||
"fp64",
|
||||
"fp64",
|
||||
80,
|
||||
1,
|
||||
0,
|
||||
0,
|
||||
8,
|
||||
CUDA_R_64F,
|
||||
CUDA_R_64F,
|
||||
CUDA_R_64F,
|
||||
CUDA_R_64F,
|
||||
CUBLAS_COMPUTE_64F,
|
||||
},
|
||||
{
|
||||
"fp32_tf32",
|
||||
"fp32",
|
||||
80,
|
||||
1,
|
||||
0,
|
||||
0,
|
||||
128,
|
||||
CUDA_R_32F,
|
||||
CUDA_R_32F,
|
||||
CUDA_R_32F,
|
||||
CUDA_R_32F,
|
||||
CUBLAS_COMPUTE_32F_FAST_TF32,
|
||||
},
|
||||
{
|
||||
"fp16_tensor",
|
||||
"fp16",
|
||||
80,
|
||||
1,
|
||||
0,
|
||||
0,
|
||||
128,
|
||||
CUDA_R_16F,
|
||||
CUDA_R_16F,
|
||||
CUDA_R_16F,
|
||||
CUDA_R_16F,
|
||||
CUBLAS_COMPUTE_32F_FAST_16F,
|
||||
},
|
||||
{
|
||||
"int8_tensor",
|
||||
"int8",
|
||||
75,
|
||||
1,
|
||||
0,
|
||||
0,
|
||||
128,
|
||||
CUDA_R_8I,
|
||||
CUDA_R_8I,
|
||||
CUDA_R_32I,
|
||||
CUDA_R_32I,
|
||||
CUBLAS_COMPUTE_32I,
|
||||
},
|
||||
{
|
||||
"fp8_e4m3",
|
||||
"fp8",
|
||||
89,
|
||||
1,
|
||||
1,
|
||||
0,
|
||||
128,
|
||||
CUDA_R_8F_E4M3,
|
||||
CUDA_R_8F_E4M3,
|
||||
CUDA_R_16BF,
|
||||
CUDA_R_16BF,
|
||||
CUBLAS_COMPUTE_32F,
|
||||
},
|
||||
{
|
||||
"fp8_e5m2",
|
||||
"fp8",
|
||||
89,
|
||||
1,
|
||||
1,
|
||||
0,
|
||||
128,
|
||||
CUDA_R_8F_E5M2,
|
||||
CUDA_R_8F_E5M2,
|
||||
CUDA_R_16BF,
|
||||
CUDA_R_16BF,
|
||||
CUBLAS_COMPUTE_32F,
|
||||
},
|
||||
#if defined(CUDA_R_4F_E2M1) && defined(CUBLASLT_MATMUL_MATRIX_SCALE_VEC16_UE4M3)
|
||||
{
|
||||
"fp4_e2m1",
|
||||
"fp4",
|
||||
100,
|
||||
1,
|
||||
0,
|
||||
1,
|
||||
128,
|
||||
CUDA_R_4F_E2M1,
|
||||
CUDA_R_4F_E2M1,
|
||||
CUDA_R_16BF,
|
||||
CUDA_R_16BF,
|
||||
CUBLAS_COMPUTE_32F,
|
||||
},
|
||||
#endif
|
||||
};
|
||||
|
||||
#define PROFILE_COUNT ((int)(sizeof(k_profiles) / sizeof(k_profiles[0])))
|
||||
|
||||
static int profile_allowed_for_run(const struct profile_desc *desc, int cc, const char *precision_filter) {
|
||||
if (!(desc->enabled && cc >= desc->min_cc)) {
|
||||
return 0;
|
||||
}
|
||||
if (precision_filter != NULL) {
|
||||
return strcmp(desc->block_label, precision_filter) == 0;
|
||||
}
|
||||
/* Mixed/all phases intentionally exclude fp64/fp4 for now: both paths are
|
||||
* unstable on the current benchmark fleet and can abort the whole mixed
|
||||
* pass after earlier phases already collected useful telemetry. */
|
||||
return strcmp(desc->block_label, "fp64") != 0 && strcmp(desc->block_label, "fp4") != 0;
|
||||
}
|
||||
|
||||
static int load_cublaslt(struct cublaslt_api *api) {
|
||||
memset(api, 0, sizeof(*api));
|
||||
api->lib = dlopen("libcublasLt.so.13", RTLD_NOW | RTLD_LOCAL);
|
||||
if (!api->lib) {
|
||||
api->lib = dlopen("libcublasLt.so", RTLD_NOW | RTLD_LOCAL);
|
||||
}
|
||||
if (!api->lib) {
|
||||
return 0;
|
||||
}
|
||||
return
|
||||
load_symbol(api->lib, "cublasLtCreate", (void **)&api->cublasLtCreate) &&
|
||||
load_symbol(api->lib, "cublasLtDestroy", (void **)&api->cublasLtDestroy) &&
|
||||
load_symbol(api->lib, "cublasLtMatmulDescCreate", (void **)&api->cublasLtMatmulDescCreate) &&
|
||||
load_symbol(api->lib, "cublasLtMatmulDescDestroy", (void **)&api->cublasLtMatmulDescDestroy) &&
|
||||
load_symbol(api->lib,
|
||||
"cublasLtMatmulDescSetAttribute",
|
||||
(void **)&api->cublasLtMatmulDescSetAttribute) &&
|
||||
load_symbol(api->lib, "cublasLtMatrixLayoutCreate", (void **)&api->cublasLtMatrixLayoutCreate) &&
|
||||
load_symbol(api->lib, "cublasLtMatrixLayoutDestroy", (void **)&api->cublasLtMatrixLayoutDestroy) &&
|
||||
load_symbol(api->lib,
|
||||
"cublasLtMatmulPreferenceCreate",
|
||||
(void **)&api->cublasLtMatmulPreferenceCreate) &&
|
||||
load_symbol(api->lib,
|
||||
"cublasLtMatmulPreferenceDestroy",
|
||||
(void **)&api->cublasLtMatmulPreferenceDestroy) &&
|
||||
load_symbol(api->lib,
|
||||
"cublasLtMatmulPreferenceSetAttribute",
|
||||
(void **)&api->cublasLtMatmulPreferenceSetAttribute) &&
|
||||
load_symbol(api->lib,
|
||||
"cublasLtMatmulAlgoGetHeuristic",
|
||||
(void **)&api->cublasLtMatmulAlgoGetHeuristic) &&
|
||||
load_symbol(api->lib, "cublasLtMatmul", (void **)&api->cublasLtMatmul);
|
||||
}
|
||||
|
||||
static const char *cublas_status_text(cublasStatus_t status) {
|
||||
switch (status) {
|
||||
case CUBLAS_STATUS_SUCCESS:
|
||||
return "CUBLAS_STATUS_SUCCESS";
|
||||
case CUBLAS_STATUS_NOT_INITIALIZED:
|
||||
return "CUBLAS_STATUS_NOT_INITIALIZED";
|
||||
case CUBLAS_STATUS_ALLOC_FAILED:
|
||||
return "CUBLAS_STATUS_ALLOC_FAILED";
|
||||
case CUBLAS_STATUS_INVALID_VALUE:
|
||||
return "CUBLAS_STATUS_INVALID_VALUE";
|
||||
case CUBLAS_STATUS_ARCH_MISMATCH:
|
||||
return "CUBLAS_STATUS_ARCH_MISMATCH";
|
||||
case CUBLAS_STATUS_MAPPING_ERROR:
|
||||
return "CUBLAS_STATUS_MAPPING_ERROR";
|
||||
case CUBLAS_STATUS_EXECUTION_FAILED:
|
||||
return "CUBLAS_STATUS_EXECUTION_FAILED";
|
||||
case CUBLAS_STATUS_INTERNAL_ERROR:
|
||||
return "CUBLAS_STATUS_INTERNAL_ERROR";
|
||||
case CUBLAS_STATUS_NOT_SUPPORTED:
|
||||
return "CUBLAS_STATUS_NOT_SUPPORTED";
|
||||
default:
|
||||
return "CUBLAS_STATUS_UNKNOWN";
|
||||
}
|
||||
}
|
||||
|
||||
static int check_cublas(const char *step, cublasStatus_t status) {
|
||||
if (status == CUBLAS_STATUS_SUCCESS) {
|
||||
return 1;
|
||||
}
|
||||
fprintf(stderr, "%s failed: %s (%d)\n", step, cublas_status_text(status), (int)status);
|
||||
return 0;
|
||||
}
|
||||
|
||||
static size_t bytes_for_elements(cudaDataType_t type, uint64_t elements) {
|
||||
switch (type) {
|
||||
case CUDA_R_32F:
|
||||
case CUDA_R_32I:
|
||||
return (size_t)(elements * 4u);
|
||||
case CUDA_R_16F:
|
||||
case CUDA_R_16BF:
|
||||
return (size_t)(elements * 2u);
|
||||
case CUDA_R_8I:
|
||||
case CUDA_R_8F_E4M3:
|
||||
case CUDA_R_8F_E5M2:
|
||||
return (size_t)(elements);
|
||||
#if defined(CUDA_R_4F_E2M1)
|
||||
case CUDA_R_4F_E2M1:
|
||||
return (size_t)((elements + 1u) / 2u);
|
||||
#endif
|
||||
default:
|
||||
return (size_t)(elements * 4u);
|
||||
}
|
||||
}
|
||||
|
||||
static cudaDataType_t matmul_scale_type(const struct profile_desc *desc) {
|
||||
if (desc->compute_type == CUBLAS_COMPUTE_32I) {
|
||||
return CUDA_R_32I;
|
||||
}
|
||||
if (desc->compute_type == CUBLAS_COMPUTE_64F) {
|
||||
return CUDA_R_64F;
|
||||
}
|
||||
return CUDA_R_32F;
|
||||
}
|
||||
|
||||
static size_t fp4_scale_bytes(uint64_t rows, uint64_t cols) {
|
||||
uint64_t row_tiles = (rows + 127u) / 128u;
|
||||
uint64_t col_tiles = (cols + 63u) / 64u;
|
||||
return (size_t)(row_tiles * col_tiles * 128u);
|
||||
}
|
||||
|
||||
static uint64_t choose_square_dim(size_t budget_bytes, size_t bytes_per_cell, int multiple) {
|
||||
double approx = sqrt((double)budget_bytes / (double)bytes_per_cell);
|
||||
uint64_t dim = (uint64_t)approx;
|
||||
if (dim < (uint64_t)multiple) {
|
||||
dim = (uint64_t)multiple;
|
||||
}
|
||||
dim = (uint64_t)round_down_size((size_t)dim, (size_t)multiple);
|
||||
if (dim < (uint64_t)multiple) {
|
||||
dim = (uint64_t)multiple;
|
||||
}
|
||||
if (dim > 65536u) {
|
||||
dim = 65536u;
|
||||
}
|
||||
return dim;
|
||||
}
|
||||
|
||||
static int device_upload(struct cuda_api *cuda, CUdeviceptr dev, const void *src, size_t bytes) {
|
||||
return check_rc(cuda, "cuMemcpyHtoD", cuda->cuMemcpyHtoD(dev, src, bytes));
|
||||
}
|
||||
|
||||
static int alloc_filled(struct cuda_api *cuda, CUdeviceptr *ptr, size_t bytes, unsigned char pattern) {
|
||||
if (!check_rc(cuda, "cuMemAlloc", cuda->cuMemAlloc(ptr, bytes))) {
|
||||
return 0;
|
||||
}
|
||||
if (!check_rc(cuda, "cuMemsetD8", cuda->cuMemsetD8(*ptr, pattern, bytes))) {
|
||||
cuda->cuMemFree(*ptr);
|
||||
*ptr = 0;
|
||||
return 0;
|
||||
}
|
||||
return 1;
|
||||
}
|
||||
|
||||
static size_t profile_scale_bytes(const struct profile_desc *desc, uint64_t m, uint64_t n, uint64_t k) {
|
||||
size_t bytes = 0;
|
||||
if (desc->needs_scalar_scale) {
|
||||
bytes += 2u * sizeof(float);
|
||||
}
|
||||
#if defined(CUBLASLT_MATMUL_MATRIX_SCALE_VEC16_UE4M3)
|
||||
if (desc->needs_block_scale) {
|
||||
bytes += fp4_scale_bytes(k, m);
|
||||
bytes += fp4_scale_bytes(k, n);
|
||||
}
|
||||
#else
|
||||
(void)m;
|
||||
(void)n;
|
||||
(void)k;
|
||||
#endif
|
||||
return bytes;
|
||||
}
|
||||
|
||||
static void destroy_profile(struct cublaslt_api *cublas, struct cuda_api *cuda, struct prepared_profile *profile) {
|
||||
if (profile->workspace_dev) {
|
||||
cuda->cuMemFree(profile->workspace_dev);
|
||||
}
|
||||
if (profile->a_scale_dev) {
|
||||
cuda->cuMemFree(profile->a_scale_dev);
|
||||
}
|
||||
if (profile->b_scale_dev) {
|
||||
cuda->cuMemFree(profile->b_scale_dev);
|
||||
}
|
||||
if (profile->d_dev) {
|
||||
cuda->cuMemFree(profile->d_dev);
|
||||
}
|
||||
if (profile->c_dev) {
|
||||
cuda->cuMemFree(profile->c_dev);
|
||||
}
|
||||
if (profile->b_dev) {
|
||||
cuda->cuMemFree(profile->b_dev);
|
||||
}
|
||||
if (profile->a_dev) {
|
||||
cuda->cuMemFree(profile->a_dev);
|
||||
}
|
||||
if (profile->preference) {
|
||||
cublas->cublasLtMatmulPreferenceDestroy(profile->preference);
|
||||
}
|
||||
if (profile->d_layout) {
|
||||
cublas->cublasLtMatrixLayoutDestroy(profile->d_layout);
|
||||
}
|
||||
if (profile->c_layout) {
|
||||
cublas->cublasLtMatrixLayoutDestroy(profile->c_layout);
|
||||
}
|
||||
if (profile->b_layout) {
|
||||
cublas->cublasLtMatrixLayoutDestroy(profile->b_layout);
|
||||
}
|
||||
if (profile->a_layout) {
|
||||
cublas->cublasLtMatrixLayoutDestroy(profile->a_layout);
|
||||
}
|
||||
if (profile->op_desc) {
|
||||
cublas->cublasLtMatmulDescDestroy(profile->op_desc);
|
||||
}
|
||||
memset(profile, 0, sizeof(*profile));
|
||||
}
|
||||
|
||||
static int prepare_profile(struct cublaslt_api *cublas,
|
||||
cublasLtHandle_t handle,
|
||||
struct cuda_api *cuda,
|
||||
const struct profile_desc *desc,
|
||||
CUstream stream,
|
||||
size_t profile_budget_bytes,
|
||||
struct prepared_profile *out) {
|
||||
size_t bytes_per_cell = 0;
|
||||
size_t attempt_budget = profile_budget_bytes;
|
||||
|
||||
bytes_per_cell += bytes_for_elements(desc->a_type, 1);
|
||||
bytes_per_cell += bytes_for_elements(desc->b_type, 1);
|
||||
bytes_per_cell += bytes_for_elements(desc->c_type, 1);
|
||||
bytes_per_cell += bytes_for_elements(desc->d_type, 1);
|
||||
if (bytes_per_cell == 0) {
|
||||
return 0;
|
||||
}
|
||||
|
||||
while (attempt_budget >= MIN_PROFILE_BUDGET_BYTES) {
|
||||
memset(out, 0, sizeof(*out));
|
||||
out->desc = *desc;
|
||||
out->stream = stream;
|
||||
|
||||
uint64_t dim = choose_square_dim(attempt_budget, bytes_per_cell, desc->min_multiple);
|
||||
out->m = dim;
|
||||
out->n = dim;
|
||||
out->k = dim;
|
||||
|
||||
size_t desired_workspace = attempt_budget / 8u;
|
||||
if (desired_workspace > 32u * 1024u * 1024u) {
|
||||
desired_workspace = 32u * 1024u * 1024u;
|
||||
}
|
||||
desired_workspace = round_down_size(desired_workspace, 256u);
|
||||
|
||||
size_t a_bytes = 0;
|
||||
size_t b_bytes = 0;
|
||||
size_t c_bytes = 0;
|
||||
size_t d_bytes = 0;
|
||||
size_t scale_bytes = 0;
|
||||
while (1) {
|
||||
a_bytes = bytes_for_elements(desc->a_type, out->k * out->m);
|
||||
b_bytes = bytes_for_elements(desc->b_type, out->k * out->n);
|
||||
c_bytes = bytes_for_elements(desc->c_type, out->m * out->n);
|
||||
d_bytes = bytes_for_elements(desc->d_type, out->m * out->n);
|
||||
scale_bytes = profile_scale_bytes(desc, out->m, out->n, out->k);
|
||||
|
||||
size_t matrix_bytes = a_bytes + b_bytes + c_bytes + d_bytes + scale_bytes;
|
||||
if (matrix_bytes <= attempt_budget) {
|
||||
size_t remaining = attempt_budget - matrix_bytes;
|
||||
out->workspace_size = desired_workspace;
|
||||
if (out->workspace_size > remaining) {
|
||||
out->workspace_size = round_down_size(remaining, 256u);
|
||||
}
|
||||
break;
|
||||
}
|
||||
|
||||
if (out->m <= (uint64_t)desc->min_multiple) {
|
||||
break;
|
||||
}
|
||||
out->m -= (uint64_t)desc->min_multiple;
|
||||
out->n = out->m;
|
||||
out->k = out->m;
|
||||
}
|
||||
if (out->m < (uint64_t)desc->min_multiple) {
|
||||
attempt_budget /= 2u;
|
||||
continue;
|
||||
}
|
||||
|
||||
if (!alloc_filled(cuda, &out->a_dev, a_bytes, 0x11) ||
|
||||
!alloc_filled(cuda, &out->b_dev, b_bytes, 0x11) ||
|
||||
!alloc_filled(cuda, &out->c_dev, c_bytes, 0x00) ||
|
||||
!alloc_filled(cuda, &out->d_dev, d_bytes, 0x00)) {
|
||||
destroy_profile(cublas, cuda, out);
|
||||
return 0;
|
||||
}
|
||||
|
||||
cudaDataType_t scale_type = matmul_scale_type(desc);
|
||||
if (!check_cublas("cublasLtMatmulDescCreate",
|
||||
cublas->cublasLtMatmulDescCreate(&out->op_desc, desc->compute_type, scale_type))) {
|
||||
destroy_profile(cublas, cuda, out);
|
||||
return 0;
|
||||
}
|
||||
|
||||
cublasOperation_t transa = CUBLAS_OP_T;
|
||||
cublasOperation_t transb = CUBLAS_OP_N;
|
||||
if (!check_cublas("set TRANSA",
|
||||
cublas->cublasLtMatmulDescSetAttribute(out->op_desc,
|
||||
CUBLASLT_MATMUL_DESC_TRANSA,
|
||||
&transa,
|
||||
sizeof(transa))) ||
|
||||
!check_cublas("set TRANSB",
|
||||
cublas->cublasLtMatmulDescSetAttribute(out->op_desc,
|
||||
CUBLASLT_MATMUL_DESC_TRANSB,
|
||||
&transb,
|
||||
sizeof(transb)))) {
|
||||
destroy_profile(cublas, cuda, out);
|
||||
return 0;
|
||||
}
|
||||
|
||||
if (desc->needs_scalar_scale) {
|
||||
float one = 1.0f;
|
||||
if (!alloc_filled(cuda, &out->a_scale_dev, sizeof(one), 0x00) ||
|
||||
!alloc_filled(cuda, &out->b_scale_dev, sizeof(one), 0x00)) {
|
||||
destroy_profile(cublas, cuda, out);
|
||||
return 0;
|
||||
}
|
||||
if (!device_upload(cuda, out->a_scale_dev, &one, sizeof(one)) ||
|
||||
!device_upload(cuda, out->b_scale_dev, &one, sizeof(one))) {
|
||||
destroy_profile(cublas, cuda, out);
|
||||
return 0;
|
||||
}
|
||||
void *a_scale_ptr = (void *)(uintptr_t)out->a_scale_dev;
|
||||
void *b_scale_ptr = (void *)(uintptr_t)out->b_scale_dev;
|
||||
if (!check_cublas("set A scale ptr",
|
||||
cublas->cublasLtMatmulDescSetAttribute(out->op_desc,
|
||||
CUBLASLT_MATMUL_DESC_A_SCALE_POINTER,
|
||||
&a_scale_ptr,
|
||||
sizeof(a_scale_ptr))) ||
|
||||
!check_cublas("set B scale ptr",
|
||||
cublas->cublasLtMatmulDescSetAttribute(out->op_desc,
|
||||
CUBLASLT_MATMUL_DESC_B_SCALE_POINTER,
|
||||
&b_scale_ptr,
|
||||
sizeof(b_scale_ptr)))) {
|
||||
destroy_profile(cublas, cuda, out);
|
||||
return 0;
|
||||
}
|
||||
}
|
||||
|
||||
#if defined(CUBLASLT_MATMUL_MATRIX_SCALE_VEC16_UE4M3)
|
||||
if (desc->needs_block_scale) {
|
||||
size_t a_scale_bytes = fp4_scale_bytes(out->k, out->m);
|
||||
size_t b_scale_bytes = fp4_scale_bytes(out->k, out->n);
|
||||
if (!alloc_filled(cuda, &out->a_scale_dev, a_scale_bytes, 0x11) ||
|
||||
!alloc_filled(cuda, &out->b_scale_dev, b_scale_bytes, 0x11)) {
|
||||
destroy_profile(cublas, cuda, out);
|
||||
return 0;
|
||||
}
|
||||
cublasLtMatmulMatrixScale_t scale_mode = CUBLASLT_MATMUL_MATRIX_SCALE_VEC16_UE4M3;
|
||||
void *a_scale_ptr = (void *)(uintptr_t)out->a_scale_dev;
|
||||
void *b_scale_ptr = (void *)(uintptr_t)out->b_scale_dev;
|
||||
if (!check_cublas("set A scale mode",
|
||||
cublas->cublasLtMatmulDescSetAttribute(out->op_desc,
|
||||
CUBLASLT_MATMUL_DESC_A_SCALE_MODE,
|
||||
&scale_mode,
|
||||
sizeof(scale_mode))) ||
|
||||
!check_cublas("set B scale mode",
|
||||
cublas->cublasLtMatmulDescSetAttribute(out->op_desc,
|
||||
CUBLASLT_MATMUL_DESC_B_SCALE_MODE,
|
||||
&scale_mode,
|
||||
sizeof(scale_mode))) ||
|
||||
!check_cublas("set A block scale ptr",
|
||||
cublas->cublasLtMatmulDescSetAttribute(out->op_desc,
|
||||
CUBLASLT_MATMUL_DESC_A_SCALE_POINTER,
|
||||
&a_scale_ptr,
|
||||
sizeof(a_scale_ptr))) ||
|
||||
!check_cublas("set B block scale ptr",
|
||||
cublas->cublasLtMatmulDescSetAttribute(out->op_desc,
|
||||
CUBLASLT_MATMUL_DESC_B_SCALE_POINTER,
|
||||
&b_scale_ptr,
|
||||
sizeof(b_scale_ptr)))) {
|
||||
destroy_profile(cublas, cuda, out);
|
||||
return 0;
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
if (!check_cublas("create A layout",
|
||||
cublas->cublasLtMatrixLayoutCreate(&out->a_layout, desc->a_type, out->k, out->m, out->k)) ||
|
||||
!check_cublas("create B layout",
|
||||
cublas->cublasLtMatrixLayoutCreate(&out->b_layout, desc->b_type, out->k, out->n, out->k)) ||
|
||||
!check_cublas("create C layout",
|
||||
cublas->cublasLtMatrixLayoutCreate(&out->c_layout, desc->c_type, out->m, out->n, out->m)) ||
|
||||
!check_cublas("create D layout",
|
||||
cublas->cublasLtMatrixLayoutCreate(&out->d_layout, desc->d_type, out->m, out->n, out->m))) {
|
||||
destroy_profile(cublas, cuda, out);
|
||||
return 0;
|
||||
}
|
||||
|
||||
if (!check_cublas("create preference", cublas->cublasLtMatmulPreferenceCreate(&out->preference))) {
|
||||
destroy_profile(cublas, cuda, out);
|
||||
return 0;
|
||||
}
|
||||
|
||||
if (out->workspace_size > 0) {
|
||||
if (!alloc_filled(cuda, &out->workspace_dev, out->workspace_size, 0x00)) {
|
||||
destroy_profile(cublas, cuda, out);
|
||||
return 0;
|
||||
}
|
||||
}
|
||||
|
||||
if (!check_cublas("set workspace",
|
||||
cublas->cublasLtMatmulPreferenceSetAttribute(
|
||||
out->preference,
|
||||
CUBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES,
|
||||
&out->workspace_size,
|
||||
sizeof(out->workspace_size)))) {
|
||||
destroy_profile(cublas, cuda, out);
|
||||
return 0;
|
||||
}
|
||||
|
||||
int found = 0;
|
||||
if (check_cublas("heuristic",
|
||||
cublas->cublasLtMatmulAlgoGetHeuristic(handle,
|
||||
out->op_desc,
|
||||
out->a_layout,
|
||||
out->b_layout,
|
||||
out->c_layout,
|
||||
out->d_layout,
|
||||
out->preference,
|
||||
1,
|
||||
&out->heuristic,
|
||||
&found)) &&
|
||||
found > 0) {
|
||||
out->ready = 1;
|
||||
return 1;
|
||||
}
|
||||
|
||||
destroy_profile(cublas, cuda, out);
|
||||
attempt_budget = round_down_size(attempt_budget * 3u / 4u, 256u);
|
||||
if (attempt_budget < MIN_PROFILE_BUDGET_BYTES) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
static int run_cublas_profile(cublasLtHandle_t handle,
|
||||
struct cublaslt_api *cublas,
|
||||
struct prepared_profile *profile) {
|
||||
int32_t alpha_i32 = 1;
|
||||
int32_t beta_i32 = 0;
|
||||
double alpha_f64 = 1.0;
|
||||
double beta_f64 = 0.0;
|
||||
float alpha = 1.0f;
|
||||
float beta = 0.0f;
|
||||
const void *alpha_ptr = α
|
||||
const void *beta_ptr = β
|
||||
if (profile->desc.compute_type == CUBLAS_COMPUTE_32I) {
|
||||
alpha_ptr = &alpha_i32;
|
||||
beta_ptr = &beta_i32;
|
||||
} else if (profile->desc.compute_type == CUBLAS_COMPUTE_64F) {
|
||||
alpha_ptr = &alpha_f64;
|
||||
beta_ptr = &beta_f64;
|
||||
}
|
||||
return check_cublas(profile->desc.name,
|
||||
cublas->cublasLtMatmul(handle,
|
||||
profile->op_desc,
|
||||
alpha_ptr,
|
||||
(const void *)(uintptr_t)profile->a_dev,
|
||||
profile->a_layout,
|
||||
(const void *)(uintptr_t)profile->b_dev,
|
||||
profile->b_layout,
|
||||
beta_ptr,
|
||||
(const void *)(uintptr_t)profile->c_dev,
|
||||
profile->c_layout,
|
||||
(void *)(uintptr_t)profile->d_dev,
|
||||
profile->d_layout,
|
||||
&profile->heuristic.algo,
|
||||
(void *)(uintptr_t)profile->workspace_dev,
|
||||
profile->workspace_size,
|
||||
profile->stream));
|
||||
}
|
||||
|
||||
static int run_cublaslt_stress(struct cuda_api *cuda,
|
||||
CUdevice dev,
|
||||
const char *device_name,
|
||||
int cc_major,
|
||||
int cc_minor,
|
||||
int seconds,
|
||||
int size_mb,
|
||||
const char *precision_filter,
|
||||
struct stress_report *report) {
|
||||
struct cublaslt_api cublas;
|
||||
struct prepared_profile prepared[MAX_STRESS_STREAMS * PROFILE_COUNT];
|
||||
cublasLtHandle_t handle = NULL;
|
||||
CUcontext ctx = NULL;
|
||||
CUstream streams[MAX_STRESS_STREAMS] = {0};
|
||||
uint16_t sample[256];
|
||||
int cc = cc_major * 10 + cc_minor;
|
||||
int planned = 0;
|
||||
int active = 0;
|
||||
int mp_count = 0;
|
||||
int stream_count = 1;
|
||||
int profile_count = PROFILE_COUNT;
|
||||
int prepared_count = 0;
|
||||
size_t requested_budget = 0;
|
||||
size_t total_budget = 0;
|
||||
size_t per_profile_budget = 0;
|
||||
int budget_profiles = 0;
|
||||
|
||||
memset(report, 0, sizeof(*report));
|
||||
snprintf(report->backend, sizeof(report->backend), "cublasLt");
|
||||
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 (!load_cublaslt(&cublas)) {
|
||||
snprintf(report->details, sizeof(report->details), "cublasLt=unavailable\n");
|
||||
return 0;
|
||||
}
|
||||
if (!check_rc(cuda, "cuCtxCreate", cuda->cuCtxCreate(&ctx, 0, dev))) {
|
||||
return 0;
|
||||
}
|
||||
if (!check_cublas("cublasLtCreate", cublas.cublasLtCreate(&handle))) {
|
||||
cuda->cuCtxDestroy(ctx);
|
||||
return 0;
|
||||
}
|
||||
|
||||
/* Count profiles matching the filter (for deciding what to run). */
|
||||
for (size_t i = 0; i < sizeof(k_profiles) / sizeof(k_profiles[0]); i++) {
|
||||
if (profile_allowed_for_run(&k_profiles[i], cc, precision_filter)) {
|
||||
planned++;
|
||||
}
|
||||
}
|
||||
if (planned <= 0) {
|
||||
snprintf(report->details, sizeof(report->details), "cublasLt_profiles=unsupported\n");
|
||||
cublas.cublasLtDestroy(handle);
|
||||
cuda->cuCtxDestroy(ctx);
|
||||
return 0;
|
||||
}
|
||||
|
||||
/* Count all profiles active on this GPU regardless of filter.
|
||||
* Mixed phases still divide budget across the full precision set, while
|
||||
* single-precision benchmark phases dedicate budget only to active
|
||||
* profiles matching precision_filter. */
|
||||
int planned_total = 0;
|
||||
for (size_t i = 0; i < sizeof(k_profiles) / sizeof(k_profiles[0]); i++) {
|
||||
if (profile_allowed_for_run(&k_profiles[i], cc, precision_filter)) {
|
||||
planned_total++;
|
||||
}
|
||||
}
|
||||
if (planned_total < planned) {
|
||||
planned_total = planned;
|
||||
}
|
||||
budget_profiles = planned_total;
|
||||
if (precision_filter != NULL) {
|
||||
budget_profiles = planned;
|
||||
}
|
||||
if (budget_profiles <= 0) {
|
||||
budget_profiles = planned_total;
|
||||
}
|
||||
|
||||
requested_budget = (size_t)size_mb * 1024u * 1024u;
|
||||
if (requested_budget < (size_t)budget_profiles * MIN_PROFILE_BUDGET_BYTES) {
|
||||
requested_budget = (size_t)budget_profiles * MIN_PROFILE_BUDGET_BYTES;
|
||||
}
|
||||
total_budget = clamp_budget_to_free_memory(cuda, requested_budget);
|
||||
if (total_budget < (size_t)budget_profiles * MIN_PROFILE_BUDGET_BYTES) {
|
||||
total_budget = (size_t)budget_profiles * MIN_PROFILE_BUDGET_BYTES;
|
||||
}
|
||||
if (query_multiprocessor_count(cuda, dev, &mp_count) &&
|
||||
cuda->cuStreamCreate &&
|
||||
cuda->cuStreamDestroy) {
|
||||
stream_count = choose_stream_count(mp_count, budget_profiles, total_budget, 1);
|
||||
}
|
||||
if (precision_filter != NULL && stream_count > MAX_SINGLE_PRECISION_STREAMS) {
|
||||
stream_count = MAX_SINGLE_PRECISION_STREAMS;
|
||||
}
|
||||
if (stream_count > 1) {
|
||||
int created = 0;
|
||||
for (; created < stream_count; created++) {
|
||||
if (!check_rc(cuda, "cuStreamCreate", cuda->cuStreamCreate(&streams[created], 0))) {
|
||||
destroy_streams(cuda, streams, created);
|
||||
stream_count = 1;
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
report->stream_count = stream_count;
|
||||
per_profile_budget = total_budget / ((size_t)budget_profiles * (size_t)stream_count);
|
||||
if (per_profile_budget < MIN_PROFILE_BUDGET_BYTES) {
|
||||
per_profile_budget = MIN_PROFILE_BUDGET_BYTES;
|
||||
}
|
||||
if (precision_filter != NULL) {
|
||||
per_profile_budget = clamp_single_precision_profile_budget(per_profile_budget);
|
||||
}
|
||||
report->buffer_mb = (int)(total_budget / (1024u * 1024u));
|
||||
append_detail(report->details,
|
||||
sizeof(report->details),
|
||||
"requested_mb=%d actual_mb=%d streams=%d mp_count=%d budget_profiles=%d per_worker_mb=%zu\n",
|
||||
size_mb,
|
||||
report->buffer_mb,
|
||||
report->stream_count,
|
||||
mp_count,
|
||||
budget_profiles,
|
||||
per_profile_budget / (1024u * 1024u));
|
||||
|
||||
for (int i = 0; i < profile_count; i++) {
|
||||
const struct profile_desc *desc = &k_profiles[i];
|
||||
if (!(desc->enabled && cc >= desc->min_cc)) {
|
||||
append_detail(report->details,
|
||||
sizeof(report->details),
|
||||
"%s=SKIPPED cc<%d\n",
|
||||
desc->name,
|
||||
desc->min_cc);
|
||||
continue;
|
||||
}
|
||||
if (!profile_allowed_for_run(desc, cc, precision_filter)) {
|
||||
append_detail(report->details,
|
||||
sizeof(report->details),
|
||||
"%s=SKIPPED benchmark_disabled\n",
|
||||
desc->name);
|
||||
continue;
|
||||
}
|
||||
for (int lane = 0; lane < stream_count; lane++) {
|
||||
CUstream stream = streams[lane];
|
||||
if (prepared_count >= (int)(sizeof(prepared) / sizeof(prepared[0]))) {
|
||||
break;
|
||||
}
|
||||
if (prepare_profile(&cublas, handle, cuda, desc, stream, per_profile_budget, &prepared[prepared_count])) {
|
||||
active++;
|
||||
append_detail(report->details,
|
||||
sizeof(report->details),
|
||||
"%s[%d]=READY dim=%llux%llux%llu block=%s stream=%d\n",
|
||||
desc->name,
|
||||
lane,
|
||||
(unsigned long long)prepared[prepared_count].m,
|
||||
(unsigned long long)prepared[prepared_count].n,
|
||||
(unsigned long long)prepared[prepared_count].k,
|
||||
desc->block_label,
|
||||
lane);
|
||||
prepared_count++;
|
||||
} else {
|
||||
append_detail(report->details,
|
||||
sizeof(report->details),
|
||||
"%s[%d]=SKIPPED unsupported\n",
|
||||
desc->name,
|
||||
lane);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
if (active <= 0) {
|
||||
cublas.cublasLtDestroy(handle);
|
||||
destroy_streams(cuda, streams, stream_count);
|
||||
cuda->cuCtxDestroy(ctx);
|
||||
return 0;
|
||||
}
|
||||
|
||||
/* Keep the GPU queue continuously full by submitting kernels without
|
||||
* synchronizing after every wave. A sync barrier after each small batch
|
||||
* creates CPU-to-GPU ping-pong gaps that prevent full TDP utilisation,
|
||||
* especially when individual kernels are short. Instead we sync at most
|
||||
* once per second (for error detection) and once at the very end. */
|
||||
double deadline = now_seconds() + (double)seconds;
|
||||
double next_sync = now_seconds() + 1.0;
|
||||
while (now_seconds() < deadline) {
|
||||
int launched = 0;
|
||||
for (int i = 0; i < prepared_count; i++) {
|
||||
if (!prepared[i].ready) {
|
||||
continue;
|
||||
}
|
||||
if (!run_cublas_profile(handle, &cublas, &prepared[i])) {
|
||||
append_detail(report->details,
|
||||
sizeof(report->details),
|
||||
"%s=FAILED runtime\n",
|
||||
prepared[i].desc.name);
|
||||
for (int j = 0; j < prepared_count; j++) {
|
||||
destroy_profile(&cublas, cuda, &prepared[j]);
|
||||
}
|
||||
cublas.cublasLtDestroy(handle);
|
||||
destroy_streams(cuda, streams, stream_count);
|
||||
cuda->cuCtxDestroy(ctx);
|
||||
return 0;
|
||||
}
|
||||
prepared[i].iterations++;
|
||||
report->iterations++;
|
||||
launched++;
|
||||
}
|
||||
if (launched <= 0) {
|
||||
break;
|
||||
}
|
||||
double now = now_seconds();
|
||||
if (now >= next_sync || now >= deadline) {
|
||||
if (!check_rc(cuda, "cuCtxSynchronize", cuda->cuCtxSynchronize())) {
|
||||
for (int i = 0; i < prepared_count; i++) {
|
||||
destroy_profile(&cublas, cuda, &prepared[i]);
|
||||
}
|
||||
cublas.cublasLtDestroy(handle);
|
||||
destroy_streams(cuda, streams, stream_count);
|
||||
cuda->cuCtxDestroy(ctx);
|
||||
return 0;
|
||||
}
|
||||
next_sync = now + 1.0;
|
||||
}
|
||||
}
|
||||
/* Final drain: ensure all queued work finishes before we read results. */
|
||||
cuda->cuCtxSynchronize();
|
||||
|
||||
for (int i = 0; i < prepared_count; i++) {
|
||||
if (!prepared[i].ready) {
|
||||
continue;
|
||||
}
|
||||
append_detail(report->details,
|
||||
sizeof(report->details),
|
||||
"%s_iterations=%lu\n",
|
||||
prepared[i].desc.name,
|
||||
prepared[i].iterations);
|
||||
}
|
||||
|
||||
for (int i = 0; i < prepared_count; i++) {
|
||||
if (prepared[i].ready) {
|
||||
if (check_rc(cuda, "cuMemcpyDtoH", cuda->cuMemcpyDtoH(sample, prepared[i].d_dev, sizeof(sample)))) {
|
||||
for (size_t j = 0; j < sizeof(sample) / sizeof(sample[0]); j++) {
|
||||
report->checksum += sample[j];
|
||||
}
|
||||
}
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < prepared_count; i++) {
|
||||
destroy_profile(&cublas, cuda, &prepared[i]);
|
||||
}
|
||||
cublas.cublasLtDestroy(handle);
|
||||
destroy_streams(cuda, streams, stream_count);
|
||||
cuda->cuCtxDestroy(ctx);
|
||||
return 1;
|
||||
}
|
||||
#endif
|
||||
@@ -0,0 +1,356 @@
|
||||
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;
|
||||
}
|
||||
@@ -0,0 +1,202 @@
|
||||
static void print_stress_report(const struct stress_report *report, int device_index, int seconds) {
|
||||
printf("device=%s\n", report->device);
|
||||
printf("device_index=%d\n", device_index);
|
||||
printf("compute_capability=%d.%d\n", report->cc_major, report->cc_minor);
|
||||
printf("backend=%s\n", report->backend);
|
||||
printf("duration_s=%d\n", seconds);
|
||||
printf("buffer_mb=%d\n", report->buffer_mb);
|
||||
printf("streams=%d\n", report->stream_count);
|
||||
printf("iterations=%lu\n", report->iterations);
|
||||
printf("checksum=%llu\n", (unsigned long long)report->checksum);
|
||||
if (report->details[0] != '\0') {
|
||||
printf("%s", report->details);
|
||||
}
|
||||
printf("status=OK\n");
|
||||
}
|
||||
|
||||
int main(int argc, char **argv) {
|
||||
int seconds = 5;
|
||||
int size_mb = 64;
|
||||
int device_index = 0;
|
||||
const char *precision_filter = NULL; /* NULL = all; else block_label to match */
|
||||
#if HAVE_CUBLASLT_HEADERS
|
||||
const char *precision_plan = NULL;
|
||||
const char *precision_plan_seconds = NULL;
|
||||
#endif
|
||||
for (int i = 1; i < argc; i++) {
|
||||
if ((strcmp(argv[i], "--seconds") == 0 || strcmp(argv[i], "-t") == 0) && i + 1 < argc) {
|
||||
seconds = atoi(argv[++i]);
|
||||
} else if ((strcmp(argv[i], "--size-mb") == 0 || strcmp(argv[i], "-m") == 0) && i + 1 < argc) {
|
||||
size_mb = atoi(argv[++i]);
|
||||
} else if ((strcmp(argv[i], "--device") == 0 || strcmp(argv[i], "-d") == 0) && i + 1 < argc) {
|
||||
device_index = atoi(argv[++i]);
|
||||
} else if (strcmp(argv[i], "--precision") == 0 && i + 1 < argc) {
|
||||
precision_filter = argv[++i];
|
||||
} else if (strcmp(argv[i], "--precision-plan") == 0 && i + 1 < argc) {
|
||||
#if HAVE_CUBLASLT_HEADERS
|
||||
precision_plan = argv[++i];
|
||||
#else
|
||||
fprintf(stderr, "--precision-plan requires a build with cuBLASLt headers\n");
|
||||
return 2;
|
||||
#endif
|
||||
} else if (strcmp(argv[i], "--precision-plan-seconds") == 0 && i + 1 < argc) {
|
||||
#if HAVE_CUBLASLT_HEADERS
|
||||
precision_plan_seconds = argv[++i];
|
||||
#else
|
||||
fprintf(stderr, "--precision-plan-seconds requires a build with cuBLASLt headers\n");
|
||||
return 2;
|
||||
#endif
|
||||
} else {
|
||||
fprintf(stderr,
|
||||
"usage: %s [--seconds N] [--size-mb N] [--device N] [--precision int8|fp8|fp16|fp32|fp64|fp4] [--precision-plan p1,p2,...,mixed] [--precision-plan-seconds s1,s2,...]\n",
|
||||
argv[0]);
|
||||
return 2;
|
||||
}
|
||||
}
|
||||
if (seconds <= 0) {
|
||||
seconds = 5;
|
||||
}
|
||||
if (size_mb <= 0) {
|
||||
size_mb = 64;
|
||||
}
|
||||
if (device_index < 0) {
|
||||
device_index = 0;
|
||||
}
|
||||
|
||||
struct cuda_api cuda;
|
||||
if (!load_cuda(&cuda)) {
|
||||
fprintf(stderr, "failed to load libcuda.so.1 or required Driver API symbols\n");
|
||||
return 1;
|
||||
}
|
||||
load_symbol(cuda.lib, "cuGetErrorName", (void **)&cuda.cuGetErrorName);
|
||||
load_symbol(cuda.lib, "cuGetErrorString", (void **)&cuda.cuGetErrorString);
|
||||
|
||||
if (!check_rc(&cuda, "cuInit", cuda.cuInit(0))) {
|
||||
return 1;
|
||||
}
|
||||
|
||||
int count = 0;
|
||||
if (!check_rc(&cuda, "cuDeviceGetCount", cuda.cuDeviceGetCount(&count))) {
|
||||
return 1;
|
||||
}
|
||||
if (count <= 0) {
|
||||
fprintf(stderr, "no CUDA devices found\n");
|
||||
return 1;
|
||||
}
|
||||
|
||||
if (device_index >= count) {
|
||||
fprintf(stderr, "device index %d out of range (found %d CUDA device(s))\n", device_index, count);
|
||||
return 1;
|
||||
}
|
||||
|
||||
CUdevice dev = 0;
|
||||
if (!check_rc(&cuda, "cuDeviceGet", cuda.cuDeviceGet(&dev, device_index))) {
|
||||
return 1;
|
||||
}
|
||||
|
||||
char name[128] = {0};
|
||||
if (!check_rc(&cuda, "cuDeviceGetName", cuda.cuDeviceGetName(name, (int)sizeof(name), dev))) {
|
||||
return 1;
|
||||
}
|
||||
|
||||
int cc_major = 0;
|
||||
int cc_minor = 0;
|
||||
if (!query_compute_capability(&cuda, dev, &cc_major, &cc_minor)) {
|
||||
return 1;
|
||||
}
|
||||
|
||||
struct stress_report report;
|
||||
int ok = 0;
|
||||
|
||||
#if HAVE_CUBLASLT_HEADERS
|
||||
if (precision_plan != NULL && precision_plan[0] != '\0') {
|
||||
char *plan_copy = strdup(precision_plan);
|
||||
char *plan_seconds_copy = NULL;
|
||||
int phase_seconds[32] = {0};
|
||||
int phase_seconds_count = 0;
|
||||
int phase_ok = 0;
|
||||
if (plan_copy == NULL) {
|
||||
fprintf(stderr, "failed to allocate precision plan buffer\n");
|
||||
return 1;
|
||||
}
|
||||
if (precision_plan_seconds != NULL && precision_plan_seconds[0] != '\0') {
|
||||
plan_seconds_copy = strdup(precision_plan_seconds);
|
||||
if (plan_seconds_copy == NULL) {
|
||||
free(plan_copy);
|
||||
fprintf(stderr, "failed to allocate precision plan seconds buffer\n");
|
||||
return 1;
|
||||
}
|
||||
for (char *sec_token = strtok(plan_seconds_copy, ",");
|
||||
sec_token != NULL && phase_seconds_count < (int)(sizeof(phase_seconds) / sizeof(phase_seconds[0]));
|
||||
sec_token = strtok(NULL, ",")) {
|
||||
while (*sec_token == ' ' || *sec_token == '\t') {
|
||||
sec_token++;
|
||||
}
|
||||
if (*sec_token == '\0') {
|
||||
continue;
|
||||
}
|
||||
phase_seconds[phase_seconds_count++] = atoi(sec_token);
|
||||
}
|
||||
}
|
||||
int phase_idx = 0;
|
||||
for (char *token = strtok(plan_copy, ","); token != NULL; token = strtok(NULL, ","), phase_idx++) {
|
||||
while (*token == ' ' || *token == '\t') {
|
||||
token++;
|
||||
}
|
||||
if (*token == '\0') {
|
||||
continue;
|
||||
}
|
||||
const char *phase_name = token;
|
||||
const char *phase_filter = token;
|
||||
if (strcmp(token, "mixed") == 0 || strcmp(token, "all") == 0) {
|
||||
phase_filter = NULL;
|
||||
}
|
||||
int phase_duration = seconds;
|
||||
if (phase_idx < phase_seconds_count && phase_seconds[phase_idx] > 0) {
|
||||
phase_duration = phase_seconds[phase_idx];
|
||||
}
|
||||
printf("phase_begin=%s\n", phase_name);
|
||||
fflush(stdout);
|
||||
memset(&report, 0, sizeof(report));
|
||||
ok = run_cublaslt_stress(&cuda, dev, name, cc_major, cc_minor, phase_duration, size_mb, phase_filter, &report);
|
||||
if (ok) {
|
||||
print_stress_report(&report, device_index, phase_duration);
|
||||
phase_ok = 1;
|
||||
} else {
|
||||
printf("phase_error=%s\n", phase_name);
|
||||
if (report.details[0] != '\0') {
|
||||
printf("%s", report.details);
|
||||
if (report.details[strlen(report.details) - 1] != '\n') {
|
||||
printf("\n");
|
||||
}
|
||||
}
|
||||
printf("status=FAILED\n");
|
||||
}
|
||||
printf("phase_end=%s\n", phase_name);
|
||||
fflush(stdout);
|
||||
}
|
||||
free(plan_seconds_copy);
|
||||
free(plan_copy);
|
||||
return phase_ok ? 0 : 1;
|
||||
}
|
||||
ok = run_cublaslt_stress(&cuda, dev, name, cc_major, cc_minor, seconds, size_mb, precision_filter, &report);
|
||||
#endif
|
||||
if (!ok) {
|
||||
if (precision_filter != NULL) {
|
||||
fprintf(stderr,
|
||||
"requested precision path unavailable: precision=%s device=%s cc=%d.%d\n",
|
||||
precision_filter,
|
||||
name,
|
||||
cc_major,
|
||||
cc_minor);
|
||||
return 1;
|
||||
}
|
||||
int ptx_mb = size_mb;
|
||||
if (!run_ptx_fallback(&cuda, dev, name, cc_major, cc_minor, seconds, ptx_mb, &report)) {
|
||||
return 1;
|
||||
}
|
||||
}
|
||||
|
||||
print_stress_report(&report, device_index, seconds);
|
||||
return 0;
|
||||
}
|
||||
+5
-1490
File diff suppressed because it is too large
Load Diff
+69
-1270
File diff suppressed because it is too large
Load Diff
@@ -1,16 +1,14 @@
|
||||
#!/bin/sh
|
||||
# Ensure memtest is present in the final ISO even if live-build's built-in
|
||||
# memtest stage does not copy the binaries or expose menu entries.
|
||||
# Ensure memtest binaries are present even if live-build's built-in memtest
|
||||
# stage does not copy them. Boot menu entries come from the canonical templates
|
||||
# enforced by build.sh after live-build finishes.
|
||||
set -e
|
||||
|
||||
: "${BEE_REQUIRE_MEMTEST:=0}"
|
||||
|
||||
# memtest86+ 6.x uses memtest86+.bin (no x64 suffix) for the BIOS binary,
|
||||
# while 5.x used memtest86+x64.bin. We normalise both to x64 names in the ISO.
|
||||
# Debian Bookworm's pinned memtest86+ package installs these exact paths.
|
||||
MEMTEST_FILES="memtest86+x64.bin memtest86+x64.efi"
|
||||
BINARY_BOOT_DIR="binary/boot"
|
||||
GRUB_CFG="binary/boot/grub/grub.cfg"
|
||||
ISOLINUX_CFG="binary/isolinux/live.cfg"
|
||||
|
||||
log() {
|
||||
echo "memtest hook: $*"
|
||||
@@ -26,14 +24,6 @@ fail_or_warn() {
|
||||
return 0
|
||||
}
|
||||
|
||||
# grub.cfg and live.cfg may not exist yet when binary hooks run — live-build
|
||||
# creates them after this hook (lb binary_grub-efi / lb binary_syslinux).
|
||||
# The template already has memtest entries hardcoded, so a missing config file
|
||||
# here is not an error; validate_iso_memtest() checks the final ISO instead.
|
||||
warn_only() {
|
||||
log "WARNING: $1"
|
||||
}
|
||||
|
||||
copy_memtest_file() {
|
||||
src="$1"
|
||||
dst_name="${2:-$(basename "$src")}"
|
||||
@@ -52,16 +42,12 @@ extract_memtest_from_deb() {
|
||||
log "extracting memtest payload from ${deb}"
|
||||
dpkg-deb -x "$deb" "$tmpdir"
|
||||
|
||||
# EFI binary: both 5.x and 6.x use memtest86+x64.efi
|
||||
if [ -f "${tmpdir}/boot/memtest86+x64.efi" ]; then
|
||||
copy_memtest_file "${tmpdir}/boot/memtest86+x64.efi"
|
||||
fi
|
||||
|
||||
# BIOS binary: 5.x = memtest86+x64.bin, 6.x = memtest86+.bin
|
||||
if [ -f "${tmpdir}/boot/memtest86+x64.bin" ]; then
|
||||
copy_memtest_file "${tmpdir}/boot/memtest86+x64.bin"
|
||||
elif [ -f "${tmpdir}/boot/memtest86+.bin" ]; then
|
||||
copy_memtest_file "${tmpdir}/boot/memtest86+.bin" "memtest86+x64.bin"
|
||||
fi
|
||||
|
||||
rm -rf "$tmpdir"
|
||||
@@ -101,10 +87,6 @@ ensure_memtest_binaries() {
|
||||
for f in ${MEMTEST_FILES}; do
|
||||
[ -f "${BINARY_BOOT_DIR}/${f}" ] || copy_memtest_file "${root}/${f}" || true
|
||||
done
|
||||
# 6.x BIOS binary may lack x64 in name — copy with normalised name
|
||||
if [ ! -f "${BINARY_BOOT_DIR}/memtest86+x64.bin" ]; then
|
||||
copy_memtest_file "${root}/memtest86+.bin" "memtest86+x64.bin" || true
|
||||
fi
|
||||
done
|
||||
|
||||
missing=0
|
||||
@@ -141,54 +123,6 @@ ensure_memtest_binaries() {
|
||||
[ "$missing" -eq 0 ] || return 0
|
||||
}
|
||||
|
||||
ensure_grub_entry() {
|
||||
[ -f "$GRUB_CFG" ] || {
|
||||
warn_only "missing ${GRUB_CFG} (will be created by lb binary_grub-efi from template)"
|
||||
return 0
|
||||
}
|
||||
|
||||
grep -q '### BEE MEMTEST ###' "$GRUB_CFG" && return 0
|
||||
|
||||
cat >> "$GRUB_CFG" <<'EOF'
|
||||
|
||||
### BEE MEMTEST ###
|
||||
if [ "${grub_platform}" = "efi" ]; then
|
||||
menuentry "Memory Test (memtest86+)" {
|
||||
chainloader /boot/memtest86+x64.efi
|
||||
}
|
||||
else
|
||||
menuentry "Memory Test (memtest86+)" {
|
||||
linux16 /boot/memtest86+x64.bin
|
||||
}
|
||||
fi
|
||||
### /BEE MEMTEST ###
|
||||
EOF
|
||||
|
||||
log "appended memtest entry to ${GRUB_CFG}"
|
||||
}
|
||||
|
||||
ensure_isolinux_entry() {
|
||||
[ -f "$ISOLINUX_CFG" ] || {
|
||||
warn_only "missing ${ISOLINUX_CFG} (will be created by lb binary_syslinux from template)"
|
||||
return 0
|
||||
}
|
||||
|
||||
grep -q '### BEE MEMTEST ###' "$ISOLINUX_CFG" && return 0
|
||||
|
||||
cat >> "$ISOLINUX_CFG" <<'EOF'
|
||||
|
||||
# ### BEE MEMTEST ###
|
||||
label memtest
|
||||
menu label ^Memory Test (memtest86+)
|
||||
linux /boot/memtest86+x64.bin
|
||||
# ### /BEE MEMTEST ###
|
||||
EOF
|
||||
|
||||
log "appended memtest entry to ${ISOLINUX_CFG}"
|
||||
}
|
||||
|
||||
log "ensuring memtest binaries and menu entries in binary image"
|
||||
log "ensuring memtest binaries in binary image"
|
||||
ensure_memtest_binaries
|
||||
ensure_grub_entry
|
||||
ensure_isolinux_entry
|
||||
log "memtest assets ready"
|
||||
|
||||
Executable
+146
@@ -0,0 +1,146 @@
|
||||
extract_live_grub_entry() {
|
||||
cfg="$1"
|
||||
live_linux="$(awk '/^[[:space:]]*linux[[:space:]]+\/live\// { print; exit }' "$cfg")"
|
||||
live_initrd="$(awk '/^[[:space:]]*initrd[[:space:]]+\/live\// { print; exit }' "$cfg")"
|
||||
[ -n "$live_linux" ] || return 1
|
||||
[ -n "$live_initrd" ] || return 1
|
||||
|
||||
grub_kernel="$(printf '%s\n' "$live_linux" | awk '{print $2}')"
|
||||
grub_append="$(printf '%s\n' "$live_linux" | cut -d' ' -f3-)"
|
||||
grub_initrd="$(printf '%s\n' "$live_initrd" | awk '{print $2}')"
|
||||
[ -n "$grub_kernel" ] || return 1
|
||||
[ -n "$grub_append" ] || return 1
|
||||
[ -n "$grub_initrd" ] || return 1
|
||||
return 0
|
||||
}
|
||||
|
||||
load_live_build_append() {
|
||||
lb_dir="$1"
|
||||
binary_cfg="$lb_dir/config/binary"
|
||||
[ -f "$binary_cfg" ] || return 1
|
||||
|
||||
# config/binary is generated by live-build and contains shell variable
|
||||
# assignments such as LB_BOOTAPPEND_LIVE="boot=live ...".
|
||||
# shellcheck disable=SC1090
|
||||
. "$binary_cfg"
|
||||
|
||||
[ -n "${LB_BOOTAPPEND_LIVE:-}" ] || return 1
|
||||
live_build_append="$LB_BOOTAPPEND_LIVE"
|
||||
return 0
|
||||
}
|
||||
|
||||
extract_live_isolinux_entry() {
|
||||
cfg="$1"
|
||||
isolinux_linux="$(awk '/^[[:space:]]*linux[[:space:]]+\/live\// { print; exit }' "$cfg")"
|
||||
isolinux_initrd="$(awk '/^[[:space:]]*initrd[[:space:]]+\/live\// { print; exit }' "$cfg")"
|
||||
isolinux_append="$(awk '/^[[:space:]]*append[[:space:]]+/ { sub(/^[[:space:]]*append[[:space:]]+/, ""); print; exit }' "$cfg")"
|
||||
[ -n "$isolinux_linux" ] || return 1
|
||||
[ -n "$isolinux_initrd" ] || return 1
|
||||
[ -n "$isolinux_append" ] || return 1
|
||||
|
||||
isolinux_kernel="$(printf '%s\n' "$isolinux_linux" | awk '{print $2}')"
|
||||
isolinux_initrd_path="$(printf '%s\n' "$isolinux_initrd" | awk '{print $2}')"
|
||||
[ -n "$isolinux_kernel" ] || return 1
|
||||
[ -n "$isolinux_initrd_path" ] || return 1
|
||||
return 0
|
||||
}
|
||||
|
||||
write_canonical_grub_cfg() {
|
||||
cfg="$1"
|
||||
kernel="$2"
|
||||
append_live="$3"
|
||||
initrd="$4"
|
||||
version_label="${PROJECT_VERSION_EFFECTIVE}"
|
||||
|
||||
template="${BUILDER_DIR}/config/bootloaders/grub-efi/grub.cfg"
|
||||
tmp_cfg="${cfg}.new"
|
||||
render_bootloader_template "$template" "$tmp_cfg" \
|
||||
"$version_label" "@KERNEL_LIVE@" "$kernel" "$append_live" "@INITRD_LIVE@" "$initrd"
|
||||
mv "$tmp_cfg" "$cfg"
|
||||
}
|
||||
|
||||
write_canonical_isolinux_cfg() {
|
||||
cfg="$1"
|
||||
kernel="$2"
|
||||
initrd="$3"
|
||||
append_live="$4"
|
||||
version_label="${PROJECT_VERSION_EFFECTIVE}"
|
||||
|
||||
template="${BUILDER_DIR}/config/bootloaders/isolinux/live.cfg.in"
|
||||
tmp_cfg="${cfg}.new"
|
||||
render_bootloader_template "$template" "$tmp_cfg" \
|
||||
"$version_label" "@LINUX@" "$kernel" "$append_live" "@INITRD@" "$initrd"
|
||||
mv "$tmp_cfg" "$cfg"
|
||||
}
|
||||
|
||||
# Render literal placeholders without sed replacement-string semantics. Values
|
||||
# containing '&', backslashes, or the sed delimiter must be copied unchanged.
|
||||
render_bootloader_template() {
|
||||
template="$1"
|
||||
output="$2"
|
||||
version="$3"
|
||||
kernel_placeholder="$4"
|
||||
kernel="$5"
|
||||
append_live="$6"
|
||||
initrd_placeholder="$7"
|
||||
initrd="$8"
|
||||
|
||||
BEE_RENDER_VERSION="$version" \
|
||||
BEE_RENDER_KERNEL_PLACEHOLDER="$kernel_placeholder" \
|
||||
BEE_RENDER_KERNEL="$kernel" \
|
||||
BEE_RENDER_APPEND="$append_live" \
|
||||
BEE_RENDER_INITRD_PLACEHOLDER="$initrd_placeholder" \
|
||||
BEE_RENDER_INITRD="$initrd" \
|
||||
awk '
|
||||
function replace_literal(text, needle, replacement, at) {
|
||||
while ((at = index(text, needle)) != 0) {
|
||||
text = substr(text, 1, at - 1) replacement substr(text, at + length(needle))
|
||||
}
|
||||
return text
|
||||
}
|
||||
{
|
||||
version = ENVIRON["BEE_RENDER_VERSION"]
|
||||
kernel_placeholder = ENVIRON["BEE_RENDER_KERNEL_PLACEHOLDER"]
|
||||
kernel = ENVIRON["BEE_RENDER_KERNEL"]
|
||||
append_live = ENVIRON["BEE_RENDER_APPEND"]
|
||||
initrd_placeholder = ENVIRON["BEE_RENDER_INITRD_PLACEHOLDER"]
|
||||
initrd = ENVIRON["BEE_RENDER_INITRD"]
|
||||
line = replace_literal($0, "@VERSION@", version)
|
||||
line = replace_literal(line, kernel_placeholder, kernel)
|
||||
line = replace_literal(line, "@APPEND_LIVE@", append_live)
|
||||
line = replace_literal(line, initrd_placeholder, initrd)
|
||||
print line
|
||||
}
|
||||
' "$template" > "$output"
|
||||
}
|
||||
|
||||
enforce_live_build_bootloader_assets() {
|
||||
lb_dir="$1"
|
||||
grub_cfg="$lb_dir/binary/boot/grub/grub.cfg"
|
||||
grub_dir="$lb_dir/binary/boot/grub"
|
||||
isolinux_cfg="$lb_dir/binary/isolinux/live.cfg"
|
||||
|
||||
if ! load_live_build_append "$lb_dir"; then
|
||||
echo "bootloader sync: WARNING: could not load LB_BOOTAPPEND_LIVE from $lb_dir/config/binary" >&2
|
||||
live_build_append=""
|
||||
fi
|
||||
|
||||
if [ -f "$grub_cfg" ]; then
|
||||
if extract_live_grub_entry "$grub_cfg"; then
|
||||
cp "${BUILDER_DIR}/config/bootloaders/grub-efi/config.cfg" "$grub_dir/config.cfg"
|
||||
write_canonical_grub_cfg "$grub_cfg" "$grub_kernel" "${live_build_append:-$grub_append}" "$grub_initrd"
|
||||
echo "bootloader sync: rewrote binary/boot/grub/grub.cfg with canonical EASY-BEE menu"
|
||||
else
|
||||
echo "bootloader sync: WARNING: could not extract live entry from $grub_cfg" >&2
|
||||
fi
|
||||
fi
|
||||
|
||||
if [ -f "$isolinux_cfg" ]; then
|
||||
if extract_live_isolinux_entry "$isolinux_cfg"; then
|
||||
write_canonical_isolinux_cfg "$isolinux_cfg" "$isolinux_kernel" "$isolinux_initrd_path" "${live_build_append:-$isolinux_append}"
|
||||
echo "bootloader sync: rewrote binary/isolinux/live.cfg with canonical EASY-BEE menu"
|
||||
else
|
||||
echo "bootloader sync: WARNING: could not extract live entry from $isolinux_cfg" >&2
|
||||
fi
|
||||
fi
|
||||
}
|
||||
Executable
+137
@@ -0,0 +1,137 @@
|
||||
cleanup_build_log() {
|
||||
status="${1:-$?}"
|
||||
trap - EXIT INT TERM HUP
|
||||
|
||||
if [ "${STEP_LOG_ACTIVE:-0}" = "1" ]; then
|
||||
cleanup_step_log "${status}" || true
|
||||
fi
|
||||
|
||||
if [ "${BUILD_LOG_ACTIVE:-0}" = "1" ]; then
|
||||
BUILD_LOG_ACTIVE=0
|
||||
exec 1>&3 2>&4
|
||||
exec 3>&- 4>&-
|
||||
if [ -n "${BUILD_TEE_PID:-}" ]; then
|
||||
wait "${BUILD_TEE_PID}" 2>/dev/null || true
|
||||
fi
|
||||
rm -rf "${BUILD_LOG_TMPDIR}"
|
||||
fi
|
||||
|
||||
if [ -n "${LOG_DIR:-}" ] && [ -d "${LOG_DIR}" ] && command -v tar >/dev/null 2>&1; then
|
||||
rm -f "${LOG_ARCHIVE}"
|
||||
tar -czf "${LOG_ARCHIVE}" -C "$(dirname "${LOG_DIR}")" "$(basename "${LOG_DIR}")" 2>/dev/null || true
|
||||
rm -rf "${LOG_DIR}"
|
||||
fi
|
||||
|
||||
exit "${status}"
|
||||
}
|
||||
|
||||
start_build_log() {
|
||||
command -v tee >/dev/null 2>&1 || {
|
||||
echo "ERROR: tee is required for build logging" >&2
|
||||
exit 1
|
||||
}
|
||||
|
||||
rm -rf "${LOG_DIR}"
|
||||
rm -f "${LOG_ARCHIVE}"
|
||||
mkdir -p "${LOG_DIR}"
|
||||
BUILD_LOG_TMPDIR="$(mktemp -d "${TMPDIR:-/tmp}/bee-build-log.XXXXXX")"
|
||||
BUILD_LOG_PIPE="${BUILD_LOG_TMPDIR}/pipe"
|
||||
mkfifo "${BUILD_LOG_PIPE}"
|
||||
|
||||
exec 3>&1 4>&2
|
||||
tee "${LOG_OUT}" < "${BUILD_LOG_PIPE}" &
|
||||
BUILD_TEE_PID=$!
|
||||
exec > "${BUILD_LOG_PIPE}" 2>&1
|
||||
BUILD_LOG_ACTIVE=1
|
||||
|
||||
trap 'cleanup_build_log "$?"' EXIT INT TERM HUP
|
||||
|
||||
echo "=== build log dir: ${LOG_DIR} ==="
|
||||
echo "=== build log: ${LOG_OUT} ==="
|
||||
echo "=== build log archive: ${LOG_ARCHIVE} ==="
|
||||
}
|
||||
|
||||
cleanup_step_log() {
|
||||
status="${1:-$?}"
|
||||
|
||||
if [ "${STEP_LOG_ACTIVE:-0}" = "1" ]; then
|
||||
STEP_LOG_ACTIVE=0
|
||||
exec 1>&5 2>&6
|
||||
exec 5>&- 6>&-
|
||||
if [ -n "${STEP_TEE_PID:-}" ]; then
|
||||
wait "${STEP_TEE_PID}" 2>/dev/null || true
|
||||
fi
|
||||
rm -rf "${STEP_LOG_TMPDIR}"
|
||||
fi
|
||||
|
||||
return "${status}"
|
||||
}
|
||||
|
||||
run_step() {
|
||||
step_name="$1"
|
||||
step_slug="$2"
|
||||
shift 2
|
||||
|
||||
step_log="${LOG_DIR}/${step_slug}.log"
|
||||
echo ""
|
||||
echo "=== step: ${step_name} ==="
|
||||
echo "=== step log: ${step_log} ==="
|
||||
|
||||
STEP_LOG_TMPDIR="$(mktemp -d "${TMPDIR:-/tmp}/bee-step-log.XXXXXX")"
|
||||
STEP_LOG_PIPE="${STEP_LOG_TMPDIR}/pipe"
|
||||
mkfifo "${STEP_LOG_PIPE}"
|
||||
|
||||
exec 5>&1 6>&2
|
||||
tee "${step_log}" < "${STEP_LOG_PIPE}" >&5 &
|
||||
STEP_TEE_PID=$!
|
||||
exec > "${STEP_LOG_PIPE}" 2>&1
|
||||
STEP_LOG_ACTIVE=1
|
||||
|
||||
set +e
|
||||
"$@"
|
||||
step_status=$?
|
||||
set -e
|
||||
|
||||
cleanup_step_log "${step_status}"
|
||||
if [ "${step_status}" -ne 0 ]; then
|
||||
echo "ERROR: step failed: ${step_name} (see ${step_log})" >&2
|
||||
exit "${step_status}"
|
||||
fi
|
||||
|
||||
echo "=== step OK: ${step_name} ==="
|
||||
}
|
||||
|
||||
run_step_sh() {
|
||||
step_name="$1"
|
||||
step_slug="$2"
|
||||
step_script="$3"
|
||||
|
||||
run_step "${step_name}" "${step_slug}" sh -c "${step_script}"
|
||||
}
|
||||
|
||||
run_optional_step_sh() {
|
||||
step_name="$1"
|
||||
step_slug="$2"
|
||||
step_script="$3"
|
||||
|
||||
if [ "${BEE_REQUIRE_MEMTEST:-0}" = "1" ]; then
|
||||
run_step_sh "${step_name}" "${step_slug}" "${step_script}"
|
||||
return 0
|
||||
fi
|
||||
|
||||
mkdir -p "${LOG_DIR}" 2>/dev/null || true
|
||||
step_log="${LOG_DIR}/${step_slug}.log"
|
||||
echo ""
|
||||
echo "=== optional step: ${step_name} ==="
|
||||
echo "=== optional step log: ${step_log} ==="
|
||||
set +e
|
||||
sh -c "${step_script}" > "${step_log}" 2>&1
|
||||
step_status=$?
|
||||
set -e
|
||||
cat "${step_log}"
|
||||
if [ "${step_status}" -ne 0 ]; then
|
||||
echo "WARNING: optional step failed: ${step_name} (see ${step_log})" >&2
|
||||
else
|
||||
echo "=== optional step OK: ${step_name} ==="
|
||||
fi
|
||||
}
|
||||
Executable
+206
@@ -0,0 +1,206 @@
|
||||
copy_memtest_from_deb() {
|
||||
deb="$1"
|
||||
dst_boot="$2"
|
||||
tmpdir="$(mktemp -d)"
|
||||
|
||||
dpkg-deb -x "$deb" "$tmpdir"
|
||||
for f in memtest86+x64.bin memtest86+x64.efi; do
|
||||
if [ -f "$tmpdir/boot/$f" ]; then
|
||||
cp "$tmpdir/boot/$f" "$dst_boot/$f"
|
||||
fi
|
||||
done
|
||||
rm -rf "$tmpdir"
|
||||
}
|
||||
|
||||
reset_live_build_stage() {
|
||||
lb_dir="$1"
|
||||
stage="$2"
|
||||
|
||||
for root in \
|
||||
"$lb_dir/.build" \
|
||||
"$lb_dir/.stage" \
|
||||
"$lb_dir/auto"; do
|
||||
[ -d "$root" ] || continue
|
||||
find "$root" -maxdepth 1 \( -name "${stage}" -o -name "${stage}.*" -o -name "*${stage}*" \) -exec rm -rf {} + 2>/dev/null || true
|
||||
done
|
||||
}
|
||||
|
||||
# State written after every successful full lb build for this variant. Keep it
|
||||
# outside the rsync-managed live-build workdir so source synchronization cannot
|
||||
# delete the state that decides whether the fast path is safe.
|
||||
FULL_BUILD_STATE_DIR="${CACHE_ROOT}/full-build-state-${BUILD_VARIANT}"
|
||||
mkdir -p "${FULL_BUILD_STATE_DIR}"
|
||||
FULL_BUILD_MARKER="${FULL_BUILD_STATE_DIR}/complete"
|
||||
FULL_BUILD_HASH_FILE="${FULL_BUILD_STATE_DIR}/heavy-config.sha256"
|
||||
FULL_BUILD_ABI_FILE="${FULL_BUILD_STATE_DIR}/kernel-abi"
|
||||
FULL_BUILD_OVERLAY_MANIFEST="${FULL_BUILD_STATE_DIR}/overlay.manifest"
|
||||
|
||||
# Hashes the content of every "heavy" config input (VERSIONS, package lists,
|
||||
# hooks, archives, auto/config, Dockerfile). Bootloader templates are excluded:
|
||||
# the fast path regenerates the complete outer ISO layer from them. Deliberately content-
|
||||
# based rather than mtime-based: mtimes get reset by git checkouts, rsync, and
|
||||
# retried builds in ways that don't track "did this content actually change
|
||||
# since the last full build", which previously let needs_full_build() silently
|
||||
# take the fast path (reusing an old squashfs built against different package
|
||||
# pins) with no error.
|
||||
hash_heavy_config() {
|
||||
(
|
||||
cd "${BUILDER_DIR}"
|
||||
find \
|
||||
VERSIONS auto/config Dockerfile \
|
||||
config/package-lists config/hooks config/archives \
|
||||
-type f -print0 2>/dev/null |
|
||||
sort -z |
|
||||
xargs -0 -r sha256sum
|
||||
) | sha256sum | awk '{print $1}'
|
||||
}
|
||||
|
||||
write_overlay_manifest() {
|
||||
out_path="$1"
|
||||
(
|
||||
cd "${OVERLAY_STAGE_DIR}"
|
||||
find . -mindepth 1 -printf '%y %P\n' | sort
|
||||
) > "$out_path"
|
||||
}
|
||||
|
||||
overlay_paths_were_removed() {
|
||||
[ -f "${FULL_BUILD_OVERLAY_MANIFEST}" ] || return 0
|
||||
current_manifest="$(mktemp)"
|
||||
write_overlay_manifest "$current_manifest"
|
||||
if comm -23 "${FULL_BUILD_OVERLAY_MANIFEST}" "$current_manifest" | grep -q .; then
|
||||
rm -f "$current_manifest"
|
||||
return 0
|
||||
fi
|
||||
rm -f "$current_manifest"
|
||||
return 1
|
||||
}
|
||||
|
||||
# Returns 0 if full lb build is needed, 1 if fast-path is safe.
|
||||
# Fast-path is safe when only light files changed since the last full build
|
||||
# (Go source, overlay scripts/configs). Heavy changes (VERSIONS, package lists,
|
||||
# hooks, archives, Dockerfile, auto/config) require a full lb build.
|
||||
needs_full_build() {
|
||||
[ -f "${FULL_BUILD_MARKER}" ] || return 0
|
||||
[ -f "${FULL_BUILD_HASH_FILE}" ] || return 0
|
||||
[ -f "${FULL_BUILD_ABI_FILE}" ] || return 0
|
||||
[ -f "${FULL_BUILD_OVERLAY_MANIFEST}" ] || return 0
|
||||
[ -f "${BUILD_WORK_DIR}/live-image-amd64.hybrid.iso" ] || return 0
|
||||
# Accept any versioned squashfs (filesystem-v*.squashfs or legacy filesystem.squashfs)
|
||||
_any_sq=$(find "${BUILD_WORK_DIR}/binary/live" -maxdepth 1 \
|
||||
-name 'filesystem*.squashfs' 2>/dev/null | head -1)
|
||||
[ -n "$_any_sq" ] || return 0
|
||||
|
||||
_old_abi="$(cat "${FULL_BUILD_ABI_FILE}" 2>/dev/null)"
|
||||
if [ "${DEBIAN_KERNEL_ABI}" != "$_old_abi" ]; then
|
||||
echo "=== full build required: kernel ABI changed (${_old_abi:-unknown} -> ${DEBIAN_KERNEL_ABI}) ==="
|
||||
return 0
|
||||
fi
|
||||
|
||||
if overlay_paths_were_removed; then
|
||||
echo "=== full build required: overlay paths were removed or changed type ==="
|
||||
return 0
|
||||
fi
|
||||
|
||||
_new_hash="$(hash_heavy_config)"
|
||||
_old_hash="$(cat "${FULL_BUILD_HASH_FILE}" 2>/dev/null)"
|
||||
|
||||
if [ "$_new_hash" != "$_old_hash" ]; then
|
||||
echo "=== full build required: heavy config content changed since last full build ==="
|
||||
return 0
|
||||
fi
|
||||
|
||||
return 1
|
||||
}
|
||||
|
||||
# Fast path: unsquash existing filesystem, rsync overlay on top, repack.
|
||||
# CACHE_ROOT must have enough free space for the extracted root filesystem.
|
||||
fast_path_repack_squashfs() (
|
||||
_old_sq=$(find "${BUILD_WORK_DIR}/binary/live" -maxdepth 1 \
|
||||
-name 'filesystem*.squashfs' | sort | head -1)
|
||||
_sq="${BUILD_WORK_DIR}/binary/live/${SQUASHFS_FILENAME}"
|
||||
_tmp_parent="$(mktemp -d "${CACHE_ROOT}/fast-unsquash-${BUILD_VARIANT}.XXXXXX")"
|
||||
_tmp="${_tmp_parent}/root"
|
||||
trap 'rm -rf "$_tmp_parent"' EXIT
|
||||
echo "=== fast-path: unsquash $(basename "$_old_sq") ($(du -sh "$_old_sq" | cut -f1) compressed) ==="
|
||||
unsquashfs -d "$_tmp" "$_old_sq"
|
||||
echo "=== fast-path: syncing overlay stage ==="
|
||||
rsync -a --checksum "${OVERLAY_STAGE_DIR}/" "$_tmp/"
|
||||
echo "=== fast-path: repacking as ${SQUASHFS_FILENAME} ==="
|
||||
_sq_new="${_sq}.new"
|
||||
rm -f "$_sq_new"
|
||||
mksquashfs "$_tmp" "$_sq_new" -comp zstd -b 1048576 -noappend -no-progress -no-xattrs
|
||||
mv "$_sq_new" "$_sq"
|
||||
rm -rf "$_tmp_parent"
|
||||
for _candidate in "${BUILD_WORK_DIR}/binary/live/"filesystem*.squashfs; do
|
||||
[ -e "$_candidate" ] || continue
|
||||
[ "$_candidate" = "$_sq" ] || rm -f "$_candidate"
|
||||
done
|
||||
echo "=== fast-path: squashfs repacked ($(du -sh "$_sq" | cut -f1)) ==="
|
||||
)
|
||||
|
||||
# Fast-path: rebuild ISO replacing the squashfs via xorriso.
|
||||
# Boot structure (El Torito, EFI, MBR hybrid) is replayed from the prior ISO.
|
||||
recover_iso_memtest() {
|
||||
lb_dir="$1"
|
||||
iso_path="$2"
|
||||
binary_boot="$lb_dir/binary/boot"
|
||||
|
||||
echo "=== attempting memtest recovery in binary tree ==="
|
||||
|
||||
mkdir -p "$binary_boot"
|
||||
|
||||
for root in \
|
||||
"$lb_dir/chroot/boot" \
|
||||
"/boot"; do
|
||||
for f in memtest86+x64.bin memtest86+x64.efi; do
|
||||
if [ ! -f "$binary_boot/$f" ] && [ -f "$root/$f" ]; then
|
||||
cp "$root/$f" "$binary_boot/$f"
|
||||
echo "memtest recovery: copied $f from $root"
|
||||
fi
|
||||
done
|
||||
done
|
||||
|
||||
if [ ! -f "$binary_boot/memtest86+x64.bin" ] || [ ! -f "$binary_boot/memtest86+x64.efi" ]; then
|
||||
for dir in \
|
||||
"$lb_dir/cache/packages.binary" \
|
||||
"$lb_dir/cache/packages.chroot" \
|
||||
"$lb_dir/chroot/var/cache/apt/archives" \
|
||||
"${BEE_CACHE_DIR:-${DIST_DIR}/cache}/lb-packages" \
|
||||
"/var/cache/apt/archives"; do
|
||||
[ -d "$dir" ] || continue
|
||||
deb="$(find "$dir" -maxdepth 1 -type f -name 'memtest86+*.deb' 2>/dev/null | head -1)"
|
||||
[ -n "$deb" ] || continue
|
||||
echo "memtest recovery: extracting payload from $deb"
|
||||
copy_memtest_from_deb "$deb" "$binary_boot"
|
||||
break
|
||||
done
|
||||
fi
|
||||
|
||||
if [ ! -f "$binary_boot/memtest86+x64.bin" ] || [ ! -f "$binary_boot/memtest86+x64.efi" ]; then
|
||||
tmpdl="$(mktemp -d)"
|
||||
if (
|
||||
cd "$tmpdl" && apt-get download memtest86+ >/dev/null 2>&1
|
||||
); then
|
||||
deb="$(find "$tmpdl" -maxdepth 1 -type f -name 'memtest86+*.deb' 2>/dev/null | head -1)"
|
||||
if [ -n "$deb" ]; then
|
||||
echo "memtest recovery: downloaded $deb"
|
||||
copy_memtest_from_deb "$deb" "$binary_boot"
|
||||
fi
|
||||
fi
|
||||
rm -rf "$tmpdl"
|
||||
fi
|
||||
|
||||
enforce_live_build_bootloader_assets "$lb_dir"
|
||||
|
||||
reset_live_build_stage "$lb_dir" "binary_checksums"
|
||||
reset_live_build_stage "$lb_dir" "binary_iso"
|
||||
reset_live_build_stage "$lb_dir" "binary_zsync"
|
||||
|
||||
run_optional_step_sh "rebuild live-build checksums after memtest recovery" "91-lb-checksums" "lb binary_checksums 2>&1"
|
||||
run_optional_step_sh "rebuild ISO after memtest recovery" "92-lb-binary-iso" "rm -f '$iso_path' && lb binary_iso 2>&1"
|
||||
run_optional_step_sh "rebuild zsync after memtest recovery" "93-lb-zsync" "lb binary_zsync 2>&1"
|
||||
|
||||
if [ ! -f "$iso_path" ]; then
|
||||
memtest_fail "ISO rebuild was skipped or failed after memtest recovery: $iso_path" "$iso_path"
|
||||
fi
|
||||
}
|
||||
Executable
+636
@@ -0,0 +1,636 @@
|
||||
iso_list_files() {
|
||||
iso_path="$1"
|
||||
|
||||
if command -v bsdtar >/dev/null 2>&1; then
|
||||
bsdtar -tf "$iso_path"
|
||||
return $?
|
||||
fi
|
||||
|
||||
if command -v xorriso >/dev/null 2>&1; then
|
||||
xorriso -indev "$iso_path" -find / -type f -print 2>/dev/null | sed 's#^/##'
|
||||
return $?
|
||||
fi
|
||||
|
||||
return 127
|
||||
}
|
||||
|
||||
iso_extract_file() {
|
||||
iso_path="$1"
|
||||
iso_member="$2"
|
||||
|
||||
if command -v bsdtar >/dev/null 2>&1; then
|
||||
bsdtar -xOf "$iso_path" "$iso_member"
|
||||
return $?
|
||||
fi
|
||||
|
||||
if command -v xorriso >/dev/null 2>&1; then
|
||||
xorriso -osirrox on -indev "$iso_path" -cat "/$iso_member" 2>/dev/null
|
||||
return $?
|
||||
fi
|
||||
|
||||
return 127
|
||||
}
|
||||
|
||||
iso_read_file_list() {
|
||||
iso_path="$1"
|
||||
out_path="$2"
|
||||
|
||||
iso_list_files "$iso_path" > "$out_path" || return 1
|
||||
[ -s "$out_path" ] || return 1
|
||||
return 0
|
||||
}
|
||||
|
||||
iso_read_member() {
|
||||
iso_path="$1"
|
||||
iso_member="$2"
|
||||
out_path="$3"
|
||||
|
||||
iso_extract_file "$iso_path" "$iso_member" > "$out_path" || return 1
|
||||
[ -s "$out_path" ] || return 1
|
||||
return 0
|
||||
}
|
||||
|
||||
require_iso_reader() {
|
||||
command -v bsdtar >/dev/null 2>&1 && return 0
|
||||
command -v xorriso >/dev/null 2>&1 && return 0
|
||||
memtest_fail "ISO reader is required for validation/debug (expected bsdtar or xorriso)" "${1:-}"
|
||||
}
|
||||
|
||||
dump_memtest_debug() {
|
||||
phase="$1"
|
||||
lb_dir="${2:-}"
|
||||
iso_path="${3:-}"
|
||||
phase_slug="$(printf '%s' "${phase}" | tr ' /' '__')"
|
||||
memtest_log="${LOG_DIR:-}/memtest-${phase_slug}.log"
|
||||
|
||||
(
|
||||
echo "=== memtest debug: ${phase} ==="
|
||||
|
||||
echo "-- auto/config --"
|
||||
if [ -f "${BUILDER_DIR}/auto/config" ]; then
|
||||
grep -n -- '--memtest' "${BUILDER_DIR}/auto/config" || echo " (no --memtest line found)"
|
||||
else
|
||||
echo " (missing ${BUILDER_DIR}/auto/config)"
|
||||
fi
|
||||
|
||||
echo "-- source bootloader templates --"
|
||||
for cfg in \
|
||||
"${BUILDER_DIR}/config/bootloaders/grub-efi/grub.cfg" \
|
||||
"${BUILDER_DIR}/config/bootloaders/isolinux/live.cfg.in"; do
|
||||
if [ -f "$cfg" ]; then
|
||||
echo " file: $cfg"
|
||||
grep -n 'Memory Test\|memtest' "$cfg" || echo " (no memtest lines)"
|
||||
fi
|
||||
done
|
||||
|
||||
echo "-- source binary hooks --"
|
||||
for hook in \
|
||||
"${BUILDER_DIR}/config/hooks/normal/9100-memtest.hook.binary"; do
|
||||
if [ -f "$hook" ]; then
|
||||
echo " hook: $hook"
|
||||
else
|
||||
echo " (missing $hook)"
|
||||
fi
|
||||
done
|
||||
|
||||
if [ -n "$lb_dir" ] && [ -d "$lb_dir" ]; then
|
||||
echo "-- live-build workdir package lists --"
|
||||
for pkg in \
|
||||
"$lb_dir/config/package-lists/bee.list.chroot" \
|
||||
"$lb_dir/config/package-lists/bee-gpu.list.chroot" \
|
||||
"$lb_dir/config/package-lists/bee-nvidia.list.chroot"; do
|
||||
if [ -f "$pkg" ]; then
|
||||
echo " file: $pkg"
|
||||
grep -n 'memtest' "$pkg" || echo " (no memtest lines)"
|
||||
fi
|
||||
done
|
||||
|
||||
echo "-- live-build chroot/boot --"
|
||||
if [ -d "$lb_dir/chroot/boot" ]; then
|
||||
find "$lb_dir/chroot/boot" -maxdepth 1 -name 'memtest*' -print | sed 's/^/ /' || true
|
||||
else
|
||||
echo " (missing $lb_dir/chroot/boot)"
|
||||
fi
|
||||
|
||||
echo "-- live-build binary/boot --"
|
||||
if [ -d "$lb_dir/binary/boot" ]; then
|
||||
find "$lb_dir/binary/boot" -maxdepth 1 -name 'memtest*' -print | sed 's/^/ /' || true
|
||||
else
|
||||
echo " (missing $lb_dir/binary/boot)"
|
||||
fi
|
||||
|
||||
echo "-- live-build binary grub cfg --"
|
||||
if [ -f "$lb_dir/binary/boot/grub/grub.cfg" ]; then
|
||||
grep -n 'Memory Test\|memtest' "$lb_dir/binary/boot/grub/grub.cfg" || echo " (no memtest lines)"
|
||||
else
|
||||
echo " (missing $lb_dir/binary/boot/grub/grub.cfg)"
|
||||
fi
|
||||
|
||||
echo "-- live-build binary isolinux cfg --"
|
||||
if [ -f "$lb_dir/binary/isolinux/live.cfg" ]; then
|
||||
grep -n 'Memory Test\|memtest' "$lb_dir/binary/isolinux/live.cfg" || echo " (no memtest lines)"
|
||||
else
|
||||
echo " (missing $lb_dir/binary/isolinux/live.cfg)"
|
||||
fi
|
||||
|
||||
echo "-- live-build package cache --"
|
||||
if [ -d "$lb_dir/cache/packages.chroot" ]; then
|
||||
find "$lb_dir/cache/packages.chroot" -maxdepth 1 -name 'memtest86+*.deb' -print | sed 's/^/ /' || true
|
||||
else
|
||||
echo " (missing $lb_dir/cache/packages.chroot)"
|
||||
fi
|
||||
fi
|
||||
|
||||
if [ -n "$iso_path" ] && [ -f "$iso_path" ]; then
|
||||
iso_files="$(mktemp)"
|
||||
iso_grub_cfg="$(mktemp)"
|
||||
iso_isolinux_cfg="$(mktemp)"
|
||||
|
||||
echo "-- ISO memtest files --"
|
||||
if iso_read_file_list "$iso_path" "$iso_files"; then
|
||||
grep 'memtest' "$iso_files" | sed 's/^/ /' || echo " (no memtest files in ISO)"
|
||||
else
|
||||
echo " (failed to list ISO contents)"
|
||||
fi
|
||||
|
||||
echo "-- ISO GRUB memtest lines --"
|
||||
if iso_read_member "$iso_path" boot/grub/grub.cfg "$iso_grub_cfg"; then
|
||||
grep -n 'Memory Test\|memtest' "$iso_grub_cfg" || echo " (no memtest lines in boot/grub/grub.cfg)"
|
||||
else
|
||||
echo " (failed to read boot/grub/grub.cfg from ISO)"
|
||||
fi
|
||||
|
||||
echo "-- ISO isolinux memtest lines --"
|
||||
if iso_read_member "$iso_path" isolinux/live.cfg "$iso_isolinux_cfg"; then
|
||||
grep -n 'Memory Test\|memtest' "$iso_isolinux_cfg" || echo " (no memtest lines in isolinux/live.cfg)"
|
||||
else
|
||||
echo " (failed to read isolinux/live.cfg from ISO)"
|
||||
fi
|
||||
|
||||
rm -f "$iso_files" "$iso_grub_cfg" "$iso_isolinux_cfg"
|
||||
fi
|
||||
|
||||
echo "=== end memtest debug: ${phase} ==="
|
||||
) | {
|
||||
if [ -n "${LOG_DIR:-}" ] && [ -d "${LOG_DIR}" ]; then
|
||||
tee "${memtest_log}"
|
||||
else
|
||||
cat
|
||||
fi
|
||||
}
|
||||
}
|
||||
|
||||
memtest_fail() {
|
||||
msg="$1"
|
||||
iso_path="${2:-}"
|
||||
level="WARNING"
|
||||
if [ "${BEE_REQUIRE_MEMTEST:-0}" = "1" ]; then
|
||||
level="ERROR"
|
||||
fi
|
||||
echo "${level}: ${msg}" >&2
|
||||
dump_memtest_debug "failure" "${LB_DIR:-}" "$iso_path" >&2
|
||||
if [ "${BEE_REQUIRE_MEMTEST:-0}" = "1" ]; then
|
||||
exit 1
|
||||
fi
|
||||
return 0
|
||||
}
|
||||
|
||||
nvidia_runtime_fail() {
|
||||
msg="$1"
|
||||
echo "ERROR: ${msg}" >&2
|
||||
exit 1
|
||||
}
|
||||
|
||||
iso_memtest_present() {
|
||||
iso_path="$1"
|
||||
iso_files="$(mktemp)"
|
||||
|
||||
[ -f "$iso_path" ] || return 1
|
||||
|
||||
if command -v bsdtar >/dev/null 2>&1; then
|
||||
:
|
||||
elif command -v xorriso >/dev/null 2>&1; then
|
||||
:
|
||||
else
|
||||
return 2
|
||||
fi
|
||||
|
||||
iso_read_file_list "$iso_path" "$iso_files" || {
|
||||
rm -f "$iso_files"
|
||||
return 2
|
||||
}
|
||||
|
||||
grep -q '^boot/memtest86+x64\.bin$' "$iso_files" || {
|
||||
rm -f "$iso_files"
|
||||
return 1
|
||||
}
|
||||
grep -q '^boot/memtest86+x64\.efi$' "$iso_files" || {
|
||||
rm -f "$iso_files"
|
||||
return 1
|
||||
}
|
||||
|
||||
grub_cfg="$(mktemp)"
|
||||
isolinux_cfg="$(mktemp)"
|
||||
|
||||
iso_read_member "$iso_path" boot/grub/grub.cfg "$grub_cfg" || {
|
||||
rm -f "$iso_files" "$grub_cfg" "$isolinux_cfg"
|
||||
return 2
|
||||
}
|
||||
iso_read_member "$iso_path" isolinux/live.cfg "$isolinux_cfg" || {
|
||||
rm -f "$iso_files" "$grub_cfg" "$isolinux_cfg"
|
||||
return 2
|
||||
}
|
||||
|
||||
grep -q 'Memory Test (memtest86+)' "$grub_cfg" || {
|
||||
rm -f "$iso_files" "$grub_cfg" "$isolinux_cfg"
|
||||
return 1
|
||||
}
|
||||
grep -q '/boot/memtest86+x64\.efi' "$grub_cfg" || {
|
||||
rm -f "$iso_files" "$grub_cfg" "$isolinux_cfg"
|
||||
return 1
|
||||
}
|
||||
grep -q '/boot/memtest86+x64\.bin' "$grub_cfg" || {
|
||||
rm -f "$iso_files" "$grub_cfg" "$isolinux_cfg"
|
||||
return 1
|
||||
}
|
||||
grep -q 'Memory Test (memtest86+)' "$isolinux_cfg" || {
|
||||
rm -f "$iso_files" "$grub_cfg" "$isolinux_cfg"
|
||||
return 1
|
||||
}
|
||||
grep -q '/boot/memtest86+x64\.bin' "$isolinux_cfg" || {
|
||||
rm -f "$iso_files" "$grub_cfg" "$isolinux_cfg"
|
||||
return 1
|
||||
}
|
||||
|
||||
rm -f "$iso_files" "$grub_cfg" "$isolinux_cfg"
|
||||
return 0
|
||||
}
|
||||
|
||||
validate_iso_memtest() {
|
||||
iso_path="$1"
|
||||
echo "=== validating memtest in ISO ==="
|
||||
|
||||
[ -f "$iso_path" ] || {
|
||||
memtest_fail "ISO not found for validation: $iso_path" "$iso_path"
|
||||
return 0
|
||||
}
|
||||
require_iso_reader "$iso_path" || return 0
|
||||
|
||||
iso_files="$(mktemp)"
|
||||
iso_read_file_list "$iso_path" "$iso_files" || {
|
||||
memtest_fail "failed to list ISO contents while validating memtest" "$iso_path"
|
||||
rm -f "$iso_files"
|
||||
return 0
|
||||
}
|
||||
|
||||
grep -q '^boot/memtest86+x64\.bin$' "$iso_files" || {
|
||||
memtest_fail "memtest BIOS binary missing in ISO: boot/memtest86+x64.bin" "$iso_path"
|
||||
rm -f "$iso_files"
|
||||
return 0
|
||||
}
|
||||
grep -q '^boot/memtest86+x64\.efi$' "$iso_files" || {
|
||||
memtest_fail "memtest EFI binary missing in ISO: boot/memtest86+x64.efi" "$iso_path"
|
||||
rm -f "$iso_files"
|
||||
return 0
|
||||
}
|
||||
|
||||
grub_cfg="$(mktemp)"
|
||||
isolinux_cfg="$(mktemp)"
|
||||
|
||||
iso_read_member "$iso_path" boot/grub/grub.cfg "$grub_cfg" || {
|
||||
memtest_fail "failed to read boot/grub/grub.cfg from ISO" "$iso_path"
|
||||
rm -f "$iso_files" "$grub_cfg" "$isolinux_cfg"
|
||||
return 0
|
||||
}
|
||||
iso_read_member "$iso_path" isolinux/live.cfg "$isolinux_cfg" || {
|
||||
memtest_fail "failed to read isolinux/live.cfg from ISO" "$iso_path"
|
||||
rm -f "$iso_files" "$grub_cfg" "$isolinux_cfg"
|
||||
return 0
|
||||
}
|
||||
|
||||
grep -q 'Memory Test (memtest86+)' "$grub_cfg" || {
|
||||
memtest_fail "GRUB menu entry for memtest is missing" "$iso_path"
|
||||
rm -f "$iso_files" "$grub_cfg" "$isolinux_cfg"
|
||||
return 0
|
||||
}
|
||||
grep -q '/boot/memtest86+x64\.efi' "$grub_cfg" || {
|
||||
memtest_fail "GRUB memtest EFI path is missing" "$iso_path"
|
||||
rm -f "$iso_files" "$grub_cfg" "$isolinux_cfg"
|
||||
return 0
|
||||
}
|
||||
grep -q '/boot/memtest86+x64\.bin' "$grub_cfg" || {
|
||||
memtest_fail "GRUB memtest BIOS path is missing" "$iso_path"
|
||||
rm -f "$iso_files" "$grub_cfg" "$isolinux_cfg"
|
||||
return 0
|
||||
}
|
||||
grep -q 'Memory Test (memtest86+)' "$isolinux_cfg" || {
|
||||
memtest_fail "isolinux menu entry for memtest is missing" "$iso_path"
|
||||
rm -f "$iso_files" "$grub_cfg" "$isolinux_cfg"
|
||||
return 0
|
||||
}
|
||||
grep -q '/boot/memtest86+x64\.bin' "$isolinux_cfg" || {
|
||||
memtest_fail "isolinux memtest path is missing" "$iso_path"
|
||||
rm -f "$iso_files" "$grub_cfg" "$isolinux_cfg"
|
||||
return 0
|
||||
}
|
||||
|
||||
rm -f "$iso_files" "$grub_cfg" "$isolinux_cfg"
|
||||
echo "=== memtest validation OK ==="
|
||||
}
|
||||
|
||||
validate_live_cmdline_params() {
|
||||
cfg="$1"
|
||||
command="$2"
|
||||
bootloader="$3"
|
||||
expected_label="$4"
|
||||
|
||||
awk -v command="$command" -v bootloader="$bootloader" -v expected_label="$expected_label" '
|
||||
function has(token, i) {
|
||||
for (i = 1; i <= NF; i++) {
|
||||
if ($i == token) {
|
||||
return 1
|
||||
}
|
||||
}
|
||||
return 0
|
||||
}
|
||||
|
||||
function reject(message) {
|
||||
printf "ERROR: %s live entry at %s:%d: %s\n", bootloader, FILENAME, NR, message
|
||||
bad = 1
|
||||
}
|
||||
|
||||
$1 == command && has("boot=live") {
|
||||
live_entries++
|
||||
|
||||
if (!has("udev.children_max=1")) {
|
||||
reject("missing udev.children_max=1")
|
||||
}
|
||||
if (!has("intel_iommu=on")) {
|
||||
reject("missing intel_iommu=on")
|
||||
}
|
||||
if (!has("iommu.passthrough=0")) {
|
||||
reject("missing iommu.passthrough=0")
|
||||
}
|
||||
if (!has("efi=disable_early_pci_dma")) {
|
||||
reject("missing efi=disable_early_pci_dma")
|
||||
}
|
||||
if (!has("live-media-label=" expected_label)) {
|
||||
reject("missing expected live-media-label=" expected_label)
|
||||
}
|
||||
if (has("iommu=pt")) {
|
||||
reject("contains forbidden iommu=pt")
|
||||
}
|
||||
|
||||
if (has("pci=realloc")) {
|
||||
failsafe_entries++
|
||||
if (!has("iommu.strict=1")) {
|
||||
reject("pci=realloc entry is missing iommu.strict=1")
|
||||
}
|
||||
} else if (has("iommu.strict=1")) {
|
||||
reject("iommu.strict=1 is allowed only in the pci=realloc fail-safe entry")
|
||||
}
|
||||
}
|
||||
|
||||
END {
|
||||
if (live_entries == 0) {
|
||||
printf "ERROR: %s config has no live boot entries\n", bootloader
|
||||
bad = 1
|
||||
}
|
||||
if (failsafe_entries != 1) {
|
||||
printf "ERROR: %s config has %d pci=realloc fail-safe entries, expected 1\n", bootloader, failsafe_entries
|
||||
bad = 1
|
||||
}
|
||||
exit bad ? 1 : 0
|
||||
}
|
||||
' "$cfg"
|
||||
}
|
||||
|
||||
validate_iso_live_boot_entries() {
|
||||
iso_path="$1"
|
||||
echo "=== validating live boot entries in ISO ==="
|
||||
|
||||
[ -f "$iso_path" ] || {
|
||||
echo "ERROR: ISO not found for live boot validation: $iso_path" >&2
|
||||
exit 1
|
||||
}
|
||||
require_iso_reader "$iso_path" >/dev/null 2>&1 || {
|
||||
echo "ERROR: ISO reader unavailable for live boot validation" >&2
|
||||
exit 1
|
||||
}
|
||||
|
||||
grub_cfg="$(mktemp)"
|
||||
isolinux_cfg="$(mktemp)"
|
||||
|
||||
iso_read_member "$iso_path" boot/grub/grub.cfg "$grub_cfg" || {
|
||||
echo "ERROR: failed to read boot/grub/grub.cfg from ISO" >&2
|
||||
rm -f "$grub_cfg" "$isolinux_cfg"
|
||||
exit 1
|
||||
}
|
||||
iso_read_member "$iso_path" isolinux/live.cfg "$isolinux_cfg" || {
|
||||
echo "ERROR: failed to read isolinux/live.cfg from ISO" >&2
|
||||
rm -f "$grub_cfg" "$isolinux_cfg"
|
||||
exit 1
|
||||
}
|
||||
|
||||
if grep -q '@APPEND_LIVE@\|@KERNEL_LIVE@\|@INITRD_LIVE@' "$grub_cfg" "$isolinux_cfg"; then
|
||||
echo "ERROR: unresolved live-build placeholders remain in ISO bootloader config" >&2
|
||||
rm -f "$grub_cfg" "$isolinux_cfg"
|
||||
exit 1
|
||||
fi
|
||||
|
||||
if grep -q 'iommu=pt' "$grub_cfg" "$isolinux_cfg"; then
|
||||
echo "ERROR: forbidden iommu=pt remains in ISO bootloader config" >&2
|
||||
rm -f "$grub_cfg" "$isolinux_cfg"
|
||||
exit 1
|
||||
fi
|
||||
|
||||
if ! validate_live_cmdline_params "$grub_cfg" linux GRUB "${BEE_ISO_VOLUME}"; then
|
||||
rm -f "$grub_cfg" "$isolinux_cfg"
|
||||
exit 1
|
||||
fi
|
||||
if ! validate_live_cmdline_params "$isolinux_cfg" append isolinux "${BEE_ISO_VOLUME}"; then
|
||||
rm -f "$grub_cfg" "$isolinux_cfg"
|
||||
exit 1
|
||||
fi
|
||||
|
||||
grep -q 'menuentry "EASY-BEE v' "$grub_cfg" || {
|
||||
echo "ERROR: GRUB default EASY-BEE entry is missing" >&2
|
||||
rm -f "$grub_cfg" "$isolinux_cfg"
|
||||
exit 1
|
||||
}
|
||||
grep -Fq "menuentry \"EASY-BEE v${PROJECT_VERSION_EFFECTIVE}\"" "$grub_cfg" || {
|
||||
echo "ERROR: GRUB version does not match ${PROJECT_VERSION_EFFECTIVE}" >&2
|
||||
rm -f "$grub_cfg" "$isolinux_cfg"
|
||||
exit 1
|
||||
}
|
||||
grep -Fq "menu label ^EASY-BEE v${PROJECT_VERSION_EFFECTIVE}" "$isolinux_cfg" || {
|
||||
echo "ERROR: isolinux version does not match ${PROJECT_VERSION_EFFECTIVE}" >&2
|
||||
rm -f "$grub_cfg" "$isolinux_cfg"
|
||||
exit 1
|
||||
}
|
||||
|
||||
if command -v xorriso >/dev/null 2>&1; then
|
||||
iso_volume="$(xorriso -indev "$iso_path" -pvd_info 2>/dev/null | awk -F: '/Volume [Ii]d/ { sub(/^[[:space:]'\''"]+/, "", $2); sub(/[[:space:]'\''"]+$/, "", $2); print $2; exit }')"
|
||||
if [ "$iso_volume" != "${BEE_ISO_VOLUME}" ]; then
|
||||
echo "ERROR: ISO volume ID is ${iso_volume:-unknown}, expected ${BEE_ISO_VOLUME}" >&2
|
||||
rm -f "$grub_cfg" "$isolinux_cfg"
|
||||
exit 1
|
||||
fi
|
||||
fi
|
||||
grep -q 'menuentry "EASY-BEE v.* -- load to RAM (toram)"' "$grub_cfg" || {
|
||||
echo "ERROR: GRUB toram entry is missing" >&2
|
||||
rm -f "$grub_cfg" "$isolinux_cfg"
|
||||
exit 1
|
||||
}
|
||||
grep -q 'linux .*boot=live ' "$grub_cfg" || {
|
||||
echo "ERROR: GRUB live entry is missing boot=live" >&2
|
||||
rm -f "$grub_cfg" "$isolinux_cfg"
|
||||
exit 1
|
||||
}
|
||||
grep -q 'linux .*boot=live .*toram ' "$grub_cfg" || {
|
||||
echo "ERROR: GRUB toram entry is missing boot=live or toram" >&2
|
||||
rm -f "$grub_cfg" "$isolinux_cfg"
|
||||
exit 1
|
||||
}
|
||||
grep -q 'linux .*live-media-label=EASY_BEE_' "$grub_cfg" || {
|
||||
echo "ERROR: GRUB live entry is missing live-media-label pinning" >&2
|
||||
rm -f "$grub_cfg" "$isolinux_cfg"
|
||||
exit 1
|
||||
}
|
||||
|
||||
grep -q 'append .*boot=live ' "$isolinux_cfg" || {
|
||||
echo "ERROR: isolinux live entry is missing boot=live" >&2
|
||||
rm -f "$grub_cfg" "$isolinux_cfg"
|
||||
exit 1
|
||||
}
|
||||
grep -q 'append .*boot=live .*toram ' "$isolinux_cfg" || {
|
||||
echo "ERROR: isolinux toram entry is missing boot=live or toram" >&2
|
||||
rm -f "$grub_cfg" "$isolinux_cfg"
|
||||
exit 1
|
||||
}
|
||||
grep -q 'append .*live-media-label=EASY_BEE_' "$isolinux_cfg" || {
|
||||
echo "ERROR: isolinux live entry is missing live-media-label pinning" >&2
|
||||
rm -f "$grub_cfg" "$isolinux_cfg"
|
||||
exit 1
|
||||
}
|
||||
|
||||
rm -f "$grub_cfg" "$isolinux_cfg"
|
||||
echo "=== live boot validation OK ==="
|
||||
}
|
||||
|
||||
validate_iso_grub_assets() {
|
||||
iso_path="$1"
|
||||
echo "=== validating GRUB assets in ISO ==="
|
||||
|
||||
[ -f "$iso_path" ] || {
|
||||
echo "ERROR: ISO not found for GRUB asset validation: $iso_path" >&2
|
||||
exit 1
|
||||
}
|
||||
require_iso_reader "$iso_path" >/dev/null 2>&1 || {
|
||||
echo "ERROR: ISO reader unavailable for GRUB asset validation" >&2
|
||||
exit 1
|
||||
}
|
||||
|
||||
iso_files="$(mktemp)"
|
||||
iso_list_files "$iso_path" > "$iso_files" || {
|
||||
echo "ERROR: failed to list ISO files for GRUB asset validation" >&2
|
||||
rm -f "$iso_files"
|
||||
exit 1
|
||||
}
|
||||
|
||||
for required in \
|
||||
boot/grub/config.cfg \
|
||||
boot/grub/grub.cfg; do
|
||||
grep -q "^${required}$" "$iso_files" || {
|
||||
echo "ERROR: missing GRUB asset in ISO: ${required}" >&2
|
||||
rm -f "$iso_files"
|
||||
exit 1
|
||||
}
|
||||
done
|
||||
|
||||
rm -f "$iso_files"
|
||||
echo "=== GRUB asset validation OK ==="
|
||||
}
|
||||
|
||||
validate_iso_nvidia_runtime() {
|
||||
iso_path="$1"
|
||||
[ "$BEE_GPU_VENDOR" = "nvidia" ] || return 0
|
||||
|
||||
echo "=== validating NVIDIA runtime in ISO ==="
|
||||
|
||||
[ -f "$iso_path" ] || nvidia_runtime_fail "ISO not found for NVIDIA runtime validation: $iso_path"
|
||||
require_iso_reader "$iso_path" >/dev/null 2>&1 || nvidia_runtime_fail "ISO reader unavailable for NVIDIA runtime validation"
|
||||
command -v unsquashfs >/dev/null 2>&1 || nvidia_runtime_fail "unsquashfs is required for NVIDIA runtime validation"
|
||||
|
||||
squashfs_tmp="$(mktemp)"
|
||||
squashfs_list="$(mktemp)"
|
||||
iso_files="$(mktemp)"
|
||||
dpkg_status_dir="$(mktemp -d)"
|
||||
iso_list_files "$iso_path" > "$iso_files" || {
|
||||
rm -f "$squashfs_tmp" "$squashfs_list" "$iso_files"
|
||||
rm -rf "$dpkg_status_dir"
|
||||
nvidia_runtime_fail "failed to list ISO files for NVIDIA runtime validation"
|
||||
}
|
||||
grep '^live/.*\.squashfs$' "$iso_files" | while IFS= read -r squashfs_member; do
|
||||
iso_read_member "$iso_path" "$squashfs_member" "$squashfs_tmp" || {
|
||||
rm -f "$squashfs_tmp" "$squashfs_list" "$iso_files"
|
||||
rm -rf "$dpkg_status_dir"
|
||||
nvidia_runtime_fail "failed to extract $squashfs_member from ISO"
|
||||
}
|
||||
unsquashfs -ll "$squashfs_tmp" >> "$squashfs_list" 2>/dev/null || {
|
||||
rm -f "$squashfs_tmp" "$squashfs_list" "$iso_files"
|
||||
rm -rf "$dpkg_status_dir"
|
||||
nvidia_runtime_fail "failed to inspect $squashfs_member from ISO"
|
||||
}
|
||||
# var/lib/dpkg/status lives in whichever squashfs layer has the base
|
||||
# rootfs (not the usr/firmware split-off layers); harmless no-op on
|
||||
# the others.
|
||||
unsquashfs -d "${dpkg_status_dir}/extract" -f "$squashfs_tmp" var/lib/dpkg/status >/dev/null 2>&1 || true
|
||||
: > "$squashfs_tmp"
|
||||
done
|
||||
|
||||
grep -Eq 'usr/bin/dcgmi$' "$squashfs_list" || {
|
||||
rm -f "$squashfs_tmp" "$squashfs_list" "$iso_files"
|
||||
rm -rf "$dpkg_status_dir"
|
||||
nvidia_runtime_fail "dcgmi missing from final NVIDIA ISO"
|
||||
}
|
||||
grep -Eq 'usr/bin/nv-hostengine$' "$squashfs_list" || {
|
||||
rm -f "$squashfs_tmp" "$squashfs_list" "$iso_files"
|
||||
rm -rf "$dpkg_status_dir"
|
||||
nvidia_runtime_fail "nv-hostengine missing from final NVIDIA ISO"
|
||||
}
|
||||
grep -Eq 'usr/bin/dcgmproftester([0-9]+)?$' "$squashfs_list" || {
|
||||
rm -f "$squashfs_tmp" "$squashfs_list" "$iso_files"
|
||||
rm -rf "$dpkg_status_dir"
|
||||
nvidia_runtime_fail "dcgmproftester missing from final NVIDIA ISO"
|
||||
}
|
||||
|
||||
# Cross-check the DCGM package version actually baked into the squashfs
|
||||
# against VERSIONS. dcgmi/nv-hostengine/dcgmproftester presence alone
|
||||
# doesn't catch a stale squashfs served by a mis-detected fast-path build
|
||||
# (dcgmi stays present across DCGM versions); this does.
|
||||
dpkg_status_file="${dpkg_status_dir}/extract/var/lib/dpkg/status"
|
||||
if [ -f "$dpkg_status_file" ]; then
|
||||
_installed_dcgm_version="$(awk '
|
||||
/^Package: datacenter-gpu-manager-4-core$/ { in_pkg=1; next }
|
||||
/^Package: / { in_pkg=0 }
|
||||
in_pkg && /^Version: / { sub(/^Version: /, ""); print; exit }
|
||||
' "$dpkg_status_file")"
|
||||
if [ -z "$_installed_dcgm_version" ]; then
|
||||
echo "=== WARNING: datacenter-gpu-manager-4-core not found in ISO dpkg status; skipping DCGM version check ==="
|
||||
else
|
||||
_installed_dcgm_no_epoch="${_installed_dcgm_version#*:}"
|
||||
if [ "$_installed_dcgm_no_epoch" != "${DCGM_VERSION}" ]; then
|
||||
rm -f "$squashfs_tmp" "$squashfs_list" "$iso_files"
|
||||
rm -rf "$dpkg_status_dir"
|
||||
nvidia_runtime_fail "DCGM version mismatch: VERSIONS pins ${DCGM_VERSION} but ISO has ${_installed_dcgm_version} (stale squashfs; retry with --clean-build)"
|
||||
fi
|
||||
fi
|
||||
else
|
||||
echo "=== WARNING: could not read dpkg status from ISO; skipping DCGM version check ==="
|
||||
fi
|
||||
|
||||
rm -f "$squashfs_tmp" "$squashfs_list" "$iso_files"
|
||||
rm -rf "$dpkg_status_dir"
|
||||
echo "=== NVIDIA runtime validation OK ==="
|
||||
}
|
||||
Executable
+98
@@ -0,0 +1,98 @@
|
||||
#!/bin/sh
|
||||
set -eu
|
||||
|
||||
BUILDER_DIR="$(CDPATH= cd -- "$(dirname "$0")" && pwd)"
|
||||
TEST_ROOT="$(mktemp -d)"
|
||||
trap 'rm -rf "$TEST_ROOT"' EXIT INT TERM HUP
|
||||
|
||||
PROJECT_VERSION_EFFECTIVE="13.0-test"
|
||||
BEE_ISO_VOLUME="EASY_BEE_TEST"
|
||||
LOG_DIR="$TEST_ROOT/log"
|
||||
mkdir -p "$LOG_DIR"
|
||||
|
||||
. "$BUILDER_DIR/lib/iso-validation.sh"
|
||||
. "$BUILDER_DIR/lib/bootloader.sh"
|
||||
|
||||
literal_template="$TEST_ROOT/literal-template"
|
||||
literal_output="$TEST_ROOT/literal-output"
|
||||
printf '%s\n' '@VERSION@ @KERNEL@ @APPEND_LIVE@ @INITRD@' > "$literal_template"
|
||||
render_bootloader_template "$literal_template" "$literal_output" \
|
||||
'13&0#test' '@KERNEL@' '/live/vmlinuz\\literal' 'boot=live marker=a&b#c\\d' '@INITRD@' '/live/initrd.img'
|
||||
literal_expected='13&0#test /live/vmlinuz\\literal boot=live marker=a&b#c\\d /live/initrd.img'
|
||||
if [ "$(cat "$literal_output")" != "$literal_expected" ]; then
|
||||
echo "ERROR: bootloader renderer changed literal replacement characters" >&2
|
||||
exit 1
|
||||
fi
|
||||
|
||||
rendered_grub="$TEST_ROOT/grub.cfg"
|
||||
rendered_isolinux="$TEST_ROOT/live.cfg"
|
||||
sed \
|
||||
-e 's#@APPEND_LIVE@#boot=live live-media-label=EASY_BEE_TEST#g' \
|
||||
-e 's#@KERNEL_LIVE@#/live/vmlinuz#g' \
|
||||
-e 's#@INITRD_LIVE@#/live/initrd.img#g' \
|
||||
-e 's#@VERSION@#13.0-test#g' \
|
||||
"$BUILDER_DIR/config/bootloaders/grub-efi/grub.cfg" > "$rendered_grub"
|
||||
sed \
|
||||
-e 's#@APPEND_LIVE@#boot=live live-media-label=EASY_BEE_TEST#g' \
|
||||
-e 's#@LINUX@#/live/vmlinuz#g' \
|
||||
-e 's#@INITRD@#/live/initrd.img#g' \
|
||||
-e 's#@VERSION@#13.0-test#g' \
|
||||
"$BUILDER_DIR/config/bootloaders/isolinux/live.cfg.in" > "$rendered_isolinux"
|
||||
|
||||
validate_live_cmdline_params "$rendered_grub" linux GRUB "$BEE_ISO_VOLUME"
|
||||
validate_live_cmdline_params "$rendered_isolinux" append isolinux "$BEE_ISO_VOLUME"
|
||||
|
||||
missing_serialization="$TEST_ROOT/missing-serialization.cfg"
|
||||
sed 's/ udev\.children_max=1//' "$rendered_grub" > "$missing_serialization"
|
||||
if validate_live_cmdline_params "$missing_serialization" linux GRUB "$BEE_ISO_VOLUME" >/dev/null 2>&1; then
|
||||
echo "ERROR: validator accepted a live entry without udev.children_max=1" >&2
|
||||
exit 1
|
||||
fi
|
||||
|
||||
CACHE_ROOT="$TEST_ROOT/cache"
|
||||
BUILD_VARIANT="test"
|
||||
BUILD_WORK_DIR="$TEST_ROOT/work"
|
||||
OVERLAY_STAGE_DIR="$TEST_ROOT/overlay"
|
||||
DEBIAN_KERNEL_ABI="6.1.0-test"
|
||||
SQUASHFS_FILENAME="filesystem-v13.0-test.squashfs"
|
||||
DIST_DIR="$TEST_ROOT/dist"
|
||||
mkdir -p "$BUILD_WORK_DIR/binary/live" "$OVERLAY_STAGE_DIR"
|
||||
touch "$BUILD_WORK_DIR/live-image-amd64.hybrid.iso"
|
||||
touch "$BUILD_WORK_DIR/binary/live/$SQUASHFS_FILENAME"
|
||||
|
||||
. "$BUILDER_DIR/lib/fast-path.sh"
|
||||
|
||||
# Isolate the decision test from repository content and GNU find extensions.
|
||||
hash_heavy_config() { printf '%s\n' test-heavy-hash; }
|
||||
write_overlay_manifest() { find "$OVERLAY_STAGE_DIR" -mindepth 1 -print | sed "s#^$OVERLAY_STAGE_DIR/##" | sort > "$1"; }
|
||||
|
||||
printf '%s\n' test-heavy-hash > "$FULL_BUILD_HASH_FILE"
|
||||
printf '%s\n' "$DEBIAN_KERNEL_ABI" > "$FULL_BUILD_ABI_FILE"
|
||||
write_overlay_manifest "$FULL_BUILD_OVERLAY_MANIFEST"
|
||||
touch "$FULL_BUILD_MARKER"
|
||||
|
||||
if needs_full_build; then
|
||||
echo "ERROR: fast path was rejected for unchanged valid state" >&2
|
||||
exit 1
|
||||
fi
|
||||
|
||||
touch "$OVERLAY_STAGE_DIR/added-by-fast-path"
|
||||
if needs_full_build; then
|
||||
echo "ERROR: an additive overlay change unnecessarily forced a full build" >&2
|
||||
exit 1
|
||||
fi
|
||||
write_overlay_manifest "$FULL_BUILD_OVERLAY_MANIFEST"
|
||||
rm "$OVERLAY_STAGE_DIR/added-by-fast-path"
|
||||
if ! needs_full_build >/dev/null; then
|
||||
echo "ERROR: removal from the latest fast-path overlay was not detected" >&2
|
||||
exit 1
|
||||
fi
|
||||
touch "$OVERLAY_STAGE_DIR/added-by-fast-path"
|
||||
|
||||
printf '%s\n' different-abi > "$FULL_BUILD_ABI_FILE"
|
||||
if ! needs_full_build >/dev/null; then
|
||||
echo "ERROR: kernel ABI change did not force a full build" >&2
|
||||
exit 1
|
||||
fi
|
||||
|
||||
echo "build library tests: OK"
|
||||
@@ -1,5 +1,5 @@
|
||||
#!/bin/sh
|
||||
# bee-nvidia-load — load NVIDIA kernel modules and create device nodes
|
||||
# bee-nvidia-load - load NVIDIA kernel modules and create device nodes
|
||||
# Called by bee-nvidia.service at boot.
|
||||
|
||||
NVIDIA_KO_DIR="/usr/local/lib/nvidia"
|
||||
@@ -28,7 +28,7 @@ have_nvidia_gpu() {
|
||||
}
|
||||
|
||||
if ! have_nvidia_gpu; then
|
||||
log "no NVIDIA GPU detected — skipping module load"
|
||||
log "no NVIDIA GPU detected - skipping module load"
|
||||
exit 0
|
||||
fi
|
||||
|
||||
@@ -65,7 +65,8 @@ load_module() {
|
||||
mod="$1"
|
||||
shift
|
||||
ko="$NVIDIA_KO_DIR/${mod}.ko"
|
||||
[ -f "$ko" ] || ko="$NVIDIA_KO_DIR/${mod//-/_}.ko"
|
||||
mod_file="$(printf '%s' "$mod" | tr '-' '_')"
|
||||
[ -f "$ko" ] || ko="$NVIDIA_KO_DIR/${mod_file}.ko"
|
||||
if [ ! -f "$ko" ]; then
|
||||
log "WARN: not found: $ko"
|
||||
return 1
|
||||
@@ -90,7 +91,7 @@ load_module_with_gsp_fallback() {
|
||||
return 1
|
||||
fi
|
||||
|
||||
# Run insmod in background — on some converted SXM→PCIe cards GSP enters an
|
||||
# Run insmod in background. On some converted SXM-to-PCIe cards GSP enters an
|
||||
# infinite crash/reload loop and insmod never returns. We check for successful
|
||||
# initialization by polling /proc/devices for nvidiactl instead of waiting for
|
||||
# insmod to exit.
|
||||
@@ -114,29 +115,29 @@ load_module_with_gsp_fallback() {
|
||||
dmesg | tail -n 10 | sed 's/^/ dmesg: /' || true
|
||||
return 1
|
||||
fi
|
||||
# insmod exited 0 but nvidiactl not yet in /proc/devices — give it a moment
|
||||
# insmod exited 0 but nvidiactl is not yet in /proc/devices; give it a moment
|
||||
sleep 2
|
||||
if nvidia_is_functional; then
|
||||
log "loaded: nvidia (GSP enabled, ${_waited}s)"
|
||||
return 0
|
||||
fi
|
||||
log "insmod exited 0 but nvidiactl missing — treating as failure"
|
||||
log "insmod exited 0 but nvidiactl missing - treating as failure"
|
||||
return 1
|
||||
fi
|
||||
sleep 1
|
||||
_waited=$((_waited + 1))
|
||||
done
|
||||
|
||||
# GSP init timed out — kill the hanging insmod and attempt gsp-off fallback
|
||||
# GSP init timed out; kill the hanging insmod and attempt gsp-off fallback.
|
||||
log "nvidia GSP init timed out after 90s"
|
||||
kill "$_insmod_pid" 2>/dev/null || true
|
||||
wait "$_insmod_pid" 2>/dev/null || true
|
||||
|
||||
# Attempt to unload the partially-initialized module
|
||||
if ! rmmod nvidia 2>/dev/null; then
|
||||
# Module is stuck in the kernel — cannot reload with different params.
|
||||
# Module is stuck in the kernel; cannot reload with different params.
|
||||
# User must reboot and select bee.nvidia.mode=gsp-off at boot menu.
|
||||
log "ERROR: rmmod nvidia failed (EBUSY) — module stuck in kernel"
|
||||
log "ERROR: rmmod nvidia failed (EBUSY) - module stuck in kernel"
|
||||
log "ERROR: reboot and select 'EASY-BEE (advanced) -> GSP=off' in boot menu"
|
||||
echo "gsp-stuck" > /run/bee-nvidia-mode
|
||||
return 1
|
||||
@@ -144,7 +145,7 @@ load_module_with_gsp_fallback() {
|
||||
|
||||
sleep 2
|
||||
log "retrying with NVreg_EnableGpuFirmware=0"
|
||||
log "WARNING: GSP disabled — power management will run via CPU path, not GPU firmware"
|
||||
log "WARNING: GSP disabled - power management will run via CPU path, not GPU firmware"
|
||||
|
||||
if insmod "$ko" NVreg_EnableGpuFirmware=0; then
|
||||
if nvidia_is_functional; then
|
||||
@@ -208,7 +209,7 @@ else
|
||||
log "GSP-off mode: skipping nvidia-modeset and nvidia-uvm during boot"
|
||||
;;
|
||||
nomsi|*)
|
||||
# nomsi: disable MSI-X/MSI interrupts — use when RmInitAdapter fails with
|
||||
# nomsi: disable MSI-X/MSI interrupts; use when RmInitAdapter fails with
|
||||
# "Failed to enable MSI-X" on one or more GPUs (IOMMU group interrupt limits).
|
||||
# NVreg_EnableMSI=0 forces legacy INTx interrupts for all GPUs.
|
||||
if ! load_module nvidia NVreg_EnableGpuFirmware=0 NVreg_EnableMSI=0; then
|
||||
@@ -230,7 +231,7 @@ if [ -n "$nvidia_major" ]; then
|
||||
done
|
||||
log "created /dev/nvidia{0-7}"
|
||||
else
|
||||
log "WARN: nvidiactl not in /proc/devices — no GPU hardware present?"
|
||||
log "WARN: nvidiactl not in /proc/devices - no GPU hardware present?"
|
||||
fi
|
||||
|
||||
uvm_major=$(grep -m1 ' nvidia-uvm$' /proc/devices | awk '{print $1}')
|
||||
@@ -255,60 +256,40 @@ if command -v nvidia-smi >/dev/null 2>&1; then
|
||||
log "WARN: failed to enable NVIDIA persistence mode"
|
||||
fi
|
||||
else
|
||||
log "WARN: nvidia-smi not found — cannot enable persistence mode"
|
||||
log "WARN: nvidia-smi not found - cannot enable persistence mode"
|
||||
fi
|
||||
|
||||
# Bound every systemctl call below: a unit whose ExecStart/ExecCondition hangs
|
||||
# (e.g. fabricmanager stuck training a bad NVSwitch fabric) must not be able to
|
||||
# wedge bee-nvidia.service forever — that would keep nvidia-dcgm.service from
|
||||
# ever starting, since it's ordered After= this one. 60s comfortably covers a
|
||||
# normal fabricmanager/dcgm startup without blocking boot indefinitely.
|
||||
SYSTEMCTL_TIMEOUT=60
|
||||
timeout_systemctl() {
|
||||
timeout "${SYSTEMCTL_TIMEOUT}" systemctl "$@"
|
||||
# Refresh nvidia-fabricmanager and nvidia-dcgm so they (re)enumerate against
|
||||
# the device nodes we just created.
|
||||
#
|
||||
# These MUST NOT block. bee-nvidia.service is Type=oneshot and ordered
|
||||
# Before=nvidia-fabricmanager.service nvidia-dcgm.service, so systemd will not
|
||||
# run either unit until this script returns. A synchronous "systemctl restart"
|
||||
# here therefore deadlocks against our own ordering and only unwedges when its
|
||||
# timeout fires. "systemctl --no-block try-restart" queues a restart only for
|
||||
# an active unit and returns without waiting. An inactive enabled unit remains
|
||||
# in the normal boot transaction and can start after bee-nvidia.service exits.
|
||||
nvidia_refresh_unit() {
|
||||
unit="$1"
|
||||
if ! command -v systemctl >/dev/null 2>&1 ||
|
||||
! systemctl list-unit-files --no-legend 2>/dev/null | awk -v wanted="$unit" '$1 == wanted { found=1 } END { exit(found ? 0 : 1) }'; then
|
||||
log "WARN: ${unit} not installed"
|
||||
return
|
||||
fi
|
||||
if systemctl --no-block try-restart "$unit" >/dev/null 2>&1; then
|
||||
log "queued refresh of ${unit} (non-blocking)"
|
||||
else
|
||||
log "WARN: could not queue refresh of ${unit}"
|
||||
fi
|
||||
}
|
||||
|
||||
# Start or refresh Fabric Manager after the NVIDIA stack is ready. On NVSwitch
|
||||
# systems CUDA/DCGM can report "system not yet initialized" until fabric
|
||||
# training completes under nvidia-fabricmanager.
|
||||
if command -v systemctl >/dev/null 2>&1 && systemctl list-unit-files --no-legend 2>/dev/null | grep -q '^nvidia-fabricmanager\.service'; then
|
||||
log "restarting nvidia-fabricmanager.service (timeout ${SYSTEMCTL_TIMEOUT}s)"
|
||||
if timeout_systemctl restart nvidia-fabricmanager.service >/dev/null 2>&1; then
|
||||
log "nvidia-fabricmanager restarted"
|
||||
elif [ $? -eq 124 ]; then
|
||||
log "WARN: systemctl restart nvidia-fabricmanager.service timed out after ${SYSTEMCTL_TIMEOUT}s"
|
||||
elif timeout_systemctl start nvidia-fabricmanager.service >/dev/null 2>&1; then
|
||||
log "nvidia-fabricmanager started"
|
||||
else
|
||||
log "WARN: failed to start nvidia-fabricmanager.service"
|
||||
systemctl status nvidia-fabricmanager.service --no-pager 2>&1 | sed 's/^/ fabricmanager: /' || true
|
||||
fi
|
||||
else
|
||||
log "WARN: nvidia-fabricmanager.service not installed"
|
||||
fi
|
||||
# On NVSwitch systems CUDA/DCGM can report "system not yet initialized" until
|
||||
# fabric training completes under nvidia-fabricmanager; on non-NVSwitch boxes
|
||||
# the unit's ExecCondition skips it and this is a no-op.
|
||||
nvidia_refresh_unit nvidia-fabricmanager.service
|
||||
|
||||
# Restart the DCGM host engine so dcgmi can discover GPUs. nv-hostengine
|
||||
# enumerates GPUs once at startup and never rescans; bee-nvidia.service now
|
||||
# orders itself Before=nvidia-dcgm.service so systemd shouldn't start it until
|
||||
# modules/device nodes exist, but restart here too in case the unit was
|
||||
# already active from a previous boot/reload with a stale empty inventory.
|
||||
# Use systemctl (not a raw nv-hostengine invocation) so systemd's own
|
||||
# supervision of nvidia-dcgm.service stays authoritative and we don't end up
|
||||
# with two host engines racing for the same port.
|
||||
if command -v systemctl >/dev/null 2>&1 && systemctl list-unit-files --no-legend 2>/dev/null | grep -q '^nvidia-dcgm\.service'; then
|
||||
log "restarting nvidia-dcgm.service (timeout ${SYSTEMCTL_TIMEOUT}s)"
|
||||
if timeout_systemctl restart nvidia-dcgm.service >/dev/null 2>&1; then
|
||||
log "nvidia-dcgm restarted"
|
||||
elif [ $? -eq 124 ]; then
|
||||
log "WARN: systemctl restart nvidia-dcgm.service timed out after ${SYSTEMCTL_TIMEOUT}s"
|
||||
elif timeout_systemctl start nvidia-dcgm.service >/dev/null 2>&1; then
|
||||
log "nvidia-dcgm started"
|
||||
else
|
||||
log "WARN: failed to start nvidia-dcgm.service"
|
||||
systemctl status nvidia-dcgm.service --no-pager 2>&1 | sed 's/^/ nvidia-dcgm: /' || true
|
||||
fi
|
||||
else
|
||||
log "WARN: nvidia-dcgm.service not installed"
|
||||
fi
|
||||
# If nvidia-dcgm is already active, restart it after the device nodes exist so
|
||||
# its hostengine refreshes its device view. Otherwise normal boot starts it.
|
||||
nvidia_refresh_unit nvidia-dcgm.service
|
||||
|
||||
log "done"
|
||||
|
||||
Reference in New Issue
Block a user