SWDEV-479260,SWDEV-483599 - Check griddim Y,Z <= 65536

Gfx12 has 16 bits for grid dim Y/Z. Detect gfxIp and return error if dim y/z > 16 bits

Change-Id: I43dd14affc9e4073d0b1232e7523967f0180fa31
This commit is contained in:
pghafari
2024-09-06 19:04:21 -04:00
zatwierdzone przez Payam Ghafari
rodzic f8beeede22
commit 0a918c8f96
+7
Wyświetl plik
@@ -619,6 +619,13 @@ hipError_t ihipLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDi
return hipErrorInvalidDeviceFunction;
}
}
constexpr auto gridDimYZmax = static_cast<uint64_t>(std::numeric_limits<uint16_t>::max()) + 1;
const auto& isa = g_devices[deviceId]->devices()[0]->isa().versionMajor();
if (isa >= 12 && (gridDim.y > gridDimYZmax || gridDim.z > gridDimYZmax)) {
return hipErrorInvalidConfiguration;
}
size_t globalWorkSizeX = static_cast<size_t>(gridDim.x) * blockDim.x;
size_t globalWorkSizeY = static_cast<size_t>(gridDim.y) * blockDim.y;
size_t globalWorkSizeZ = static_cast<size_t>(gridDim.z) * blockDim.z;