22

How do I allocate and transfer(to and from Host) 2D arrays in device memory in Cuda?

4 Answers 4

20

I found a solution to this problem. I didn't have to flatten the array.

The inbuilt cudaMallocPitch() function did the job. And I could transfer the array to and from device using cudaMemcpy2D() function.

For example

cudaMallocPitch((void**) &array, &pitch, a*sizeof(float), b);

This creates a 2D array of size a*b with the pitch as passed in as parameter.

The following code creates a 2D array and loops over the elements. It compiles readily, you may use it.

#include<stdio.h>
#include<cuda.h>
#define height 50
#define width 50

// Device code
__global__ void kernel(float* devPtr, int pitch)
{
    for (int r = 0; r < height; ++r) {
        float* row = (float*)((char*)devPtr + r * pitch);
        for (int c = 0; c < width; ++c) {
             float element = row[c];
        }
    }
}

//Host Code
int main()
{    
    float* devPtr;
    size_t pitch;
    cudaMallocPitch((void**)&devPtr, &pitch, width * sizeof(float), height);
    kernel<<<100, 512>>>(devPtr, pitch);
    return 0;
}
1
  • is it possible to allocate a new row for the array later on?
    – scatman
    Apr 12, 2011 at 6:08
4

Flatten it: make it one-dimensional. See how it's done here

3

Your device code could be faster. Try utilizing the threads more.

__global__ void kernel(float* devPtr, int pitch)
{
    int r = threadIdx.x;

    float* row = (float*)((char*)devPtr + r * pitch);
    for (int c = 0; c < width; ++c) {
         float element = row[c];
    }
}

Then you calculate the blocks and threads allocation appropriate so that each thread deals with a single element.

1
  • The code Gitmo posted is a useless sample from the docs. Yes, your version is faster, but how do you do this in parallel for rows and columns? Strictly speaking you could have a mess in your hands because you don't check if r is less than the actual number of rows
    – darda
    Jun 19, 2014 at 0:25
1
+50
#include <stdio.h>
#include <stdlib.h>
#include <sys/time.h>
#include <cuda.h>
#define MAX_ITER 1000000
#define MAX 100 //maximum value of the matrix element
#define TOL 0.000001

// Generate a random float number with the maximum value of max
float rand_float(int max){
  return ((float)rand()/(float)(RAND_MAX)) * max;
}

__global__ void kernel(float **device_2Darray1, float **device_2Darray2, float **device_2Darray3, int rows, int cols) {

  // Calculate the row index
  int row = blockIdx.y * blockDim.y + threadIdx.y;

  // Calculate the column index
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  
  // Check if the thread is within the array bounds
  if (row < rows && col < cols) {
    // Perform the computation
    device_2Darray3[row][col] = device_2Darray1[row][col] + device_2Darray2[row][col];
  }
}


int main(int argc, char *argv[]){
  float **host_2Darray;
  float **device_2Darray;

  int rows = 10; // or whatever value you want
  int cols = 10; // or whatever value you want

  // allocate memory for the host
  host_2Darray = (float**)malloc(rows * sizeof(float*));
  for(int i = 0; i < rows; i++){
    host_2Darray[i] = (float*)malloc(cols * sizeof(float));
    for(int j = 0; j < cols; j++){
      host_2Darray[i][j] = rand_float(MAX);
    }
  }

  // allocate memory for the device
  cudaMalloc((void***)&device_2Darray, rows * sizeof(float*));
  for(int i = 0; i < rows; i++){
    cudaMalloc((void**)&device_2Darray[i], cols * sizeof(float));
  }

  // copy host memory to device
  for(int i = 0; i < rows; i++){
    cudaMemcpy(device_2Darray[i], host_2Darray[i], cols * sizeof(float), cudaMemcpyHostToDevice);
  }

  // call the kernel
  dim3 threadsPerBlock(16, 16);
  dim3 blocksPerGrid((rows + threadsPerBlock.x - 1) / threadsPerBlock.x, 
                     (cols + threadsPerBlock.y - 1) / threadsPerBlock.y);
  kernel<<<blocksPerGrid, threadsPerBlock>>>(device_2Darray, rows, cols);

  // copy device memory back to host
  for(int i = 0; i < rows; i++){
    cudaMemcpy(host_2Darray[i], device_2Darray[i], cols * sizeof(float), cudaMemcpyDeviceToHost);
  }

  // free device memory
  for(int i = 0; i < rows; i++){
    cudaFree(device_2Darray[i]);
  }
  cudaFree(device_2Darray);

  // free host memory
  for(int i = 0; i < rows; i++){
    free(host_2Darray[i]);
  }
  free(host_2Darray);

  return 0;
}

The method of creating 2D arrays in CUDA is more complicated than that of a 1D array because the device (GPU) memory is linear. We actually create an array of pointers (each pointer pointing to a 1D array), hence the double pointers.

As for the global void kernel function: Simple operation like adding two 2D arrays elementwise. You would use blockIdx and blockDim along with threadIdx to compute the global index of the thread, and then use this index to specify which element of the array this thread is responsible for. In this example, blockIdx.y and blockIdx.x gives the index of the current thread within the block in y and x directions. You can think of it as a 3-level hierarchy: grid -> block -> thread.

You would need to adjust the rest of your code to allocate and initialize the second array, allocate the output array, copy the output back to the host, etc.

Don't forget to modify the kernel invocation as well to pass the correct arguments. You have to be expected to do some of the work yourself, like integrating in the concepts, as no one has access to your full code and how it all works together as a system.

2
  • 1
    I already awarded 50 points to your answer. However, now I see that some parts of the source code are missing. Have you tested your source code?
    – user366312
    Apr 30, 2023 at 16:54
  • I updated the code, sorry about that. May 26, 2023 at 16:00

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service and acknowledge you have read our privacy policy.

Not the answer you're looking for? Browse other questions tagged or ask your own question.