1

I have been reading a few threads on stackoverflow about 2D arrays and cudaMallocPitch and I have tried to use cudaMallocPitch with the small documentation I have found. However I'm now facing a problem.

I need to go through an array and do something similar : 

 for(int k=0; k<100; ++k){
     for(i=SID; i<SID+stride; ++i){
        while(-1 < j && Driver[k][j] != Road[i]){
            j = Pilot[j][k];

        }
        ++j;
     }
  }

I was thus wondering, how should I adapt this code to make it work with the pitch, because I have read that I had to update the pointer to the beginning of the row. Of course my kernel receives the following :

__global__ void driving(char *Driver, size_t pitch_driver, 
                        char *Road, int *Pilot, size_t pitch_pilot)

And I'm not really sure how to make things working, I've been reading and trying, but it seems not working at the moment.

Thank you.

Edit 1: I have been reading this thread in particular :How to use 2D Arrays in CUDA? and came across the lines : 

for (int row = 0; row < rowCount; row++)  
 {  
     // update the pointer to point to the beginning of the next row  
    float* rowData = (float*)(((char*)d_array) + (row * pitch));  
    for (int column = 0; column < columnCount; column++)  
     {  
       rowData[column] = 123.0; // make every value in the array 123.0  
       destinationArray[(row*columnCount) + column] = rowData[column];  
      }  
 }

Which is updating the pointer of the next row, I'am not sure how to use to make my 2 for loops and while working such as in the previous code.

At the moment I can only access one dimension of my array but not the other one.

it returns the value 2, but when I try my multiple comparisons, it only returns 0, or even comparing two values do not work.

Community
  • 1
  • 1
Anoracx
  • 438
  • 7
  • 24
  • 2
    What, exactly is "not working"? What part of the process don't you understand? – talonmies Feb 28 '13 at 06:15
  • It's not working when I try to go through the array. Based on something similar to what i edited. – Anoracx Feb 28 '13 at 12:07
  • I think your questions boils down to this: True 2D arrays cause problems in C, and even more problems in CUDA. Instead, `cudaMallocPitch()` represents the 2D array as a 1D array that you simply access as: `array[row*pitch + column];`. That's it. If you really need a pointer to the start of a row, you can use `ptr = &array[row*pitch];` – Peter Feb 28 '13 at 14:27
  • 1
    @Peter: Your access code isn't correct. There is no guarantee that the pitch that cudaMallocPitch returns will be a nice multiple of the word size of the array type. The pitch is selected purely on GPU memory controller page size and map requirements. The API doesn't know anything about the type being allocated, only the *minimum* byte width of each row to be allocated. – talonmies Feb 28 '13 at 14:56
  • I actually tried what @Peter said and it's not working, I still cant access the second row. – Anoracx Feb 28 '13 at 14:59
  • Sorry about that, you're right. pitch is reported in bytes. The correct access pattern is given in the docs for cudaMallocPitch. The OP also has it listed fairly clearly in the rowData line of the sample code. @Anoracx, you'll need to be more specific about what isn't working. – Peter Feb 28 '13 at 16:59
  • @talonmies and peter, I updated the post a second time. It seems that when I want to compare 2 chars it doesn't work. However 2 int seems to work. – Anoracx Feb 28 '13 at 22:51

1 Answers1

2

In the CUDA Reference Manual it says:

5.8.2.17 cudaError_t cudaMallocPitch (void devPtr, size_t pitch, size_t width, size_t height)

[...]

Given the row and column of an array element of type T, the address is computed as:

T* pElement = (T*)((char*)BaseAddress + Row * pitch) + Column;

So you need to cast your pointer first to char*, do the math and then cast it back to your type.

Hugo Maxwell
  • 723
  • 5
  • 13