| 243 | } |
| 244 | |
| 245 | void SparseFrameBuffer::setTiles(const std::vector<uint32_t> &_tileIDs) |
| 246 | { |
| 247 | RKCOMMON_IF_TRACING_ENABLED({ |
| 248 | rkcommon::tracing::beginEvent("SparseFB::setTiles", "ospray"); |
| 249 | rkcommon::tracing::setCounter("SparseFB::numTiles", _tileIDs.size()); |
| 250 | }); |
| 251 | // (Re-)configure the sparse framebuffer based on the tileIDs we're passed |
| 252 | tileIDs = _tileIDs; |
| 253 | numRenderTasks = |
| 254 | vec2i(tileIDs.size() * TILE_SIZE, TILE_SIZE) / getRenderTaskSize(); |
| 255 | |
| 256 | if (hasVarianceBuffer && !tileIDs.empty()) { |
| 257 | taskErrorBuffer = devicert::make_buffer_shared_unique<float>( |
| 258 | getISPCDevice().getDRTDevice(), numRenderTasks.long_product()); |
| 259 | std::fill(taskErrorBuffer->begin(), taskErrorBuffer->end(), inf); |
| 260 | } else { |
| 261 | taskErrorBuffer = nullptr; |
| 262 | } |
| 263 | |
| 264 | if (!tileIDs.empty()) { |
| 265 | tiles = devicert::make_buffer_device_shadowed_unique<Tile>( |
| 266 | getISPCDevice().getDRTDevice(), tileIDs.size()); |
| 267 | const vec2f rcpSize = rcp(vec2f(size)); |
| 268 | #ifndef OSPRAY_TARGET_SYCL |
| 269 | rkcommon::tasking::parallel_for(tiles->size(), [&](size_t i) { |
| 270 | Tile &t = (*tiles)[i]; |
| 271 | t.fbSize = size; |
| 272 | t.rcp_fbSize = rcpSize; |
| 273 | t.region = getTileRegion(tileIDs[i]); |
| 274 | t.accumID = 0; |
| 275 | }); |
| 276 | #else |
| 277 | // TODO: refactor and simplify this, we don't need so many copies of TileIDs |
| 278 | devicert::BufferDeviceShadowedImpl<uint32_t> deviceTileIDs( |
| 279 | device.getDRTDevice(), tileIDs); |
| 280 | deviceTileIDs.copyToDevice(); |
| 281 | |
| 282 | // In SYCL, populate the tiles on the device. |
| 283 | // TODO: Best to unify the codepaths more here and do the ISPC device |
| 284 | // tile population in ISPC so it also runs on the "device" |
| 285 | const uint32_t *deviceTileIDsPtr = deviceTileIDs.devicePtr(); |
| 286 | const size_t numTasks = tiles->size(); |
| 287 | const vec2i fbSize = size; |
| 288 | const vec2i fbTotalTiles = totalTiles; |
| 289 | Tile *tilesDevice = tiles->devicePtr(); |
| 290 | sycl::queue *queue = |
| 291 | static_cast<sycl::queue *>(device.getDRTDevice().getSyclQueuePtr()); |
| 292 | queue |
| 293 | ->submit([&](sycl::handler &cgh) { |
| 294 | const sycl::nd_range<1> dispatchRange = |
| 295 | device.computeDispatchRange(numTasks, 16); |
| 296 | cgh.parallel_for(dispatchRange, [=](sycl::nd_item<1> taskIndex) { |
| 297 | if (taskIndex.get_global_id(0) < numTasks) { |
| 298 | const size_t tid = taskIndex.get_global_id(0); |
| 299 | tilesDevice[tid].fbSize = fbSize; |
| 300 | tilesDevice[tid].rcp_fbSize = rcpSize; |
| 301 | tilesDevice[tid].region = ospray::getTileRegion( |
| 302 | deviceTileIDsPtr[tid], fbSize, fbTotalTiles); |
no test coverage detected