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