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
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
| //*********************************************************************
//
// Multiply a vector by a float data
//
//*********************************************************************
// common SDK header for standard utilities and system libs
#include <oclUtils.h>
#include "oclVectorMul.h"
namespace liboclvecmul {
// Name of the file with the source code for the computation kernel
// *********************************************************************
const char* cSourceFile = "VectorMul.cl";
// Host buffers for demo
// *********************************************************************
void *src, *dst; // Host buffers for OpenCL test
// OpenCL Vars
cl_context cxGPUContext; // OpenCL context
cl_command_queue cqCommandQue; // OpenCL command que
cl_device_id* cdDevices; // OpenCL device list
cl_program cpProgram; // OpenCL program
cl_kernel ckKernel; // OpenCL kernel
cl_mem cmDevSrc; // OpenCL device source buffer
cl_mem cmDevDst; // OpenCL device destination buffer
size_t szGlobalWorkSize; // 1D var for Total # of work items
size_t szLocalWorkSize; // 1D var for # of work items in the work group
size_t szParmDataBytes; // Byte size of context information
size_t szKernelLength; // Byte size of kernel code
cl_int ciErr1, ciErr2; // Error code var
char* cPathAndName = NULL; // var for full paths to data, src, etc.
char* cSourceCL = NULL; // Buffer to hold source for compilation
// demo config vars
int iNumElements = 8388608;//16777216;//11444777; // Length of float arrays to process (odd # for illustration)
float value = 10.0f;
shrBOOL bQuickTest = shrFALSE;
// Forward Declarations
// *********************************************************************
void Cleanup (int iExitCode);
// Main function
// *********************************************************************
int multiplie(int argc, char **argv)
{
// get command line arg for quick test, if provided
bQuickTest = shrCheckCmdLineFlag(argc, (const char**)argv, "noprompt");
// start logs
shrSetLogFileName ("oclVectorMul.txt");
// shrLog(LOGBOTH, 0.0, "%s Starting...\n\n# of float elements per Array \t= %u\n", argv[0], iNumElements);
// set and log Global and Local work size dimensions
szLocalWorkSize = 256;
szGlobalWorkSize = shrRoundUp((int)szLocalWorkSize, iNumElements); // rounded up to the nearest multiple of the LocalWorkSize
// shrLog(LOGBOTH, 0.0, "Global Work Size \t\t= %u\nLocal Work Size \t\t= %u\n# of Work Groups \t\t= %u\n\n",
// szGlobalWorkSize, szLocalWorkSize, (szGlobalWorkSize % szLocalWorkSize + szGlobalWorkSize/szLocalWorkSize));
// Allocate and initialize host arrays
src = (void *)malloc(sizeof(cl_float) * szGlobalWorkSize);
dst = (void *)malloc(sizeof(cl_float) * szGlobalWorkSize);
shrFillArray((float*)src, iNumElements);
// shrLog(LOGBOTH, 0.0, "Allocate and Init Host Mem...\n");
// Create the OpenCL context on a GPU device
cxGPUContext = clCreateContextFromType(0, CL_DEVICE_TYPE_GPU, NULL, NULL, &ciErr1);
// shrLog(LOGBOTH, 0.0, "clCreateContextFromType...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog(LOGBOTH, 0.0, "Error in clCreateContextFromType, near Line %u in file %u", __LINE__, __FILE__);
Cleanup(EXIT_FAILURE);
}
// Get the list of GPU devices associated with context
ciErr1 = clGetContextInfo(cxGPUContext, CL_CONTEXT_DEVICES, 0, NULL, &szParmDataBytes);
cdDevices = (cl_device_id*)malloc(szParmDataBytes);
ciErr1 |= clGetContextInfo(cxGPUContext, CL_CONTEXT_DEVICES, szParmDataBytes, cdDevices, NULL);
// shrLog(LOGBOTH, 0.0, "clGetContextInfo...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog(LOGBOTH, 0.0, "Error in clGetContextInfo, near Line %u in file %u", __LINE__, __FILE__);
Cleanup(EXIT_FAILURE);
}
// Create a command-queue
cqCommandQue = clCreateCommandQueue(cxGPUContext, cdDevices[0], 0, &ciErr1);
// shrLog(LOGBOTH, 0.0, "clCreateCommandQueue...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog(LOGBOTH, 0.0, "Error in clCreateCommandQueue, near Line %u in file %u", __LINE__, __FILE__);
Cleanup(EXIT_FAILURE);
}
// Allocate the OpenCL source and result buffer memory objects on the device GMEM, and copy the data to the device
shrDeltaT(0);
cmDevSrc = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_float) * szGlobalWorkSize, src, &ciErr1);
double memCpyHDTime = shrDeltaT(0);
cmDevDst = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, sizeof(cl_float) * szGlobalWorkSize, NULL, &ciErr2);
ciErr1 |= ciErr2;
// shrLog(LOGBOTH, 0.0, "clCreateBuffer...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog(LOGBOTH, 0.0, "Error in clCreateBuffer, near Line %u in file %u", __LINE__, __FILE__);
Cleanup(EXIT_FAILURE);
}
// Read the OpenCL kernel in from source file
cPathAndName = shrFindFilePath(cSourceFile, argv[0]);
cSourceCL = oclLoadProgSource(cPathAndName, "", &szKernelLength);
// shrLog(LOGBOTH, 0.0, "oclLoadProgSource (%s)...\n", cSourceFile);
// Create the program
cpProgram = clCreateProgramWithSource(cxGPUContext, 1, (const char **)&cSourceCL, &szKernelLength, &ciErr1);
// shrLog(LOGBOTH, 0.0, "clCreateProgramWithSource...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog(LOGBOTH, 0.0, "Error in clCreateProgramWithSource, near Line %u in file %u", __LINE__, __FILE__);
Cleanup(EXIT_FAILURE);
}
// Build the program
ciErr1 = clBuildProgram(cpProgram, 0, NULL, NULL, NULL, NULL);
// shrLog(LOGBOTH, 0.0, "clBuildProgram...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog(LOGBOTH, 0.0, "Error in clBuildProgram, near Line %u in file %u", __LINE__, __FILE__);
Cleanup(EXIT_FAILURE);
}
// Create the kernel
ckKernel = clCreateKernel(cpProgram, "VectorMul", &ciErr1);
// shrLog(LOGBOTH, 0.0, "clCreateKernel...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog(LOGBOTH, 0.0, "Error in clCreateKernel, near Line %u in file %u", __LINE__, __FILE__);
Cleanup(EXIT_FAILURE);
}
// Set the Argument values
ciErr1 = clSetKernelArg(ckKernel, 0, sizeof(cl_mem), (void*)&cmDevSrc);
ciErr1 |= clSetKernelArg(ckKernel, 1, sizeof(cl_float), (void*)&value);
ciErr1 |= clSetKernelArg(ckKernel, 2, sizeof(cl_mem), (void*)&cmDevDst);
ciErr1 |= clSetKernelArg(ckKernel, 3, sizeof(cl_int), (void*)&iNumElements);
// shrLog(LOGBOTH, 0.0, "clSetKernelArg...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog(LOGBOTH, 0.0, "Error in clSetKernelArg, near Line %u in file %u", __LINE__, __FILE__);
Cleanup(EXIT_FAILURE);
}
// Launch kernel
shrDeltaT(0);
ciErr1 = clEnqueueNDRangeKernel(cqCommandQue, ckKernel, 1, NULL, &szGlobalWorkSize, &szLocalWorkSize, 0, NULL, NULL);
double computeTime = shrDeltaT(0);
// shrLog(LOGBOTH, 0.0, "clEnqueueNDRangeKernel...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog(LOGBOTH, 0.0, "Error in clEnqueueNDRangeKernel, near Line %u in file %u", __LINE__, __FILE__);
Cleanup(EXIT_FAILURE);
}
// Synchronous/blocking read of results, and check accumulated errors
shrDeltaT(0);
ciErr1 = clEnqueueReadBuffer(cqCommandQue, cmDevDst, CL_TRUE, 0, sizeof(cl_float) * szGlobalWorkSize, dst, 0, NULL, NULL);
double memCpyDHTime = shrDeltaT(0);
// shrLog(LOGBOTH, 0.0, "clEnqueueReadBuffer...\n\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog(LOGBOTH, 0.0, "Error in clEnqueueReadBuffer, near Line %u in file %u", __LINE__, __FILE__);
Cleanup(EXIT_FAILURE);
}
printf("\nMultiplication\n\n\t*%d datas\n\t*%d Global Work Size\n\t*%d Local Work Size\n", iNumElements, (int)szGlobalWorkSize, (int)szLocalWorkSize);
printf("\n=============================================\n");
printf("Time to copy datas HOST -> DEVICE : %.3f ms\n", (memCpyHDTime*1000));
printf("Time to compute : %.3f ms\n", (computeTime*1000));
printf("Time to copy datas DEVICE -> HOST : %.3f ms\n", (memCpyDHTime*1000));
printf("--------------------------------------------\n");
printf("Total time : %.3f ms\n", ((memCpyHDTime+computeTime+memCpyDHTime)*1000));
printf("=============================================\n\n");
float *srcf, *dstf;
srcf = (float*)src;
dstf = (float*)dst;
for (int i=0; i<iNumElements; i++)
{
if (srcf[i]*10 != dstf[i])
printf("Error at indice %d\n", i);
}
// Cleanup and leave
Cleanup (EXIT_SUCCESS);
}
void Cleanup (int iExitCode)
{
// Cleanup allocated objects
shrLog(LOGBOTH, 0.0, "\nStarting Cleanup...\n\n");
if(cdDevices)free(cdDevices);
if(cPathAndName)free(cPathAndName);
if(cSourceCL)free(cSourceCL);
if(ckKernel)clReleaseKernel(ckKernel);
if(cpProgram)clReleaseProgram(cpProgram);
if(cqCommandQue)clReleaseCommandQueue(cqCommandQue);
if(cxGPUContext)clReleaseContext(cxGPUContext);
if(cmDevSrc)clReleaseMemObject(cmDevSrc);
if(cmDevDst)clReleaseMemObject(cmDevDst);
// Free host memory
free(src);
free (dst);
// finalize logs and leave
if (bQuickTest)
{
shrLog(LOGBOTH | CLOSELOG, 0.0, "oclVectorMul Ending...\n");
}
else
{
shrLog(LOGBOTH | CLOSELOG, 0.0, "oclVectorMul Ending...\nPress Enter to Exit\n");
getchar();
}
exit (iExitCode);
}
} // namespace |
OpenCL et cmake
Répondre avec citation
Partager