Fix issue 3542: fixed opencl, oneapi, and cuda sparse coo2dense launc… · arrayfire/arrayfire@989b71b · GitHub

1. arrayfire
2. arrayfire

Repository navigation

• Code
• Issues294 (294)
• Pull requests43 (43)
• Actions
• Projects
• Wiki
• Security and quality
• Insights

Commit 989b71b

Browse filesBrowse the repository at this point in the historyBrowse files

edwinsolisf

authored

Fix issue 3542: fixed opencl, oneapi, and cuda sparse coo2dense launch dimensions (#3583)

1 parent 90b27ac commit 989b71b
Copy full SHA for 989b71b

5 files changed

+21-28Lines changed: 21 additions & 28 deletions

File tree

Expand file treeCollapse file tree

Open diff view settings

Filter options

• src/backend
  • cuda/kernel
    • sparse.cuh
    • sparse.hpp
  • oneapi/kernel
    • sparse.hpp
  • opencl/kernel
    • coo2dense.cl
    • sparse.hpp

Expand file treeCollapse file tree

Open diff view settings

Collapse file

‎src/backend/cuda/kernel/sparse.cuh‎

Copy file name to clipboardExpand all lines: src/backend/cuda/kernel/sparse.cuh

+6-8Lines changed: 6 additions & 8 deletions

Original file line number	Diff line number	Diff line change
@@ -17,15 +17,13 @@ namespace cuda {
17	17	template<typename T>
18	18	__global__ void coo2Dense(Param<T> output, CParam<T> values, CParam<int> rowIdx,
19	19	CParam<int> colIdx) {
20		- int id = blockIdx.x * blockDim.x * reps + threadIdx.x;
21		- if (id >= values.dims[0]) return;
	20	+ for (int i = threadIdx.x; i < reps * blockDim.x; i += blockDim.x) {
	21	+ int id = i + blockIdx.x * blockDim.x * reps;
	22	+ if (id >= values.dims[0]) return;
22	23
23		- for (int i = threadIdx.x; i <= reps * blockDim.x; i += blockDim.x) {
24		- if (i >= values.dims[0]) return;
25		-
26		- T v = values.ptr[i];
27		- int r = rowIdx.ptr[i];
28		- int c = colIdx.ptr[i];
	24	+ T v = values.ptr[id];
	25	+ int r = rowIdx.ptr[id];
	26	+ int c = colIdx.ptr[id];
29	27
30	28	int offset = r + c * output.strides[1];
31	29

Collapse file

‎src/backend/cuda/kernel/sparse.hpp‎

Copy file name to clipboardExpand all lines: src/backend/cuda/kernel/sparse.hpp

+1-1Lines changed: 1 addition & 1 deletion

Original file line number	Diff line number	Diff line change
@@ -30,7 +30,7 @@ void coo2dense(Param<T> output, CParam<T> values, CParam<int> rowIdx,
30	30
31	31	dim3 threads(256, 1, 1);
32	32
33		- dim3 blocks(divup(output.dims[0], threads.x * reps), 1, 1);
	33	+ dim3 blocks(divup(values.dims[0], threads.x * reps), 1, 1);
34	34
35	35	EnqueueArgs qArgs(blocks, threads, getActiveStream());
36	36

Collapse file

‎src/backend/oneapi/kernel/sparse.hpp‎

Copy file name to clipboardExpand all lines: src/backend/oneapi/kernel/sparse.hpp

+7-10Lines changed: 7 additions & 10 deletions

Original file line number	Diff line number	Diff line change
@@ -47,19 +47,16 @@ class coo2DenseCreateKernel {
47	47	void operator()(sycl::nd_item<2> it) const {
48	48	sycl::group g = it.get_group();
49	49
50		- const int id = g.get_group_id(0) * g.get_local_range(0) * REPEAT +
51		- it.get_local_id(0);
52		-
53		- if (id >= values_.dims[0]) return;
54		-
55	50	const int dimSize = g.get_local_range(0);
56	51
57	52	for (int i = it.get_local_id(0); i < REPEAT * dimSize; i += dimSize) {
58		- if (i >= values_.dims[0]) return;
	53	+ const int id =
	54	+ g.get_group_id(0) * g.get_local_range(0) * REPEAT + i;
	55	+ if (id >= values_.dims[0]) return;
59	56
60		- T v = vPtr_[i];
61		- int r = rPtr_[i];
62		- int c = cPtr_[i];
	57	+ T v = vPtr_[id];
	58	+ int r = rPtr_[id];
	59	+ int c = cPtr_[id];
63	60
64	61	int offset = r + c * output_.strides[1];
65	62
@@ -83,7 +80,7 @@ void coo2dense(Param<T> out, const Param<T> values, const Param<int> rowIdx,
83	80	const Param<int> colIdx) {
84	81	auto local = sycl::range(THREADS_PER_BLOCK, 1);
85	82	auto global = sycl::range(
86		- divup(out.info.dims[0], local[0] * REPEAT) * THREADS_PER_BLOCK, 1);
	83	+ divup(values.info.dims[0], local[0] * REPEAT) * THREADS_PER_BLOCK, 1);
87	84
88	85	getQueue().submit([&](auto &h) {
89	86	sycl::accessor d_rowIdx{*rowIdx.data, h, sycl::read_only};

Collapse file

‎src/backend/opencl/kernel/coo2dense.cl‎

Copy file name to clipboardExpand all lines: src/backend/opencl/kernel/coo2dense.cl

+5-8Lines changed: 5 additions & 8 deletions

Original file line number	Diff line number	Diff line change
@@ -11,18 +11,15 @@ kernel void coo2Dense(global T *oPtr, const KParam output, global const T *vPtr,
11	11	const KParam values, global const int *rPtr,
12	12	const KParam rowIdx, global const int *cPtr,
13	13	const KParam colIdx) {
14		- const int id = get_group_id(0) * get_local_size(0) * reps + get_local_id(0);
15		-
16		- if (id >= values.dims[0]) return;
17		-
18	14	const int dimSize = get_local_size(0);
19	15
20	16	for (int i = get_local_id(0); i < reps * dimSize; i += dimSize) {
21		- if (i >= values.dims[0]) return;
	17	+ const int id = i + get_group_id(0) * dimSize * reps;
	18	+ if (id >= values.dims[0]) return;
22	19
23		- T v = vPtr[i];
24		- int r = rPtr[i];
25		- int c = cPtr[i];
	20	+ T v = vPtr[id];
	21	+ int r = rPtr[id];
	22	+ int c = cPtr[id];
26	23
27	24	int offset = r + c * output.strides[1];
28	25

Collapse file

‎src/backend/opencl/kernel/sparse.hpp‎

Copy file name to clipboardExpand all lines: src/backend/opencl/kernel/sparse.hpp

+2-1Lines changed: 2 additions & 1 deletion

Original file line number	Diff line number	Diff line change
@@ -49,7 +49,8 @@ void coo2dense(Param out, const Param values, const Param rowIdx,
49	49	cl::NDRange local(THREADS_PER_GROUP, 1, 1);
50	50
51	51	cl::NDRange global(
52		- divup(out.info.dims[0], local[0] * REPEAT) * THREADS_PER_GROUP, 1, 1);
	52	+ divup(values.info.dims[0], local[0] * REPEAT) * THREADS_PER_GROUP, 1,
	53	+ 1);
53	54
54	55	coo2dense(cl::EnqueueArgs(getQueue(), global, local), *out.data, out.info,
55	56	*values.data, values.info, *rowIdx.data, rowIdx.info,

0 commit comments

Comments

0 (0)