Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Already on GitHub? Sign in to your account

CUDA's reduce_by_key should fall back to the generic path if the size of the data type is too large #318

Closed
jaredhoberock opened this Issue Dec 21, 2012 · 3 comments

Comments

Projects
None yet
2 participants
Owner

jaredhoberock commented Dec 21, 2012

I don't believe there is a significant limitation on the generic path.

What's the data size limit of reduce_by_key? I have a dataset of 180million items, it works if i passed in 1million, but if fails if 10million or more.
code snippet is:
size_t num_runs = thrust::reduce_by_key
(d_col_ptr.begin(), d_col_ptr.begin() + testsize, // input key sequence
thrust::constant_iterator(1), // input value sequence
output.begin(), // output key sequence
lengths.begin() // output value sequence
).first - output.begin(); // compute the output size
it's on Mac OS X 10.8.
gzhangtest/thrust]$ gcc --version
i686-apple-darwin11-llvm-gcc-4.2 (GCC) 4.2.1 (Based on Apple Inc. build 5658) (LLVM build 2336.11.00)
gzhangtest/thrust]$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2012 NVIDIA Corporation
Built on Fri_Sep_28_16:10:16_PDT_2012
Cuda compilation tools, release 5.0, V0.2.1221

gzhangtest/thrust]$ nvcc -m64 compress.cu -o compress
gzhangtest/thrust]$ ./compress 1000000
Datafile /dbdata/gra/func_area_sk, size is: 364767580, number of data items: 182383790
run-length encoded output:
(136,1)(-9,1)(128,1)(-9,2)(131,1)(-9,1)(136,1)(-9,1)(129,1)(136,3)(-9,4)(136,1)(-9,1)(119,1)(129,1)(-9,2)(163,1)(-9,1)(129,1)(-9,2)
gzhangtest/thrust]$ ./compress 10000000
Datafile /dbdata/gra/func_area_sk, size is: 364767580, number of data items: 182383790
libc++abi.dylib: terminate called throwing an exception
Abort trap: 6

thanks,
George

Owner

jaredhoberock commented Apr 4, 2013

Hi George,
The only data size limit is the limit of temporary memory which can be allocated by the algorithm. We usually recommend folks use the thrust-users mailing list for this sort of question.

If you can post a self-contained program which reproduces the problem, someone should be able to help you out.

Owner

jaredhoberock commented Sep 18, 2013

Relying on shmalloc will essentially address this problem.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment