| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
| 8 个月前 | ||
[MLIR][NVVM] Add clusterlaunchcontrol Ops (#156585) This change adds the clusterlaunchcontrol.try.cancel and clusterlaunchcontrol.query.cancel Ops to the NVVM dialect. Tests are added in clusterlaunchcontrol.mlir. PTX Reference: https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-clusterlaunchcontrol-try-cancel | 10 个月前 | |
[MLIR][NVVM] Add support for converting fp4/6/8 to fp16x2 (#162439) This change adds the following NVVM dialect Ops for converting fp4/6/8 to fp16x2: - convert.f4x2.to.f16x2 - convert.f6x2.to.f16x2 - convert.f8x2.to.f16x2 - convert.f8x2.to.bf16x2 Tests are added in convert_fp4x2.mlir, convert_fp6x2.mlir, and convert_fp8x2.mlir. PTX Reference: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cvt | 9 个月前 | |
[MLIR][NVVM] Add support for converting fp4/6/8 to fp16x2 (#162439) This change adds the following NVVM dialect Ops for converting fp4/6/8 to fp16x2: - convert.f4x2.to.f16x2 - convert.f6x2.to.f16x2 - convert.f8x2.to.f16x2 - convert.f8x2.to.bf16x2 Tests are added in convert_fp4x2.mlir, convert_fp6x2.mlir, and convert_fp8x2.mlir. PTX Reference: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cvt | 9 个月前 | |
[MLIR][NVVM] Add support for converting fp4/6/8 to fp16x2 (#162439) This change adds the following NVVM dialect Ops for converting fp4/6/8 to fp16x2: - convert.f4x2.to.f16x2 - convert.f6x2.to.f16x2 - convert.f8x2.to.f16x2 - convert.f8x2.to.bf16x2 Tests are added in convert_fp4x2.mlir, convert_fp6x2.mlir, and convert_fp8x2.mlir. PTX Reference: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cvt | 9 个月前 | |
[MLIR][NVVM] Add support for Convert Ops with rs-rounding mode (#165736) Added NVVM dialect operations for stochastic rounding (.rs) conversions from F32 to various packed floating-point formats. These operations map to existing PTX instructions and LLVM intrinsics. Supported conversions: - F32x2 to F16x2/BF16x2 (with optional relu and satfinite modifiers) - F32x4 to packed F8 formats (E4M3, E5M2) - F32x4 to packed F6 formats (E2M3, E3M2) - F32x4 to packed F4 format (E2M1) All operations support stochastic rounding with randomness provided via an rbits parameter, and optional relu and saturation modifiers. | 9 个月前 | |
[MLIR][NVVM] Rename cvt Ops to convert (#140868) This patch renames the cvt Ops and related structures to convert in the NVVM dialect to be more descriptive. | 1 年前 | |
[MLIR][NVVM] Add an explicit mask operand to elect.sync (#145509) This patch adds a mask operand to elect.sync explicitly. When provided, this overrides the default value of 0xffffffff. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 1 年前 | |
[MLIR][NVVM] Fix the lowering of mbarrier.test.wait (#166555) PR #165993 accidentally broke the lowering of the test.wait Op. This patch fixes the issue and adds tests to verify the lowering to intrinsics for all mbarrier Ops, ensuring similar regressions are caught in the future. Additionally, the cp-async-mbarrier test is moved to the mbarriers.mlir test file to keep all related tests together. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 9 个月前 | |
[mlir][NVVM] Add nvvm.membar operation (#166698) Add nvvm.membar operation with level as defined in https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-membar This will be used to replace direct intrinsic call in CUDA Fortran for threadfence(), threadfence_block and thread fence_system() currently lowered here: https://github.com/llvm/llvm-project/blob/e700f157026bf8b4d58f936c5db8f152e269d77f/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp#L1310 The nvvm membar intrsinsic are also used in CUDA C/C++ (https://github.com/llvm/llvm-project/blob/49f55f4991227f3c7a2b8161bbf45c74b7023944/clang/lib/Headers/__clang_cuda_device_functions.h#L528) | 9 个月前 | |
[MLIR][NVVM][NVGPU] Combine prefetch and prefetch.tensormap (#153134) This PR combines the prefetch and prefetch.tensormap NVVM Ops to one prefetch Op. The tensormap variant is lowered through the newly added intrinsics. The lowering of the NVGPU tma.prefetch.descriptor Op is changed from lowering to the prefetch.tensormap Op to prefetch. PTX Spec Reference: https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-prefetch-prefetchu | 11 个月前 | |
[MLIR][NVVM] Update redux.sync op (#166125) This change: - Updates the redux.sync NVVM Op input and output type constraints. - Adds a verifier for the Op to prevent stack dumps and hitting llvm_unreachable in certain invalid usage scenarios. Instead, we gracefully error out with an informative message now. | 9 个月前 | |
[MLIR][NVVM] Update Op verifiers to prevent ungraceful exits (#165677) Updates the following Ops to prevent ungraceful exits with a stack-dump in certain cases of incorrect usages, and instead gracefully error out with a more informative error message: - tcgen05.ld - shfl.sync | 9 个月前 | |
[MLIR][NVVM] [NFC] Rename Tcgen05GroupKind to CTAGroupKind (#156448) ...as the cta_group::1/2 are used in non-tcgen05 Ops like TMA Loads also. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 11 个月前 | |
[MLIR][NVVM] [NFC] Rename Tcgen05GroupKind to CTAGroupKind (#156448) ...as the cta_group::1/2 are used in non-tcgen05 Ops like TMA Loads also. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 11 个月前 | |
[MLIR][NVVM] [NFC] Rename Tcgen05GroupKind to CTAGroupKind (#156448) ...as the cta_group::1/2 are used in non-tcgen05 Ops like TMA Loads also. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 11 个月前 | |
[MLIR][NVVM] [NFC] Update test cmd-lines and doc links (#128207) For the NVVM Dialect tests under Target/LLVMIR/nvvm/ dir, we verify the lowering to the intrinsics using mlir-translate. Remove the -verify-diagnostics option from the cmd-line for these tests since all the verifier checks are tested through the nvvmir-invalid.mlir file. Similarly, remove the split-input-file option which is not relevant here. Update a few remaining links in the NVVMOps.td file. All the reference links follow the same style now. Rename the tcgen05-barriers.mlir file to tcgen05-commit.mlir and move the wait/fence tests to a separate file. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 1 年前 | |
[MLIR][NVVM] Update Op verifiers to prevent ungraceful exits (#165677) Updates the following Ops to prevent ungraceful exits with a stack-dump in certain cases of incorrect usages, and instead gracefully error out with a more informative error message: - tcgen05.ld - shfl.sync | 9 个月前 | |
[MLIR][NVVM] Add support for tcgen05.{ld, st} (#130728) This commit adds support for tcgen05.{ld, st} to the NVVM Dialect with tests under tcgen05-ld.mlir and tcgen05-st.mlir respectively | 1 年前 | |
Reland "[MLIR][NVVM] Add tcgen05.mma MLIR Ops (#164356)" (#168638) Reland commit fb829bf11feeb53f815a3abf539e63ec3a23ed3d with additional fixes relating to post-merge CI failure /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp: In function ‘constexpr llvm::nvvm::CTAGroupKind getNVVMCtaGroupKind(mlir::NVVM::CTAGroupKind)’: /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/llvm/include/llvm/Support/ErrorHandling.h:165:36: error: call to non-constexpr function ‘void llvm::llvm_unreachable_internal(const char*, const char*, unsigned int)’ ::llvm::llvm_unreachable_internal(msg, __FILE__, __LINE__) ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~ /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp:73:3: note: in expansion of macro ‘llvm_unreachable’ llvm_unreachable("unsupported cta_group value"); ^ | 8 个月前 | |
Reland "[MLIR][NVVM] Add tcgen05.mma MLIR Ops (#164356)" (#168638) Reland commit fb829bf11feeb53f815a3abf539e63ec3a23ed3d with additional fixes relating to post-merge CI failure /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp: In function ‘constexpr llvm::nvvm::CTAGroupKind getNVVMCtaGroupKind(mlir::NVVM::CTAGroupKind)’: /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/llvm/include/llvm/Support/ErrorHandling.h:165:36: error: call to non-constexpr function ‘void llvm::llvm_unreachable_internal(const char*, const char*, unsigned int)’ ::llvm::llvm_unreachable_internal(msg, __FILE__, __LINE__) ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~ /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp:73:3: note: in expansion of macro ‘llvm_unreachable’ llvm_unreachable("unsupported cta_group value"); ^ | 8 个月前 | |
Reland "[MLIR][NVVM] Add tcgen05.mma MLIR Ops (#164356)" (#168638) Reland commit fb829bf11feeb53f815a3abf539e63ec3a23ed3d with additional fixes relating to post-merge CI failure /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp: In function ‘constexpr llvm::nvvm::CTAGroupKind getNVVMCtaGroupKind(mlir::NVVM::CTAGroupKind)’: /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/llvm/include/llvm/Support/ErrorHandling.h:165:36: error: call to non-constexpr function ‘void llvm::llvm_unreachable_internal(const char*, const char*, unsigned int)’ ::llvm::llvm_unreachable_internal(msg, __FILE__, __LINE__) ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~ /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp:73:3: note: in expansion of macro ‘llvm_unreachable’ llvm_unreachable("unsupported cta_group value"); ^ | 8 个月前 | |
Reland "[MLIR][NVVM] Add tcgen05.mma MLIR Ops (#164356)" (#168638) Reland commit fb829bf11feeb53f815a3abf539e63ec3a23ed3d with additional fixes relating to post-merge CI failure /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp: In function ‘constexpr llvm::nvvm::CTAGroupKind getNVVMCtaGroupKind(mlir::NVVM::CTAGroupKind)’: /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/llvm/include/llvm/Support/ErrorHandling.h:165:36: error: call to non-constexpr function ‘void llvm::llvm_unreachable_internal(const char*, const char*, unsigned int)’ ::llvm::llvm_unreachable_internal(msg, __FILE__, __LINE__) ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~ /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp:73:3: note: in expansion of macro ‘llvm_unreachable’ llvm_unreachable("unsupported cta_group value"); ^ | 8 个月前 | |
Reland "[MLIR][NVVM] Add tcgen05.mma MLIR Ops (#164356)" (#168638) Reland commit fb829bf11feeb53f815a3abf539e63ec3a23ed3d with additional fixes relating to post-merge CI failure /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp: In function ‘constexpr llvm::nvvm::CTAGroupKind getNVVMCtaGroupKind(mlir::NVVM::CTAGroupKind)’: /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/llvm/include/llvm/Support/ErrorHandling.h:165:36: error: call to non-constexpr function ‘void llvm::llvm_unreachable_internal(const char*, const char*, unsigned int)’ ::llvm::llvm_unreachable_internal(msg, __FILE__, __LINE__) ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~ /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp:73:3: note: in expansion of macro ‘llvm_unreachable’ llvm_unreachable("unsupported cta_group value"); ^ | 8 个月前 | |
Reland "[MLIR][NVVM] Add tcgen05.mma MLIR Ops (#164356)" (#168638) Reland commit fb829bf11feeb53f815a3abf539e63ec3a23ed3d with additional fixes relating to post-merge CI failure /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp: In function ‘constexpr llvm::nvvm::CTAGroupKind getNVVMCtaGroupKind(mlir::NVVM::CTAGroupKind)’: /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/llvm/include/llvm/Support/ErrorHandling.h:165:36: error: call to non-constexpr function ‘void llvm::llvm_unreachable_internal(const char*, const char*, unsigned int)’ ::llvm::llvm_unreachable_internal(msg, __FILE__, __LINE__) ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~ /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp:73:3: note: in expansion of macro ‘llvm_unreachable’ llvm_unreachable("unsupported cta_group value"); ^ | 8 个月前 | |
Reland "[MLIR][NVVM] Add tcgen05.mma MLIR Ops (#164356)" (#168638) Reland commit fb829bf11feeb53f815a3abf539e63ec3a23ed3d with additional fixes relating to post-merge CI failure /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp: In function ‘constexpr llvm::nvvm::CTAGroupKind getNVVMCtaGroupKind(mlir::NVVM::CTAGroupKind)’: /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/llvm/include/llvm/Support/ErrorHandling.h:165:36: error: call to non-constexpr function ‘void llvm::llvm_unreachable_internal(const char*, const char*, unsigned int)’ ::llvm::llvm_unreachable_internal(msg, __FILE__, __LINE__) ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~ /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp:73:3: note: in expansion of macro ‘llvm_unreachable’ llvm_unreachable("unsupported cta_group value"); ^ | 8 个月前 | |
Reland "[MLIR][NVVM] Add tcgen05.mma MLIR Ops (#164356)" (#168638) Reland commit fb829bf11feeb53f815a3abf539e63ec3a23ed3d with additional fixes relating to post-merge CI failure /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp: In function ‘constexpr llvm::nvvm::CTAGroupKind getNVVMCtaGroupKind(mlir::NVVM::CTAGroupKind)’: /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/llvm/include/llvm/Support/ErrorHandling.h:165:36: error: call to non-constexpr function ‘void llvm::llvm_unreachable_internal(const char*, const char*, unsigned int)’ ::llvm::llvm_unreachable_internal(msg, __FILE__, __LINE__) ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~ /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp:73:3: note: in expansion of macro ‘llvm_unreachable’ llvm_unreachable("unsupported cta_group value"); ^ | 8 个月前 | |
Reland "[MLIR][NVVM] Add tcgen05.mma MLIR Ops (#164356)" (#168638) Reland commit fb829bf11feeb53f815a3abf539e63ec3a23ed3d with additional fixes relating to post-merge CI failure /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp: In function ‘constexpr llvm::nvvm::CTAGroupKind getNVVMCtaGroupKind(mlir::NVVM::CTAGroupKind)’: /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/llvm/include/llvm/Support/ErrorHandling.h:165:36: error: call to non-constexpr function ‘void llvm::llvm_unreachable_internal(const char*, const char*, unsigned int)’ ::llvm::llvm_unreachable_internal(msg, __FILE__, __LINE__) ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~ /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp:73:3: note: in expansion of macro ‘llvm_unreachable’ llvm_unreachable("unsupported cta_group value"); ^ | 8 个月前 | |
Reland "[MLIR][NVVM] Add tcgen05.mma MLIR Ops (#164356)" (#168638) Reland commit fb829bf11feeb53f815a3abf539e63ec3a23ed3d with additional fixes relating to post-merge CI failure /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp: In function ‘constexpr llvm::nvvm::CTAGroupKind getNVVMCtaGroupKind(mlir::NVVM::CTAGroupKind)’: /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/llvm/include/llvm/Support/ErrorHandling.h:165:36: error: call to non-constexpr function ‘void llvm::llvm_unreachable_internal(const char*, const char*, unsigned int)’ ::llvm::llvm_unreachable_internal(msg, __FILE__, __LINE__) ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~ /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp:73:3: note: in expansion of macro ‘llvm_unreachable’ llvm_unreachable("unsupported cta_group value"); ^ | 8 个月前 | |
Reland "[MLIR][NVVM] Add tcgen05.mma MLIR Ops (#164356)" (#168638) Reland commit fb829bf11feeb53f815a3abf539e63ec3a23ed3d with additional fixes relating to post-merge CI failure /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp: In function ‘constexpr llvm::nvvm::CTAGroupKind getNVVMCtaGroupKind(mlir::NVVM::CTAGroupKind)’: /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/llvm/include/llvm/Support/ErrorHandling.h:165:36: error: call to non-constexpr function ‘void llvm::llvm_unreachable_internal(const char*, const char*, unsigned int)’ ::llvm::llvm_unreachable_internal(msg, __FILE__, __LINE__) ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~ /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp:73:3: note: in expansion of macro ‘llvm_unreachable’ llvm_unreachable("unsupported cta_group value"); ^ | 8 个月前 | |
Reland "[MLIR][NVVM] Add tcgen05.mma MLIR Ops (#164356)" (#168638) Reland commit fb829bf11feeb53f815a3abf539e63ec3a23ed3d with additional fixes relating to post-merge CI failure /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp: In function ‘constexpr llvm::nvvm::CTAGroupKind getNVVMCtaGroupKind(mlir::NVVM::CTAGroupKind)’: /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/llvm/include/llvm/Support/ErrorHandling.h:165:36: error: call to non-constexpr function ‘void llvm::llvm_unreachable_internal(const char*, const char*, unsigned int)’ ::llvm::llvm_unreachable_internal(msg, __FILE__, __LINE__) ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~ /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp:73:3: note: in expansion of macro ‘llvm_unreachable’ llvm_unreachable("unsupported cta_group value"); ^ | 8 个月前 | |
Reland "[MLIR][NVVM] Add tcgen05.mma MLIR Ops (#164356)" (#168638) Reland commit fb829bf11feeb53f815a3abf539e63ec3a23ed3d with additional fixes relating to post-merge CI failure /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp: In function ‘constexpr llvm::nvvm::CTAGroupKind getNVVMCtaGroupKind(mlir::NVVM::CTAGroupKind)’: /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/llvm/include/llvm/Support/ErrorHandling.h:165:36: error: call to non-constexpr function ‘void llvm::llvm_unreachable_internal(const char*, const char*, unsigned int)’ ::llvm::llvm_unreachable_internal(msg, __FILE__, __LINE__) ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~ /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp:73:3: note: in expansion of macro ‘llvm_unreachable’ llvm_unreachable("unsupported cta_group value"); ^ | 8 个月前 | |
[MLIR][NVVM] [NFC] Rename Tcgen05GroupKind to CTAGroupKind (#156448) ...as the cta_group::1/2 are used in non-tcgen05 Ops like TMA Loads also. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 11 个月前 | |
[MLIR][NVVM] Add Op to create tcgen05-mma smem descriptor (#141651) This patch adds an Op to create the shared-memory descriptor for Tcgen05 MMA. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 1 年前 | |
[MLIR][NVVM] Add support for tcgen05.{ld, st} (#130728) This commit adds support for tcgen05.{ld, st} to the NVVM Dialect with tests under tcgen05-ld.mlir and tcgen05-st.mlir respectively | 1 年前 | |
[MLIR][NVVM] Add support for shared::cta destination (#168056) This patch adds support for shared::cta as destination space in the TMA non-tensor copy Op (from global to shared::cta). * Appropriate verifier checks are added. * Unit tests are added to verify the lowering. The related intrinsic changes were merged through PR #167508. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 8 个月前 | |
[MLIR][NVVM] Add support for shared::cta destination (#168056) This patch adds support for shared::cta as destination space in the TMA non-tensor copy Op (from global to shared::cta). * Appropriate verifier checks are added. * Unit tests are added to verify the lowering. The related intrinsic changes were merged through PR #167508. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 8 个月前 | |
[MLIR][NVVM] Update TMA Load Op (#156347) This patch includes im2col and gather mode support for the TMA Load Op. The lowering is also updated to intrinsics except when a Predicate is given. This completes the Blackwell additions on this Op. * NVVM Dialect has support for Shared::Cluster address-space now. So, this patch also updates the Op to use AS(7) instead of AS(3). The corresponding inline-ptx based unit tests are also updated. * lit tests are added for all combinations. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 10 个月前 | |
[MLIR][NVVM] Update TMA Load Op (#156347) This patch includes im2col and gather mode support for the TMA Load Op. The lowering is also updated to intrinsics except when a Predicate is given. This completes the Blackwell additions on this Op. * NVVM Dialect has support for Shared::Cluster address-space now. So, this patch also updates the Op to use AS(7) instead of AS(3). The corresponding inline-ptx based unit tests are also updated. * lit tests are added for all combinations. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 10 个月前 | |
[MLIR][NVVM] Update TMA Load Op (#156347) This patch includes im2col and gather mode support for the TMA Load Op. The lowering is also updated to intrinsics except when a Predicate is given. This completes the Blackwell additions on this Op. * NVVM Dialect has support for Shared::Cluster address-space now. So, this patch also updates the Op to use AS(7) instead of AS(3). The corresponding inline-ptx based unit tests are also updated. * lit tests are added for all combinations. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 10 个月前 | |
[MLIR][NVVM] Update TMA Load Op (#156347) This patch includes im2col and gather mode support for the TMA Load Op. The lowering is also updated to intrinsics except when a Predicate is given. This completes the Blackwell additions on this Op. * NVVM Dialect has support for Shared::Cluster address-space now. So, this patch also updates the Op to use AS(7) instead of AS(3). The corresponding inline-ptx based unit tests are also updated. * lit tests are added for all combinations. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 10 个月前 | |
[MLIR][NVVM] Update TMA Load Op (#156347) This patch includes im2col and gather mode support for the TMA Load Op. The lowering is also updated to intrinsics except when a Predicate is given. This completes the Blackwell additions on this Op. * NVVM Dialect has support for Shared::Cluster address-space now. So, this patch also updates the Op to use AS(7) instead of AS(3). The corresponding inline-ptx based unit tests are also updated. * lit tests are added for all combinations. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 10 个月前 | |
[MLIR][NVVM] Update TMA tensor prefetch Op (#153464) This patch updates the TMA Tensor prefetch Op to add support for im2col_w/w128 and tile_gather4 modes. This completes support for all modes available in Blackwell. * lit tests are added for all possible combinations. * The invalid tests are moved to a separate file with more coverage. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 11 个月前 | |
[MLIR][NVVM] Update TMA tensor prefetch Op (#153464) This patch updates the TMA Tensor prefetch Op to add support for im2col_w/w128 and tile_gather4 modes. This completes support for all modes available in Blackwell. * lit tests are added for all possible combinations. * The invalid tests are moved to a separate file with more coverage. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 11 个月前 | |
[MLIR][NVVM] Update TMA Store Op (#155435) This patch includes im2col and scatter mode support to the TMA Store Op. The lowering is also updated to intrinsics except when Predicate is given. This completes the Blackwell additions on this Op. * lit tests are added for all combinations. * Move the TMA reduce invalid tests to their own file. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 11 个月前 | |
[MLIR][NVVM] Update TMA Store Op (#155435) This patch includes im2col and scatter mode support to the TMA Store Op. The lowering is also updated to intrinsics except when Predicate is given. This completes the Blackwell additions on this Op. * lit tests are added for all combinations. * Move the TMA reduce invalid tests to their own file. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 11 个月前 | |
[MLIR][NVVM] Fix undef in cp.async.bulk.tensor.reduce Op (#157423) This change: - Moves the LLVMIR lowering code of the NVVM dialect cp.async.bulk.tensor.reduce Op to NVVMDialect.cpp. - Fixes the usage of undef in the lowering since it is now deprecated. - Removes macros to use a table to look up intrinsics instead. The tests are updated accordingly. | 10 个月前 | |
[MLIR][NVVM] Update TMA Store Op (#155435) This patch includes im2col and scatter mode support to the TMA Store Op. The lowering is also updated to intrinsics except when Predicate is given. This completes the Blackwell additions on this Op. * lit tests are added for all combinations. * Move the TMA reduce invalid tests to their own file. Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 11 个月前 |
| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
| 8 个月前 | ||
| 10 个月前 | ||
| 9 个月前 | ||
| 9 个月前 | ||
| 9 个月前 | ||
| 9 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 9 个月前 | ||
| 9 个月前 | ||
| 11 个月前 | ||
| 9 个月前 | ||
| 9 个月前 | ||
| 11 个月前 | ||
| 11 个月前 | ||
| 11 个月前 | ||
| 1 年前 | ||
| 9 个月前 | ||
| 1 年前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 11 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 10 个月前 | ||
| 10 个月前 | ||
| 10 个月前 | ||
| 10 个月前 | ||
| 10 个月前 | ||
| 11 个月前 | ||
| 11 个月前 | ||
| 11 个月前 | ||
| 11 个月前 | ||
| 10 个月前 | ||
| 11 个月前 |