index of max_element with raw pointers

Question

I've been looking at Thrust and I stumbled upon a question that almost (but not quite) answered mine: Finding the maximum element value AND its position using CUDA Thrust

The example posted in there in the answer works fine, but how to do the same thing with raw pointers? Let us assume this code which I believe to be correct (ignore the kernel configuration, it's for simplicity):

float* d_A;
const unsigned int noElems = 10;
cudaMalloc(&d_A, noElems * sizeof(float));
initDeviceVector<<<1, noElems>>>(d_A);

thrust::device_ptr<float> d_ptr = thrust::device_pointer_cast(d_A);     
thrust::device_vector<float>::iterator iter = 
    thrust::max_element(d_ptr, d_ptr + noElems);

I can't quite figure out how to extract the position using iter and raw pointers.

Thank you for your time.

True, I think I understand what I'm supposed to do (just subtract iter from d_ptr), what I can't figure out is how to do it in Thrust. Obviously, iter - d_ptr won't work simply because of different datatypes. Is there a way, for example, to get an integer value out of both iter and d_ptr? I checked the documentation and was unable to extract the information from there. — Intri, Jun 04 '14 at 12:07
If I understand correctly your question, you need two operations: 1) convert a raw pointer to a `device_ptr` so that you can feed `max_element` with raw pointers; 2) convert an iterator to a pointer so that you can point to the max element value afterwards. Would these two posts help you: [Converting thrust::iterators to and from raw pointers](http://stackoverflow.com/questions/12201446/converting-thrustiterators-to-and-from-raw-pointers) and [Converting an STL vector iterator to a raw pointer](http://artins.org/ben/software/converting-an-stl-vector-iterator-to-a-raw-pointer). — Vitality, Jun 04 '14 at 14:06

score 2 · Accepted Answer · answered Jun 04 '14 at 14:09

There's probably a number of ways to do this. However working directly from your code, we can compare the value of iter to a device_ptr if we convert it to a suitable device pointer first.

The following fully worked example demonstrates this:

$ cat t436.cu
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/extrema.h>
#include <stdio.h>


__global__ void initDeviceVector(float *data){
  int idx = threadIdx.x+blockDim.x*blockIdx.x;
  data[idx] = idx%7;
}

  int main(){

  float* d_A;
  const unsigned int noElems = 10;
  cudaMalloc(&d_A, noElems * sizeof(float));
  initDeviceVector<<<1, noElems>>>(d_A);

  thrust::device_ptr<float> d_ptr = thrust::device_pointer_cast(d_A);
  thrust::device_vector<float>::iterator iter = thrust::max_element(d_ptr, d_ptr + noElems);

  int pos = thrust::device_pointer_cast(&(iter[0])) - d_ptr;

  printf("pos = %d\n", pos);
  return 0;
}

$ nvcc -arch=sm_20 -o t436 t436.cu
$ ./t436
pos = 6
$

Thank you, this is it. One more smaller question, if I may: in my experiments, I stored the result of the max_element function directly to thrust::device_ptr and it seems to work as well. Does Thrust do the pointer casting implicitly? — Intri, Jun 04 '14 at 14:29
Yes, it does, and doing that considerably simplifies the code. — Robert Crovella, Jun 04 '14 at 18:51

index of max_element with raw pointers

1 Answers1