Only the first device result aligns with the host’s computation

[mengllm]

@and-ivanov @benrothen

Hi,

Regardless of whether I use generated.bin or extracted.bin, and whether I use checksum_kernel or checksum_kernel_from_data, the device’s checksum result changes with each execution of cuLaunchKernel. Only the first device result aligns with the host’s computation. Is this phenomenon reasonable? If so, how can verification between the device and host be achieved?

add print msg in code: int warmup_repeats = warmup ? 10 : 0; State device_result_temp; for (int iters = 0; iters < warmup_repeats + repeats; iters++){ Time t1 = timer::now();
CUDA_DRV_CHECK(cuLaunchKernel(checksum_kernel,
/* grid size */ gridSize, 1, 1, /* block size */ blockSize, 1, 1,
/* shared mem */ 0, /* stream */ nullptr, args, 0));
CUDA_DRV_CHECK(cuCtxSynchronize()); // wait kernel to stop
Time t2 = timer::now(); if (iters >= warmup_repeats) { runtime += seconds(t2 - t1);
}
CUDA_DRV_CHECK(cuMemcpyDtoH(&device_result_temp, device_state, sizeof(State)));
printf("execution %d： device_result_temp.c: %" PRIx32 " \n", iters, device_result_temp.c);

logs: execution 0： device_result_temp.c: 8442a62d base address 0x7f0d9c000000 Data pointer alignment is good! execution 1： device_result_temp.c: 8387062d base address 0x7f0d9c000000 Data pointer alignment is good! execution 2： device_result_temp.c: 203062d base address 0x7f0d9c000000 Data pointer alignment is good! execution 3： device_result_temp.c: 26ab062d base address 0x7f0d9c000000 Data pointer alignment is good! execution 4： device_result_temp.c: f224462d base address 0x7f0d9c000000 Data pointer alignment is good! execution 5： device_result_temp.c: 6a21462d base address 0x7f0d9c000000 Data pointer alignment is good! execution 6： device_result_temp.c: edec462d base address 0x7f0d9c000000 Data pointer alignment is good! execution 7： device_result_temp.c: 22efe62d base address 0x7f0d9c000000 Data pointer alignment is good! execution 8： device_result_temp.c: c0d4062d base address 0x7f0d9c000000 Data pointer alignment is good! execution 9： device_result_temp.c: 81fa862d base address 0x7f0d9c000000 Data pointer alignment is good! execution 10： device_result_temp.c: 7129e62d checksum Runtime: 1.1418 s result DtoH took 2.1416e-05 s size:104 GPU clocks: 1609585358 Optimal clocks 684800000 Observed 43 % of peak performance Computing checksum on host... 100% Verification on host took 376.599 s verification FAILED! dev: 7129e62d host: 8442a62d

Meanwhile,regardless of how many times the host computes the checksum,the result remains unchanged: Computing checksum on host... 100% host verification result: 8442a62d Computing checksum on host... 100% host verification result: 8442a62d Computing checksum on host... 100% host verification result: 8442a62d Computing checksum on host... 100% host verification result: 8442a62d Computing checksum on host... 100% host verification result: 8442a62d Computing checksum on host... 100% host verification result: 8442a62d Computing checksum on host... 100% host verification result: 8442a62d Computing checksum on host... 100% host verification result: 8442a62d Computing checksum on host... 100% host verification result: 8442a62d Computing checksum on host... 100% host verification result: 8442a62d Computing checksum on host... 100% host verification result: 8442a62d