As already indicated by kangshiyin, the improvements arising from the use of cudaMallocPitch
depend on the compute capability and are expected to be more significant for older ones. However, for most recent compute capabilities, pitched memory allocation does not seem to lead to a relevant speedup.
The code below provides a performance testbench between the uses of non-pitched and pitched memories. In particular, the code performs the summation between three (non-pitched or pitched) matrices. The reason for dealing with three matrices is the need to highlight memory transactions as compared to computation, so to highlight the differences between non-pitched and pitched allocations. Below are the timing results for a GTX 960
card and a GT 920M
cards.
GTX 960
Non-pitched - Time = 3.242208; Memory = 65320000 bytes
Pitched - Time = 3.150944; Memory = 65433600 bytes
GT 920M
Non-pitched - Time = 20.496799; Memory = 65320000 bytes
Pitched - Time = 20.418560; Memory = 65433600 bytes
As it can be seen, there is not much difference in the two implementations for the two cards. The above results also show the increase in memory occupancy due to the use of pitched memory allocation.
Here is the code:
#include<stdio.h>
#include<cuda.h>
#include<cuda_runtime.h>
#include<device_launch_parameters.h>
#include "Utilities.cuh"
#include "TimingGPU.cuh"
#define BLOCKSIZE_x 16
#define BLOCKSIZE_y 16
/******************/
/* TEST KERNEL 2D */
/******************/
__global__ void test_kernel_2D(float * __restrict__ devPtrA, float * __restrict__ devPtrB, float * __restrict__ devPtrC, const int Nrows, const int Ncols)
{
int tidx = blockIdx.x * blockDim.x + threadIdx.x;
int tidy = blockIdx.y * blockDim.y + threadIdx.y;
if ((tidx < Ncols) && (tidy < Nrows)) {
devPtrA[tidy * Ncols + tidx] = devPtrA[tidy * Ncols + tidx] + devPtrB[tidy * Ncols + tidx] + devPtrC[tidy * Ncols + tidx];
}
}
/**************************/
/* TEST KERNEL PITCHED 2D */
/**************************/
__global__ void test_kernel_Pitched_2D(float * __restrict__ devPtrA, float * __restrict__ devPtrB, float * __restrict__ devPtrC, const size_t pitchA, const size_t pitchB, const size_t pitchC, const int Nrows, const int Ncols)
{
int tidx = blockIdx.x * blockDim.x + threadIdx.x;
int tidy = blockIdx.y * blockDim.y + threadIdx.y;
if ((tidx < Ncols) && (tidy < Nrows))
{
float *row_a = (float *)((char*)devPtrA + tidy * pitchA);
float *row_b = (float *)((char*)devPtrB + tidy * pitchB);
float *row_c = (float *)((char*)devPtrC + tidy * pitchC);
row_a[tidx] = row_a[tidx] + row_b[tidx] + row_c[tidx];
}
}
/********/
/* MAIN */
/********/
int main()
{
const int Nrows = 7100;
const int Ncols = 2300;
TimingGPU timerGPU;
float *hostPtrA = (float *)malloc(Nrows * Ncols * sizeof(float));
float *hostPtrB = (float *)malloc(Nrows * Ncols * sizeof(float));
float *hostPtrC = (float *)malloc(Nrows * Ncols * sizeof(float));
float *devPtrA, *devPtrPitchedA;
float *devPtrB, *devPtrPitchedB;
float *devPtrC, *devPtrPitchedC;
size_t pitchA, pitchB, pitchC;
for (int i = 0; i < Nrows; i++)
for (int j = 0; j < Ncols; j++) {
hostPtrA[i * Ncols + j] = 1.f;
hostPtrB[i * Ncols + j] = 2.f;
hostPtrC[i * Ncols + j] = 3.f;
//printf("row %i column %i value %f \n", i, j, hostPtr[i][j]);
}
// --- 2D non-pitched allocation and host->device memcopy
gpuErrchk(cudaMalloc(&devPtrA, Nrows * Ncols * sizeof(float)));
gpuErrchk(cudaMalloc(&devPtrB, Nrows * Ncols * sizeof(float)));
gpuErrchk(cudaMalloc(&devPtrC, Nrows * Ncols * sizeof(float)));
gpuErrchk(cudaMemcpy(devPtrA, hostPtrA, Nrows * Ncols * sizeof(float), cudaMemcpyHostToDevice));
gpuErrchk(cudaMemcpy(devPtrB, hostPtrB, Nrows * Ncols * sizeof(float), cudaMemcpyHostToDevice));
gpuErrchk(cudaMemcpy(devPtrC, hostPtrC, Nrows * Ncols * sizeof(float), cudaMemcpyHostToDevice));
// --- 2D pitched allocation and host->device memcopy
gpuErrchk(cudaMallocPitch(&devPtrPitchedA, &pitchA, Ncols * sizeof(float), Nrows));
gpuErrchk(cudaMallocPitch(&devPtrPitchedB, &pitchB, Ncols * sizeof(float), Nrows));
gpuErrchk(cudaMallocPitch(&devPtrPitchedC, &pitchC, Ncols * sizeof(float), Nrows));
gpuErrchk(cudaMemcpy2D(devPtrPitchedA, pitchA, hostPtrA, Ncols * sizeof(float), Ncols*sizeof(float), Nrows, cudaMemcpyHostToDevice));
gpuErrchk(cudaMemcpy2D(devPtrPitchedB, pitchB, hostPtrB, Ncols * sizeof(float), Ncols*sizeof(float), Nrows, cudaMemcpyHostToDevice));
gpuErrchk(cudaMemcpy2D(devPtrPitchedC, pitchC, hostPtrC, Ncols * sizeof(float), Ncols*sizeof(float), Nrows, cudaMemcpyHostToDevice));
dim3 gridSize(iDivUp(Ncols, BLOCKSIZE_x), iDivUp(Nrows, BLOCKSIZE_y));
dim3 blockSize(BLOCKSIZE_y, BLOCKSIZE_x);
timerGPU.StartCounter();
test_kernel_2D << <gridSize, blockSize >> >(devPtrA, devPtrB, devPtrC, Nrows, Ncols);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
printf("Non-pitched - Time = %f; Memory = %i bytes \n", timerGPU.GetCounter(), Nrows * Ncols * sizeof(float));
timerGPU.StartCounter();
test_kernel_Pitched_2D << <gridSize, blockSize >> >(devPtrPitchedA, devPtrPitchedB, devPtrPitchedC, pitchA, pitchB, pitchC, Nrows, Ncols);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
printf("Pitched - Time = %f; Memory = %i bytes \n", timerGPU.GetCounter(), Nrows * pitchA);
//gpuErrchk(cudaMemcpy2D(hostPtr, Ncols * sizeof(float), devPtrPitched, pitch, Ncols * sizeof(float), Nrows, cudaMemcpyDeviceToHost));
gpuErrchk(cudaMemcpy(hostPtrA, devPtrA, Nrows * Ncols * sizeof(float), cudaMemcpyDeviceToHost));
gpuErrchk(cudaMemcpy(hostPtrB, devPtrB, Nrows * Ncols * sizeof(float), cudaMemcpyDeviceToHost));
gpuErrchk(cudaMemcpy(hostPtrC, devPtrC, Nrows * Ncols * sizeof(float), cudaMemcpyDeviceToHost));
//for (int i = 0; i < Nrows; i++)
// for (int j = 0; j < Ncols; j++)
// printf("row %i column %i value %f \n", i, j, hostPtr[i * Ncols + j]);
return 0;
}