Come menzionato da sgarizvi, CUDA supporta solo quattro, modalità non personalizzabili di indirizzo, vale a dire, morsetto, confine, involucro e specchio, che sono descritti nella sezione 3.2.11.1. della guida alla programmazione CUDA.
I primi due lavorano in entrambe le coordinate non normalizzate e normalizzate, mentre le ultime due solo in coordinate normalizzate.
Per descrivere i primi due, consideriamo il caso di coordinate non normalizzate e consideriamo i segnali 1D, per semplicità. In questo caso, la sequenza di input è c[k]
, con k=0,...,M-1
.
cudaAddressModeClamp
Il segnale c[k]
prosegue all'esterno k=0,...,M-1
modo che c[k] = c[0]
per k < 0
e c[k] = c[M-1]
per k >= M
.
cudaAddressModeBorder
Il segnale c[k]
prosegue all'esterno k=0,...,M-1
modo che c[k] = 0
per k < 0
e per k >= M
.
Ora, per descrivere le ultime due modalità di indirizzo, siamo costretti a considerare le coordinate normalizzate, in modo che i campioni del segnale di ingresso 1D siano assunti da c[k/M]
, con k=0,...,M-1
.
cudaAddressModeWrap
Il segnale c[k/M]
prosegue all'esterno k=0,...,M-1
modo che sia periodico con periodo pari a M
. In altre parole, c[(k + p * M)/M] = c[k/M]
per qualsiasi numero intero (positivo, negativo o negativo) p
.
cudaAddressModeMirror
Il segnale c[k/M]
prosegue all'esterno k=0,...,M-1
modo che sia periodico con periodo pari a 2 * M - 2
.In altre parole, c[l/M] = c[k/M]
per qualsiasi l
e k
tale che (l + k)mod(2 * M - 2) = 0
.
Il codice seguente illustra tutte le quattro modalità di indirizzi disponibili
#include <stdio.h>
texture<float, 1, cudaReadModeElementType> texture_clamp;
texture<float, 1, cudaReadModeElementType> texture_border;
texture<float, 1, cudaReadModeElementType> texture_wrap;
texture<float, 1, cudaReadModeElementType> texture_mirror;
/********************/
/* CUDA ERROR CHECK */
/********************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
if (code != cudaSuccess)
{
fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
}
/******************************/
/* CUDA ADDRESS MODE CLAMPING */
/******************************/
__global__ void Test_texture_clamping(const int M) {
printf("Texture clamping - i = %i; value = %f\n", -threadIdx.x, tex1D(texture_clamp, -(float)threadIdx.x));
printf("Texture clamping - i = %i; value = %f\n", M + threadIdx.x, tex1D(texture_clamp, (float)(M + threadIdx.x)));
}
/****************************/
/* CUDA ADDRESS MODE BORDER */
/****************************/
__global__ void Test_texture_border(const int M) {
printf("Texture border - i = %i; value = %f\n", -threadIdx.x, tex1D(texture_border, -(float)threadIdx.x));
printf("Texture border - i = %i; value = %f\n", M + threadIdx.x, tex1D(texture_border, (float)(M + threadIdx.x)));
}
/**************************/
/* CUDA ADDRESS MODE WRAP */
/**************************/
__global__ void Test_texture_wrap(const int M) {
printf("Texture wrap - i = %i; value = %f\n", -threadIdx.x, tex1D(texture_wrap, -(float)threadIdx.x/(float)M));
printf("Texture wrap - i = %i; value = %f\n", M + threadIdx.x, tex1D(texture_wrap, (float)(M + threadIdx.x)/(float)M));
}
/****************************/
/* CUDA ADDRESS MODE MIRROR */
/****************************/
__global__ void Test_texture_mirror(const int M) {
printf("Texture mirror - i = %i; value = %f\n", -threadIdx.x, tex1D(texture_mirror, -(float)threadIdx.x/(float)M));
printf("Texture mirror - i = %i; value = %f\n", M + threadIdx.x, tex1D(texture_mirror, (float)(M + threadIdx.x)/(float)M));
}
/********/
/* MAIN */
/********/
void main(){
const int M = 4;
// --- Host side memory allocation and initialization
float *h_data = (float*)malloc(M * sizeof(float));
for (int i=0; i<M; i++) h_data[i] = (float)i;
// --- Texture clamping
cudaArray* d_data_clamping = NULL; gpuErrchk(cudaMallocArray(&d_data_clamping, &texture_clamp.channelDesc, M, 1));
gpuErrchk(cudaMemcpyToArray(d_data_clamping, 0, 0, h_data, M * sizeof(float), cudaMemcpyHostToDevice));
cudaBindTextureToArray(texture_clamp, d_data_clamping);
texture_clamp.normalized = false;
texture_clamp.addressMode[0] = cudaAddressModeClamp;
dim3 dimBlock(2 * M, 1); dim3 dimGrid(1, 1);
Test_texture_clamping<<<dimGrid,dimBlock>>>(M);
printf("\n\n\n");
// --- Texture border
cudaArray* d_data_border = NULL; gpuErrchk(cudaMallocArray(&d_data_border, &texture_border.channelDesc, M, 1));
gpuErrchk(cudaMemcpyToArray(d_data_border, 0, 0, h_data, M * sizeof(float), cudaMemcpyHostToDevice));
cudaBindTextureToArray(texture_border, d_data_border);
texture_border.normalized = false;
texture_border.addressMode[0] = cudaAddressModeBorder;
Test_texture_border<<<dimGrid,dimBlock>>>(M);
printf("\n\n\n");
// --- Texture wrap
cudaArray* d_data_wrap = NULL; gpuErrchk(cudaMallocArray(&d_data_wrap, &texture_wrap.channelDesc, M, 1));
gpuErrchk(cudaMemcpyToArray(d_data_wrap, 0, 0, h_data, M * sizeof(float), cudaMemcpyHostToDevice));
cudaBindTextureToArray(texture_wrap, d_data_wrap);
texture_wrap.normalized = true;
texture_wrap.addressMode[0] = cudaAddressModeWrap;
Test_texture_wrap<<<dimGrid,dimBlock>>>(M);
printf("\n\n\n");
// --- Texture mirror
cudaArray* d_data_mirror = NULL; gpuErrchk(cudaMallocArray(&d_data_mirror, &texture_mirror.channelDesc, M, 1));
gpuErrchk(cudaMemcpyToArray(d_data_mirror, 0, 0, h_data, M * sizeof(float), cudaMemcpyHostToDevice));
cudaBindTextureToArray(texture_mirror, d_data_mirror);
texture_mirror.normalized = true ;
texture_mirror.addressMode[0] = cudaAddressModeMirror;
Test_texture_mirror<<<dimGrid,dimBlock>>>(M);
printf("\n\n\n");
}
Queste sono le uscite
index -7 -6 -5 -4 -3 -2 -1 0 1 2 3 4 5 6 7 8 9 10 11
clamp 0 0 0 0 0 0 0 0 1 2 3 3 3 3 3 3 3 3 3
border 0 0 0 0 0 0 0 0 1 2 3 0 0 0 0 0 0 0 0
wrap 1 2 3 0 1 2 3 0 1 2 3 0 1 2 3 0 1 2 3
mirror 1 2 3 3 2 1 0 0 1 2 3 3 2 1 0 0 1 2 3
l'hardware supporta l'impostazione del colore, ma che non è esposta in CUDA. Probabilmente perché nessuna delle classiche modalità di indirizzamento richiede parametri aggiuntivi. NVIDIA lo ha registrato come funzione richiesta. Per ovviare al problema, è possibile disegnare un bordo di 1 pixel del colore necessario intorno alla trama e utilizzare la modalità di indirizzamento della pinza insieme alle coordinate ottimizzate. –
@RogerDahl Avevo indovinato che questo è solo un problema dell'API CUDA. Perché il colore del bordo può essere impostato in DirectX per lo stesso hardware. In ogni caso, non posso modificare la trama in questo caso particolare, quindi nessuna soluzione per me :-) –