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
THANKS FOR SHARING SUCH A GREAT INFORMATION.. IT REALLY HELPFUL TO ME..I ALWAYS SEARCH TO READ THE QUALITY CONTENT AND FINALLY I FOUND THIS IN YOU POST. KEEP IT UP!
ReplyDelete섯다
LOVE TO READ YOUR BLOG! THANKS FOR POSTING SOME VALUABLE POST.
ReplyDelete토토사이트
Malatya
ReplyDeleteKırıkkale
Aksaray
Bitlis
Manisa
NVYR6
kars
ReplyDeletesinop
sakarya
ankara
çorum
82T1
https://titandijital.com.tr/
ReplyDeletenevşehir parça eşya taşıma
bolu parça eşya taşıma
batman parça eşya taşıma
bayburt parça eşya taşıma
DS84
ankara parça eşya taşıma
ReplyDeletetakipçi satın al
antalya rent a car
antalya rent a car
ankara parça eşya taşıma
22LJOV
antalya evden eve nakliyat
ReplyDeleteankara evden eve nakliyat
bursa evden eve nakliyat
yalova evden eve nakliyat
gümüşhane evden eve nakliyat
1YX
5CAA7
ReplyDeleteErzincan Lojistik
Ardahan Evden Eve Nakliyat
Kırıkkale Parça Eşya Taşıma
Antalya Parça Eşya Taşıma
Muğla Evden Eve Nakliyat
76267
ReplyDeleteerzurum canlı sohbet odaları
isparta sesli sohbet sitesi
bartın sesli sohbet siteler
sinop görüntülü canlı sohbet
eskişehir sesli görüntülü sohbet
adıyaman görüntülü sohbet sitesi
mobil sesli sohbet
sivas rastgele sohbet siteleri
bursa sohbet sitesi
1D6AB
ReplyDeleteosmaniye tamamen ücretsiz sohbet siteleri
canli goruntulu sohbet siteleri
mardin görüntülü sohbet yabancı
Antalya En İyi Rastgele Görüntülü Sohbet
aydın canlı sohbet
Gümüşhane Canlı Görüntülü Sohbet Siteleri
bitlis sesli mobil sohbet
isparta sohbet chat
aksaray yabancı görüntülü sohbet siteleri
3DA11
ReplyDeleteHakkari Tamamen Ücretsiz Sohbet Siteleri
karaman görüntülü sohbet kızlarla
kocaeli canli sohbet bedava
istanbul canlı sohbet et
ücretsiz sohbet odaları
telefonda sohbet
istanbul telefonda kadınlarla sohbet
rastgele sohbet
Şırnak Canli Goruntulu Sohbet Siteleri
626F9
ReplyDeleteKripto Para Çıkarma
Gate io Borsası Güvenilir mi
MEME Coin Hangi Borsada
Kripto Para Kazma
Parasız Görüntülü Sohbet
Loop Network Coin Hangi Borsada
Kwai Beğeni Hilesi
Bitcoin Nasıl Kazanılır
Binance Borsası Güvenilir mi