/docs/MyDocs

To get this branch, use:
bzr branch http://darksoft.org/webbzr/docs/MyDocs

« back to all changes in this revision

Viewing changes to Development/libraries/cuda/examples/multiply/matrixMul.cu

  • Committer: Suren A. Chilingaryan
  • Date: 2009-04-09 03:21:08 UTC
  • Revision ID: csa@dside.dyndns.org-20090409032108-w4edamdh4adrgdu3
import

Show diffs side-by-side

added added

removed removed

Lines of Context:
Development/libraries/cuda/examples/multiply/matrixMul.cu
 
1
/*
 
2
 * Copyright 1993-2007 NVIDIA Corporation.  All rights reserved.
 
3
 *
 
4
 * NOTICE TO USER:
 
5
 *
 
6
 * This source code is subject to NVIDIA ownership rights under U.S. and
 
7
 * international Copyright laws.  Users and possessors of this source code
 
8
 * are hereby granted a nonexclusive, royalty-free license to use this code
 
9
 * in individual and commercial software.
 
10
 *
 
11
 * NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE
 
12
 * CODE FOR ANY PURPOSE.  IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR
 
13
 * IMPLIED WARRANTY OF ANY KIND.  NVIDIA DISCLAIMS ALL WARRANTIES WITH
 
14
 * REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF
 
15
 * MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
 
16
 * IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL,
 
17
 * OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS
 
18
 * OF USE, DATA OR PROFITS,  WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE
 
19
 * OR OTHER TORTIOUS ACTION,  ARISING OUT OF OR IN CONNECTION WITH THE USE
 
20
 * OR PERFORMANCE OF THIS SOURCE CODE.
 
21
 *
 
22
 * U.S. Government End Users.   This source code is a "commercial item" as
 
23
 * that term is defined at  48 C.F.R. 2.101 (OCT 1995), consisting  of
 
24
 * "commercial computer  software"  and "commercial computer software
 
25
 * documentation" as such terms are  used in 48 C.F.R. 12.212 (SEPT 1995)
 
26
 * and is provided to the U.S. Government only as a commercial end item.
 
27
 * Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through
 
28
 * 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the
 
29
 * source code with only those rights set forth herein.
 
30
 *
 
31
 * Any use of this source code in individual and commercial software must
 
32
 * include, in the user documentation and internal comments to the code,
 
33
 * the above Disclaimer and U.S. Government End Users Notice.
 
34
 */
 
35
 
 
36
/* Matrix multiplication: C = A * B.
 
37
 * Host code.
 
38
 *
 
39
 * This sample implements matrix multiplication and is exactly the same as
 
40
 * Chapter 7 of the programming guide.
 
41
 * It has been written for clarity of exposition to illustrate various CUDA
 
42
 * programming principles, not with the goal of providing the most
 
43
 * performant generic kernel for matrix multiplication.
 
44
 *
 
45
 * CUBLAS provides high-performance matrix multiplication.
 
46
 */
 
47
 
 
48
// includes, system
 
49
#include <stdlib.h>
 
50
#include <stdio.h>
 
51
#include <string.h>
 
52
#include <math.h>
 
53
 
 
54
// includes, project
 
55
#include <cutil.h>
 
56
 
 
57
// includes, kernels
 
58
#include <matrixMul_kernel.cu>
 
59
 
 
60
#define iters 1
 
61
 
 
62
////////////////////////////////////////////////////////////////////////////////
 
63
// declaration, forward
 
64
void runTest(int argc, char** argv);
 
65
void randomInit(float*, int, int, int);
 
66
void printDiff(float*, float*, int, int);
 
67
 
 
68
extern "C"
 
69
void computeGold(float*, const float*, const float*, unsigned int, unsigned int, unsigned int);
 
70
 
 
71
////////////////////////////////////////////////////////////////////////////////
 
72
// Program main
 
73
////////////////////////////////////////////////////////////////////////////////
 
74
int
 
75
main(int argc, char** argv)
 
76
{
 
77
    runTest(argc, argv);
 
78
 
 
79
//    CUT_EXIT(argc, argv);
 
80
}
 
81
 
 
82
////////////////////////////////////////////////////////////////////////////////
 
83
//! Run a simple test for CUDA
 
84
////////////////////////////////////////////////////////////////////////////////
 
85
 
 
86
 
 
87
void
 
88
runTest(int argc, char** argv)
 
