Peter, from the context all the questions you are asking are superfluous. The tags include "GPU Computing" and "CUDA Programming". Everyone in this area knows the terms "device memory" and "host". You are guessing right that this is actually a question about CUDA. "Thrust" is a higher level library which provides a nice C++ interface and executes kernels in the background, e.g. for matrix or vector operations. This way one does not have to write CUDA device code directly.

Unfortunately, I am not familiar with the Thrust library myself. But, the overall question is, what you are trying to do? What a zip operation would do is take the values stored in the three arrays you provided in the example code. These would be copied to an array of structs which you also need to allocate with cudaMalloc before. Why would you want to do that? The extra copy of the zip operation takes time and afterwards you would loose performance when accessing memory. On the GPU it is always (at least 99.9% of the time) more efficient to store a struct of arrays (like you have done) than to store an array of struct (the struct would have an int, a float, and an int). So, why do you want to use a zip_iterator? What specific problem are you trying to solve?

Hi, Peter: Nice try! Pretty much you are right about the zip_iterator's potentially laborous work and the GPU's powerful paralell (concurrent) processing capability. Yes, driver has to be in, however it is specific to a GPU card. That driver is installed seperately from CUDA, which includes functions, cudaMalloc() as an example. Thanks for answering my question!

@Peter: In the context of GPU programming a kernel is a function to be executed on the GPU. The meaning has not relation to an OS kernel. The compiled kernel (compilation is done during runtime) is send to the GPU together with the input parameters (i.e. your data) which is then executed in parallel on the GPU. Usually, you would hand in an array to the kernel. Each instance of the kernel will have its own index. Knowing the index of your instance and the size of the array you can compute on which part of the array your instance of the kernel will do the computation. Since each instance is only processing a chunk (or maybe just one element) of the array all instances can run in parallel. When invoking the kernel you tell the number of kernel instances you want to start in parallel.

The GPU does have its own memory and access to this memory is a lot faster than access to regular RAM. Only the most recent version of CUDA allows to access CPU memory directly. cudaMalloc is the equivalent of malloc to allocate memory on the GPU. cudaCopy is then used to copy data from CPU memory to GPU memory. (Actually, the driver is smarter than this and I think that cudaCopy only moves the data to an area in CPU memory where you cannot mess with it anymore. It is then automatically copied to the GPU when a kernel needs to access it.)

Back to the original question. Zipping is a term from functional programming. What the zip_iterator does is "it takes multiple input sequences and yields a sequence of tuples" (taken from the CUDA docs). After reading the docs I have to correct my previous statement. zip_iterator does not copy the data into a single iterator, but only provides a virtual array to the existing values. Hence, performance will not decrease as I stated before. You will find an example of how it is done in the official documentation: http://thrust.github.io/doc/classthrust_1_1zip__iterator.html

Hi, Simon: The doc you linked is among the most helpful I have ever read about the usage  of zip_iterator in Thrust. However it just can not be used directly, because in order to refer to the raw pointer in device memory, thrust asks for the wrapping of it and thus we get device_ptr defined in Thrust. Thrust operates on the data pointed to by raw pointer via device_ptr. I got the question because zip_iterator can zip device_vector but device_ptr. Why do I need to zip the three arrays? The short reply is in my project , I need to get the two elements of the two int arrays associated with the smallest element of the float array. The size of the array is huge and the arrays are already in device memory, why Thrust can not be given a try? 

I'm not sure if I am understanding completely what you are trying to do. Maybe this link http://docs.nvidia.com/cuda/thrust/#zip-iterator will also help a little bit in finding your answer (I am not a real CUDA expert, but have a good understanding of GPGPU, mostly through OpenCL).

I am not sure what minimum() would do on a tuple. It might do another thing than you expect. Here is how I would go about this: With min_element I would extract the element from the float array. The subtraction (min_element - first element) gives you the index which can then be used on the two int arrays.

Hi, Simon, what you have suggested works perfectly! Thanks.