4298 case Intrinsic::nvvm_match_all_sync_i32p:
4299 case Intrinsic::nvvm_match_all_sync_i64p:
4304 Info.memVT = MVT::i1;
4310 case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_col:
4311 case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_row:
4312 case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_col_stride:
4313 case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_row_stride:
4314 case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_col:
4315 case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_row:
4316 case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_col_stride:
4317 case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_row_stride:
4318 case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_col:
4319 case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_row:
4320 case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_col_stride:
4321 case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_row_stride:
4322 case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_col:
4323 case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_row:
4324 case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_col_stride:
4325 case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_row_stride:
4326 case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_col:
4327 case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_row:
4328 case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_col_stride:
4329 case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_row_stride:
4330 case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_col:
4331 case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_row:
4332 case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_col_stride:
4333 case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_row_stride: {
4335 Info.memVT = MVT::v8f16;
4336 Info.ptrVal =
I.getArgOperand(0);
4339 Info.align =
Align(16);
4343 case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_col:
4344 case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_col_stride:
4345 case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_col_stride:
4346 case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_col:
4347 case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_row:
4348 case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_row_stride:
4349 case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_row_stride:
4350 case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_row:
4351 case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_col:
4352 case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_col_stride:
4353 case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_row:
4354 case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_row_stride:
4355 case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_col:
4356 case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_col_stride:
4357 case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_col_stride:
4358 case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_col:
4359 case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_row:
4360 case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_row_stride:
4361 case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_row_stride:
4362 case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_row:
4363 case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_col:
4364 case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_col_stride:
4365 case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_row:
4366 case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_row_stride: {
4368 Info.memVT = MVT::v2i32;
4369 Info.ptrVal =
I.getArgOperand(0);
4372 Info.align =
Align(8);
4377 case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_col:
4378 case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_col_stride:
4379 case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_col_stride:
4380 case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_col:
4381 case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_row:
4382 case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_row_stride:
4383 case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_row_stride:
4384 case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_row:
4385 case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_col:
4386 case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_col_stride:
4387 case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_row:
4388 case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_row_stride:
4389 case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_col:
4390 case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_col_stride:
4391 case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_row:
4392 case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_row_stride:
4394 case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_col:
4395 case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_col_stride:
4396 case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_col_stride:
4397 case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_col:
4398 case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_row:
4399 case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_row_stride:
4400 case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_row_stride:
4401 case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_row:
4402 case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_col:
4403 case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_col_stride:
4404 case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_row:
4405 case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_row_stride:
4406 case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_col:
4407 case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_col_stride:
4408 case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_row:
4409 case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_row_stride:
4410 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x4_b16:
4411 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x4_trans_b16:
4412 case Intrinsic::nvvm_ldmatrix_sync_aligned_m16n16_x2_trans_b8:
4413 case Intrinsic::nvvm_ldmatrix_sync_aligned_m16n16_x2_trans_b8x16_b4x16_p64:
4414 case Intrinsic::nvvm_ldmatrix_sync_aligned_m16n16_x2_trans_b8x16_b6x16_p32:
4415 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n16_x4_b8x16_b4x16_p64:
4416 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n16_x4_b8x16_b6x16_p32: {
4418 Info.memVT = MVT::v4i32;
4419 Info.ptrVal =
I.getArgOperand(0);
4422 Info.align =
Align(16);
4427 case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_col:
4428 case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_col_stride:
4429 case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_col_stride:
4430 case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_col:
4431 case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_row:
4432 case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_row_stride:
4433 case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_row_stride:
4434 case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_row:
4436 case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_col:
4437 case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_col_stride:
4438 case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_col_stride:
4439 case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_col:
4440 case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_row:
4441 case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_row_stride:
4442 case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_row_stride:
4443 case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_row:
4444 case Intrinsic::nvvm_wmma_m8n8k128_load_a_b1_row:
4445 case Intrinsic::nvvm_wmma_m8n8k128_load_a_b1_row_stride:
4446 case Intrinsic::nvvm_wmma_m8n8k128_load_b_b1_col:
4447 case Intrinsic::nvvm_wmma_m8n8k128_load_b_b1_col_stride:
4448 case Intrinsic::nvvm_wmma_m8n8k32_load_a_s4_row:
4449 case Intrinsic::nvvm_wmma_m8n8k32_load_a_s4_row_stride:
4450 case Intrinsic::nvvm_wmma_m8n8k32_load_a_u4_row_stride:
4451 case Intrinsic::nvvm_wmma_m8n8k32_load_a_u4_row:
4452 case Intrinsic::nvvm_wmma_m8n8k32_load_b_s4_col:
4453 case Intrinsic::nvvm_wmma_m8n8k32_load_b_s4_col_stride:
4454 case Intrinsic::nvvm_wmma_m8n8k32_load_b_u4_col_stride:
4455 case Intrinsic::nvvm_wmma_m8n8k32_load_b_u4_col:
4456 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x1_b16:
4457 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x1_trans_b16:
4458 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n16_x1_b8x16_b4x16_p64:
4459 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n16_x1_b8x16_b6x16_p32: {
4461 Info.memVT = MVT::i32;
4462 Info.ptrVal =
I.getArgOperand(0);
4465 Info.align =
Align(4);
4470 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_col:
4471 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_row:
4472 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_col_stride:
4473 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_row_stride:
4474 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_col:
4475 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_row:
4476 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_col_stride:
4477 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_row_stride:
4478 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_col:
4479 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_row:
4480 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_col_stride:
4481 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_row_stride: {
4483 Info.memVT = MVT::v4f16;
4484 Info.ptrVal =
I.getArgOperand(0);
4487 Info.align =
Align(16);
4492 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_col:
4493 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_row:
4494 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_col_stride:
4495 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_row_stride:
4496 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_col:
4497 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_row:
4498 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_col_stride:
4499 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_row_stride:
4500 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_col:
4501 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_row:
4502 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_col_stride:
4503 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_row_stride:
4504 case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_col:
4505 case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_row:
4506 case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_col_stride:
4507 case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_row_stride: {
4509 Info.memVT = MVT::v8f32;
4510 Info.ptrVal =
I.getArgOperand(0);
4513 Info.align =
Align(16);
4518 case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_col:
4519 case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_col_stride:
4520 case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_row:
4521 case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_row_stride:
4523 case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_col:
4524 case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_col_stride:
4525 case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_row:
4526 case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_row_stride:
4528 case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_col:
4529 case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_col_stride:
4530 case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_row:
4531 case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_row_stride:
4532 case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_col:
4533 case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_col_stride:
4534 case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_row:
4535 case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_row_stride:
4536 case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_col:
4537 case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_col_stride:
4538 case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_row:
4539 case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_row_stride: {
4541 Info.memVT = MVT::v8i32;
4542 Info.ptrVal =
I.getArgOperand(0);
4545 Info.align =
Align(16);
4550 case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_col:
4551 case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_col_stride:
4552 case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_row:
4553 case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_row_stride:
4554 case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_col:
4555 case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_col_stride:
4556 case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_row:
4557 case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_row_stride:
4558 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x2_b16:
4559 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x2_trans_b16:
4560 case Intrinsic::nvvm_ldmatrix_sync_aligned_m16n16_x1_trans_b8:
4561 case Intrinsic::nvvm_ldmatrix_sync_aligned_m16n16_x1_trans_b8x16_b4x16_p64:
4562 case Intrinsic::nvvm_ldmatrix_sync_aligned_m16n16_x1_trans_b8x16_b6x16_p32:
4563 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n16_x2_b8x16_b4x16_p64:
4564 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n16_x2_b8x16_b6x16_p32: {
4566 Info.memVT = MVT::v2i32;
4567 Info.ptrVal =
I.getArgOperand(0);
4570 Info.align =
Align(8);
4575 case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_col:
4576 case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_col_stride:
4577 case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_row:
4578 case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_row_stride:
4580 case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_col:
4581 case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_col_stride:
4582 case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_row:
4583 case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_row_stride: {
4585 Info.memVT = MVT::f64;
4586 Info.ptrVal =
I.getArgOperand(0);
4589 Info.align =
Align(8);
4594 case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_col:
4595 case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_col_stride:
4596 case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_row:
4597 case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_row_stride: {
4599 Info.memVT = MVT::v2f64;
4600 Info.ptrVal =
I.getArgOperand(0);
4603 Info.align =
Align(16);
4608 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_col:
4609 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_row:
4610 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_col_stride:
4611 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_row_stride:
4612 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_col:
4613 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_row:
4614 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_col_stride:
4615 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_row_stride:
4616 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_col:
4617 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_row:
4618 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_col_stride:
4619 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_row_stride: {
4621 Info.memVT = MVT::v4f16;
4622 Info.ptrVal =
I.getArgOperand(0);
4625 Info.align =
Align(16);
4630 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_col:
4631 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_row:
4632 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_col_stride:
4633 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_row_stride:
4634 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_col:
4635 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_row:
4636 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_col_stride:
4637 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_row_stride:
4638 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_col:
4639 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_row:
4640 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_col_stride:
4641 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_row_stride:
4642 case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_col:
4643 case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_row:
4644 case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_col_stride:
4645 case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_row_stride: {
4647 Info.memVT = MVT::v8f32;
4648 Info.ptrVal =
I.getArgOperand(0);
4651 Info.align =
Align(16);
4656 case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_col:
4657 case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_col_stride:
4658 case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_row:
4659 case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_row_stride:
4660 case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_col:
4661 case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_col_stride:
4662 case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_row:
4663 case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_row_stride:
4664 case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_col:
4665 case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_col_stride:
4666 case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_row:
4667 case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_row_stride: {
4669 Info.memVT = MVT::v8i32;
4670 Info.ptrVal =
I.getArgOperand(0);
4673 Info.align =
Align(16);
4678 case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_col:
4679 case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_col_stride:
4680 case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_row:
4681 case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_row_stride:
4682 case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_col:
4683 case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_col_stride:
4684 case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_row:
4685 case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_row_stride:
4686 case Intrinsic::nvvm_stmatrix_sync_aligned_m8n8_x2_b16:
4687 case Intrinsic::nvvm_stmatrix_sync_aligned_m8n8_x2_trans_b16:
4688 case Intrinsic::nvvm_stmatrix_sync_aligned_m16n8_x2_trans_b8: {
4690 Info.memVT = MVT::v2i32;
4691 Info.ptrVal =
I.getArgOperand(0);
4694 Info.align =
Align(8);
4699 case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_col:
4700 case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_col_stride:
4701 case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_row:
4702 case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_row_stride: {
4704 Info.memVT = MVT::v2f64;
4705 Info.ptrVal =
I.getArgOperand(0);
4708 Info.align =
Align(16);
4713 case Intrinsic::nvvm_stmatrix_sync_aligned_m8n8_x1_b16:
4714 case Intrinsic::nvvm_stmatrix_sync_aligned_m8n8_x1_trans_b16:
4715 case Intrinsic::nvvm_stmatrix_sync_aligned_m16n8_x1_trans_b8: {
4717 Info.memVT = MVT::i32;
4718 Info.ptrVal =
I.getArgOperand(0);
4721 Info.align =
Align(4);
4726 case Intrinsic::nvvm_stmatrix_sync_aligned_m8n8_x4_b16:
4727 case Intrinsic::nvvm_stmatrix_sync_aligned_m8n8_x4_trans_b16:
4728 case Intrinsic::nvvm_stmatrix_sync_aligned_m16n8_x4_trans_b8: {
4730 Info.memVT = MVT::v4i32;
4731 Info.ptrVal =
I.getArgOperand(0);
4734 Info.align =
Align(16);
4739 case Intrinsic::nvvm_prefetch_tensormap: {
4740 auto &
DL =
I.getDataLayout();
4743 Info.ptrVal =
I.getArgOperand(0);
4752 case Intrinsic::nvvm_tensormap_replace_global_address:
4753 case Intrinsic::nvvm_tensormap_replace_global_stride: {
4755 Info.memVT = MVT::i64;
4756 Info.ptrVal =
I.getArgOperand(0);
4764 case Intrinsic::nvvm_tensormap_replace_rank:
4765 case Intrinsic::nvvm_tensormap_replace_box_dim:
4766 case Intrinsic::nvvm_tensormap_replace_global_dim:
4767 case Intrinsic::nvvm_tensormap_replace_element_stride:
4768 case Intrinsic::nvvm_tensormap_replace_elemtype:
4769 case Intrinsic::nvvm_tensormap_replace_interleave_layout:
4770 case Intrinsic::nvvm_tensormap_replace_swizzle_mode:
4771 case Intrinsic::nvvm_tensormap_replace_swizzle_atomicity:
4772 case Intrinsic::nvvm_tensormap_replace_fill_mode: {
4774 Info.memVT = MVT::i32;
4775 Info.ptrVal =
I.getArgOperand(0);
4783 case Intrinsic::nvvm_ldu_global_i:
4784 case Intrinsic::nvvm_ldu_global_f:
4785 case Intrinsic::nvvm_ldu_global_p: {
4788 Info.ptrVal =
I.getArgOperand(0);
4796 case Intrinsic::nvvm_tex_1d_v4f32_s32:
4797 case Intrinsic::nvvm_tex_1d_v4f32_f32:
4798 case Intrinsic::nvvm_tex_1d_level_v4f32_f32:
4799 case Intrinsic::nvvm_tex_1d_grad_v4f32_f32:
4800 case Intrinsic::nvvm_tex_1d_array_v4f32_s32:
4801 case Intrinsic::nvvm_tex_1d_array_v4f32_f32:
4802 case Intrinsic::nvvm_tex_1d_array_level_v4f32_f32:
4803 case Intrinsic::nvvm_tex_1d_array_grad_v4f32_f32:
4804 case Intrinsic::nvvm_tex_2d_v4f32_s32:
4805 case Intrinsic::nvvm_tex_2d_v4f32_f32:
4806 case Intrinsic::nvvm_tex_2d_level_v4f32_f32:
4807 case Intrinsic::nvvm_tex_2d_grad_v4f32_f32:
4808 case Intrinsic::nvvm_tex_2d_array_v4f32_s32:
4809 case Intrinsic::nvvm_tex_2d_array_v4f32_f32:
4810 case Intrinsic::nvvm_tex_2d_array_level_v4f32_f32:
4811 case Intrinsic::nvvm_tex_2d_array_grad_v4f32_f32:
4812 case Intrinsic::nvvm_tex_3d_v4f32_s32:
4813 case Intrinsic::nvvm_tex_3d_v4f32_f32:
4814 case Intrinsic::nvvm_tex_3d_level_v4f32_f32:
4815 case Intrinsic::nvvm_tex_3d_grad_v4f32_f32:
4816 case Intrinsic::nvvm_tex_cube_v4f32_f32:
4817 case Intrinsic::nvvm_tex_cube_level_v4f32_f32:
4818 case Intrinsic::nvvm_tex_cube_array_v4f32_f32:
4819 case Intrinsic::nvvm_tex_cube_array_level_v4f32_f32:
4820 case Intrinsic::nvvm_tld4_r_2d_v4f32_f32:
4821 case Intrinsic::nvvm_tld4_g_2d_v4f32_f32:
4822 case Intrinsic::nvvm_tld4_b_2d_v4f32_f32:
4823 case Intrinsic::nvvm_tld4_a_2d_v4f32_f32:
4824 case Intrinsic::nvvm_tex_unified_1d_v4f32_s32:
4825 case Intrinsic::nvvm_tex_unified_1d_v4f32_f32:
4826 case Intrinsic::nvvm_tex_unified_1d_level_v4f32_f32:
4827 case Intrinsic::nvvm_tex_unified_1d_grad_v4f32_f32:
4828 case Intrinsic::nvvm_tex_unified_1d_array_v4f32_s32:
4829 case Intrinsic::nvvm_tex_unified_1d_array_v4f32_f32:
4830 case Intrinsic::nvvm_tex_unified_1d_array_level_v4f32_f32:
4831 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4f32_f32:
4832 case Intrinsic::nvvm_tex_unified_2d_v4f32_s32:
4833 case Intrinsic::nvvm_tex_unified_2d_v4f32_f32:
4834 case Intrinsic::nvvm_tex_unified_2d_level_v4f32_f32:
4835 case Intrinsic::nvvm_tex_unified_2d_grad_v4f32_f32:
4836 case Intrinsic::nvvm_tex_unified_2d_array_v4f32_s32:
4837 case Intrinsic::nvvm_tex_unified_2d_array_v4f32_f32:
4838 case Intrinsic::nvvm_tex_unified_2d_array_level_v4f32_f32:
4839 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4f32_f32:
4840 case Intrinsic::nvvm_tex_unified_3d_v4f32_s32:
4841 case Intrinsic::nvvm_tex_unified_3d_v4f32_f32:
4842 case Intrinsic::nvvm_tex_unified_3d_level_v4f32_f32:
4843 case Intrinsic::nvvm_tex_unified_3d_grad_v4f32_f32:
4844 case Intrinsic::nvvm_tex_unified_cube_v4f32_f32:
4845 case Intrinsic::nvvm_tex_unified_cube_level_v4f32_f32:
4846 case Intrinsic::nvvm_tex_unified_cube_array_v4f32_f32:
4847 case Intrinsic::nvvm_tex_unified_cube_array_level_v4f32_f32:
4848 case Intrinsic::nvvm_tex_unified_cube_grad_v4f32_f32:
4849 case Intrinsic::nvvm_tex_unified_cube_array_grad_v4f32_f32:
4850 case Intrinsic::nvvm_tld4_unified_r_2d_v4f32_f32:
4851 case Intrinsic::nvvm_tld4_unified_g_2d_v4f32_f32:
4852 case Intrinsic::nvvm_tld4_unified_b_2d_v4f32_f32:
4853 case Intrinsic::nvvm_tld4_unified_a_2d_v4f32_f32:
4855 Info.memVT = MVT::v4f32;
4856 Info.ptrVal =
nullptr;
4859 Info.align =
Align(16);
4863 case Intrinsic::nvvm_tex_1d_v4s32_s32:
4864 case Intrinsic::nvvm_tex_1d_v4s32_f32:
4865 case Intrinsic::nvvm_tex_1d_level_v4s32_f32:
4866 case Intrinsic::nvvm_tex_1d_grad_v4s32_f32:
4867 case Intrinsic::nvvm_tex_1d_array_v4s32_s32:
4868 case Intrinsic::nvvm_tex_1d_array_v4s32_f32:
4869 case Intrinsic::nvvm_tex_1d_array_level_v4s32_f32:
4870 case Intrinsic::nvvm_tex_1d_array_grad_v4s32_f32:
4871 case Intrinsic::nvvm_tex_2d_v4s32_s32:
4872 case Intrinsic::nvvm_tex_2d_v4s32_f32:
4873 case Intrinsic::nvvm_tex_2d_level_v4s32_f32:
4874 case Intrinsic::nvvm_tex_2d_grad_v4s32_f32:
4875 case Intrinsic::nvvm_tex_2d_array_v4s32_s32:
4876 case Intrinsic::nvvm_tex_2d_array_v4s32_f32:
4877 case Intrinsic::nvvm_tex_2d_array_level_v4s32_f32:
4878 case Intrinsic::nvvm_tex_2d_array_grad_v4s32_f32:
4879 case Intrinsic::nvvm_tex_3d_v4s32_s32:
4880 case Intrinsic::nvvm_tex_3d_v4s32_f32:
4881 case Intrinsic::nvvm_tex_3d_level_v4s32_f32:
4882 case Intrinsic::nvvm_tex_3d_grad_v4s32_f32:
4883 case Intrinsic::nvvm_tex_cube_v4s32_f32:
4884 case Intrinsic::nvvm_tex_cube_level_v4s32_f32:
4885 case Intrinsic::nvvm_tex_cube_array_v4s32_f32:
4886 case Intrinsic::nvvm_tex_cube_array_level_v4s32_f32:
4887 case Intrinsic::nvvm_tex_cube_v4u32_f32:
4888 case Intrinsic::nvvm_tex_cube_level_v4u32_f32:
4889 case Intrinsic::nvvm_tex_cube_array_v4u32_f32:
4890 case Intrinsic::nvvm_tex_cube_array_level_v4u32_f32:
4891 case Intrinsic::nvvm_tex_1d_v4u32_s32:
4892 case Intrinsic::nvvm_tex_1d_v4u32_f32:
4893 case Intrinsic::nvvm_tex_1d_level_v4u32_f32:
4894 case Intrinsic::nvvm_tex_1d_grad_v4u32_f32:
4895 case Intrinsic::nvvm_tex_1d_array_v4u32_s32:
4896 case Intrinsic::nvvm_tex_1d_array_v4u32_f32:
4897 case Intrinsic::nvvm_tex_1d_array_level_v4u32_f32:
4898 case Intrinsic::nvvm_tex_1d_array_grad_v4u32_f32:
4899 case Intrinsic::nvvm_tex_2d_v4u32_s32:
4900 case Intrinsic::nvvm_tex_2d_v4u32_f32:
4901 case Intrinsic::nvvm_tex_2d_level_v4u32_f32:
4902 case Intrinsic::nvvm_tex_2d_grad_v4u32_f32:
4903 case Intrinsic::nvvm_tex_2d_array_v4u32_s32:
4904 case Intrinsic::nvvm_tex_2d_array_v4u32_f32:
4905 case Intrinsic::nvvm_tex_2d_array_level_v4u32_f32:
4906 case Intrinsic::nvvm_tex_2d_array_grad_v4u32_f32:
4907 case Intrinsic::nvvm_tex_3d_v4u32_s32:
4908 case Intrinsic::nvvm_tex_3d_v4u32_f32:
4909 case Intrinsic::nvvm_tex_3d_level_v4u32_f32:
4910 case Intrinsic::nvvm_tex_3d_grad_v4u32_f32:
4911 case Intrinsic::nvvm_tld4_r_2d_v4s32_f32:
4912 case Intrinsic::nvvm_tld4_g_2d_v4s32_f32:
4913 case Intrinsic::nvvm_tld4_b_2d_v4s32_f32:
4914 case Intrinsic::nvvm_tld4_a_2d_v4s32_f32:
4915 case Intrinsic::nvvm_tld4_r_2d_v4u32_f32:
4916 case Intrinsic::nvvm_tld4_g_2d_v4u32_f32:
4917 case Intrinsic::nvvm_tld4_b_2d_v4u32_f32:
4918 case Intrinsic::nvvm_tld4_a_2d_v4u32_f32:
4919 case Intrinsic::nvvm_tex_unified_1d_v4s32_s32:
4920 case Intrinsic::nvvm_tex_unified_1d_v4s32_f32:
4921 case Intrinsic::nvvm_tex_unified_1d_level_v4s32_f32:
4922 case Intrinsic::nvvm_tex_unified_1d_grad_v4s32_f32:
4923 case Intrinsic::nvvm_tex_unified_1d_array_v4s32_s32:
4924 case Intrinsic::nvvm_tex_unified_1d_array_v4s32_f32:
4925 case Intrinsic::nvvm_tex_unified_1d_array_level_v4s32_f32:
4926 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4s32_f32:
4927 case Intrinsic::nvvm_tex_unified_2d_v4s32_s32:
4928 case Intrinsic::nvvm_tex_unified_2d_v4s32_f32:
4929 case Intrinsic::nvvm_tex_unified_2d_level_v4s32_f32:
4930 case Intrinsic::nvvm_tex_unified_2d_grad_v4s32_f32:
4931 case Intrinsic::nvvm_tex_unified_2d_array_v4s32_s32:
4932 case Intrinsic::nvvm_tex_unified_2d_array_v4s32_f32:
4933 case Intrinsic::nvvm_tex_unified_2d_array_level_v4s32_f32:
4934 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4s32_f32:
4935 case Intrinsic::nvvm_tex_unified_3d_v4s32_s32:
4936 case Intrinsic::nvvm_tex_unified_3d_v4s32_f32:
4937 case Intrinsic::nvvm_tex_unified_3d_level_v4s32_f32:
4938 case Intrinsic::nvvm_tex_unified_3d_grad_v4s32_f32:
4939 case Intrinsic::nvvm_tex_unified_1d_v4u32_s32:
4940 case Intrinsic::nvvm_tex_unified_1d_v4u32_f32:
4941 case Intrinsic::nvvm_tex_unified_1d_level_v4u32_f32:
4942 case Intrinsic::nvvm_tex_unified_1d_grad_v4u32_f32:
4943 case Intrinsic::nvvm_tex_unified_1d_array_v4u32_s32:
4944 case Intrinsic::nvvm_tex_unified_1d_array_v4u32_f32:
4945 case Intrinsic::nvvm_tex_unified_1d_array_level_v4u32_f32:
4946 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4u32_f32:
4947 case Intrinsic::nvvm_tex_unified_2d_v4u32_s32:
4948 case Intrinsic::nvvm_tex_unified_2d_v4u32_f32:
4949 case Intrinsic::nvvm_tex_unified_2d_level_v4u32_f32:
4950 case Intrinsic::nvvm_tex_unified_2d_grad_v4u32_f32:
4951 case Intrinsic::nvvm_tex_unified_2d_array_v4u32_s32:
4952 case Intrinsic::nvvm_tex_unified_2d_array_v4u32_f32:
4953 case Intrinsic::nvvm_tex_unified_2d_array_level_v4u32_f32:
4954 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4u32_f32:
4955 case Intrinsic::nvvm_tex_unified_3d_v4u32_s32:
4956 case Intrinsic::nvvm_tex_unified_3d_v4u32_f32:
4957 case Intrinsic::nvvm_tex_unified_3d_level_v4u32_f32:
4958 case Intrinsic::nvvm_tex_unified_3d_grad_v4u32_f32:
4959 case Intrinsic::nvvm_tex_unified_cube_v4s32_f32:
4960 case Intrinsic::nvvm_tex_unified_cube_level_v4s32_f32:
4961 case Intrinsic::nvvm_tex_unified_cube_array_v4s32_f32:
4962 case Intrinsic::nvvm_tex_unified_cube_array_level_v4s32_f32:
4963 case Intrinsic::nvvm_tex_unified_cube_v4u32_f32:
4964 case Intrinsic::nvvm_tex_unified_cube_level_v4u32_f32:
4965 case Intrinsic::nvvm_tex_unified_cube_array_v4u32_f32:
4966 case Intrinsic::nvvm_tex_unified_cube_array_level_v4u32_f32:
4967 case Intrinsic::nvvm_tex_unified_cube_grad_v4s32_f32:
4968 case Intrinsic::nvvm_tex_unified_cube_grad_v4u32_f32:
4969 case Intrinsic::nvvm_tex_unified_cube_array_grad_v4s32_f32:
4970 case Intrinsic::nvvm_tex_unified_cube_array_grad_v4u32_f32:
4971 case Intrinsic::nvvm_tld4_unified_r_2d_v4s32_f32:
4972 case Intrinsic::nvvm_tld4_unified_g_2d_v4s32_f32:
4973 case Intrinsic::nvvm_tld4_unified_b_2d_v4s32_f32:
4974 case Intrinsic::nvvm_tld4_unified_a_2d_v4s32_f32:
4975 case Intrinsic::nvvm_tld4_unified_r_2d_v4u32_f32:
4976 case Intrinsic::nvvm_tld4_unified_g_2d_v4u32_f32:
4977 case Intrinsic::nvvm_tld4_unified_b_2d_v4u32_f32:
4978 case Intrinsic::nvvm_tld4_unified_a_2d_v4u32_f32:
4980 Info.memVT = MVT::v4i32;
4981 Info.ptrVal =
nullptr;
4984 Info.align =
Align(16);
4988 case Intrinsic::nvvm_suld_1d_i8_clamp:
4989 case Intrinsic::nvvm_suld_1d_v2i8_clamp:
4990 case Intrinsic::nvvm_suld_1d_v4i8_clamp:
4991 case Intrinsic::nvvm_suld_1d_array_i8_clamp:
4992 case Intrinsic::nvvm_suld_1d_array_v2i8_clamp:
4993 case Intrinsic::nvvm_suld_1d_array_v4i8_clamp:
4994 case Intrinsic::nvvm_suld_2d_i8_clamp:
4995 case Intrinsic::nvvm_suld_2d_v2i8_clamp:
4996 case Intrinsic::nvvm_suld_2d_v4i8_clamp:
4997 case Intrinsic::nvvm_suld_2d_array_i8_clamp:
4998 case Intrinsic::nvvm_suld_2d_array_v2i8_clamp:
4999 case Intrinsic::nvvm_suld_2d_array_v4i8_clamp:
5000 case Intrinsic::nvvm_suld_3d_i8_clamp:
5001 case Intrinsic::nvvm_suld_3d_v2i8_clamp:
5002 case Intrinsic::nvvm_suld_3d_v4i8_clamp:
5003 case Intrinsic::nvvm_suld_1d_i8_trap:
5004 case Intrinsic::nvvm_suld_1d_v2i8_trap:
5005 case Intrinsic::nvvm_suld_1d_v4i8_trap:
5006 case Intrinsic::nvvm_suld_1d_array_i8_trap:
5007 case Intrinsic::nvvm_suld_1d_array_v2i8_trap:
5008 case Intrinsic::nvvm_suld_1d_array_v4i8_trap:
5009 case Intrinsic::nvvm_suld_2d_i8_trap:
5010 case Intrinsic::nvvm_suld_2d_v2i8_trap:
5011 case Intrinsic::nvvm_suld_2d_v4i8_trap:
5012 case Intrinsic::nvvm_suld_2d_array_i8_trap:
5013 case Intrinsic::nvvm_suld_2d_array_v2i8_trap:
5014 case Intrinsic::nvvm_suld_2d_array_v4i8_trap:
5015 case Intrinsic::nvvm_suld_3d_i8_trap:
5016 case Intrinsic::nvvm_suld_3d_v2i8_trap:
5017 case Intrinsic::nvvm_suld_3d_v4i8_trap:
5018 case Intrinsic::nvvm_suld_1d_i8_zero:
5019 case Intrinsic::nvvm_suld_1d_v2i8_zero:
5020 case Intrinsic::nvvm_suld_1d_v4i8_zero:
5021 case Intrinsic::nvvm_suld_1d_array_i8_zero:
5022 case Intrinsic::nvvm_suld_1d_array_v2i8_zero:
5023 case Intrinsic::nvvm_suld_1d_array_v4i8_zero:
5024 case Intrinsic::nvvm_suld_2d_i8_zero:
5025 case Intrinsic::nvvm_suld_2d_v2i8_zero:
5026 case Intrinsic::nvvm_suld_2d_v4i8_zero:
5027 case Intrinsic::nvvm_suld_2d_array_i8_zero:
5028 case Intrinsic::nvvm_suld_2d_array_v2i8_zero:
5029 case Intrinsic::nvvm_suld_2d_array_v4i8_zero:
5030 case Intrinsic::nvvm_suld_3d_i8_zero:
5031 case Intrinsic::nvvm_suld_3d_v2i8_zero:
5032 case Intrinsic::nvvm_suld_3d_v4i8_zero:
5034 Info.memVT = MVT::i8;
5035 Info.ptrVal =
nullptr;
5038 Info.align =
Align(16);
5042 case Intrinsic::nvvm_suld_1d_i16_clamp:
5043 case Intrinsic::nvvm_suld_1d_v2i16_clamp:
5044 case Intrinsic::nvvm_suld_1d_v4i16_clamp:
5045 case Intrinsic::nvvm_suld_1d_array_i16_clamp:
5046 case Intrinsic::nvvm_suld_1d_array_v2i16_clamp:
5047 case Intrinsic::nvvm_suld_1d_array_v4i16_clamp:
5048 case Intrinsic::nvvm_suld_2d_i16_clamp:
5049 case Intrinsic::nvvm_suld_2d_v2i16_clamp:
5050 case Intrinsic::nvvm_suld_2d_v4i16_clamp:
5051 case Intrinsic::nvvm_suld_2d_array_i16_clamp:
5052 case Intrinsic::nvvm_suld_2d_array_v2i16_clamp:
5053 case Intrinsic::nvvm_suld_2d_array_v4i16_clamp:
5054 case Intrinsic::nvvm_suld_3d_i16_clamp:
5055 case Intrinsic::nvvm_suld_3d_v2i16_clamp:
5056 case Intrinsic::nvvm_suld_3d_v4i16_clamp:
5057 case Intrinsic::nvvm_suld_1d_i16_trap:
5058 case Intrinsic::nvvm_suld_1d_v2i16_trap:
5059 case Intrinsic::nvvm_suld_1d_v4i16_trap:
5060 case Intrinsic::nvvm_suld_1d_array_i16_trap:
5061 case Intrinsic::nvvm_suld_1d_array_v2i16_trap:
5062 case Intrinsic::nvvm_suld_1d_array_v4i16_trap:
5063 case Intrinsic::nvvm_suld_2d_i16_trap:
5064 case Intrinsic::nvvm_suld_2d_v2i16_trap:
5065 case Intrinsic::nvvm_suld_2d_v4i16_trap:
5066 case Intrinsic::nvvm_suld_2d_array_i16_trap:
5067 case Intrinsic::nvvm_suld_2d_array_v2i16_trap:
5068 case Intrinsic::nvvm_suld_2d_array_v4i16_trap:
5069 case Intrinsic::nvvm_suld_3d_i16_trap:
5070 case Intrinsic::nvvm_suld_3d_v2i16_trap:
5071 case Intrinsic::nvvm_suld_3d_v4i16_trap:
5072 case Intrinsic::nvvm_suld_1d_i16_zero:
5073 case Intrinsic::nvvm_suld_1d_v2i16_zero:
5074 case Intrinsic::nvvm_suld_1d_v4i16_zero:
5075 case Intrinsic::nvvm_suld_1d_array_i16_zero:
5076 case Intrinsic::nvvm_suld_1d_array_v2i16_zero:
5077 case Intrinsic::nvvm_suld_1d_array_v4i16_zero:
5078 case Intrinsic::nvvm_suld_2d_i16_zero:
5079 case Intrinsic::nvvm_suld_2d_v2i16_zero:
5080 case Intrinsic::nvvm_suld_2d_v4i16_zero:
5081 case Intrinsic::nvvm_suld_2d_array_i16_zero:
5082 case Intrinsic::nvvm_suld_2d_array_v2i16_zero:
5083 case Intrinsic::nvvm_suld_2d_array_v4i16_zero:
5084 case Intrinsic::nvvm_suld_3d_i16_zero:
5085 case Intrinsic::nvvm_suld_3d_v2i16_zero:
5086 case Intrinsic::nvvm_suld_3d_v4i16_zero:
5088 Info.memVT = MVT::i16;
5089 Info.ptrVal =
nullptr;
5092 Info.align =
Align(16);
5096 case Intrinsic::nvvm_suld_1d_i32_clamp:
5097 case Intrinsic::nvvm_suld_1d_v2i32_clamp:
5098 case Intrinsic::nvvm_suld_1d_v4i32_clamp:
5099 case Intrinsic::nvvm_suld_1d_array_i32_clamp:
5100 case Intrinsic::nvvm_suld_1d_array_v2i32_clamp:
5101 case Intrinsic::nvvm_suld_1d_array_v4i32_clamp:
5102 case Intrinsic::nvvm_suld_2d_i32_clamp:
5103 case Intrinsic::nvvm_suld_2d_v2i32_clamp:
5104 case Intrinsic::nvvm_suld_2d_v4i32_clamp:
5105 case Intrinsic::nvvm_suld_2d_array_i32_clamp:
5106 case Intrinsic::nvvm_suld_2d_array_v2i32_clamp:
5107 case Intrinsic::nvvm_suld_2d_array_v4i32_clamp:
5108 case Intrinsic::nvvm_suld_3d_i32_clamp:
5109 case Intrinsic::nvvm_suld_3d_v2i32_clamp:
5110 case Intrinsic::nvvm_suld_3d_v4i32_clamp:
5111 case Intrinsic::nvvm_suld_1d_i32_trap:
5112 case Intrinsic::nvvm_suld_1d_v2i32_trap:
5113 case Intrinsic::nvvm_suld_1d_v4i32_trap:
5114 case Intrinsic::nvvm_suld_1d_array_i32_trap:
5115 case Intrinsic::nvvm_suld_1d_array_v2i32_trap:
5116 case Intrinsic::nvvm_suld_1d_array_v4i32_trap:
5117 case Intrinsic::nvvm_suld_2d_i32_trap:
5118 case Intrinsic::nvvm_suld_2d_v2i32_trap:
5119 case Intrinsic::nvvm_suld_2d_v4i32_trap:
5120 case Intrinsic::nvvm_suld_2d_array_i32_trap:
5121 case Intrinsic::nvvm_suld_2d_array_v2i32_trap:
5122 case Intrinsic::nvvm_suld_2d_array_v4i32_trap:
5123 case Intrinsic::nvvm_suld_3d_i32_trap:
5124 case Intrinsic::nvvm_suld_3d_v2i32_trap:
5125 case Intrinsic::nvvm_suld_3d_v4i32_trap:
5126 case Intrinsic::nvvm_suld_1d_i32_zero:
5127 case Intrinsic::nvvm_suld_1d_v2i32_zero:
5128 case Intrinsic::nvvm_suld_1d_v4i32_zero:
5129 case Intrinsic::nvvm_suld_1d_array_i32_zero:
5130 case Intrinsic::nvvm_suld_1d_array_v2i32_zero:
5131 case Intrinsic::nvvm_suld_1d_array_v4i32_zero:
5132 case Intrinsic::nvvm_suld_2d_i32_zero:
5133 case Intrinsic::nvvm_suld_2d_v2i32_zero:
5134 case Intrinsic::nvvm_suld_2d_v4i32_zero:
5135 case Intrinsic::nvvm_suld_2d_array_i32_zero:
5136 case Intrinsic::nvvm_suld_2d_array_v2i32_zero:
5137 case Intrinsic::nvvm_suld_2d_array_v4i32_zero:
5138 case Intrinsic::nvvm_suld_3d_i32_zero:
5139 case Intrinsic::nvvm_suld_3d_v2i32_zero:
5140 case Intrinsic::nvvm_suld_3d_v4i32_zero:
5142 Info.memVT = MVT::i32;
5143 Info.ptrVal =
nullptr;
5146 Info.align =
Align(16);
5150 case Intrinsic::nvvm_suld_1d_i64_clamp:
5151 case Intrinsic::nvvm_suld_1d_v2i64_clamp:
5152 case Intrinsic::nvvm_suld_1d_array_i64_clamp:
5153 case Intrinsic::nvvm_suld_1d_array_v2i64_clamp:
5154 case Intrinsic::nvvm_suld_2d_i64_clamp:
5155 case Intrinsic::nvvm_suld_2d_v2i64_clamp:
5156 case Intrinsic::nvvm_suld_2d_array_i64_clamp:
5157 case Intrinsic::nvvm_suld_2d_array_v2i64_clamp:
5158 case Intrinsic::nvvm_suld_3d_i64_clamp:
5159 case Intrinsic::nvvm_suld_3d_v2i64_clamp:
5160 case Intrinsic::nvvm_suld_1d_i64_trap:
5161 case Intrinsic::nvvm_suld_1d_v2i64_trap:
5162 case Intrinsic::nvvm_suld_1d_array_i64_trap:
5163 case Intrinsic::nvvm_suld_1d_array_v2i64_trap:
5164 case Intrinsic::nvvm_suld_2d_i64_trap:
5165 case Intrinsic::nvvm_suld_2d_v2i64_trap:
5166 case Intrinsic::nvvm_suld_2d_array_i64_trap:
5167 case Intrinsic::nvvm_suld_2d_array_v2i64_trap:
5168 case Intrinsic::nvvm_suld_3d_i64_trap:
5169 case Intrinsic::nvvm_suld_3d_v2i64_trap:
5170 case Intrinsic::nvvm_suld_1d_i64_zero:
5171 case Intrinsic::nvvm_suld_1d_v2i64_zero:
5172 case Intrinsic::nvvm_suld_1d_array_i64_zero:
5173 case Intrinsic::nvvm_suld_1d_array_v2i64_zero:
5174 case Intrinsic::nvvm_suld_2d_i64_zero:
5175 case Intrinsic::nvvm_suld_2d_v2i64_zero:
5176 case Intrinsic::nvvm_suld_2d_array_i64_zero:
5177 case Intrinsic::nvvm_suld_2d_array_v2i64_zero:
5178 case Intrinsic::nvvm_suld_3d_i64_zero:
5179 case Intrinsic::nvvm_suld_3d_v2i64_zero:
5181 Info.memVT = MVT::i64;
5182 Info.ptrVal =
nullptr;
5185 Info.align =
Align(16);
5189 case Intrinsic::nvvm_tcgen05_ld_16x64b_x1:
5190 case Intrinsic::nvvm_tcgen05_ld_32x32b_x1:
5191 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x1: {
5193 Info.memVT = MVT::v1i32;
5194 Info.ptrVal =
I.getArgOperand(0);
5202 case Intrinsic::nvvm_tcgen05_ld_16x64b_x2:
5203 case Intrinsic::nvvm_tcgen05_ld_16x128b_x1:
5204 case Intrinsic::nvvm_tcgen05_ld_32x32b_x2:
5205 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x2:
5206 case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x2_i32:
5207 case Intrinsic::nvvm_tcgen05_ld_red_16x32bx2_x2_i32: {
5209 Info.memVT = MVT::v2i32;
5210 Info.ptrVal =
I.getArgOperand(0);
5218 case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x2_f32:
5219 case Intrinsic::nvvm_tcgen05_ld_red_16x32bx2_x2_f32: {
5221 Info.memVT = MVT::v2f32;
5222 Info.ptrVal =
I.getArgOperand(0);
5230 case Intrinsic::nvvm_tcgen05_ld_16x64b_x4:
5231 case Intrinsic::nvvm_tcgen05_ld_16x128b_x2:
5232 case Intrinsic::nvvm_tcgen05_ld_32x32b_x4:
5233 case Intrinsic::nvvm_tcgen05_ld_16x256b_x1:
5234 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x4:
5235 case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x4_i32:
5236 case Intrinsic::nvvm_tcgen05_ld_red_16x32bx2_x4_i32: {
5238 Info.memVT = MVT::v4i32;
5239 Info.ptrVal =
I.getArgOperand(0);
5247 case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x4_f32:
5248 case Intrinsic::nvvm_tcgen05_ld_red_16x32bx2_x4_f32: {
5250 Info.memVT = MVT::v4f32;
5251 Info.ptrVal =
I.getArgOperand(0);
5259 case Intrinsic::nvvm_tcgen05_ld_16x64b_x8:
5260 case Intrinsic::nvvm_tcgen05_ld_16x128b_x4:
5261 case Intrinsic::nvvm_tcgen05_ld_16x256b_x2:
5262 case Intrinsic::nvvm_tcgen05_ld_32x32b_x8:
5263 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x8:
5264 case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x8_i32:
5265 case Intrinsic::nvvm_tcgen05_ld_red_16x32bx2_x8_i32: {
5267 Info.memVT = MVT::v8i32;
5268 Info.ptrVal =
I.getArgOperand(0);
5276 case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x8_f32:
5277 case Intrinsic::nvvm_tcgen05_ld_red_16x32bx2_x8_f32: {
5279 Info.memVT = MVT::v8f32;
5280 Info.ptrVal =
I.getArgOperand(0);
5288 case Intrinsic::nvvm_tcgen05_ld_16x64b_x16:
5289 case Intrinsic::nvvm_tcgen05_ld_16x128b_x8:
5290 case Intrinsic::nvvm_tcgen05_ld_16x256b_x4:
5291 case Intrinsic::nvvm_tcgen05_ld_32x32b_x16:
5292 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x16:
5293 case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x16_i32:
5294 case Intrinsic::nvvm_tcgen05_ld_red_16x32bx2_x16_i32: {
5296 Info.memVT = MVT::v16i32;
5297 Info.ptrVal =
I.getArgOperand(0);
5305 case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x16_f32:
5306 case Intrinsic::nvvm_tcgen05_ld_red_16x32bx2_x16_f32: {
5308 Info.memVT = MVT::v16f32;
5309 Info.ptrVal =
I.getArgOperand(0);
5317 case Intrinsic::nvvm_tcgen05_ld_16x64b_x32:
5318 case Intrinsic::nvvm_tcgen05_ld_16x128b_x16:
5319 case Intrinsic::nvvm_tcgen05_ld_16x256b_x8:
5320 case Intrinsic::nvvm_tcgen05_ld_32x32b_x32:
5321 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x32:
5322 case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x32_i32:
5323 case Intrinsic::nvvm_tcgen05_ld_red_16x32bx2_x32_i32: {
5325 Info.memVT = MVT::v32i32;
5326 Info.ptrVal =
I.getArgOperand(0);
5334 case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x32_f32:
5335 case Intrinsic::nvvm_tcgen05_ld_red_16x32bx2_x32_f32: {
5337 Info.memVT = MVT::v32f32;
5338 Info.ptrVal =
I.getArgOperand(0);
5346 case Intrinsic::nvvm_tcgen05_ld_16x64b_x64:
5347 case Intrinsic::nvvm_tcgen05_ld_16x128b_x32:
5348 case Intrinsic::nvvm_tcgen05_ld_16x256b_x16:
5349 case Intrinsic::nvvm_tcgen05_ld_32x32b_x64:
5350 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x64:
5351 case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x64_i32:
5352 case Intrinsic::nvvm_tcgen05_ld_red_16x32bx2_x64_i32: {
5354 Info.memVT = MVT::v64i32;
5355 Info.ptrVal =
I.getArgOperand(0);
5363 case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x64_f32:
5364 case Intrinsic::nvvm_tcgen05_ld_red_16x32bx2_x64_f32: {
5366 Info.memVT = MVT::v64f32;
5367 Info.ptrVal =
I.getArgOperand(0);
5375 case Intrinsic::nvvm_tcgen05_ld_16x64b_x128:
5376 case Intrinsic::nvvm_tcgen05_ld_16x128b_x64:
5377 case Intrinsic::nvvm_tcgen05_ld_16x256b_x32:
5378 case Intrinsic::nvvm_tcgen05_ld_32x32b_x128:
5379 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x128:
5380 case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x128_i32:
5381 case Intrinsic::nvvm_tcgen05_ld_red_16x32bx2_x128_i32: {
5383 Info.memVT = MVT::v128i32;
5384 Info.ptrVal =
I.getArgOperand(0);
5392 case Intrinsic::nvvm_tcgen05_ld_red_32x32b_x128_f32:
5393 case Intrinsic::nvvm_tcgen05_ld_red_16x32bx2_x128_f32: {
5395 Info.memVT = MVT::v128f32;
5396 Info.ptrVal =
I.getArgOperand(0);
5404 case Intrinsic::nvvm_tcgen05_st_16x64b_x1:
5405 case Intrinsic::nvvm_tcgen05_st_32x32b_x1:
5406 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x1: {
5408 Info.memVT = MVT::v1i32;
5409 Info.ptrVal =
I.getArgOperand(0);
5417 case Intrinsic::nvvm_tcgen05_st_16x64b_x2:
5418 case Intrinsic::nvvm_tcgen05_st_16x128b_x1:
5419 case Intrinsic::nvvm_tcgen05_st_32x32b_x2:
5420 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x2: {
5422 Info.memVT = MVT::v2i32;
5423 Info.ptrVal =
I.getArgOperand(0);
5431 case Intrinsic::nvvm_tcgen05_st_16x64b_x4:
5432 case Intrinsic::nvvm_tcgen05_st_16x128b_x2:
5433 case Intrinsic::nvvm_tcgen05_st_16x256b_x1:
5434 case Intrinsic::nvvm_tcgen05_st_32x32b_x4:
5435 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x4: {
5437 Info.memVT = MVT::v4i32;
5438 Info.ptrVal =
I.getArgOperand(0);
5446 case Intrinsic::nvvm_tcgen05_st_16x64b_x8:
5447 case Intrinsic::nvvm_tcgen05_st_16x128b_x4:
5448 case Intrinsic::nvvm_tcgen05_st_16x256b_x2:
5449 case Intrinsic::nvvm_tcgen05_st_32x32b_x8:
5450 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x8: {
5452 Info.memVT = MVT::v8i32;
5453 Info.ptrVal =
I.getArgOperand(0);
5461 case Intrinsic::nvvm_tcgen05_st_16x64b_x16:
5462 case Intrinsic::nvvm_tcgen05_st_16x128b_x8:
5463 case Intrinsic::nvvm_tcgen05_st_16x256b_x4:
5464 case Intrinsic::nvvm_tcgen05_st_32x32b_x16:
5465 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x16: {
5467 Info.memVT = MVT::v16i32;
5468 Info.ptrVal =
I.getArgOperand(0);
5476 case Intrinsic::nvvm_tcgen05_st_16x64b_x32:
5477 case Intrinsic::nvvm_tcgen05_st_16x128b_x16:
5478 case Intrinsic::nvvm_tcgen05_st_16x256b_x8:
5479 case Intrinsic::nvvm_tcgen05_st_32x32b_x32:
5480 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x32: {
5482 Info.memVT = MVT::v32i32;
5483 Info.ptrVal =
I.getArgOperand(0);
5491 case Intrinsic::nvvm_tcgen05_st_16x64b_x64:
5492 case Intrinsic::nvvm_tcgen05_st_16x128b_x32:
5493 case Intrinsic::nvvm_tcgen05_st_16x256b_x16:
5494 case Intrinsic::nvvm_tcgen05_st_32x32b_x64:
5495 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x64: {
5497 Info.memVT = MVT::v64i32;
5498 Info.ptrVal =
I.getArgOperand(0);
5506 case Intrinsic::nvvm_tcgen05_st_16x64b_x128:
5507 case Intrinsic::nvvm_tcgen05_st_16x128b_x64:
5508 case Intrinsic::nvvm_tcgen05_st_16x256b_x32:
5509 case Intrinsic::nvvm_tcgen05_st_32x32b_x128:
5510 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x128: {
5512 Info.memVT = MVT::v128i32;
5513 Info.ptrVal =
I.getArgOperand(0);
5521 nvvm_tcgen05_mma_shared_f8f6f4_disable_output_lane_cg1_decompress_b:
5523 nvvm_tcgen05_mma_tensor_f8f6f4_disable_output_lane_cg1_decompress_b:
5524 case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg1:
5525 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg1:
5526 case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg1:
5527 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg1:
5528 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1:
5529 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1:
5530 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1_ashift:
5532 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1_ashift:
5533 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1:
5534 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1:
5535 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1_ashift:
5537 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1_ashift: {
5540 Info.memVT = MVT::v4i32;
5541 Info.ptrVal =
I.getArgOperand(0);
5544 Info.align =
Align(16);
5550 nvvm_tcgen05_mma_shared_f8f6f4_disable_output_lane_cg2_decompress_b:
5552 nvvm_tcgen05_mma_tensor_f8f6f4_disable_output_lane_cg2_decompress_b:
5553 case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg2:
5554 case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg2:
5555 case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg2:
5556 case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg2:
5557 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2:
5558 case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2:
5559 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2:
5560 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2:
5561 case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2_ashift:
5563 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2_ashift:
5564 case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2_ashift:
5566 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2_ashift: {
5569 Info.memVT = MVT::v8i32;
5570 Info.ptrVal =
I.getArgOperand(0);
5573 Info.align =
Align(16);
5577 case Intrinsic::nvvm_tcgen05_alloc_cg1:
5578 case Intrinsic::nvvm_tcgen05_alloc_cg2:
5580 Info.memVT = MVT::i32;
5581 Info.ptrVal =
I.getArgOperand(0);
5584 Info.align =
Align(4);