1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
|
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <cuda.h>
#define TEXSIZE 4096
texture<float4, 1, cudaReadModeElementType> texData;
void cuda_errcheck(const char *msg) {
cudaError_t err;
if ((err = cudaGetLastError()) != cudaSuccess) {
char errmsg[1024];
sprintf(errmsg,"CUDA error %s: %s", msg, cudaGetErrorString(err));
}
}
void cuda_bind_texture_data(const float4 *t) {
static cudaArray *ct;
if ( ! ct ) {
cudaMallocArray(&ct, &texData.channelDesc, TEXSIZE, 1);
cuda_errcheck("allocating texDataArray");
} cudaMemcpyToArray(ct, 0, 0, t, TEXSIZE*sizeof(float4), cudaMemcpyHostToDevice);
cuda_errcheck("memcpy to texDataArray");
texData.normalized = true;
texData.addressMode[0] = cudaAddressModeClamp;
texData.addressMode[1] = cudaAddressModeClamp;
texData.filterMode = cudaFilterModeLinear;
cudaBindTextureToArray(texData, ct);
cuda_errcheck("binding texDataArray to texture");
}
__global__ void testKernel(float4 *data, float *coord)
{
int gsize = blockDim.x * gridDim.x;
int gid = threadIdx.x + blockDim.x * blockIdx.x;
float norm_coord = (float)(gid + 32) / gsize;
data[gid] = tex1D(texData, norm_coord);
coord[gid] = norm_coord;
}
#define NUMTHREADS (TEXSIZE * 4)
int main(int argc, char *argv[])
{
float4 h_texData[TEXSIZE];
for (int t = 0; t < TEXSIZE; t++) {
h_texData[t].x = t + 10000;
h_texData[t].y = t + 20000;
h_texData[t].z = t + 30000;
h_texData[t].w = t + 40000;
}
float4 *h_output = new float4[NUMTHREADS];
float4 *d_output;
cudaMalloc(&d_output, NUMTHREADS * sizeof(float4));
cuda_errcheck("output malloc");
float *h_coord = new float[NUMTHREADS];
float *d_coord;
cudaMalloc(&d_coord, NUMTHREADS * sizeof(float));
cuda_errcheck("coord malloc");
cudaSetDevice(0);
cuda_errcheck("device init");
cuda_bind_texture_data(h_texData);
dim3 dimBlock(256, 1, 1);
dim3 dimGrid(NUMTHREADS / 256, 1, 1);
testKernel<<<dimGrid, dimBlock>>> (d_output, d_coord);
cuda_errcheck("kernel launch");
cudaThreadSynchronize();
cudaMemcpy(h_output, d_output, NUMTHREADS * sizeof(float4), cudaMemcpyDeviceToHost);
cuda_errcheck("output copy");
cudaMemcpy(h_coord, d_coord, NUMTHREADS * sizeof(float), cudaMemcpyDeviceToHost);
cuda_errcheck("coord copy");
for (int t = 0; t < NUMTHREADS; t++) {
printf("output[%d] (%.06f) = (%5.3f, %5.3f, %5.3f, %5.3f)\n", t, h_coord[t],
h_output[t].x, h_output[t].y, h_output[t].z, h_output[t].w);
}
return 0;
}
|
