portelli/Grid
1
0
Fork 0
mirror of https://github.com/paboyle/Grid.git synced 2026-06-12 23:23:10 +01:00
Code Releases Activity

tests/debug: test plan-before-malloc vs malloc-before-plan ordering

Browse Source 
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
This commit is contained in:
Peter Boyle
2026-05-19 21:40:17 -04:00
parent ad9d03fd85
commit e5996b440d
1 changed files with 40 additions and 56 deletions
Download Patch File Download Diff File Expand all files Collapse all files
tests/debug/Test_hipfft_minimal.cc
+40 -56
View File
@@ -35,8 +35,11 @@ static const char *hipfftResultString(hipfftResult r) {
} }
} }
// Plan creation + execution for (G, howmany) using hipfftCreate+hipfftMakePlanMany. // Plan creation + execution for (G, howmany).
// This is the path Grid's FFT.h now uses. // Tests two orderings to isolate whether a prior hipMalloc poisons hipfft
// plan creation for small G on ROCm 7:
// A) plan BEFORE hipMalloc — hypothesis: succeeds
// B) hipMalloc BEFORE plan — hypothesis: fails for G < 32
static void tryPlanAndExec(int G, long howmany) { static void tryPlanAndExec(int G, long howmany) {
int n[] = {G}; int n[] = {G};
long nelems = (long)G * howmany; long nelems = (long)G * howmany;
@@ -44,68 +47,49 @@ static void tryPlanAndExec(int G, long howmany) {
printf("--- G=%-4d howmany=%-10ld total_elems=%-12ld ---\n", printf("--- G=%-4d howmany=%-10ld total_elems=%-12ld ---\n",
G, howmany, nelems); G, howmany, nelems);
// Allocate device buffer (hipfftDoubleComplex = 16 bytes each) // --- A: create plan first, allocate buffer afterwards ---
hipfftDoubleComplex *dbuf = nullptr;
hipError_t herr = hipMalloc(&dbuf, nelems * sizeof(hipfftDoubleComplex));
if (herr != hipSuccess) {
printf(" hipMalloc failed (%d) for %ld elems — skipping\n\n", (int)herr, nelems);
return;
}
hipMemset(dbuf, 0, nelems * sizeof(hipfftDoubleComplex));
// 1. hipfftPlanMany (one-step, nullptr embed) — current Grid path
{
hipfftHandle p;
hipfftResult rv = hipfftPlanMany(&p, 1, n,
nullptr, 1, G,
nullptr, 1, G,
HIPFFT_Z2Z, (int)howmany);
printf(" hipfftPlanMany create : %d (%s)\n", (int)rv, hipfftResultString(rv));
if (rv == HIPFFT_SUCCESS) {
rv = hipfftExecZ2Z(p, dbuf, dbuf, HIPFFT_FORWARD);
hipDeviceSynchronize();
printf(" hipfftPlanMany execFwd: %d (%s)\n", (int)rv, hipfftResultString(rv));
hipfftDestroy(p);
}
}
// 2. hipfftCreate + hipfftMakePlanMany (two-step) — also current Grid path
{ {
hipfftHandle p; hipfftHandle p;
size_t workSize = 0; size_t workSize = 0;
hipfftResult rc = hipfftCreate(&p); hipfftCreate(&p);
if (rc == HIPFFT_SUCCESS) { hipfftResult rv = hipfftMakePlanMany(p, 1, n,
hipfftResult rv = hipfftMakePlanMany(p, 1, n, nullptr, 1, G, nullptr, 1, G,
nullptr, 1, G, HIPFFT_Z2Z, (int)howmany, &workSize);
nullptr, 1, G, printf(" plan-first create : %d (%s)\n", (int)rv, hipfftResultString(rv));
HIPFFT_Z2Z, (int)howmany, &workSize);
printf(" hipfftMakePlanMany : %d (%s) workSize=%zu\n",
(int)rv, hipfftResultString(rv), workSize);
if (rv == HIPFFT_SUCCESS) {
rv = hipfftExecZ2Z(p, dbuf, dbuf, HIPFFT_FORWARD);
hipDeviceSynchronize();
printf(" hipfftMakePlanMany exec : %d (%s)\n", (int)rv, hipfftResultString(rv));
}
hipfftDestroy(p);
} else {
printf(" hipfftCreate : %d (%s)\n", (int)rc, hipfftResultString(rc));
}
}
// 3. hipfftPlan1d (simplest API, batch = howmany)
{
hipfftHandle p;
hipfftResult rv = hipfftPlan1d(&p, G, HIPFFT_Z2Z, (int)howmany);
printf(" hipfftPlan1d create : %d (%s)\n", (int)rv, hipfftResultString(rv));
if (rv == HIPFFT_SUCCESS) { if (rv == HIPFFT_SUCCESS) {
rv = hipfftExecZ2Z(p, dbuf, dbuf, HIPFFT_FORWARD); hipfftDoubleComplex *buf = nullptr;
hipMalloc(&buf, nelems * sizeof(hipfftDoubleComplex));
hipMemset(buf, 0, nelems * sizeof(hipfftDoubleComplex));
rv = hipfftExecZ2Z(p, buf, buf, HIPFFT_FORWARD);
hipDeviceSynchronize(); hipDeviceSynchronize();
printf(" hipfftPlan1d execFwd: %d (%s)\n", (int)rv, hipfftResultString(rv)); printf(" plan-first execFwd: %d (%s)\n", (int)rv, hipfftResultString(rv));
hipfftDestroy(p); hipFree(buf);
} }
hipfftDestroy(p);
}
// --- B: hipMalloc first, create plan afterwards ---
{
hipfftDoubleComplex *buf = nullptr;
hipMalloc(&buf, nelems * sizeof(hipfftDoubleComplex));
hipMemset(buf, 0, nelems * sizeof(hipfftDoubleComplex));
hipfftHandle p;
size_t workSize = 0;
hipfftCreate(&p);
hipfftResult rv = hipfftMakePlanMany(p, 1, n,
nullptr, 1, G, nullptr, 1, G,
HIPFFT_Z2Z, (int)howmany, &workSize);
printf(" malloc-first create : %d (%s)\n", (int)rv, hipfftResultString(rv));
if (rv == HIPFFT_SUCCESS) {
rv = hipfftExecZ2Z(p, buf, buf, HIPFFT_FORWARD);
hipDeviceSynchronize();
printf(" malloc-first execFwd: %d (%s)\n", (int)rv, hipfftResultString(rv));
}
hipfftDestroy(p);
hipFree(buf);
} }
hipFree(dbuf);
printf("\n"); printf("\n");
} }