Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
milvus
milvus
提交
75ca0d96
M
milvus
项目概览
milvus
/
milvus
12 个月 前同步成功
通知
261
Star
22476
Fork
2472
代码
文件
提交
分支
Tags
贡献者
分支图
Diff
Issue
0
列表
看板
标记
里程碑
合并请求
0
DevOps
流水线
流水线任务
计划
Wiki
0
Wiki
分析
仓库
DevOps
项目成员
Pages
M
milvus
项目概览
项目概览
详情
发布
仓库
仓库
文件
提交
分支
标签
贡献者
分支图
比较
Issue
0
Issue
0
列表
看板
标记
里程碑
合并请求
0
合并请求
0
Pages
DevOps
DevOps
流水线
流水线任务
计划
分析
分析
仓库分析
DevOps
Wiki
0
Wiki
成员
成员
收起侧边栏
关闭侧边栏
动态
分支图
创建新Issue
流水线任务
提交
Issue看板
体验新版 GitCode,发现更多精彩内容 >>
未验证
提交
75ca0d96
编写于
4月 21, 2020
作者:
M
Morry Niu
提交者:
GitHub
4月 21, 2020
浏览文件
操作
浏览文件
下载
电子邮件补丁
差异文件
keep bitset naming consistent (#2004)
Signed-off-by:
N
Xinyao Niu
<
morryniu123@gmail.com
>
上级
55cc6be6
变更
2
隐藏空白更改
内联
并排
Showing
2 changed file
with
11 addition
and
11 deletion
+11
-11
core/src/index/thirdparty/faiss/gpu/impl/L2Select.cu
core/src/index/thirdparty/faiss/gpu/impl/L2Select.cu
+5
-5
core/src/index/thirdparty/faiss/gpu/utils/BlockSelectKernel.cuh
...rc/index/thirdparty/faiss/gpu/utils/BlockSelectKernel.cuh
+6
-6
未找到文件。
core/src/index/thirdparty/faiss/gpu/impl/L2Select.cu
浏览文件 @
75ca0d96
...
...
@@ -45,7 +45,7 @@ __global__ void l2SelectMin1(Tensor<T, 2, true> productDistances,
// FIXME: if we have exact multiples, don't need this
bool
endRow
=
(
blockIdx
.
x
==
gridDim
.
x
-
1
);
bool
bitset
Is
Empty
=
(
bitset
.
getSize
(
0
)
==
0
);
bool
bitsetEmpty
=
(
bitset
.
getSize
(
0
)
==
0
);
if
(
endRow
)
{
if
(
productDistances
.
getSize
(
0
)
%
kRowsPerBlock
==
0
)
{
...
...
@@ -57,7 +57,7 @@ __global__ void l2SelectMin1(Tensor<T, 2, true> productDistances,
for
(
int
row
=
rowStart
;
row
<
productDistances
.
getSize
(
0
);
++
row
)
{
for
(
int
col
=
threadIdx
.
x
;
col
<
productDistances
.
getSize
(
1
);
col
+=
blockDim
.
x
)
{
if
(
bitset
Is
Empty
||
(
!
(
bitset
[
col
>>
3
]
&
(
0x1
<<
(
col
&
0x7
)))))
{
if
(
bitsetEmpty
||
(
!
(
bitset
[
col
>>
3
]
&
(
0x1
<<
(
col
&
0x7
)))))
{
distance
[
0
]
=
Math
<
T
>::
add
(
centroidDistances
[
col
],
productDistances
[
row
][
col
]);
}
else
{
...
...
@@ -149,11 +149,11 @@ __global__ void l2SelectMinK(Tensor<T, 2, true> productDistances,
int
limit
=
utils
::
roundDown
(
productDistances
.
getSize
(
1
),
kWarpSize
);
int
i
=
threadIdx
.
x
;
bool
bitset
Is
Empty
=
(
bitset
.
getSize
(
0
)
==
0
);
bool
bitsetEmpty
=
(
bitset
.
getSize
(
0
)
==
0
);
T
v
;
for
(;
i
<
limit
;
i
+=
blockDim
.
x
)
{
if
(
bitset
Is
Empty
||
(
!
(
bitset
[
i
>>
3
]
&
(
0x1
<<
(
i
&
0x7
)))))
{
if
(
bitsetEmpty
||
(
!
(
bitset
[
i
>>
3
]
&
(
0x1
<<
(
i
&
0x7
)))))
{
v
=
Math
<
T
>::
add
(
centroidDistances
[
i
],
productDistances
[
row
][
i
]);
}
else
{
...
...
@@ -164,7 +164,7 @@ __global__ void l2SelectMinK(Tensor<T, 2, true> productDistances,
}
if
(
i
<
productDistances
.
getSize
(
1
))
{
if
(
bitset
Is
Empty
||
(
!
(
bitset
[
i
>>
3
]
&
(
0x1
<<
(
i
&
0x7
)))))
{
if
(
bitsetEmpty
||
(
!
(
bitset
[
i
>>
3
]
&
(
0x1
<<
(
i
&
0x7
)))))
{
v
=
Math
<
T
>::
add
(
centroidDistances
[
i
],
productDistances
[
row
][
i
]);
}
else
{
...
...
core/src/index/thirdparty/faiss/gpu/utils/BlockSelectKernel.cuh
浏览文件 @
75ca0d96
...
...
@@ -142,10 +142,10 @@ __global__ void blockSelect(Tensor<K, 2, true> in,
// Whole warps must participate in the selection
int
limit
=
utils
::
roundDown
(
in
.
getSize
(
1
),
kWarpSize
);
bool
bitset
Is
Empty
=
(
bitset
.
getSize
(
0
)
==
0
);
bool
bitsetEmpty
=
(
bitset
.
getSize
(
0
)
==
0
);
for
(;
i
<
limit
;
i
+=
ThreadsPerBlock
)
{
if
(
bitset
Is
Empty
||
(
!
(
bitset
[
i
>>
3
]
&
(
0x1
<<
(
i
&
0x7
)))))
{
if
(
bitsetEmpty
||
(
!
(
bitset
[
i
>>
3
]
&
(
0x1
<<
(
i
&
0x7
)))))
{
heap
.
add
(
*
inStart
,
(
IndexType
)
i
);
}
else
{
heap
.
add
(
-
1.0
,
(
IndexType
)
i
);
...
...
@@ -156,7 +156,7 @@ __global__ void blockSelect(Tensor<K, 2, true> in,
// Handle last remainder fraction of a warp of elements
if
(
i
<
in
.
getSize
(
1
))
{
if
(
bitset
Is
Empty
||
(
!
(
bitset
[
i
>>
3
]
&
(
0x1
<<
(
i
&
0x7
)))))
{
if
(
bitsetEmpty
||
(
!
(
bitset
[
i
>>
3
]
&
(
0x1
<<
(
i
&
0x7
)))))
{
heap
.
addThreadQ
(
*
inStart
,
(
IndexType
)
i
);
}
else
{
heap
.
addThreadQ
(
-
1.0
,
(
IndexType
)
i
);
...
...
@@ -204,10 +204,10 @@ __global__ void blockSelectPair(Tensor<K, 2, true> inK,
// Whole warps must participate in the selection
int
limit
=
utils
::
roundDown
(
inK
.
getSize
(
1
),
kWarpSize
);
bool
bitset
Is
Empty
=
(
bitset
.
getSize
(
0
)
==
0
);
bool
bitsetEmpty
=
(
bitset
.
getSize
(
0
)
==
0
);
for
(;
i
<
limit
;
i
+=
ThreadsPerBlock
)
{
if
(
bitset
Is
Empty
||
(
!
(
bitset
[
i
>>
3
]
&
(
0x1
<<
(
i
&
0x7
)))))
{
if
(
bitsetEmpty
||
(
!
(
bitset
[
i
>>
3
]
&
(
0x1
<<
(
i
&
0x7
)))))
{
heap
.
add
(
*
inKStart
,
*
inVStart
);
}
else
{
heap
.
add
(
-
1.0
,
*
inVStart
);
...
...
@@ -219,7 +219,7 @@ __global__ void blockSelectPair(Tensor<K, 2, true> inK,
// Handle last remainder fraction of a warp of elements
if
(
i
<
inK
.
getSize
(
1
))
{
if
(
bitset
Is
Empty
||
(
!
(
bitset
[
i
>>
3
]
&
(
0x1
<<
(
i
&
0x7
)))))
{
if
(
bitsetEmpty
||
(
!
(
bitset
[
i
>>
3
]
&
(
0x1
<<
(
i
&
0x7
)))))
{
heap
.
addThreadQ
(
*
inKStart
,
*
inVStart
);
}
else
{
heap
.
addThreadQ
(
-
1.0
,
*
inVStart
);
...
...
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录