| 1260 | } |
| 1261 | |
| 1262 | void BM_WhileLoop(int num_iters) { |
| 1263 | // Benchmark a simple kernel to measure while loop overheads. |
| 1264 | tensorflow::testing::StopTiming(); |
| 1265 | |
| 1266 | se::Platform* platform = PlatformUtil::GetDefaultPlatform().ValueOrDie(); |
| 1267 | auto executors = PlatformUtil::GetStreamExecutors(platform).ValueOrDie(); |
| 1268 | se::StreamExecutorMemoryAllocator allocator(platform, executors); |
| 1269 | LocalClient* client = |
| 1270 | ClientLibrary::GetOrCreateLocalClient(platform).ValueOrDie(); |
| 1271 | |
| 1272 | const int64 seq_len = 100; |
| 1273 | Shape loop_state_shape = ShapeUtil::MakeTupleShape( |
| 1274 | {ShapeUtil::MakeShape(S32, {}), |
| 1275 | ShapeUtil::MakeShape(F32, {seq_len, 1024, 1024})}); |
| 1276 | |
| 1277 | // Create while condition computation with 'loop_limit'. |
| 1278 | const int32 loop_limit = 100; |
| 1279 | XlaComputation condition; |
| 1280 | { |
| 1281 | XlaBuilder builder("condition"); |
| 1282 | auto prev = Parameter(&builder, 0, loop_state_shape, "prev"); |
| 1283 | auto iteration = GetTupleElement(prev, 0); |
| 1284 | Lt(iteration, ConstantR0<int32>(&builder, loop_limit)); |
| 1285 | condition = builder.Build().ConsumeValueOrDie(); |
| 1286 | } |
| 1287 | |
| 1288 | // Create while body computation with unit loop increment. |
| 1289 | XlaComputation body; |
| 1290 | { |
| 1291 | XlaBuilder builder("body"); |
| 1292 | auto prev = Parameter(&builder, 0, loop_state_shape, "prev"); |
| 1293 | // TupleElement 0 |
| 1294 | auto iteration = GetTupleElement(prev, 0); |
| 1295 | auto out0 = Add(iteration, ConstantR0<int32>(&builder, 1)); |
| 1296 | // TupleElement 1 |
| 1297 | auto input = GetTupleElement(prev, 1); |
| 1298 | // Update. |
| 1299 | auto one = ConstantR0<float>(&builder, 1.0); |
| 1300 | auto update = Broadcast(one, {1, 1024, 1024}); |
| 1301 | // Starts = iteration * 2; |
| 1302 | auto zero = ConstantR0<int32>(&builder, 0); |
| 1303 | // UpdateSlice. |
| 1304 | auto out1 = DynamicUpdateSlice(input, update, {zero, zero, zero}); |
| 1305 | Tuple(&builder, {out0, out1}); |
| 1306 | body = builder.Build().ConsumeValueOrDie(); |
| 1307 | } |
| 1308 | |
| 1309 | // Create a While instruction. |
| 1310 | XlaBuilder builder("while"); |
| 1311 | auto zero = ConstantR0<float>(&builder, 0.0); |
| 1312 | auto input = Broadcast(zero, {seq_len, 1024, 1024}); |
| 1313 | auto init = Tuple(&builder, {ConstantR0<int32>(&builder, 0), input}); |
| 1314 | While(condition, body, init); |
| 1315 | auto computation = builder.Build().ConsumeValueOrDie(); |
| 1316 | |
| 1317 | TF_ASSERT_OK_AND_ASSIGN( |
| 1318 | auto executables, |
| 1319 | client->Compile(computation, {}, ExecutableBuildOptions())); |
nothing calls this directly
no test coverage detected