mochimodev/winminer

WDDM causes kernel timed out

Opened this issue · 1 comments

Sometimes it shows "unspecified launch failure" since all gpu mem and kernel got crashed before running cuda_find_peach. Mostly happens on low-end CPU rigs (i don't know why), undervolted GPUs.
Solution is:

  1. unlock WDDM: https://docs.nvidia.com/gameworks/content/developertools/desktop/timeout_detection_recovery.htm

  2. split this kernel:
    https://github.com/mochimodev/winminer/blob/master/winminer_v1/winminer/algo/peach/cuda_peach.cu#L401

here is what my code does (which fixes almost all crashed cases):

__global__ void cuda_build_map(uint32_t offset, uint8_t *g_map)
{
    const uint32_t thread = offset + blockDim.x * blockIdx.x + threadIdx.x;
    if (thread < MAP)
        cuda_gen_tile(thread, g_map);
}

gridBuildMap = get_number_SM;
for (int offset_build_map = 0; offset_build_map < MAP; offset_build_map += 256 * gridBuildMap)
                {
                    cuda_build_map << <gridBuildMap, 256, 0, ctx.queue >> >(offset_build_map, ctx.d_map);
                }

since it only runs once per block, it won't hurt the performance much.

plus, you can't check kernel errors if it's not synced yet, at this line: https://github.com/mochimodev/winminer/blob/master/winminer_v1/winminer/algo/peach/cuda_peach.cu#L403