CUDA Array in CUDA | How to use CUDA Array in CUDA | CUDA ARRAY
CUDA Array in CUDA | How to use CUDA Array in CUDA
In the previous articles [this, this] we have discussed texturememory and
cudaChannelFormatDesc ()
in CUDA. In this article we learn how to use CUDA Array in CUDA programming,
which will be very useful when you start using Texture memory and Surface
memory (will be discusses in future article).
Before start discussion please learn how to
use cudaChannelFormatDesc () in
CUDA, which also introduce CUDA Array. If you walked from use cudaChannelFormatDesc () in CUDA than read future
otherwise go through given link.
Let’s start our discussion over CUDA Array
CUDA arrays are similar like an
array in CUDA but specially they are memory areas dedicate to textures. They
are read-only for the GPU (and are only accessible through texture fetches),
and can be written to by the CPU using
cudaMemcpyToArray ().
In essence, CUDA arrays are opaque
memory layouts optimized for texture fetching. They are one-dimensional,
two-dimensional, or three-dimensional and composed of elements, each of which
has 1, 2 or 4 components that may be signed or unsigned 8-, 16- or 32-bit
integers, 16-bit floats, or 32-bit floats. CUDA arrays are only readable by kernels
through texture fetching and may only be bound to texture references
with the same number of packed components.
Allocating a CUDA array requires
the specification of a channel format description matching the one used
by the texture that needs to be mapped on it.
Note:
Devices with capability
2.0 or higher can write to CUDA arrays using surfaces, a read-write
feature similar to textures. Table
F-2 listed about the
CUDA Array in terms of device compute
capability.
Please be familiar with the following
functions;
Allocation of CUDA Array
in CUDA
|
texture<> someTex;
cudaArray *dArray;
cudaMallocArray( &dArray,
&someTex.channelDesc, width, height);
|
Using CUDA Array in CUDA
|
cudaMemcpyToArray ( dArray, wOffset,
hOffset, source, size, cudaMemcpyHostToDevice);
cudaBindTextureToArray ( someTex,
dArray );
Copies
size bytes from the memory area pointed by Source to the CUDA Array dArray
starting at upper left corner (wOffset, hOffset)
|
Releasing CUDA Array
memory in CUDA
For releasing the
memory pointer by CUDA Array we use cudaFreeArray
()
|
cudaFreeArray
( dArray );
|
Complete
Example of CUDA Array and
cudaChannelFormatDesc ()
|
#include "HEADER.h"
#include "ERROR.h"
// 2D float texture
texture<float, cudaTextureType2D,
cudaReadModeElementType> texRef;
// Simple transformation kernel
__global__
void transformKernel(float* output, int width, int height, float theta)
{
// Calculate
normalized texture coordinates
unsigned
int x =
blockIdx.x *
blockDim.x +
threadIdx.x;
unsigned
int y =
blockIdx.y *
blockDim.y +
threadIdx.y;
float u = x / (float)width;
float v = y / (float)height;
// Transform
coordinates
u -=
0.5f;
v -=
0.5f;
float tu
= u * cosf(theta) - v * sinf(theta) + 0.5f ;
float tv
= v * cosf(theta) + u * sinf(theta) + 0.5f ;
// Read from
texture and write to global memory
output[y
* width + x] = tex2D(texRef, tu, tv);
}
// Host code
int main()
{
int
width = 3 ,
height = 3 ;
float
h_data[3][3] ;
for ( int i =0 ; i<3 ; i++ )
for
( int j = 0 ; j<3; j++ )
h_data [i][j] = i*j ;
int size
= width*height*sizeof(float) ;
// Allocate
CUDA array in device memory
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0,
cudaChannelFormatKindFloat );
cudaArray*
cuArray;
cudaMallocArray(&cuArray, &channelDesc, width, height);
// Copy to
device memory some data located at address h_data
// in host
memory
cudaMemcpyToArray(cuArray, 0, 0, h_data, size,
cudaMemcpyHostToDevice);
// Set texture
parameters
texRef.addressMode[0] = cudaAddressModeWrap;
texRef.addressMode[1] = cudaAddressModeWrap;
texRef.filterMode
= cudaFilterModeLinear;
texRef.normalized
= true;
// Bind the
array to the texture reference
cudaBindTextureToArray(texRef, cuArray, channelDesc);
// Allocate
result of transformation in device memory
float* output;
cudaMalloc(&output, size );
// Invoke
kernel
dim3 dimBlock(16, 16);
dim3 dimGrid(
(width +
dimBlock.x -
1) /
dimBlock.x,
(height +
dimBlock.y -
1) /
dimBlock.y );
transformKernel<<<dimGrid, dimBlock>>>(output, width, height, 90 );
printf ("\n Original array \n");
for ( int i =0 ; i<3 ; i++ )
{
for
( int j = 0 ; j<3; j++ )
printf ("%f ", h_data [i][j] );
printf ("\n");
}
CUDA_CALL (
cudaMemcpy (h_data ,
output, size,
cudaMemcpyDeviceToHost ) );
printf ("\n After operation \n ");
for ( int i =0 ; i<3 ; i++ )
{
for
( int j = 0 ; j<3; j++ )
printf ("%f ", h_data [i][j] );
printf ("\n");
}
system ("pause");
// Free device
memory
cudaFreeArray(cuArray);
cudaFree(output);
return 0;
}
|
Sample
output
Fig.
I hope you must like this article and have learned CUDA Array and how to use CUDA Array in CUDA.
Got
Questions?
Feel
free to ask me any question because I'd be happy to walk you through step by
step!
Posts
