diff --git a/swe-paddle/tasks/PaddlePaddle__Paddle-64320/README.md b/swe-paddle/tasks/PaddlePaddle__Paddle-64320/README.md new file mode 100644 index 000000000..a35cd7bbf --- /dev/null +++ b/swe-paddle/tasks/PaddlePaddle__Paddle-64320/README.md @@ -0,0 +1,37 @@ +# PaddlePaddle__Paddle-64320 + +This directory converts Paddle PR #64320 into a SWE-Paddle community task candidate. + +## Source + +| Field | Value | +| --- | --- | +| Repo | `PaddlePaddle/Paddle` | +| PR | [64320](https://github.com/PaddlePaddle/Paddle/pull/64320) | +| PR title | 【Hackathon 6th No.17】为 Paddle 新增 sparse.mask_as API -part | +| Base commit | `605f5e20305db0e4932a20d3e0e6cf7d7d9631d8` | +| Merged at | `2024-06-07T09:10:27Z` | +| Hackathon | `6th` task `17` | +| Task type | `feature_enhancement` | +| Resource | CPU + GPU | + +## Summary + +Add `sparse.mask_as` API for Paddle, which extracts values from a dense tensor at positions indicated by a sparse mask and outputs a sparse tensor. + +## Files + +- `proposal.md`: candidate proposal for maintainer triage. +- `instruction.md`: self-contained problem statement for the coding agent. +- `solution/code.patch`: gold patch from the merged PR. +- `tests/test.patch`: test patch exposing the target behavior. +- `tests/test.sh`: minimal target test command. +- `environment/README.md`: environment notes for reproduction. + +## Verification + +```bash +bash tests/test.sh +``` + +Expected behavior: applying `tests/test.patch` to `base_commit` should fail on the target behavior; applying both `tests/test.patch` and `solution/code.patch` should pass the target tests. diff --git a/swe-paddle/tasks/PaddlePaddle__Paddle-64320/environment/README.md b/swe-paddle/tasks/PaddlePaddle__Paddle-64320/environment/README.md new file mode 100644 index 000000000..85bd29e04 --- /dev/null +++ b/swe-paddle/tasks/PaddlePaddle__Paddle-64320/environment/README.md @@ -0,0 +1,27 @@ +# Environment Notes + +This candidate is part of the SWE-Paddle task set. + +## Expected Environment + +- Repository: `PaddlePaddle/Paddle` +- Base commit: `605f5e20305db0e4932a20d3e0e6cf7d7d9631d8` +- Resource: CPU + GPU +- GPU required: yes (CUDA kernels are included) +- Build path: Paddle source checkout at the base commit. This task involves C++ CPU kernels, CUDA GPU kernels, Python API, and YAML op definitions, so source build is required. + +## Run Order + +1. Check out `PaddlePaddle/Paddle` at the base commit. +2. Apply `tests/test.patch`. +3. Run `bash tests/test.sh`; the target behavior should fail before the fix. +4. Apply `solution/code.patch`. +5. Run `bash tests/test.sh` again; the target behavior should pass after the gold patch. + +## Minimal Test Command + +```bash +bash tests/test.sh +``` + +The verifier is responsible for deriving stable F2P and P2P node IDs from repeated runs. diff --git a/swe-paddle/tasks/PaddlePaddle__Paddle-64320/instruction.md b/swe-paddle/tasks/PaddlePaddle__Paddle-64320/instruction.md new file mode 100644 index 000000000..6238ee3ad --- /dev/null +++ b/swe-paddle/tasks/PaddlePaddle__Paddle-64320/instruction.md @@ -0,0 +1,24 @@ +# 新增 sparse.mask_as API + +## 详细描述 + +为 Paddle 稀疏计算新增 `sparse.mask_as` API。该 API 根据给定的稀疏 mask（SparseCooTensor 或 SparseCsrTensor），从稠密 Tensor 中提取对应非零位置的值，输出一个与 mask 具有相同 indices 的稀疏 Tensor。 + +要求支持： +- COO 格式：支持 1-D ~ 4-D 输入 +- CSR 格式：支持 2-D 和 3-D 输入（其他维度应报错） +- 数据类型：float32, float64, int32, int64, complex64, complex128, int8, int16, float16 +- 前向计算 + 反向梯度 + +## 验收说明 + +- `paddle.sparse.mask_as(x, mask)` 应能正确根据 mask 的 indices 从稠密 Tensor x 中提取值 +- 支持 COO 和 CSR 两种稀疏格式 +- CSR 格式仅支持 2-D 和 3-D，其他维度应报错 +- 反向梯度应正确传播 + +## Acceptance Criteria + +- The behavior described above should be fixed. +- Existing valid behavior should remain unchanged. +- Do not satisfy the task by deleting tests, weakening assertions, or bypassing validation broadly. diff --git a/swe-paddle/tasks/PaddlePaddle__Paddle-64320/solution/code.patch b/swe-paddle/tasks/PaddlePaddle__Paddle-64320/solution/code.patch new file mode 100644 index 000000000..602a3c2b4 --- /dev/null +++ b/swe-paddle/tasks/PaddlePaddle__Paddle-64320/solution/code.patch @@ -0,0 +1,872 @@ +diff --git a/paddle/phi/kernels/sparse/cpu/mask_grad_kernel.cc b/paddle/phi/kernels/sparse/cpu/mask_grad_kernel.cc +new file mode 100644 +index 00000000000000..3503c88b2ef8b4 +--- /dev/null ++++ b/paddle/phi/kernels/sparse/cpu/mask_grad_kernel.cc +@@ -0,0 +1,56 @@ ++/* Copyright (c) 2022 PaddlePaddle Authors. All Rights Reserved. ++ ++Licensed under the Apache License, Version 2.0 (the "License"); ++you may not use this file except in compliance with the License. ++You may obtain a copy of the License at ++ ++ http://www.apache.org/licenses/LICENSE-2.0 ++ ++Unless required by applicable law or agreed to in writing, software ++distributed under the License is distributed on an "AS IS" BASIS, ++WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ++See the License for the specific language governing permissions and ++limitations under the License. */ ++ ++#include "paddle/phi/kernels/sparse/mask_grad_kernel.h" ++#include "paddle/phi/kernels/sparse/mask_kernel.h" ++#include "paddle/phi/kernels/sparse/sparse_utils_kernel.h" ++ ++#include "paddle/phi/backends/cpu/cpu_context.h" ++#include "paddle/phi/core/kernel_registry.h" ++ ++PD_REGISTER_KERNEL(mask_as_coo_grad, ++ CPU, ++ ALL_LAYOUT, ++ phi::sparse::MaskAsCooGradKernel, ++ float, ++ double, ++ phi::dtype::float16, ++ uint8_t, ++ int8_t, ++ int16_t, ++ int, ++ int64_t, ++ bool, ++ phi::dtype::complex, ++ phi::dtype::complex) { ++ kernel->InputAt(1).SetDataLayout(phi::DataLayout::SPARSE_COO); ++} ++ ++PD_REGISTER_KERNEL(mask_as_csr_grad, ++ CPU, ++ ALL_LAYOUT, ++ phi::sparse::MaskAsCsrGradKernel, ++ float, ++ double, ++ phi::dtype::float16, ++ uint8_t, ++ int8_t, ++ int16_t, ++ int, ++ int64_t, ++ bool, ++ phi::dtype::complex, ++ phi::dtype::complex) { ++ kernel->InputAt(1).SetDataLayout(phi::DataLayout::SPARSE_CSR); ++} +diff --git a/paddle/phi/kernels/sparse/cpu/mask_kernel.cc b/paddle/phi/kernels/sparse/cpu/mask_kernel.cc +index 5213dd44a4c07c..658a26452dafb6 100644 +--- a/paddle/phi/kernels/sparse/cpu/mask_kernel.cc ++++ b/paddle/phi/kernels/sparse/cpu/mask_kernel.cc +@@ -13,6 +13,7 @@ See the License for the specific language governing permissions and + limitations under the License. */ + + #include "paddle/phi/kernels/sparse/mask_kernel.h" ++#include "paddle/phi/kernels/sparse/sparse_utils_kernel.h" + + #include "paddle/common/ddim.h" + #include "paddle/phi/api/ext/dispatch.h" +@@ -75,16 +76,116 @@ void MaskCooCPUKernel(const CPUContext& dev_ctx, + * x and mask must have the same shape. + **/ + template +-void MaskCooKernel(const Context& dev_ctx, +- const DenseTensor& x, +- const SparseCooTensor& mask, +- SparseCooTensor* out) { ++void MaskAsCooKernel(const Context& dev_ctx, ++ const DenseTensor& x, ++ const SparseCooTensor& mask, ++ SparseCooTensor* out) { + PD_VISIT_BASE_INTEGRAL_TYPES( + mask.indices().dtype(), "MaskCooCPUKernel", ([&] { + MaskCooCPUKernel(dev_ctx, x, mask, out); + })); + } + ++template ++void MaskCsr2DCPUKernel(const CPUContext& dev_ctx, ++ const DenseTensor& x, ++ const SparseCsrTensor& mask, ++ SparseCsrTensor* out) { ++ const DenseTensor& mask_cols = mask.cols(); ++ const DenseTensor& mask_crows = mask.crows(); ++ int64_t num_non_zeros = mask.nnz(); ++ ++ DenseTensor out_cols = phi::EmptyLike(dev_ctx, mask_cols); ++ DenseTensor out_crows = phi::EmptyLike(dev_ctx, mask_crows); ++ DenseTensor out_values = phi::Empty(dev_ctx, {num_non_zeros}); ++ ++ phi::Copy(dev_ctx, mask_cols, dev_ctx.GetPlace(), false, &out_cols); ++ phi::Copy(dev_ctx, mask_crows, dev_ctx.GetPlace(), false, &out_crows); ++ ++ int64_t numel = 0; ++ for (int64_t i = 0; i < mask_crows.numel() - 1; ++i) { ++ for (int64_t j = mask_crows.data()[i]; ++ j < mask_crows.data()[i + 1]; ++ ++j) { ++ IntT col_idx = mask_cols.data()[numel]; ++ ++ out_values.data()[numel] = ++ x.data()[(i / x.dims()[0]) * x.dims()[1] + ++ (i % x.dims()[0]) * x.dims()[1] + col_idx]; ++ ++ ++numel; ++ } ++ } ++ ++ out->SetMember(out_crows, out_cols, out_values, x.dims()); ++} ++ ++template ++void MaskCsr3DCPUKernel(const CPUContext& dev_ctx, ++ const DenseTensor& x, ++ const SparseCsrTensor& mask, ++ SparseCsrTensor* out) { ++ const DenseTensor& mask_cols = mask.cols(); ++ const DenseTensor& mask_crows = mask.crows(); ++ int64_t num_non_zeros = mask.nnz(); ++ ++ DenseTensor out_cols = phi::EmptyLike(dev_ctx, mask_cols); ++ DenseTensor out_crows = phi::EmptyLike(dev_ctx, mask_crows); ++ DenseTensor out_values = phi::Empty(dev_ctx, {num_non_zeros}); ++ ++ phi::Copy(dev_ctx, mask_cols, dev_ctx.GetPlace(), false, &out_cols); ++ phi::Copy(dev_ctx, mask_crows, dev_ctx.GetPlace(), false, &out_crows); ++ ++ int64_t numel = 0; ++ for (int64_t i = 0; i < mask_crows.numel() - 1; ++i) { ++ for (int64_t j = mask_crows.data()[i]; ++ j < mask_crows.data()[i + 1]; ++ ++j) { ++ IntT col_idx = mask_cols.data()[numel]; ++ ++ out_values.data()[numel] = ++ x.data()[(i / (mask_crows.numel() / x.dims()[0])) * ++ (x.dims()[1] * x.dims()[2]) + ++ (i % (mask_crows.numel() / x.dims()[0])) * x.dims()[2] + ++ col_idx]; ++ ++ ++numel; ++ } ++ } ++ ++ out->SetMember(out_crows, out_cols, out_values, x.dims()); ++} ++ ++/** ++ * @brief Filter the DenseTensor x by the ++ * mask.crows(), mask.cols() and output a SparseCsrTensor ++ * x and mask must have the same shape. ++ **/ ++template ++void MaskAsCsrKernel(const Context& dev_ctx, ++ const DenseTensor& x, ++ const SparseCsrTensor& mask, ++ SparseCsrTensor* out) { ++ const phi::DDim& x_dims = x.dims(); ++ if (x_dims.size() == 2) { ++ PD_VISIT_BASE_INTEGRAL_TYPES( ++ mask.crows().dtype(), "MaskCsr2DCPUKernel", ([&] { ++ MaskCsr2DCPUKernel(dev_ctx, x, mask, out); ++ })); ++ } else if (x_dims.size() == 3) { ++ PD_VISIT_BASE_INTEGRAL_TYPES( ++ mask.crows().dtype(), "MaskCsr3DCPUKernel", ([&] { ++ MaskCsr3DCPUKernel(dev_ctx, x, mask, out); ++ })); ++ } else { ++ // throw exception ++ phi::errors::InvalidArgument( ++ "mask_as for Sparse CSR Tensor only support 2-D or 3-D, but got " ++ "%d-D.", ++ x_dims.size()); ++ } ++} ++ + template + void MaskHelperCooCPUKernel(const CPUContext& dev_ctx, + const SparseCooTensor& x, +@@ -157,10 +258,26 @@ void MaskHelperCooKernel(const Context& dev_ctx, + } // namespace sparse + } // namespace phi + +-PD_REGISTER_KERNEL(mask_coo, ++PD_REGISTER_KERNEL(mask_helper_coo, ++ CPU, ++ ALL_LAYOUT, ++ phi::sparse::MaskHelperCooKernel, ++ float, ++ double, ++ phi::dtype::float16, ++ uint8_t, ++ int16_t, ++ int, ++ int64_t, ++ phi::dtype::complex, ++ phi::dtype::complex) { ++ kernel->InputAt(0).SetDataLayout(phi::DataLayout::SPARSE_COO); ++} ++ ++PD_REGISTER_KERNEL(mask_as_coo, + CPU, + ALL_LAYOUT, +- phi::sparse::MaskCooKernel, ++ phi::sparse::MaskAsCooKernel, + float, + double, + uint8_t, +@@ -174,18 +291,19 @@ PD_REGISTER_KERNEL(mask_coo, + kernel->InputAt(1).SetDataLayout(phi::DataLayout::SPARSE_COO); + } + +-PD_REGISTER_KERNEL(mask_helper_coo, ++PD_REGISTER_KERNEL(mask_as_csr, + CPU, + ALL_LAYOUT, +- phi::sparse::MaskHelperCooKernel, ++ phi::sparse::MaskAsCsrKernel, + float, + double, +- phi::dtype::float16, + uint8_t, ++ int8_t, + int16_t, + int, + int64_t, ++ bool, + phi::dtype::complex, + phi::dtype::complex) { +- kernel->InputAt(0).SetDataLayout(phi::DataLayout::SPARSE_COO); ++ kernel->InputAt(1).SetDataLayout(phi::DataLayout::SPARSE_CSR); + } +diff --git a/paddle/phi/kernels/sparse/gpu/mask_grad_kernel.cu b/paddle/phi/kernels/sparse/gpu/mask_grad_kernel.cu +new file mode 100644 +index 00000000000000..1e4e3276d82e15 +--- /dev/null ++++ b/paddle/phi/kernels/sparse/gpu/mask_grad_kernel.cu +@@ -0,0 +1,56 @@ ++/* Copyright (c) 2022 PaddlePaddle Authors. All Rights Reserved. ++ ++Licensed under the Apache License, Version 2.0 (the "License"); ++you may not use this file except in compliance with the License. ++You may obtain a copy of the License at ++ ++ http://www.apache.org/licenses/LICENSE-2.0 ++ ++Unless required by applicable law or agreed to in writing, software ++distributed under the License is distributed on an "AS IS" BASIS, ++WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ++See the License for the specific language governing permissions and ++limitations under the License. */ ++ ++#include "paddle/phi/kernels/sparse/mask_grad_kernel.h" ++#include "paddle/phi/kernels/sparse/mask_kernel.h" ++#include "paddle/phi/kernels/sparse/sparse_utils_kernel.h" ++ ++#include "paddle/phi/backends/cpu/cpu_context.h" ++#include "paddle/phi/core/kernel_registry.h" ++ ++PD_REGISTER_KERNEL(mask_as_coo_grad, ++ GPU, ++ ALL_LAYOUT, ++ phi::sparse::MaskAsCooGradKernel, ++ float, ++ double, ++ phi::dtype::float16, ++ uint8_t, ++ int8_t, ++ int16_t, ++ int, ++ int64_t, ++ bool, ++ phi::dtype::complex, ++ phi::dtype::complex) { ++ kernel->InputAt(1).SetDataLayout(phi::DataLayout::SPARSE_COO); ++} ++ ++PD_REGISTER_KERNEL(mask_as_csr_grad, ++ GPU, ++ ALL_LAYOUT, ++ phi::sparse::MaskAsCsrGradKernel, ++ float, ++ double, ++ phi::dtype::float16, ++ uint8_t, ++ int8_t, ++ int16_t, ++ int, ++ int64_t, ++ bool, ++ phi::dtype::complex, ++ phi::dtype::complex) { ++ kernel->InputAt(1).SetDataLayout(phi::DataLayout::SPARSE_CSR); ++} +diff --git a/paddle/phi/kernels/sparse/gpu/mask_kernel.cu b/paddle/phi/kernels/sparse/gpu/mask_kernel.cu +index 0941ad69b0dd2d..3459f6802b8819 100644 +--- a/paddle/phi/kernels/sparse/gpu/mask_kernel.cu ++++ b/paddle/phi/kernels/sparse/gpu/mask_kernel.cu +@@ -12,7 +12,10 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + See the License for the specific language governing permissions and + limitations under the License. */ + ++#include ++ + #include "paddle/phi/kernels/sparse/mask_kernel.h" ++#include "paddle/phi/kernels/sparse/sparse_utils_kernel.h" + + #include "paddle/common/ddim.h" + #include "paddle/phi/backends/gpu/gpu_info.h" +@@ -106,22 +109,256 @@ void MaskCooGPUKernel(const GPUContext& dev_ctx, + out->SetMember(out_indices, out_values, dims, true); + } + ++template ++__global__ void ConvertCsrCrowsToCooRows(const IntT* crows_ptr, ++ const IntT* crows_offsets, ++ IntT* rows_ptr, ++ IntT* batch_ptr, ++ const int rows) { ++ const int b = blockIdx.y; ++ const int64_t offset = crows_offsets ? crows_offsets[b] : 0; ++ const int tid = threadIdx.x + blockIdx.x * blockDim.x; ++ for (int i = tid; i < rows; i += gridDim.x * blockDim.x) { ++ for (int j = crows_ptr[b * (rows + 1) + i]; ++ j < crows_ptr[b * (rows + 1) + i + 1]; ++ j++) { ++ rows_ptr[offset + j] = i; ++ if (batch_ptr) { ++ batch_ptr[offset + j] = b; ++ } ++ } ++ } ++} ++ ++template ++__global__ void GetBatchSizes(const IntT* crows, ++ const int rows, ++ const int batches, ++ IntT* batch_sizes) { ++ const int tid = threadIdx.x + blockIdx.x * blockDim.x; ++ if (tid < batches) { ++ batch_sizes[tid] = crows[tid * (rows + 1) + rows]; ++ } ++} ++ ++template ++void MaskCsr2DGPUKernel(const GPUContext& dev_ctx, ++ const DenseTensor& x, ++ const SparseCsrTensor& mask, ++ SparseCsrTensor* out) { ++ const DenseTensor& mask_cols = mask.cols(); ++ const DenseTensor& mask_crows = mask.crows(); ++ int64_t num_non_zeros = mask.nnz(); ++ ++ DenseTensor out_cols = phi::EmptyLike(dev_ctx, mask_cols); ++ DenseTensor out_crows = phi::EmptyLike(dev_ctx, mask_crows); ++ DenseTensor out_values = phi::Empty(dev_ctx, {num_non_zeros}); ++ ++ phi::Copy(dev_ctx, mask_cols, dev_ctx.GetPlace(), false, &out_cols); ++ phi::Copy(dev_ctx, mask_crows, dev_ctx.GetPlace(), false, &out_crows); ++ ++ const DDim& dims = x.dims(); ++ const int64_t non_zero_num = mask.nnz(); ++ int64_t sparse_dim = 2; ++ DenseTensor sparse_offsets = phi::Empty(dev_ctx, {sparse_dim}); ++ std::vector h_sparse_offsets(sparse_dim); ++ phi::funcs::sparse::CalcOffsetsPerDim( ++ dims, sparse_dim, h_sparse_offsets.data()); ++ ++ phi::backends::gpu::GpuMemcpyAsync(sparse_offsets.data(), ++ &h_sparse_offsets[0], ++ sizeof(int64_t) * sparse_dim, ++ gpuMemcpyHostToDevice, ++ dev_ctx.stream()); ++ ++ const auto& csr_crows = mask.crows(); ++ const auto& csr_cols = mask.cols(); ++ const IntT* csr_crows_data = csr_crows.data(); ++ const IntT* csr_cols_data = csr_cols.data(); ++ ++ const int batches = 1; ++ const int rows = dims[0]; ++ auto dims_2d = flatten_to_2d(dims, sparse_dim); ++ const int cols = dims_2d[1]; ++ ++ DenseTensor indices = phi::Empty(dev_ctx, {sparse_dim, non_zero_num}); ++ IntT* coo_indices = indices.data(); ++ IntT* batch_ptr = nullptr; ++ IntT* coo_rows_data = coo_indices; ++ IntT* coo_cols_data = coo_rows_data + non_zero_num; ++ IntT* offsets_ptr = nullptr; ++ ++ auto config = phi::backends::gpu::GetGpuLaunchConfig1D(dev_ctx, rows, 1); ++ config.block_per_grid.y = batches; ++ ConvertCsrCrowsToCooRows ++ <<>>( ++ csr_crows_data, offsets_ptr, coo_rows_data, batch_ptr, rows); ++ phi::backends::gpu::GpuMemcpyAsync(coo_cols_data, ++ csr_cols_data, ++ sizeof(IntT) * non_zero_num, ++ gpuMemcpyDeviceToDevice, ++ dev_ctx.stream()); ++ ++ const T* x_ptr = x.data(); ++ const IntT* indices_ptr = coo_indices; ++ T* out_values_ptr = out_values.data(); ++ ++ auto config_mask = ++ phi::backends::gpu::GetGpuLaunchConfig1D(dev_ctx, non_zero_num * cols, 1); ++ MaskKernel<<>>(x_ptr, ++ indices_ptr, ++ sparse_offsets.data(), ++ non_zero_num, ++ cols, ++ sparse_dim, ++ out_values_ptr); ++ ++ out->SetMember(out_crows, out_cols, out_values, x.dims()); ++} ++ ++template ++void MaskCsr3DGPUKernel(const GPUContext& dev_ctx, ++ const DenseTensor& x, ++ const SparseCsrTensor& mask, ++ SparseCsrTensor* out) { ++ const DenseTensor& mask_cols = mask.cols(); ++ const DenseTensor& mask_crows = mask.crows(); ++ int64_t num_non_zeros = mask.nnz(); ++ ++ DenseTensor out_cols = phi::EmptyLike(dev_ctx, mask_cols); ++ DenseTensor out_crows = phi::EmptyLike(dev_ctx, mask_crows); ++ DenseTensor out_values = phi::Empty(dev_ctx, {num_non_zeros}); ++ ++ phi::Copy(dev_ctx, mask_cols, dev_ctx.GetPlace(), false, &out_cols); ++ phi::Copy(dev_ctx, mask_crows, dev_ctx.GetPlace(), false, &out_crows); ++ ++ const DDim& dims = x.dims(); ++ const int64_t non_zero_num = mask.nnz(); ++ int64_t sparse_dim = 3; ++ DenseTensor sparse_offsets = phi::Empty(dev_ctx, {sparse_dim}); ++ std::vector h_sparse_offsets(sparse_dim); ++ phi::funcs::sparse::CalcOffsetsPerDim( ++ dims, sparse_dim, h_sparse_offsets.data()); ++ ++ phi::backends::gpu::GpuMemcpyAsync(sparse_offsets.data(), ++ &h_sparse_offsets[0], ++ sizeof(int64_t) * sparse_dim, ++ gpuMemcpyHostToDevice, ++ dev_ctx.stream()); ++ ++ const auto& csr_crows = mask.crows(); ++ const auto& csr_cols = mask.cols(); ++ const IntT* csr_crows_data = csr_crows.data(); ++ const IntT* csr_cols_data = csr_cols.data(); ++ ++ const int batches = dims[0]; ++ const int rows = dims[1]; ++ auto dims_2d = flatten_to_2d(dims, sparse_dim); ++ const int cols = dims_2d[1]; ++ ++ DenseTensor indices = phi::Empty(dev_ctx, {sparse_dim, non_zero_num}); ++ DenseTensor offsets = phi::Empty(dev_ctx, {batches}); ++ IntT* coo_indices = indices.data(); ++ IntT* batch_ptr = coo_indices; ++ IntT* coo_rows_data = batch_ptr + non_zero_num; ++ IntT* coo_cols_data = coo_rows_data + non_zero_num; ++ IntT* offsets_ptr = offsets.data(); ++ ++ auto config_batch = ++ phi::backends::gpu::GetGpuLaunchConfig1D(dev_ctx, batches, 1); ++ GetBatchSizes ++ <<>>( ++ csr_crows_data, rows, batches, offsets_ptr); ++ ++#ifdef PADDLE_WITH_HIP ++ thrust::exclusive_scan(thrust::hip::par.on(dev_ctx.stream()), ++#else ++ thrust::exclusive_scan(thrust::cuda::par.on(dev_ctx.stream()), ++#endif ++ offsets_ptr, ++ offsets_ptr + batches, ++ offsets_ptr); ++ ++ auto config = phi::backends::gpu::GetGpuLaunchConfig1D(dev_ctx, rows, 1); ++ config.block_per_grid.y = batches; ++ ConvertCsrCrowsToCooRows ++ <<>>( ++ csr_crows_data, offsets_ptr, coo_rows_data, batch_ptr, rows); ++ phi::backends::gpu::GpuMemcpyAsync(coo_cols_data, ++ csr_cols_data, ++ sizeof(IntT) * non_zero_num, ++ gpuMemcpyDeviceToDevice, ++ dev_ctx.stream()); ++ ++ const T* x_ptr = x.data(); ++ const IntT* indices_ptr = coo_indices; ++ T* out_values_ptr = out_values.data(); ++ ++ auto config_mask = ++ phi::backends::gpu::GetGpuLaunchConfig1D(dev_ctx, non_zero_num * cols, 1); ++ MaskKernel<<>>(x_ptr, ++ indices_ptr, ++ sparse_offsets.data(), ++ non_zero_num, ++ cols, ++ sparse_dim, ++ out_values_ptr); ++ ++ out->SetMember(out_crows, out_cols, out_values, x.dims()); ++} ++ + /** + * @brief Filter the DenseTensor x by the + * mask.indices() and output a SparseCooTensor + * x and mask must have the same shape. + **/ + template +-void MaskCooKernel(const Context& dev_ctx, +- const DenseTensor& x, +- const SparseCooTensor& mask, +- SparseCooTensor* out) { ++void MaskAsCooKernel(const Context& dev_ctx, ++ const DenseTensor& x, ++ const SparseCooTensor& mask, ++ SparseCooTensor* out) { + PD_VISIT_BASE_INTEGRAL_TYPES( + mask.indices().dtype(), "MaskCooGPUKernel", ([&] { + MaskCooGPUKernel(dev_ctx, x, mask, out); + })); + } + ++/** ++ * @brief Filter the DenseTensor x by the ++ * mask.crows(), mask.cols() and output a SparseCsrTensor ++ * x and mask must have the same shape. ++ **/ ++template ++void MaskAsCsrKernel(const Context& dev_ctx, ++ const DenseTensor& x, ++ const SparseCsrTensor& mask, ++ SparseCsrTensor* out) { ++ const phi::DDim& x_dims = x.dims(); ++ if (x_dims.size() == 2) { ++ PD_VISIT_BASE_INTEGRAL_TYPES( ++ mask.crows().dtype(), "MaskCsr2DGPUKernel", ([&] { ++ MaskCsr2DGPUKernel(dev_ctx, x, mask, out); ++ })); ++ } else if (x_dims.size() == 3) { ++ PD_VISIT_BASE_INTEGRAL_TYPES( ++ mask.crows().dtype(), "MaskCsr3DGPUKernel", ([&] { ++ MaskCsr3DGPUKernel(dev_ctx, x, mask, out); ++ })); ++ } else { ++ // throw exception ++ phi::errors::InvalidArgument( ++ "mask_as for Sparse CSR Tensor only support 2-D or 3-D, but got " ++ "%d-D.", ++ x_dims.size()); ++ } ++} ++ + template + __global__ void MaskTable(const IntT* x_indexs, + const int n, +@@ -296,10 +533,26 @@ void MaskHelperCooKernel(const Context& dev_ctx, + } // namespace sparse + } // namespace phi + +-PD_REGISTER_KERNEL(mask_coo, ++PD_REGISTER_KERNEL(mask_helper_coo, + GPU, + ALL_LAYOUT, +- phi::sparse::MaskCooKernel, ++ phi::sparse::MaskHelperCooKernel, ++ float, ++ double, ++ phi::dtype::float16, ++ uint8_t, ++ int16_t, ++ int, ++ int64_t, ++ phi::dtype::complex, ++ phi::dtype::complex) { ++ kernel->InputAt(0).SetDataLayout(phi::DataLayout::SPARSE_COO); ++} ++ ++PD_REGISTER_KERNEL(mask_as_coo, ++ GPU, ++ ALL_LAYOUT, ++ phi::sparse::MaskAsCooKernel, + float, + double, + phi::dtype::float16, +@@ -314,18 +567,20 @@ PD_REGISTER_KERNEL(mask_coo, + kernel->InputAt(1).SetDataLayout(phi::DataLayout::SPARSE_COO); + } + +-PD_REGISTER_KERNEL(mask_helper_coo, ++PD_REGISTER_KERNEL(mask_as_csr, + GPU, + ALL_LAYOUT, +- phi::sparse::MaskHelperCooKernel, ++ phi::sparse::MaskAsCsrKernel, + float, + double, + phi::dtype::float16, + uint8_t, ++ int8_t, + int16_t, + int, + int64_t, ++ bool, + phi::dtype::complex, + phi::dtype::complex) { +- kernel->InputAt(0).SetDataLayout(phi::DataLayout::SPARSE_COO); ++ kernel->InputAt(1).SetDataLayout(phi::DataLayout::SPARSE_CSR); + } +diff --git a/paddle/phi/kernels/sparse/mask_grad_kernel.h b/paddle/phi/kernels/sparse/mask_grad_kernel.h +new file mode 100644 +index 00000000000000..687562aa300d1c +--- /dev/null ++++ b/paddle/phi/kernels/sparse/mask_grad_kernel.h +@@ -0,0 +1,45 @@ ++/* Copyright (c) 2022 PaddlePaddle Authors. All Rights Reserved. ++ ++Licensed under the Apache License, Version 2.0 (the "License"); ++you may not use this file except in compliance with the License. ++You may obtain a copy of the License at ++ ++ http://www.apache.org/licenses/LICENSE-2.0 ++ ++Unless required by applicable law or agreed to in writing, software ++distributed under the License is distributed on an "AS IS" BASIS, ++WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ++See the License for the specific language governing permissions and ++limitations under the License. */ ++ ++#pragma once ++ ++#include "paddle/phi/core/dense_tensor.h" ++#include "paddle/phi/core/sparse_coo_tensor.h" ++#include "paddle/phi/core/sparse_csr_tensor.h" ++#include "paddle/phi/kernels/sparse/mask_kernel.h" ++#include "paddle/phi/kernels/sparse/sparse_utils_kernel.h" ++ ++namespace phi { ++namespace sparse { ++ ++template ++void MaskAsCooGradKernel(const Context& dev_ctx, ++ const DenseTensor& x, ++ const SparseCooTensor& mask, ++ const SparseCooTensor& out_grad, ++ DenseTensor* x_grad) { ++ CooToDenseKernel(dev_ctx, out_grad, x_grad); ++} ++ ++template ++void MaskAsCsrGradKernel(const Context& dev_ctx, ++ const DenseTensor& x, ++ const SparseCsrTensor& mask, ++ const SparseCsrTensor& out_grad, ++ DenseTensor* x_grad) { ++ CsrToDenseKernel(dev_ctx, out_grad, x_grad); ++} ++ ++} // namespace sparse ++} // namespace phi +diff --git a/paddle/phi/kernels/sparse/mask_kernel.h b/paddle/phi/kernels/sparse/mask_kernel.h +index 5ffc7fb4aa44d9..5be993e243b193 100644 +--- a/paddle/phi/kernels/sparse/mask_kernel.h ++++ b/paddle/phi/kernels/sparse/mask_kernel.h +@@ -16,21 +16,28 @@ limitations under the License. */ + + #include "paddle/phi/core/dense_tensor.h" + #include "paddle/phi/core/sparse_coo_tensor.h" ++#include "paddle/phi/core/sparse_csr_tensor.h" + + namespace phi { + namespace sparse { + +-template +-void MaskCooKernel(const Context& dev_ctx, +- const DenseTensor& x, +- const SparseCooTensor& mask, +- SparseCooTensor* out); +- + template + void MaskHelperCooKernel(const Context& dev_ctx, + const SparseCooTensor& x, + const DenseTensor& mask_indices, + DenseTensor* out); + ++template ++void MaskAsCooKernel(const Context& dev_ctx, ++ const DenseTensor& x, ++ const SparseCooTensor& mask, ++ SparseCooTensor* out); ++ ++template ++void MaskAsCsrKernel(const Context& dev_ctx, ++ const DenseTensor& x, ++ const SparseCsrTensor& mask, ++ SparseCsrTensor* out); ++ + } // namespace sparse + } // namespace phi +diff --git a/paddle/phi/kernels/sparse/sparse_utils_grad_kernel.cc b/paddle/phi/kernels/sparse/sparse_utils_grad_kernel.cc +index f5915c7acb84ce..2b802615486f42 100644 +--- a/paddle/phi/kernels/sparse/sparse_utils_grad_kernel.cc ++++ b/paddle/phi/kernels/sparse/sparse_utils_grad_kernel.cc +@@ -32,7 +32,7 @@ void CooToDenseGradKernel(const Context& dev_ctx, + const SparseCooTensor& x, + const DenseTensor& out_grad, + SparseCooTensor* x_grad) { +- MaskCooKernel(dev_ctx, out_grad, x, x_grad); ++ MaskAsCooKernel(dev_ctx, out_grad, x, x_grad); + } + + } // namespace sparse +diff --git a/paddle/phi/ops/yaml/sparse_backward.yaml b/paddle/phi/ops/yaml/sparse_backward.yaml +index 3e614b942d3019..f7734af1bf6ecd 100644 +--- a/paddle/phi/ops/yaml/sparse_backward.yaml ++++ b/paddle/phi/ops/yaml/sparse_backward.yaml +@@ -184,6 +184,17 @@ + func : log1p_coo_grad {sparse_coo, sparse_coo -> sparse_coo}, + log1p_csr_grad {sparse_csr, sparse_csr -> sparse_csr} + ++- backward_op : mask_as_grad ++ forward : mask_as(Tensor x, Tensor mask) -> Tensor(out) ++ args : (Tensor x, Tensor mask, Tensor out_grad) ++ output : Tensor(x_grad) ++ infer_meta : ++ func : UnchangedInferMeta ++ param : [x] ++ kernel : ++ func : mask_as_coo_grad {dense, sparse_coo, sparse_coo -> dense}, ++ mask_as_csr_grad {dense, sparse_csr, sparse_csr -> dense} ++ + - backward_op : masked_matmul_grad + forward : masked_matmul(Tensor x, Tensor y, Tensor mask) -> Tensor(out) + args : (Tensor x, Tensor y, Tensor out_grad) +diff --git a/paddle/phi/ops/yaml/sparse_ops.yaml b/paddle/phi/ops/yaml/sparse_ops.yaml +index ac230be485c095..80cef73a6c1f5e 100644 +--- a/paddle/phi/ops/yaml/sparse_ops.yaml ++++ b/paddle/phi/ops/yaml/sparse_ops.yaml +@@ -497,6 +497,18 @@ + func : indices_coo{sparse_coo -> dense} + layout : x + ++- op: mask_as ++ args : (Tensor x, Tensor mask) ++ output : Tensor(out) ++ infer_meta : ++ func : UnchangedInferMeta ++ param : [x] ++ kernel : ++ func : mask_as_coo{dense, sparse_coo -> sparse_coo}, ++ mask_as_csr{dense, sparse_csr -> sparse_csr} ++ layout : x ++ backward: mask_as_grad ++ + - op: masked_matmul + args : (Tensor x, Tensor y, Tensor mask) + output : Tensor(out) +diff --git a/python/paddle/sparse/__init__.py b/python/paddle/sparse/__init__.py +index 661143f12dae8f..98f5ca0b13ee54 100644 +--- a/python/paddle/sparse/__init__.py ++++ b/python/paddle/sparse/__init__.py +@@ -17,6 +17,7 @@ + add, + divide, + is_same_shape, ++ mask_as, + masked_matmul, + matmul, + multiply, +@@ -77,6 +78,7 @@ + 'expm1', + 'mv', + 'matmul', ++ 'mask_as', + 'masked_matmul', + 'addmm', + 'add', +diff --git a/python/paddle/sparse/binary.py b/python/paddle/sparse/binary.py +index 3aac3d5e7f1446..abc943ac3c1fc2 100644 +--- a/python/paddle/sparse/binary.py ++++ b/python/paddle/sparse/binary.py +@@ -452,3 +452,60 @@ def is_same_shape(x, y): + + """ + return x.is_same_shape(y) ++ ++ ++@dygraph_only ++def mask_as(x, mask, name=None): ++ r""" ++ Filter the input dense tensor `x` using the `indices` of the sparse matrix `mask`, ++ which in turn generates a sparse matrix of the corresponding format. ++ The input `x` and `mask` must have the same shape, and the sparse tensor returned has the same indices as `mask` ++ even `zero` values exist in the coresponding indices. ++ ++ Args: ++ x (Tensor): The input tensor. It should be a DenseTensor. ++ The data type can be float32, float64, int32, int64, complex64, complex128, int8, int16, float16. ++ mask (Tensor): The input tensor. It can be SparseCooTensor or SparseCsrTensor. ++ It should be 2D or 3D when the mask is SparseCsrTensor. ++ name (str, optional): Name for the operation (optional, default is None). For more information, please refer to :ref:`api_guide_Name`. ++ ++ Returns: ++ Tensor: A sparse tensor. ++ ++ Examples: ++ .. code-block:: python ++ ++ >>> import paddle ++ >>> paddle.set_device('cpu') ++ ++ >>> # csr sparse tensor ++ >>> crows = [0, 2, 3, 5] ++ >>> cols = [1, 3, 2, 0, 1] ++ >>> values = [1., 2., 3., 4., 5.] ++ >>> dense_shape = [3, 4] ++ >>> csr = paddle.sparse.sparse_csr_tensor(crows, cols, values, dense_shape) ++ >>> paddle.seed(2024) ++ >>> x = paddle.rand(dense_shape).astype(csr.dtype) ++ >>> out = paddle.sparse.mask_as(x, csr) ++ >>> print(out) ++ Tensor(shape=[3, 4], dtype=paddle.float32, place=Place(cpu), stop_gradient=True, ++ crows=[0, 2, 3, 5], ++ cols=[1, 3, 2, 0, 1], ++ values=[0.23659813, 0.08467803, 0.64152628, 0.66596609, 0.90394485]) ++ ++ >>> # coo sparse tensor ++ >>> indices = [[0, 1, 2], [1, 2, 0]] ++ >>> values = [1.0, 2.0, 3.0] ++ >>> dense_shape = [3, 3] ++ >>> coo = paddle.sparse.sparse_coo_tensor(indices, values, dense_shape) ++ >>> paddle.seed(2024) ++ >>> x = paddle.rand(dense_shape).astype(coo.dtype) ++ >>> out = paddle.sparse.mask_as(x, coo) ++ >>> print(out) ++ Tensor(shape=[3, 3], dtype=paddle.float32, place=Place(cpu), stop_gradient=True, ++ indices=[[0, 1, 2], ++ [1, 2, 0]], ++ values=[0.23659813, 0.40340215, 0.64152628]) ++ ++ """ ++ return _C_ops.sparse_mask_as(x, mask) diff --git a/swe-paddle/tasks/PaddlePaddle__Paddle-64320/tests/test.patch b/swe-paddle/tasks/PaddlePaddle__Paddle-64320/tests/test.patch new file mode 100644 index 000000000..9cbfe9c34 --- /dev/null +++ b/swe-paddle/tasks/PaddlePaddle__Paddle-64320/tests/test.patch @@ -0,0 +1,176 @@ +diff --git a/test/legacy_test/CMakeLists.txt b/test/legacy_test/CMakeLists.txt +index f84458dd494f30..33aa88c3a7c516 100644 +--- a/test/legacy_test/CMakeLists.txt ++++ b/test/legacy_test/CMakeLists.txt +@@ -1131,6 +1131,7 @@ set_pir_tests_properties() + set_tests_properties(test_nadam_op PROPERTIES TIMEOUT 100) + set_tests_properties(test_radam_op PROPERTIES TIMEOUT 100) + set_tests_properties(test_nan_inf PROPERTIES TIMEOUT 120) ++set_tests_properties(test_sparse_mask_as_op PROPERTIES TIMEOUT 120) + set_tests_properties(test_bicubic_interp_op PROPERTIES TIMEOUT 120) + set_tests_properties(test_bilinear_interp_op PROPERTIES TIMEOUT 120) + set_tests_properties(test_conv2d_op_depthwise_conv +diff --git a/test/legacy_test/test_sparse_mask_as_op.py b/test/legacy_test/test_sparse_mask_as_op.py +new file mode 100644 +index 00000000000000..f4cd639452b5d3 +--- /dev/null ++++ b/test/legacy_test/test_sparse_mask_as_op.py +@@ -0,0 +1,159 @@ ++# Copyright (c) 2024 PaddlePaddle Authors. All Rights Reserved. ++# ++# Licensed under the Apache License, Version 2.0 (the "License"); ++# you may not use this file except in compliance with the License. ++# You may obtain a copy of the License at ++# ++# http://www.apache.org/licenses/LICENSE-2.0 ++# ++# Unless required by applicable law or agreed to in writing, software ++# distributed under the License is distributed on an "AS IS" BASIS, ++# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. ++# See the License for the specific language governing permissions and ++# limitations under the License. ++ ++import unittest ++ ++import numpy as np ++ ++import paddle ++ ++ ++def generate_data(shape, dtype): ++ """ ++ Generate `data` and `mask` with the same shape and dtype. ++ """ ++ _mask = np.random.randint(0, 2, shape) ++ if np.sum(_mask) == 0: ++ _mask.flat[0] = 1 ++ mask = (np.random.randint(-100, 100, shape) * _mask).astype(dtype) ++ data = np.random.randint(-100, 100, shape).astype(dtype) ++ return data, mask ++ ++ ++class TestMaskAs(unittest.TestCase): ++ def setUp(self): ++ self.init_format() ++ self.places = [paddle.CPUPlace()] ++ if paddle.is_compiled_with_cuda(): ++ self.places.append(paddle.CUDAPlace(0)) ++ ++ def init_format(self): ++ self.format = None ++ ++ def check(self, shape, dtype, place, check_grad=True): ++ paddle.disable_static() ++ dense_data_np, dense_mask_np = generate_data(shape, dtype) ++ ++ dense_data_pd = paddle.to_tensor( ++ dense_data_np, dtype=dtype, place=place ++ ) ++ dense_data_pd.stop_gradient = False ++ ++ if self.format == 'coo': ++ sparse_mask_pd = paddle.to_tensor( ++ dense_mask_np, dtype=dtype, place=place ++ ).to_sparse_coo(len(shape)) ++ else: ++ sparse_mask_pd = paddle.to_tensor( ++ dense_mask_np, dtype=dtype, place=place ++ ).to_sparse_csr() ++ ++ sparse_out_pd = paddle.sparse.mask_as(dense_data_pd, sparse_mask_pd) ++ ++ # compare the tensor from sparse->dense with reference numpy data ++ # the result only keeps the values where mask not zero, like: ++ # dense_data_np ++ # [[ 38. 15. 76.] ++ # [-98. -75. 10.] ++ # [-52. 49. -48.]] ++ # dense_mask_np ++ # [[-70. 0. 0.] ++ # [-50. 34. 60.] ++ # [-34. 0. -18.]] ++ # dense_data_np_ref ++ # [[ 38. 0. 0.] ++ # [-98. -75. 10.] ++ # [-52. 0. -48.]] ++ dense_data_np_ref = dense_data_np * (dense_mask_np != 0) ++ np.testing.assert_allclose( ++ sparse_out_pd.to_dense().numpy(), dense_data_np_ref ++ ) ++ ++ if check_grad: ++ # with sparse_out_pd backward, we get the grad from dense_data_pd ++ sparse_out_pd.backward() ++ dense_data_grad = dense_data_pd.grad ++ ++ self.assertEqual( ++ list(dense_data_grad.shape), list(dense_data_pd.shape) ++ ) ++ self.assertEqual(dense_data_grad.dtype, dense_data_pd.dtype) ++ ++ # make a dense data to compare the grad from sparse_out_pd ++ grad_ref = np.ones_like(dense_mask_np) * (dense_mask_np != 0) ++ ++ np.testing.assert_allclose( ++ dense_data_pd.grad.numpy(), ++ grad_ref, ++ ) ++ ++ def check_with_dtypes(self, shape): ++ for place in self.places: ++ self.check(shape, 'float32', place) ++ self.check(shape, 'float64', place) ++ self.check(shape, 'int32', place) ++ self.check(shape, 'int64', place) ++ self.check(shape, 'complex64', place) ++ self.check(shape, 'complex128', place) ++ ++ # `int8`` not registered in `FullLikeCooKernel`, so skip check_grad ++ self.check(shape, 'int8', place, check_grad=False) ++ ++ # `int16` not registered in `multiply`, so skip check_grad ++ self.check(shape, 'int16', place, check_grad=False) ++ ++ if paddle.is_compiled_with_cuda(): ++ place = paddle.CUDAPlace(0) ++ self.check(shape, 'float16', place) ++ ++ ++class TestMaskAsCoo(TestMaskAs): ++ def init_format(self): ++ self.format = 'coo' ++ ++ def test_1d(self): ++ self.check_with_dtypes((5,)) ++ ++ def test_2d(self): ++ self.check_with_dtypes((5, 3)) ++ ++ def test_3d(self): ++ self.check_with_dtypes((5, 3, 4)) ++ ++ def test_4d(self): ++ self.check_with_dtypes((5, 3, 4, 2)) ++ ++ ++class TestMaskAsCsr(TestMaskAs): ++ def init_format(self): ++ self.format = 'csr' ++ ++ def test_2d(self): ++ self.check_with_dtypes((5, 3)) ++ ++ def test_3d(self): ++ self.check_with_dtypes((5, 3, 4)) ++ ++ def test_error_dimension(self): ++ # error 1d ++ with self.assertRaises(ValueError): ++ self.check_with_dtypes((5,)) ++ ++ # error 4d ++ with self.assertRaises(ValueError): ++ self.check_with_dtypes((5, 3, 4, 2)) ++ ++ ++if __name__ == "__main__": diff --git a/swe-paddle/tasks/PaddlePaddle__Paddle-64320/tests/test.sh b/swe-paddle/tasks/PaddlePaddle__Paddle-64320/tests/test.sh new file mode 100644 index 000000000..b40127d4e --- /dev/null +++ b/swe-paddle/tasks/PaddlePaddle__Paddle-64320/tests/test.sh @@ -0,0 +1,4 @@ +#!/usr/bin/env bash +set -euo pipefail + +python -m pytest test/legacy_test/test_sparse_mask_as_op.py -q