Skip to content

Commit 3e21725

Browse files
zhuyuegongchensu
zhuyue
authored and
gongchensu
committed
Fix format errors.
1 parent 2ffd9f5 commit 3e21725

File tree

9 files changed

+183
-210
lines changed

9 files changed

+183
-210
lines changed

‎src/infiniop-test/include/ops.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -44,7 +44,7 @@ DECLARE_INFINIOP_TEST(all_equal)
4444
REGISTER_INFINIOP_TEST(causal_softmax) \
4545
REGISTER_INFINIOP_TEST(rearrange) \
4646
REGISTER_INFINIOP_TEST(sub) \
47-
REGISTER_INFINIOP_TEST(all_equal) \
47+
REGISTER_INFINIOP_TEST(all_equal) \
4848
}
4949

5050
namespace infiniop_test {

‎src/infiniop-test/src/ops/all_equal.cpp‎

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -1,9 +1,9 @@
1+
#include "../../../include/infiniop/ops/all_equal.h"
12
#include "ops.hpp"
23
#include "utils.hpp"
34
#include <infinirt.h>
45
#include <iomanip>
56
#include <iostream>
6-
#include "../../../include/infiniop/ops/all_equal.h"
77

88
namespace infiniop_test::all_equal {
99
struct Test::Attributes {
@@ -41,9 +41,9 @@ std::shared_ptr<infiniop_test::Result> Test::run(
4141
auto b = _attributes->b->to(device, device_id);
4242
auto c = _attributes->c->to(device, device_id);
4343
CHECK_OR(infiniopCreateAllEqualDescriptor(handle, &op_desc,
44-
c->desc(),
45-
a->desc(),
46-
b->desc()),
44+
c->desc(),
45+
a->desc(),
46+
b->desc()),
4747
return TEST_FAILED(OP_CREATION_FAILED, "Failed to create op descriptor."));
4848
size_t workspace_size;
4949
CHECK_OR(infiniopGetAllEqualWorkspaceSize(op_desc, &workspace_size),
@@ -52,10 +52,10 @@ std::shared_ptr<infiniop_test::Result> Test::run(
5252
CHECK_OR(infinirtMalloc(&workspace, workspace_size),
5353
return TEST_FAILED(OP_CREATION_FAILED, "Failed to allocate workspace."));
5454
CHECK_OR(infiniopAllEqual(op_desc, workspace, workspace_size,
55-
c->data(),
56-
a->data(),
57-
b->data(),
58-
nullptr),
55+
c->data(),
56+
a->data(),
57+
b->data(),
58+
nullptr),
5959
return TEST_FAILED(OP_EXECUTION_FAILED, "Failed during execution."));
6060

6161
try {

‎src/infiniop/ops/all_equal/all_equal.h‎

Lines changed: 35 additions & 37 deletions
Original file line numberDiff line numberDiff line change
@@ -6,43 +6,41 @@
66
#include "../../tensor.h"
77
#include "info.h"
88

9-
#define DESCRIPTOR(NAMESPACE) \
10-
namespace op::all_equal::NAMESPACE { \
11-
class Descriptor final : public InfiniopDescriptor { \
12-
struct Opaque; \
13-
Opaque *_opaque; \
14-
op::all_equal::AllEqualInfo _info; \
15-
size_t _workspace_size; \
16-
Descriptor( \
17-
infiniDtype_t dtype, \
18-
op::all_equal::AllEqualInfo info, \
19-
size_t workspace_size_, \
20-
Opaque *opaque, \
21-
infiniDevice_t device_type, \
22-
int device_id \
23-
) : InfiniopDescriptor{device_type, device_id}, \
24-
_opaque(opaque), \
25-
_info(info), \
26-
_workspace_size(workspace_size_) {} \
27-
public: \
28-
~Descriptor(); \
29-
size_t workspaceSize() const { return _workspace_size; } \
30-
static infiniStatus_t create( \
31-
infiniopHandle_t handle, \
32-
Descriptor **desc_ptr, \
33-
infiniopTensorDescriptor_t c_desc, \
34-
infiniopTensorDescriptor_t a_desc, \
35-
infiniopTensorDescriptor_t b_desc \
36-
); \
37-
infiniStatus_t calculate( \
38-
void *workspace, \
39-
size_t workspace_size, \
40-
void * c, \
41-
const void * a, \
42-
const void * b, \
43-
void *stream \
44-
) const; \
45-
}; \
9+
#define DESCRIPTOR(NAMESPACE) \
10+
namespace op::all_equal::NAMESPACE { \
11+
class Descriptor final : public InfiniopDescriptor { \
12+
struct Opaque; \
13+
Opaque *_opaque; \
14+
op::all_equal::AllEqualInfo _info; \
15+
size_t _workspace_size; \
16+
Descriptor( \
17+
infiniDtype_t dtype, \
18+
op::all_equal::AllEqualInfo info, \
19+
size_t workspace_size_, \
20+
Opaque *opaque, \
21+
infiniDevice_t device_type, \
22+
int device_id) : InfiniopDescriptor{device_type, device_id}, \
23+
_opaque(opaque), \
24+
_info(info), \
25+
_workspace_size(workspace_size_) {} \
26+
\
27+
public: \
28+
~Descriptor(); \
29+
size_t workspaceSize() const { return _workspace_size; } \
30+
static infiniStatus_t create( \
31+
infiniopHandle_t handle, \
32+
Descriptor **desc_ptr, \
33+
infiniopTensorDescriptor_t c_desc, \
34+
infiniopTensorDescriptor_t a_desc, \
35+
infiniopTensorDescriptor_t b_desc); \
36+
infiniStatus_t calculate( \
37+
void *workspace, \
38+
size_t workspace_size, \
39+
void *c, \
40+
const void *a, \
41+
const void *b, \
42+
void *stream) const; \
43+
}; \
4644
}
4745

4846
#endif

‎src/infiniop/ops/all_equal/cpu/all_equal_cpu.cc‎

Lines changed: 20 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -12,62 +12,56 @@ infiniStatus_t Descriptor::create(
1212
Descriptor **desc_ptr,
1313
infiniopTensorDescriptor_t c_desc,
1414
infiniopTensorDescriptor_t a_desc,
15-
infiniopTensorDescriptor_t b_desc
16-
) {
15+
infiniopTensorDescriptor_t b_desc) {
1716
auto handle = reinterpret_cast<device::cpu::Handle *>(handle_);
1817

19-
// --------------------- start: check data type and calculate workspace size ----------------------
18+
// --------------------- start: check data type and calculate workspace size ----------------------
2019
auto dtype = c_desc->dtype();
2120
CHECK_DTYPE(dtype, INFINI_DTYPE_BOOL);
2221
CHECK_OR_RETURN(b_desc->dtype() == a_desc->dtype(), INFINI_STATUS_BAD_TENSOR_DTYPE);
2322
size_t WorkSpaceSize = 0;
24-
// ---------------------- end: check data type and calculate workspace size -----------------------
23+
// ---------------------- end: check data type and calculate workspace size -----------------------
2524

2625
auto result = AllEqualInfo::createAllEqualInfo(
2726
c_desc,
2827
a_desc,
29-
b_desc
30-
);
28+
b_desc);
3129
CHECK_RESULT(result);
3230
const AllEqualInfo &info = result.take();
33-
31+
3432
*desc_ptr = new Descriptor(
3533
dtype, std::move(info), WorkSpaceSize,
3634
nullptr,
37-
handle->device, handle->device_id
38-
);
35+
handle->device, handle->device_id);
3936

4037
return INFINI_STATUS_SUCCESS;
4138
}
4239

43-
4440
infiniStatus_t Descriptor::calculate(
4541
void *workspace,
4642
size_t workspace_size,
47-
void * c,
48-
const void * a,
49-
const void * b,
50-
void *stream
51-
) const {
43+
void *c,
44+
const void *a,
45+
const void *b,
46+
void *stream) const {
5247
std::vector<ptrdiff_t> contiguous_strides(_info.ndim);
53-
ptrdiff_t last_dim = 1;
48+
ptrdiff_t last_dim = 1;
5449
ptrdiff_t last_stride = 1;
55-
for(size_t d = 0; d < _info.ndim; d ++)
56-
{
57-
contiguous_strides[d] = last_dim * last_stride;
50+
for (size_t d = 0; d < _info.ndim; d++) {
51+
contiguous_strides[d] = last_dim * last_stride;
5852
last_dim = _info.a_shape[d];
5953
last_stride = contiguous_strides[d];
6054
}
6155
size_t total_size = last_dim * last_stride;
6256
size_t elem_size = infiniSizeOf(_info.dtype);
63-
auto c_ptr = reinterpret_cast<bool*>(c);
57+
auto c_ptr = reinterpret_cast<bool *>(c);
6458
*c_ptr = true;
65-
#pragma omp parallel for
66-
for(size_t i = 0; i < total_size; i ++) {
67-
auto a_ptr = reinterpret_cast<const char*>(a);
68-
auto b_ptr = reinterpret_cast<const char*>(b);
59+
#pragma omp parallel for
60+
for (size_t i = 0; i < total_size; i++) {
61+
auto a_ptr = reinterpret_cast<const char *>(a);
62+
auto b_ptr = reinterpret_cast<const char *>(b);
6963
size_t rem = i;
70-
for(int d = _info.ndim - 1; d >= 0; d --) {
64+
for (int d = _info.ndim - 1; d >= 0; d--) {
7165
size_t dim_index = rem / contiguous_strides[d];
7266
rem = rem % contiguous_strides[d];
7367
a_ptr += dim_index * _info.a_strides[d];
@@ -79,4 +73,4 @@ infiniStatus_t Descriptor::calculate(
7973
}
8074
return INFINI_STATUS_SUCCESS;
8175
}
82-
}
76+
} // namespace op::all_equal::cpu

‎src/infiniop/ops/all_equal/cpu/all_equal_cpu.h‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -5,5 +5,4 @@
55

66
DESCRIPTOR(cpu)
77

8-
98
#endif // __ALL_EQUAL_CPU_H__

‎src/infiniop/ops/all_equal/cuda/kernel.cuh‎

Lines changed: 15 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -3,48 +3,47 @@
33
// ------------------------------- start: perform operator on CUDA --------------------------------
44
template <unsigned int BLOCK_SIZE, typename Tdata>
55
__device__ void allEqualKernel(
6-
bool * c,
7-
const Tdata * a,
8-
const Tdata * b,
6+
bool *c,
7+
const Tdata *a,
8+
const Tdata *b,
99
size_t ndim,
1010
size_t total_size,
11-
ptrdiff_t* contiguous_strides,
12-
ptrdiff_t* a_strides,
13-
ptrdiff_t* b_strides
14-
) {
11+
ptrdiff_t *contiguous_strides,
12+
ptrdiff_t *a_strides,
13+
ptrdiff_t *b_strides) {
1514
// 使用共享内存来避免竞态条件
1615
__shared__ bool block_result;
17-
16+
1817
if (threadIdx.x == 0) {
1918
block_result = true;
2019
}
2120
__syncthreads();
22-
21+
2322
// 每个线程检查自己负责的元素
2423
bool thread_result = true;
25-
for(size_t i = threadIdx.x; i < total_size; i += BLOCK_SIZE) {
24+
for (size_t i = threadIdx.x; i < total_size; i += BLOCK_SIZE) {
2625
auto a_ptr = a;
2726
auto b_ptr = b;
2827
size_t rem = i;
29-
for(int d = ndim - 1; d >= 0; d --) {
28+
for (int d = ndim - 1; d >= 0; d--) {
3029
size_t dim_index = rem / contiguous_strides[d];
3130
rem = rem % contiguous_strides[d];
3231
a_ptr += dim_index * a_strides[d];
3332
b_ptr += dim_index * b_strides[d];
3433
}
3534
if (*a_ptr != *b_ptr) {
3635
thread_result = false;
37-
break; // 发现不匹配，提前退出
36+
break; // 发现不匹配，提前退出
3837
}
3938
}
40-
39+
4140
// 使用原子操作来安全地更新结果
4241
if (!thread_result) {
43-
atomicAnd((int*)&block_result, 0);
42+
atomicAnd((int *)&block_result, 0);
4443
}
45-
44+
4645
__syncthreads();
47-
46+
4847
// 只有第一个线程写入最终结果
4948
if (threadIdx.x == 0) {
5049
*c = block_result;

‎src/infiniop/ops/all_equal/info.h‎

Lines changed: 10 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -12,35 +12,34 @@ class AllEqualInfo {
1212
AllEqualInfo() = default;
1313

1414
public:
15-
// ---------------------------- start: define member variables of Info ----------------------------
16-
size_t ndim;
15+
// ---------------------------- start: define member variables of Info ----------------------------
16+
size_t ndim;
1717
infiniDtype_t dtype;
1818
std::vector<size_t> a_shape;
1919
std::vector<ptrdiff_t> a_strides;
2020
std::vector<ptrdiff_t> b_strides;
2121

22-
// ----------------------------- end: define member variables of Info -----------------------------
22+
// ----------------------------- end: define member variables of Info -----------------------------
2323

2424
static utils::Result<AllEqualInfo> createAllEqualInfo(
2525
infiniopTensorDescriptor_t c_desc,
2626
infiniopTensorDescriptor_t a_desc,
27-
infiniopTensorDescriptor_t b_desc
28-
) {
29-
// ------------------------- start: check tensor shape and input validity -------------------------
27+
infiniopTensorDescriptor_t b_desc) {
28+
// ------------------------- start: check tensor shape and input validity -------------------------
3029
CHECK_OR_RETURN(c_desc->ndim() == 1 && c_desc->dim(0) == 1, INFINI_STATUS_BAD_TENSOR_SHAPE);
3130
CHECK_SAME_SHAPE(a_desc->shape(), b_desc->shape());
32-
// -------------------------- end: check tensor shape and input validity --------------------------
31+
// -------------------------- end: check tensor shape and input validity --------------------------
3332
return utils::Result<AllEqualInfo>(AllEqualInfo{
34-
// ------------------------------ start: create an instance of Info -------------------------------
35-
a_desc->ndim(),
33+
// ------------------------------ start: create an instance of Info -------------------------------
34+
a_desc->ndim(),
3635
a_desc->dtype(),
3736
a_desc->shape(),
3837
a_desc->strides(),
3938
b_desc->strides()
40-
// ------------------------------- end: create an instance of Info --------------------------------
39+
// ------------------------------- end: create an instance of Info --------------------------------
4140
});
4241
}
4342
};
44-
}
43+
} // namespace op::all_equal
4544

4645
#endif // __EQUAL_INFO_H__

0 commit comments

Comments
 (0)