| 26 | } |
| 27 | |
| 28 | ChecksumForward::Result ChecksumForwardImpl::exec( |
| 29 | _megdnn_tensor_in data, _megdnn_workspace workspace) { |
| 30 | auto wbundle = get_wbundle(data.layout); |
| 31 | wbundle.set(workspace.raw_ptr); |
| 32 | Result result; |
| 33 | memset(&result, 0, sizeof(result)); |
| 34 | check_exec(data.layout, workspace.size); |
| 35 | auto stream = cuda_stream(handle()); |
| 36 | |
| 37 | auto ptr = static_cast<uint8_t*>(data.raw_ptr()); |
| 38 | size_t size_all = data.layout.shape[0], size_ints = size_all / sizeof(uint32_t); |
| 39 | auto last_val_size = std::min<size_t>(size_all, 4); |
| 40 | cuda_check(cudaMemcpyAsync( |
| 41 | &result.last_val, ptr + size_all - last_val_size, last_val_size, |
| 42 | cudaMemcpyDeviceToHost, stream)); |
| 43 | if (size_ints) { |
| 44 | checksum::calc( |
| 45 | static_cast<uint32_t*>(wbundle.get(1)), |
| 46 | static_cast<uint32_t*>(data.raw_ptr()), |
| 47 | static_cast<uint32_t*>(wbundle.get(0)), size_ints, stream); |
| 48 | cuda_check(cudaMemcpyAsync( |
| 49 | &result.checksum, wbundle.get(1), sizeof(result.checksum), |
| 50 | cudaMemcpyDeviceToHost, stream)); |
| 51 | } |
| 52 | cuda_check(cudaStreamSynchronize(stream)); |
| 53 | return result; |
| 54 | } |
| 55 | |
| 56 | // vim: syntax=cpp.doxygen |
nothing calls this directly
no test coverage detected