89
{
 
90
    int i;
 
91
    CUT_DEVICE_INIT(argc, argv);
 
92
 
 
93
    // set seed for rand()
 
94
    srand(2006);
 
95
 
 
96
    // allocate host memory for matrices A and B
 
97
    unsigned int size_A = WA * HA;
 
98
    unsigned int mem_size_A = sizeof(float) * size_A;
 
99
    float* h_A = (float*) malloc(mem_size_A);
 
100
    unsigned int size_B = WB * HB;
 
101
    unsigned int mem_size_B = sizeof(float) * size_B;
 
102
    float* h_B = (float*) malloc(mem_size_B);
 
103
 
 
104
    // initialize host memory
 
105
    randomInit(h_A, size_A, WA, HA);
 
106
    randomInit(h_B, size_B, WB, HB);
 
107
 
 
108
 
 
109
 
 
110
    // allocate device memory
 
111
    float* d_A;
 
112
    CUDA_SAFE_CALL(cudaMalloc((void**) &d_A, mem_size_A));
 
113
    float* d_B;
 
114
    CUDA_SAFE_CALL(cudaMalloc((void**) &d_B, mem_size_B));
 
115
 
 
116
    // copy host memory to device
 
117
    CUDA_SAFE_CALL(cudaMemcpy(d_A, h_A, mem_size_A, cudaMemcpyHostToDevice) );
 
118
    CUDA_SAFE_CALL(cudaMemcpy(d_B, h_B, mem_size_B, cudaMemcpyHostToDevice) );
 
119
 
 
120
    float *d_T;
 
121
    CUDA_SAFE_CALL(cudaMalloc((void**) &d_T, mem_size_A + mem_size_B));
 
122
 
 
123
    tex.normalized = 0;
 
124
    CUDA_SAFE_CALL(cudaBindTexture(0, tex, d_T, mem_size_A + mem_size_B));
 
125
 
 
126
//    CUDA_SAFE_CALL(cudaMemset(d_T, 0, mem_size_A + mem_size_B) );
 
127
    CUDA_SAFE_CALL(cudaMemcpy(d_T         , h_A, mem_size_A, cudaMemcpyHostToDevice) );
 
128
    CUDA_SAFE_CALL(cudaMemcpy(d_T + size_A, h_B, mem_size_B, cudaMemcpyHostToDevice) );
 
129
 
 
130
 
 
131
 
 
132
    // allocate device memory for result
 
133
    unsigned int size_C = WC * HC;
 
134
    unsigned int mem_size_C = sizeof(float) * size_C;
 
135
    float* d_C;
 
136
    CUDA_SAFE_CALL(cudaMalloc((void**) &d_C, mem_size_C));
 
137
 
 
138
    // allocate host memory for the result
 
139
    float* h_C = (float*) malloc(mem_size_C);
 
140
    
 
141
    // create and start timer
 
142
    unsigned int timer = 0;
 
143
    cutCreateTimer(&timer);
 
144
 
 
145
    // setup execution parameters
 
146
    dim3 threads(BLOCK_SIZE, BLOCK_SIZE);
 
147
    dim3 grid(WC / threads.x, HC / threads.y);
 
148
/*
 
149
 
 
150
    testTex<<< grid, threads >>>(d_C, 0, size_A, WA, WB);
 
151
 
 
152
    CUDA_SAFE_CALL(cudaMemcpy(h_C, d_C, mem_size_C, cudaMemcpyDeviceToHost) );
 
153
 
 
154
    for (int k=0;k<10;k++) {
 
155
     for (int l=0;l<10;l++)
 
156
        printf("%f ", h_C[k*WA+l]);
 
157
     printf("\n");
 
158
    }
 
159
    printf("\n");
 
160
    exit(1);
 
161
*/
 
162
 
 
163
    // execute the kernel
 
164
    matrixTexSimpleMul<<< grid, threads >>>(d_C, 0, size_A, WA, WB);
 
165
    cudaThreadSynchronize();
 
166
 
 
167
    cutResetTimer(timer);
 
168
    cutStartTimer(timer);
 
169
    for (i=0;i<iters;++i) {
 
170
        matrixTexSimpleMul<<< grid, threads >>>(d_C, 0, size_A, WA, WB);
 
171
    }
 
172
    cudaThreadSynchronize();
 
173
    cutStopTimer(timer);
 
174
    printf("Processing time (simple,tex): %f (ms) \n", cutGetTimerValue(timer)/iters);
 
175
 
 
176
    // copy result from device to host
 
177
    CUDA_SAFE_CALL(cudaMemcpy(h_C, d_C, mem_size_C, cudaMemcpyDeviceToHost) );
 
178
 
 
179
    matrixSimpleMul<<< grid, threads >>>(d_C, d_A, d_B, WA, WB);
 
180
    cudaThreadSynchronize();
 
181
 
 
182
    // execute the kernel
 
183
    cutResetTimer(timer);
 
184
    cutStartTimer(timer);
 
185
    for (i=0;i<iters;++i) {
 
186
        matrixSimpleMul<<< grid, threads >>>(d_C, d_A, d_B, WA, WB);
 
187
    }
 
188
    cudaThreadSynchronize();
 
189
    cutStopTimer(timer);
 
190
    printf("Processing time (Simple): %f (ms) \n", cutGetTimerValue(timer) / iters);
 
191
 
 
192
 
 
193
 
 
194
    // execute the kernel
 
195
    matrixTexMul<<< grid, threads >>>(d_C, 0, size_A, WA, WB);
 
196
    cudaThreadSynchronize();
 
197
 
 
198
    cutResetTimer(timer);
 
199
    cutStartTimer(timer);
 
200
    for (i=0;i<iters;++i) {
 
201
        matrixTexMul<<< grid, threads >>>(d_C, 0, size_A, WA, WB);
 
202
    }
 
203
    cudaThreadSynchronize();
 
204
    cutStopTimer(timer);
 
205
    printf("Processing time (texture): %f (ms) \n", cutGetTimerValue(timer)/iters);
 
206
 
 
207
 
 
208
 
 
209
    matrixMul<<< grid, threads >>>(d_C, d_A, d_B, WA, WB);
 
210
    cudaThreadSynchronize();
 
211
 
 
212
    // execute the kernel
 
213
    cutResetTimer(timer);
 
214
    cutStartTimer(timer);
 
215
    for (i=0;i<iters;++i) {
 
216
        matrixMul<<< grid, threads >>>(d_C, d_A, d_B, WA, WB);
 
217
    }
 
218
    cudaThreadSynchronize();
 
219
    cutStopTimer(timer);
 
220
    printf("Processing time (Device): %f (ms) \n", cutGetTimerValue(timer) / iters);
 
221
 
 
222
    // check if kernel execution generated and error
 
223
    CUT_CHECK_ERROR("Kernel execution failed");
 
224
 
 
225
 
 
226
 
 
227
    // compute reference solution
 
228
    float* reference = (float*) malloc(mem_size_C);
 
229
/*
 
230
    cutResetTimer(timer);
 
231
    cutStartTimer(timer);
 
232
    computeGold(reference, h_A, h_B, HA, WA, WB);
 
233
    cutStopTimer(timer);
 
234
    printf("Processing time (CPU): %f (ms) \n", cutGetTimerValue(timer));
 
235
 
 
236
    // check result
 
237
    CUTBoolean res = cutCompareL2fe(reference, h_C, size_C, 1e-6f);
 
238
    printf("Test %s \n", (1 == res) ? "PASSED" : "FAILED");
 
239
    if (res!=1) printDiff(reference, h_C, WC, HC);
 
240
*/
 
241
 
 
242
    // clean up memory
 
243
    free(h_A);
 
244
    free(h_B);
 
245
    free(h_C);
 
246
    free(reference);
 
247
    CUDA_SAFE_CALL(cudaFree(d_A));
 
248
    CUDA_SAFE_CALL(cudaFree(d_B));
 
249
    CUDA_SAFE_CALL(cudaFree(d_C));
 
250
    CUDA_SAFE_CALL(cudaFree(d_T));
 
251
    CUT_SAFE_CALL(cutDeleteTimer(timer));
 
252
}
 
253
 
 
254
// Allocates a matrix with random float entries.
 
255
void randomInit(float* data, int size, int w, int h)
 
256
{
 
257
 
 
258
    for (int i = 0; i < w; i++)
 
259
        for (int j = 0; j < h; j++)
 
260
            data[i + j*w] = i + 0.0001 * j;
 
261
 
 
262
 
 
263
//    for (int i = 0; i < size; ++i)
 
264
//        data[i] = rand() / (float)RAND_MAX;
 
265
 
 
266
}
 
267
 
 
268
void printDiff(float *data1, float *data2, int width, int height)
 
269
{
 
270
  int i,j,k;
 
271
  int error_count=0;
 
272
  for (j=0; j<height; j++) {
 
273
    for (i=0; i<width; i++) {
 
274
      k = j*width+i;
 
275
      if (data1[k] != data2[k]) {
 
276
         printf("diff(%d,%d) CPU=%4.4f, GPU=%4.4f \n", i,j, data1[k], data2[k]);
 
277
         error_count++;
 
278
      }
 
279
    }
 
280
  }
 
281
  printf(" nTotal Errors = %d n", error_count);
 
282
}
 
283