func main() {
var i, j cl.CL_size_t
// Rows and columns in the input image
inputFile := "test.png"
outputFile := "output.png"
refFile := "ref.png"
// Homegrown function to read a BMP from file
inputpixels, imageWidth, imageHeight, err1 := utils.Read_image_data(inputFile)
if err1 != nil {
log.Fatal(err1)
return
} else {
fmt.Printf("width=%d, height=%d (%d)\n", imageWidth, imageHeight, inputpixels[0])
}
// Output image on the host
outputpixels := make([]uint16, imageHeight*imageWidth)
inputImage := make([]float32, imageHeight*imageWidth)
outputImage := make([]float32, imageHeight*imageWidth)
refImage := make([]float32, imageHeight*imageWidth)
for i = 0; i < imageHeight*imageWidth; i++ {
inputImage[i] = float32(inputpixels[i])
}
// 45 degree motion blur
var filter = [49]float32{0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0,
0, 0, -1, 0, 1, 0, 0,
0, 0, -2, 0, 2, 0, 0,
0, 0, -1, 0, 1, 0, 0,
0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0}
// The convolution filter is 7x7
filterWidth := cl.CL_size_t(7)
filterSize := cl.CL_size_t(filterWidth * filterWidth) // Assume a square kernel
// Set up the OpenCL environment
var status cl.CL_int
// Discovery platform
var platform [1]cl.CL_platform_id
status = cl.CLGetPlatformIDs(1, platform[:], nil)
chk(status, "clGetPlatformIDs")
// Discover device
var device [1]cl.CL_device_id
cl.CLGetDeviceIDs(platform[0], cl.CL_DEVICE_TYPE_ALL, 1, device[:], nil)
chk(status, "clGetDeviceIDs")
// Create context
//var props =[3]cl.CL_context_properties{cl.CL_CONTEXT_PLATFORM,
// (cl.CL_context_properties)(unsafe.Pointer(&platform[0])), 0};
var context cl.CL_context
context = cl.CLCreateContext(nil, 1, device[:], nil, nil, &status)
chk(status, "clCreateContext")
// Create command queue
var queue cl.CL_command_queue
queue = cl.CLCreateCommandQueue(context, device[0], 0, &status)
chk(status, "clCreateCommandQueue")
// The image format describes how the data will be stored in memory
var format cl.CL_image_format
format.Image_channel_order = cl.CL_R // single channel
format.Image_channel_data_type = cl.CL_FLOAT // float data type
var desc cl.CL_image_desc
desc.Image_type = cl.CL_MEM_OBJECT_IMAGE2D
desc.Image_width = imageWidth
desc.Image_height = imageHeight
desc.Image_depth = 0
desc.Image_array_size = 0
desc.Image_row_pitch = 0
desc.Image_slice_pitch = 0
desc.Num_mip_levels = 0
desc.Num_samples = 0
desc.Buffer = cl.CL_mem{}
// Create space for the source image on the device
d_inputImage := cl.CLCreateImage(context, cl.CL_MEM_READ_ONLY, &format, &desc,
nil, &status)
chk(status, "clCreateImage")
// Create space for the output image on the device
d_outputImage := cl.CLCreateImage(context, cl.CL_MEM_WRITE_ONLY, &format, &desc,
nil, &status)
chk(status, "clCreateImage")
// Create space for the 7x7 filter on the device
d_filter := cl.CLCreateBuffer(context, 0, filterSize*cl.CL_size_t(unsafe.Sizeof(filter[0])),
nil, &status)
chk(status, "clCreateBuffer")
// Copy the source image to the device
var origin = [3]cl.CL_size_t{0, 0, 0} // Offset within the image to copy from
var region = [3]cl.CL_size_t{cl.CL_size_t(imageWidth), cl.CL_size_t(imageHeight), 1} // Elements to per dimension
status = cl.CLEnqueueWriteImage(queue, d_inputImage, cl.CL_FALSE, origin, region,
0, 0, unsafe.Pointer(&inputImage[0]), 0, nil, nil)
chk(status, "clEnqueueWriteImage")
// Copy the 7x7 filter to the device
status = cl.CLEnqueueWriteBuffer(queue, d_filter, cl.CL_FALSE, 0,
filterSize*cl.CL_size_t(unsafe.Sizeof(filter[0])), unsafe.Pointer(&filter[0]), 0, nil, nil)
chk(status, "clEnqueueWriteBuffer")
// Create the image sampler
sampler := cl.CLCreateSampler(context, cl.CL_FALSE,
cl.CL_ADDRESS_CLAMP_TO_EDGE, cl.CL_FILTER_NEAREST, &status)
chk(status, "clCreateSampler")
// Create a program object with source and build it
program := utils.Build_program(context, device[:], "convolution.cl", nil)
kernel := cl.CLCreateKernel(*program, []byte("convolution"), &status)
chk(status, "clCreateKernel")
// Set the kernel arguments
var w, h, f cl.CL_int
w = cl.CL_int(imageWidth)
h = cl.CL_int(imageHeight)
f = cl.CL_int(filterWidth)
status = cl.CLSetKernelArg(kernel, 0, cl.CL_size_t(unsafe.Sizeof(d_inputImage)), unsafe.Pointer(&d_inputImage))
status |= cl.CLSetKernelArg(kernel, 1, cl.CL_size_t(unsafe.Sizeof(d_outputImage)), unsafe.Pointer(&d_outputImage))
status |= cl.CLSetKernelArg(kernel, 2, cl.CL_size_t(unsafe.Sizeof(h)), unsafe.Pointer(&h))
status |= cl.CLSetKernelArg(kernel, 3, cl.CL_size_t(unsafe.Sizeof(w)), unsafe.Pointer(&w))
status |= cl.CLSetKernelArg(kernel, 4, cl.CL_size_t(unsafe.Sizeof(d_filter)), unsafe.Pointer(&d_filter))
status |= cl.CLSetKernelArg(kernel, 5, cl.CL_size_t(unsafe.Sizeof(f)), unsafe.Pointer(&f))
status |= cl.CLSetKernelArg(kernel, 6, cl.CL_size_t(unsafe.Sizeof(sampler)), unsafe.Pointer(&sampler))
chk(status, "clSetKernelArg")
// Set the work item dimensions
var globalSize = [2]cl.CL_size_t{imageWidth, imageHeight}
status = cl.CLEnqueueNDRangeKernel(queue, kernel, 2, nil, globalSize[:], nil, 0,
nil, nil)
chk(status, "clEnqueueNDRange")
// Read the image back to the host
status = cl.CLEnqueueReadImage(queue, d_outputImage, cl.CL_TRUE, origin,
region, 0, 0, unsafe.Pointer(&outputImage[0]), 0, nil, nil)
chk(status, "clEnqueueReadImage")
// Write the output image to file
for i = 0; i < imageHeight*imageWidth; i++ {
outputpixels[i] = uint16(outputImage[i])
}
utils.Write_image_data(outputFile, outputpixels, imageWidth, imageHeight)
// Compute the reference image
for i = 0; i < imageHeight; i++ {
for j = 0; j < imageWidth; j++ {
refImage[i*imageWidth+j] = 0
}
}
// Iterate over the rows of the source image
halfFilterWidth := filterWidth / 2
var sum float32
for i = 0; i < imageHeight; i++ {
// Iterate over the columns of the source image
for j = 0; j < imageWidth; j++ {
sum = 0 // Reset sum for new source pixel
// Apply the filter to the neighborhood
for k := -halfFilterWidth; k <= halfFilterWidth; k++ {
for l := -halfFilterWidth; l <= halfFilterWidth; l++ {
if i+k >= 0 && i+k < imageHeight &&
j+l >= 0 && j+l < imageWidth {
sum += inputImage[(i+k)*imageWidth+j+l] *
filter[(k+halfFilterWidth)*filterWidth+
l+halfFilterWidth]
} else {
i_k := i + k
j_l := j + l
if i+k < 0 {
i_k = 0
} else if i+k >= imageHeight {
i_k = imageHeight - 1
}
if j+l < 0 {
j_l = 0
} else if j+l >= imageWidth {
j_l = imageWidth - 1
}
sum += inputImage[(i_k)*imageWidth+j_l] *
filter[(k+halfFilterWidth)*filterWidth+
l+halfFilterWidth]
}
}
}
refImage[i*imageWidth+j] = sum
}
}
// Write the ref image to file
for i = 0; i < imageHeight*imageWidth; i++ {
outputpixels[i] = uint16(refImage[i])
}
utils.Write_image_data(refFile, outputpixels, imageWidth, imageHeight)
failed := 0
for i = 0; i < imageHeight; i++ {
for j = 0; j < imageWidth; j++ {
if math.Abs(float64(outputImage[i*imageWidth+j]-refImage[i*imageWidth+j])) > 0.01 {
//fmt.Printf("Results are INCORRECT\n");
//fmt.Printf("Pixel mismatch at <%d,%d> (%f vs. %f) %f\n", i, j,
// outputImage[i*imageWidth+j], refImage[i*imageWidth+j], inputImage[i*imageWidth+j]);
failed++
}
}
}
fmt.Printf("Mismatch Pixel number/Total pixel number = %d/%d\n", failed, imageWidth*imageHeight)
// Free OpenCL resources
cl.CLReleaseKernel(kernel)
cl.CLReleaseProgram(*program)
cl.CLReleaseCommandQueue(queue)
cl.CLReleaseMemObject(d_inputImage)
cl.CLReleaseMemObject(d_outputImage)
cl.CLReleaseMemObject(d_filter)
cl.CLReleaseSampler(sampler)
cl.CLReleaseContext(context)
}
func TestMatvec(t *testing.T) {
/* Host/device data structures */
var platform [1]cl.CL_platform_id
var device [1]cl.CL_device_id
var context cl.CL_context
var queue cl.CL_command_queue
var i, err cl.CL_int
/* Program/kernel data structures */
var program cl.CL_program
var program_buffer [1][]byte
var program_log interface{}
var program_size [1]cl.CL_size_t
var log_size cl.CL_size_t
var kernel cl.CL_kernel
/* Data and buffers */
var mat [16]float32
var vec, result [4]float32
var correct = [4]float32{0.0, 0.0, 0.0, 0.0}
var mat_buff, vec_buff, res_buff cl.CL_mem
/* Initialize data to be processed by the kernel */
for i = 0; i < 16; i++ {
mat[i] = float32(i) * 2.0
}
for i = 0; i < 4; i++ {
vec[i] = float32(i) * 3.0
correct[0] += mat[i] * vec[i]
correct[1] += mat[i+4] * vec[i]
correct[2] += mat[i+8] * vec[i]
correct[3] += mat[i+12] * vec[i]
}
/* Identify a platform */
err = cl.CLGetPlatformIDs(1, platform[:], nil)
if err < 0 {
t.Errorf("Couldn't find any platforms")
}
/* Access a device */
err = cl.CLGetDeviceIDs(platform[0], cl.CL_DEVICE_TYPE_GPU, 1, device[:], nil)
if err < 0 {
err = cl.CLGetDeviceIDs(platform[0], cl.CL_DEVICE_TYPE_CPU, 1, device[:], nil)
if err < 0 {
t.Errorf("Couldn't find any devices")
}
}
/* Create the context */
context = cl.CLCreateContext(nil, 1, device[:], nil, nil, &err)
if err < 0 {
t.Errorf("Couldn't create a context")
}
/* Read program file and place content into buffer */
program_handle, err1 := os.Open("matvec.cl")
if err1 != nil {
t.Errorf("Couldn't find the program file")
}
defer program_handle.Close()
fi, err2 := program_handle.Stat()
if err2 != nil {
t.Errorf("Couldn't find the program stat")
}
program_size[0] = cl.CL_size_t(fi.Size())
program_buffer[0] = make([]byte, program_size[0])
read_size, err3 := program_handle.Read(program_buffer[0])
if err3 != nil || cl.CL_size_t(read_size) != program_size[0] {
t.Errorf("read file error or file size wrong")
}
/* Create a program containing all program content */
program = cl.CLCreateProgramWithSource(context, 1,
program_buffer[:], program_size[:], &err)
if err < 0 {
t.Errorf("Couldn't create the program")
}
/* Build program */
err = cl.CLBuildProgram(program, 1, device[:], nil, nil, nil)
if err < 0 {
/* Find size of log and print to std output */
cl.CLGetProgramBuildInfo(program, device[0], cl.CL_PROGRAM_BUILD_LOG,
0, nil, &log_size)
//program_log = (char*) malloc(log_size+1);
//program_log[log_size] = '\0';
cl.CLGetProgramBuildInfo(program, device[0], cl.CL_PROGRAM_BUILD_LOG,
log_size, &program_log, nil)
t.Errorf("%s\n", program_log)
//free(program_log);
}
/* Create kernel for the mat_vec_mult function */
kernel = cl.CLCreateKernel(program, []byte("matvec_mult"), &err)
if err < 0 {
t.Errorf("Couldn't create the kernel")
return
}
/* Create CL buffers to hold input and output data */
mat_buff = cl.CLCreateBuffer(context, cl.CL_MEM_READ_ONLY|
cl.CL_MEM_COPY_HOST_PTR, cl.CL_size_t(unsafe.Sizeof(mat)), unsafe.Pointer(&mat[0]), &err)
if err < 0 {
t.Errorf("Couldn't create a buffer object")
return
}
vec_buff = cl.CLCreateBuffer(context, cl.CL_MEM_READ_ONLY|
cl.CL_MEM_COPY_HOST_PTR, cl.CL_size_t(unsafe.Sizeof(vec)), unsafe.Pointer(&vec[0]), nil)
res_buff = cl.CLCreateBuffer(context, cl.CL_MEM_WRITE_ONLY,
cl.CL_size_t(unsafe.Sizeof(result)), nil, nil)
/* Create kernel arguments from the CL buffers */
err = cl.CLSetKernelArg(kernel, 0, cl.CL_size_t(unsafe.Sizeof(mat_buff)), unsafe.Pointer(&mat_buff))
if err < 0 {
t.Errorf("Couldn't set the kernel argument")
return
}
cl.CLSetKernelArg(kernel, 1, cl.CL_size_t(unsafe.Sizeof(vec_buff)), unsafe.Pointer(&vec_buff))
cl.CLSetKernelArg(kernel, 2, cl.CL_size_t(unsafe.Sizeof(res_buff)), unsafe.Pointer(&res_buff))
/* Create a CL command queue for the device*/
queue = cl.CLCreateCommandQueue(context, device[0], 0, &err)
if err < 0 {
t.Errorf("Couldn't create the command queue, errcode=%d\n", err)
return
}
/* Enqueue the command queue to the device */
var work_units_per_kernel = [1]cl.CL_size_t{4} /* 4 work-units per kernel */
err = cl.CLEnqueueNDRangeKernel(queue, kernel, 1, nil, work_units_per_kernel[:],
nil, 0, nil, nil)
if err < 0 {
t.Errorf("Couldn't enqueue the kernel execution command, errcode=%d\n", err)
return
}
/* Read the result */
err = cl.CLEnqueueReadBuffer(queue, res_buff, cl.CL_TRUE, 0, cl.CL_size_t(unsafe.Sizeof(result)),
unsafe.Pointer(&result[0]), 0, nil, nil)
if err < 0 {
t.Errorf("Couldn't enqueue the read buffer command")
return
}
/* Test the result */
if (result[0] == correct[0]) && (result[1] == correct[1]) &&
(result[2] == correct[2]) && (result[3] == correct[3]) {
t.Logf("Matrix-vector multiplication successful.")
} else {
t.Errorf("Matrix-vector multiplication unsuccessful.")
}
/* Deallocate resources */
cl.CLReleaseMemObject(mat_buff)
cl.CLReleaseMemObject(vec_buff)
cl.CLReleaseMemObject(res_buff)
cl.CLReleaseKernel(kernel)
cl.CLReleaseCommandQueue(queue)
cl.CLReleaseProgram(program)
cl.CLReleaseContext(context)
}
func main() {
/* OpenCL data structures */
var device []cl.CL_device_id
var context cl.CL_context
var queue cl.CL_command_queue
var program *cl.CL_program
var kernel cl.CL_kernel
var err cl.CL_int
/* Data and buffers */
var select1 [4]float32
var select2 [2]cl.CL_uchar
var select1_buffer, select2_buffer cl.CL_mem
/* Create a context */
device = utils.Create_device()
context = cl.CLCreateContext(nil, 1, device[:], nil, nil, &err)
if err < 0 {
println("Couldn't create a context")
return
}
/* Create a kernel */
program = utils.Build_program(context, device[:], PROGRAM_FILE, nil)
kernel = cl.CLCreateKernel(*program, KERNEL_FUNC, &err)
if err < 0 {
println("Couldn't create a kernel")
return
}
/* Create a write-only buffer to hold the output data */
select1_buffer = cl.CLCreateBuffer(context, cl.CL_MEM_WRITE_ONLY,
cl.CL_size_t(unsafe.Sizeof(select1)), nil, &err)
if err < 0 {
println("Couldn't create a buffer")
return
}
select2_buffer = cl.CLCreateBuffer(context, cl.CL_MEM_WRITE_ONLY,
cl.CL_size_t(unsafe.Sizeof(select2)), nil, &err)
/* Create kernel argument */
err = cl.CLSetKernelArg(kernel, 0, cl.CL_size_t(unsafe.Sizeof(select1_buffer)), unsafe.Pointer(&select1_buffer))
if err < 0 {
println("Couldn't set a kernel argument")
return
}
cl.CLSetKernelArg(kernel, 1, cl.CL_size_t(unsafe.Sizeof(select2_buffer)), unsafe.Pointer(&select2_buffer))
/* Create a command queue */
queue = cl.CLCreateCommandQueue(context, device[0], 0, &err)
if err < 0 {
println("Couldn't create a command queue")
return
}
/* Enqueue kernel */
err = cl.CLEnqueueTask(queue, kernel, 0, nil, nil)
if err < 0 {
println("Couldn't enqueue the kernel")
return
}
/* Read and print the result */
err = cl.CLEnqueueReadBuffer(queue, select1_buffer, cl.CL_TRUE, 0,
cl.CL_size_t(unsafe.Sizeof(select1)), unsafe.Pointer(&select1), 0, nil, nil)
if err < 0 {
println("Couldn't read the buffer")
return
}
cl.CLEnqueueReadBuffer(queue, select2_buffer, cl.CL_TRUE, 0,
cl.CL_size_t(unsafe.Sizeof(select2)), unsafe.Pointer(&select2), 0, nil, nil)
fmt.Printf("select: ")
for i := 0; i < 3; i++ {
fmt.Printf("%.2f, ", select1[i])
}
fmt.Printf("%.2f\n", select1[3])
fmt.Printf("bitselect: %X, %X\n", select2[0], select2[1])
/* Deallocate resources */
cl.CLReleaseMemObject(select1_buffer)
cl.CLReleaseMemObject(select2_buffer)
cl.CLReleaseKernel(kernel)
cl.CLReleaseCommandQueue(queue)
cl.CLReleaseProgram(*program)
cl.CLReleaseContext(context)
}
func main() {
// This code executes on the OpenCL host
// Host data
var size cl.CL_int
var A []cl.CL_int //input array
var B []cl.CL_int //input array
var C []cl.CL_int //output array
// Elements in each array
const elements = cl.CL_size_t(2048)
// Compute the size of the data
datasize := cl.CL_size_t(unsafe.Sizeof(size)) * elements
// Allocate space for input/output data
A = make([]cl.CL_int, datasize)
B = make([]cl.CL_int, datasize)
C = make([]cl.CL_int, datasize)
// Initialize the input data
for i := cl.CL_int(0); i < cl.CL_int(elements); i++ {
A[i] = i
B[i] = i
}
// Use this to check the output of each API call
var status cl.CL_int
//-----------------------------------------------------
// STEP 1: Discover and initialize the platforms
//-----------------------------------------------------
var numPlatforms cl.CL_uint
var platforms []cl.CL_platform_id
// Use clGetPlatformIDs() to retrieve the number of
// platforms
status = cl.CLGetPlatformIDs(0, nil, &numPlatforms)
// Allocate enough space for each platform
platforms = make([]cl.CL_platform_id, numPlatforms)
// Fill in platforms with clGetPlatformIDs()
status = cl.CLGetPlatformIDs(numPlatforms, platforms, nil)
if status != cl.CL_SUCCESS {
println("CLGetPlatformIDs status!=cl.CL_SUCCESS")
return
}
//-----------------------------------------------------
// STEP 2: Discover and initialize the devices
//-----------------------------------------------------
var numDevices cl.CL_uint
var devices []cl.CL_device_id
// Use clGetDeviceIDs() to retrieve the number of
// devices present
status = cl.CLGetDeviceIDs(platforms[0],
cl.CL_DEVICE_TYPE_ALL,
0,
nil,
&numDevices)
if status != cl.CL_SUCCESS {
println("CLGetDeviceIDs status!=cl.CL_SUCCESS")
return
}
// Allocate enough space for each device
devices = make([]cl.CL_device_id, numDevices)
// Fill in devices with clGetDeviceIDs()
status = cl.CLGetDeviceIDs(platforms[0],
cl.CL_DEVICE_TYPE_ALL,
numDevices,
devices,
nil)
if status != cl.CL_SUCCESS {
println("CLGetDeviceIDs status!=cl.CL_SUCCESS")
return
}
//-----------------------------------------------------
// STEP 3: Create a context
//-----------------------------------------------------
var context cl.CL_context
// Create a context using clCreateContext() and
// associate it with the devices
context = cl.CLCreateContext(nil,
numDevices,
devices,
nil,
nil,
&status)
if status != cl.CL_SUCCESS {
println("CLCreateContext status!=cl.CL_SUCCESS")
return
}
//-----------------------------------------------------
// STEP 4: Create a command queue
//-----------------------------------------------------
var cmdQueue cl.CL_command_queue
// Create a command queue using clCreateCommandQueue(),
// and associate it with the device you want to execute
// on
cmdQueue = cl.CLCreateCommandQueue(context,
devices[0],
0,
&status)
if status != cl.CL_SUCCESS {
println("CLCreateCommandQueue status!=cl.CL_SUCCESS")
return
}
//-----------------------------------------------------
// STEP 5: Create device buffers
//-----------------------------------------------------
var bufferA cl.CL_mem // Input array on the device
var bufferB cl.CL_mem // Input array on the device
var bufferC cl.CL_mem // Output array on the device
// Use clCreateBuffer() to create a buffer object (d_A)
// that will contain the data from the host array A
bufferA = cl.CLCreateBuffer(context,
cl.CL_MEM_READ_ONLY,
datasize,
nil,
&status)
if status != cl.CL_SUCCESS {
println("CLCreateBuffer status!=cl.CL_SUCCESS")
return
}
// Use clCreateBuffer() to create a buffer object (d_B)
// that will contain the data from the host array B
bufferB = cl.CLCreateBuffer(context,
cl.CL_MEM_READ_ONLY,
datasize,
nil,
&status)
if status != cl.CL_SUCCESS {
println("CLCreateBuffer status!=cl.CL_SUCCESS")
return
}
// Use clCreateBuffer() to create a buffer object (d_C)
// with enough space to hold the output data
bufferC = cl.CLCreateBuffer(context,
cl.CL_MEM_WRITE_ONLY,
datasize,
nil,
&status)
if status != cl.CL_SUCCESS {
println("CLCreateBuffer status!=cl.CL_SUCCESS")
return
}
//-----------------------------------------------------
// STEP 6: Write host data to device buffers
//-----------------------------------------------------
// Use clEnqueueWriteBuffer() to write input array A to
// the device buffer bufferA
status = cl.CLEnqueueWriteBuffer(cmdQueue,
bufferA,
cl.CL_FALSE,
0,
datasize,
unsafe.Pointer(&A[0]),
0,
nil,
nil)
if status != cl.CL_SUCCESS {
println("CLEnqueueWriteBuffer status!=cl.CL_SUCCESS")
return
}
// Use clEnqueueWriteBuffer() to write input array B to
// the device buffer bufferB
status = cl.CLEnqueueWriteBuffer(cmdQueue,
bufferB,
cl.CL_FALSE,
0,
datasize,
unsafe.Pointer(&B[0]),
0,
nil,
nil)
if status != cl.CL_SUCCESS {
println("CLEnqueueWriteBuffer status!=cl.CL_SUCCESS")
return
}
//-----------------------------------------------------
// STEP 7: Create and compile the program
//-----------------------------------------------------
programSource, programeSize := utils.Load_programsource("chapter2.cl")
// Create a program using clCreateProgramWithSource()
program := cl.CLCreateProgramWithSource(context,
1,
programSource[:],
programeSize[:],
&status)
if status != cl.CL_SUCCESS {
println("CLCreateProgramWithSource status!=cl.CL_SUCCESS")
return
}
// Build (compile) the program for the devices with
// clBuildProgram()
status = cl.CLBuildProgram(program,
numDevices,
devices,
nil,
nil,
nil)
if status != cl.CL_SUCCESS {
println("CLBuildProgram status!=cl.CL_SUCCESS")
return
}
//-----------------------------------------------------
// STEP 8: Create the kernel
//-----------------------------------------------------
var kernel cl.CL_kernel
// Use clCreateKernel() to create a kernel from the
// vector addition function (named "vecadd")
kernel = cl.CLCreateKernel(program, []byte("vecadd"), &status)
if status != cl.CL_SUCCESS {
println("CLCreateKernel status!=cl.CL_SUCCESS")
return
}
//-----------------------------------------------------
// STEP 9: Set the kernel arguments
//-----------------------------------------------------
// Associate the input and output buffers with the
// kernel
// using clSetKernelArg()
status = cl.CLSetKernelArg(kernel,
0,
cl.CL_size_t(unsafe.Sizeof(bufferA)),
unsafe.Pointer(&bufferA))
status |= cl.CLSetKernelArg(kernel,
1,
cl.CL_size_t(unsafe.Sizeof(bufferB)),
unsafe.Pointer(&bufferB))
status |= cl.CLSetKernelArg(kernel,
2,
cl.CL_size_t(unsafe.Sizeof(bufferC)),
unsafe.Pointer(&bufferC))
if status != cl.CL_SUCCESS {
println("CLSetKernelArg status!=cl.CL_SUCCESS")
return
}
//-----------------------------------------------------
// STEP 10: Configure the work-item structure
//-----------------------------------------------------
// Define an index space (global work size) of work
// items for
// execution. A workgroup size (local work size) is not
// required,
// but can be used.
var globalWorkSize [1]cl.CL_size_t
// There are 'elements' work-items
globalWorkSize[0] = elements
//-----------------------------------------------------
// STEP 11: Enqueue the kernel for execution
//-----------------------------------------------------
// Execute the kernel by using
// clEnqueueNDRangeKernel().
// 'globalWorkSize' is the 1D dimension of the
// work-items
status = cl.CLEnqueueNDRangeKernel(cmdQueue,
kernel,
1,
nil,
globalWorkSize[:],
nil,
0,
nil,
nil)
if status != cl.CL_SUCCESS {
println("CLEnqueueNDRangeKernel status!=cl.CL_SUCCESS")
return
}
//-----------------------------------------------------
// STEP 12: Read the output buffer back to the host
//-----------------------------------------------------
// Use clEnqueueReadBuffer() to read the OpenCL output
// buffer (bufferC)
// to the host output array (C)
cl.CLEnqueueReadBuffer(cmdQueue,
bufferC,
cl.CL_TRUE,
0,
datasize,
unsafe.Pointer(&C[0]),
0,
nil,
nil)
if status != cl.CL_SUCCESS {
println("CLEnqueueReadBuffer status!=cl.CL_SUCCESS")
return
}
// Verify the output
result := true
for i := cl.CL_int(0); i < cl.CL_int(elements); i++ {
if C[i] != i+i {
result = false
break
}
}
if result {
println("Output is correct\n")
} else {
println("Output is incorrect\n")
}
//-----------------------------------------------------
// STEP 13: Release OpenCL resources
//-----------------------------------------------------
// Free OpenCL resources
cl.CLReleaseKernel(kernel)
cl.CLReleaseProgram(program)
cl.CLReleaseCommandQueue(cmdQueue)
cl.CLReleaseMemObject(bufferA)
cl.CLReleaseMemObject(bufferB)
cl.CLReleaseMemObject(bufferC)
cl.CLReleaseContext(context)
}
func main() {
/* OpenCL data structures */
var device []cl.CL_device_id
var context cl.CL_context
var queue cl.CL_command_queue
var program *cl.CL_program
var kernel cl.CL_kernel
var err cl.CL_int
/* Data and events */
var num_vectors cl.CL_int
var data [NUM_BYTES]byte
var data_buffer cl.CL_mem
var prof_event cl.CL_event
var total_time cl.CL_ulong
var time_start, time_end interface{}
var mapped_memory unsafe.Pointer
/* Create a device and context */
device = utils.Create_device()
context = cl.CLCreateContext(nil, 1, device[:], nil, nil, &err)
if err < 0 {
println("Couldn't create a context")
return
}
/* Build the program and create a kernel */
program = utils.Build_program(context, device[:], PROGRAM_FILE, nil)
kernel = cl.CLCreateKernel(*program, KERNEL_FUNC, &err)
if err < 0 {
println("Couldn't create a kernel")
return
}
/* Create a buffer to hold data */
data_buffer = cl.CLCreateBuffer(context, cl.CL_MEM_WRITE_ONLY,
cl.CL_size_t(unsafe.Sizeof(data[0]))*NUM_BYTES, nil, &err)
if err < 0 {
println("Couldn't create a buffer")
return
}
/* Create kernel argument */
err = cl.CLSetKernelArg(kernel, 0, cl.CL_size_t(unsafe.Sizeof(data_buffer)), unsafe.Pointer(&data_buffer))
if err < 0 {
println("Couldn't set a kernel argument")
return
}
/* Tell kernel number of char16 vectors */
num_vectors = NUM_BYTES / 16
cl.CLSetKernelArg(kernel, 1, cl.CL_size_t(unsafe.Sizeof(num_vectors)), unsafe.Pointer(&num_vectors))
/* Create a command queue */
queue = cl.CLCreateCommandQueue(context, device[0],
cl.CL_QUEUE_PROFILING_ENABLE, &err)
if err < 0 {
println("Couldn't create a command queue")
return
}
total_time = 0.0
for i := 0; i < NUM_ITERATIONS; i++ {
/* Enqueue kernel */
err = cl.CLEnqueueTask(queue, kernel, 0, nil, nil)
if err < 0 {
println("Couldn't enqueue the kernel")
return
}
if PROFILE_READ == 1 {
/* Read the buffer */
err = cl.CLEnqueueReadBuffer(queue, data_buffer, cl.CL_TRUE, 0,
cl.CL_size_t(unsafe.Sizeof(data[0]))*NUM_BYTES, unsafe.Pointer(&data[0]), 0, nil, &prof_event)
if err < 0 {
println("Couldn't read the buffer")
return
}
} else {
/* Create memory map */
mapped_memory = cl.CLEnqueueMapBuffer(queue, data_buffer, cl.CL_TRUE,
cl.CL_MAP_READ, 0, cl.CL_size_t(unsafe.Sizeof(data[0]))*NUM_BYTES, 0, nil, &prof_event, &err)
if err < 0 {
println("Couldn't map the buffer to host memory")
return
}
}
/* Get profiling information */
cl.CLGetEventProfilingInfo(prof_event, cl.CL_PROFILING_COMMAND_START,
cl.CL_size_t(unsafe.Sizeof(total_time)), &time_start, nil)
cl.CLGetEventProfilingInfo(prof_event, cl.CL_PROFILING_COMMAND_END,
cl.CL_size_t(unsafe.Sizeof(total_time)), &time_end, nil)
total_time += time_end.(cl.CL_ulong) - time_start.(cl.CL_ulong)
if PROFILE_READ == 0 {
/* Unmap the buffer */
err = cl.CLEnqueueUnmapMemObject(queue, data_buffer, mapped_memory,
0, nil, nil)
if err < 0 {
println("Couldn't unmap the buffer")
return
}
}
}
if PROFILE_READ == 1 {
fmt.Printf("Average read time: %v\n", total_time/NUM_ITERATIONS)
} else {
fmt.Printf("Average map time: %v\n", total_time/NUM_ITERATIONS)
}
/* Deallocate resources */
cl.CLReleaseEvent(prof_event)
cl.CLReleaseMemObject(data_buffer)
cl.CLReleaseKernel(kernel)
cl.CLReleaseCommandQueue(queue)
cl.CLReleaseProgram(*program)
cl.CLReleaseContext(context)
}
func main() {
/* OpenCL data structures */
var device []cl.CL_device_id
var context cl.CL_context
var queue cl.CL_command_queue
var program *cl.CL_program
var kernel cl.CL_kernel
var err cl.CL_int
var err1 error
var global_size [2]cl.CL_size_t
/* Image data */
var pixels []uint16
var png_format cl.CL_image_format
var input_image, output_image cl.CL_mem
var origin, region [3]cl.CL_size_t
var width, height cl.CL_size_t
/* Open input file and read image data */
pixels, width, height, err1 = utils.Read_image_data(INPUT_FILE)
if err1 != nil {
return
} else {
fmt.Printf("width=%d, height=%d", width, height)
}
/* Create a device and context */
device = utils.Create_device()
context = cl.CLCreateContext(nil, 1, device[:], nil, nil, &err)
if err < 0 {
println("Couldn't create a context")
return
}
/* Build the program and create a kernel */
program = utils.Build_program(context, device[:], PROGRAM_FILE, nil)
kernel = cl.CLCreateKernel(*program, KERNEL_FUNC, &err)
if err < 0 {
fmt.Printf("Couldn't create a kernel: %d", err)
return
}
/* Create image object */
png_format.Image_channel_order = cl.CL_LUMINANCE
png_format.Image_channel_data_type = cl.CL_UNORM_INT16
input_image = cl.CLCreateImage2D(context,
cl.CL_MEM_READ_ONLY|cl.CL_MEM_COPY_HOST_PTR,
&png_format, width, height, 0, unsafe.Pointer(&pixels[0]), &err)
output_image = cl.CLCreateImage2D(context,
cl.CL_MEM_WRITE_ONLY, &png_format, width, height, 0, nil, &err)
if err < 0 {
println("Couldn't create the image object")
return
}
/* Create kernel arguments */
err = cl.CLSetKernelArg(kernel, 0, cl.CL_size_t(unsafe.Sizeof(input_image)), unsafe.Pointer(&input_image))
err |= cl.CLSetKernelArg(kernel, 1, cl.CL_size_t(unsafe.Sizeof(output_image)), unsafe.Pointer(&output_image))
if err < 0 {
println("Couldn't set a kernel argument")
return
}
/* Create a command queue */
queue = cl.CLCreateCommandQueue(context, device[0], 0, &err)
if err < 0 {
println("Couldn't create a command queue")
return
}
/* Enqueue kernel */
global_size[0] = width
global_size[1] = height
err = cl.CLEnqueueNDRangeKernel(queue, kernel, 2, nil, global_size[:],
nil, 0, nil, nil)
if err < 0 {
println("Couldn't enqueue the kernel")
return
}
/* Read the image object */
origin[0] = 0
origin[1] = 0
origin[2] = 0
region[0] = width
region[1] = height
region[2] = 1
err = cl.CLEnqueueReadImage(queue, output_image, cl.CL_TRUE, origin,
region, 0, 0, unsafe.Pointer(&pixels[0]), 0, nil, nil)
if err < 0 {
println("Couldn't read from the image object")
return
}
/* Create output PNG file and write data */
utils.Write_image_data(OUTPUT_FILE, pixels, width, height)
/* Deallocate resources */
cl.CLReleaseMemObject(input_image)
cl.CLReleaseMemObject(output_image)
cl.CLReleaseKernel(kernel)
cl.CLReleaseCommandQueue(queue)
cl.CLReleaseProgram(*program)
cl.CLReleaseContext(context)
}
func main() {
/* OpenCL data structures */
var device []cl.CL_device_id
var context cl.CL_context
var queue cl.CL_command_queue
var program *cl.CL_program
var kernel cl.CL_kernel
var err cl.CL_int
/* Data and events */
var data []float32
var data_buffer cl.CL_mem
var user_event, kernel_event, read_event [1]cl.CL_event
/* Initialize data */
data = make([]float32, 4)
for i := 0; i < 4; i++ {
data[i] = float32(i) * 1.0
}
/* Create a device and context */
device = utils.Create_device()
context = cl.CLCreateContext(nil, 1, device[:], nil, nil, &err)
if err < 0 {
println("Couldn't create a context")
return
}
/* Build the program and create a kernel */
program = utils.Build_program(context, device[:], PROGRAM_FILE, nil)
kernel = cl.CLCreateKernel(*program, KERNEL_FUNC, &err)
if err < 0 {
println("Couldn't create a kernel")
return
}
/* Create a buffer to hold data */
data_buffer = cl.CLCreateBuffer(context,
cl.CL_MEM_READ_WRITE|cl.CL_MEM_COPY_HOST_PTR,
cl.CL_size_t(unsafe.Sizeof(data[0]))*4, unsafe.Pointer(&data[0]), &err)
if err < 0 {
println("Couldn't create a buffer")
return
}
/* Create kernel argument */
err = cl.CLSetKernelArg(kernel, 0, cl.CL_size_t(unsafe.Sizeof(data_buffer)), unsafe.Pointer(&data_buffer))
if err < 0 {
println("Couldn't set a kernel argument")
return
}
/* Create a command queue */
queue = cl.CLCreateCommandQueue(context, device[0],
cl.CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err)
if err < 0 {
println("Couldn't create a command queue")
return
}
/* Configure events */
user_event[0] = cl.CLCreateUserEvent(context, &err)
if err < 0 {
println("Couldn't enqueue the kernel")
return
}
/* Enqueue kernel */
err = cl.CLEnqueueTask(queue, kernel, 1, user_event[:], &kernel_event[0])
if err < 0 {
println("Couldn't enqueue the kernel")
return
}
/* Read the buffer */
err = cl.CLEnqueueReadBuffer(queue, data_buffer, cl.CL_FALSE, 0,
cl.CL_size_t(unsafe.Sizeof(data[0]))*4, unsafe.Pointer(&data[0]), 1, kernel_event[:], &read_event[0])
if err < 0 {
println("Couldn't read the buffer")
return
}
/* Set callback for event */
err = cl.CLSetEventCallback(read_event[0], cl.CL_COMPLETE,
read_complete, unsafe.Pointer(&data))
if err < 0 {
println("Couldn't set callback for event")
return
}
/* Sleep for a second to demonstrate the that commands haven't
started executing. Then prompt user */
time.Sleep(1)
fmt.Printf("Old data: %4.2f, %4.2f, %4.2f, %4.2f\n",
data[0], data[1], data[2], data[3])
fmt.Printf("Press ENTER to continue.\n")
//getchar();
reader := bufio.NewReader(os.Stdin)
reader.ReadString('\n')
/* Set user event to success */
cl.CLSetUserEventStatus(user_event[0], cl.CL_SUCCESS)
/* Deallocate resources */
cl.CLReleaseEvent(read_event[0])
cl.CLReleaseEvent(kernel_event[0])
cl.CLReleaseEvent(user_event[0])
cl.CLReleaseMemObject(data_buffer)
cl.CLReleaseKernel(kernel)
cl.CLReleaseCommandQueue(queue)
cl.CLReleaseProgram(*program)
cl.CLReleaseContext(context)
}
func main() {
// Use this to check the output of each API call
var status cl.CL_int
//-----------------------------------------------------
// STEP 1: Discover and initialize the platforms
//-----------------------------------------------------
var numPlatforms cl.CL_uint
var platforms []cl.CL_platform_id
// Use clGetPlatformIDs() to retrieve the number of
// platforms
status = cl.CLGetPlatformIDs(0, nil, &numPlatforms)
// Allocate enough space for each platform
platforms = make([]cl.CL_platform_id, numPlatforms)
// Fill in platforms with clGetPlatformIDs()
status = cl.CLGetPlatformIDs(numPlatforms, platforms, nil)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLGetPlatformIDs")
//-----------------------------------------------------
// STEP 2: Discover and initialize the GPU devices
//-----------------------------------------------------
var numDevices cl.CL_uint
var devices []cl.CL_device_id
// Use clGetDeviceIDs() to retrieve the number of
// devices present
status = cl.CLGetDeviceIDs(platforms[0],
cl.CL_DEVICE_TYPE_GPU,
0,
nil,
&numDevices)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLGetDeviceIDs")
// Allocate enough space for each device
devices = make([]cl.CL_device_id, numDevices)
// Fill in devices with clGetDeviceIDs()
status = cl.CLGetDeviceIDs(platforms[0],
cl.CL_DEVICE_TYPE_GPU,
numDevices,
devices,
nil)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLGetDeviceIDs")
//-----------------------------------------------------
// STEP 3: Create a context
//-----------------------------------------------------
var context cl.CL_context
// Create a context using clCreateContext() and
// associate it with the devices
context = cl.CLCreateContext(nil,
numDevices,
devices,
nil,
nil,
&status)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLCreateContext")
defer cl.CLReleaseContext(context)
//-----------------------------------------------------
// STEP 4: Create a command queue
//-----------------------------------------------------
var commandQueue [MAX_COMMAND_QUEUE]cl.CL_command_queue
// Create a command queue using clCreateCommandQueueWithProperties(),
// and associate it with the device you want to execute
for i := 0; i < MAX_COMMAND_QUEUE; i++ {
commandQueue[i] = cl.CLCreateCommandQueueWithProperties(context,
devices[0],
nil,
&status)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLCreateCommandQueueWithProperties")
defer cl.CLReleaseCommandQueue(commandQueue[i])
}
//-----------------------------------------------------
// STEP 5: Create device buffers
//-----------------------------------------------------
producerGroupSize := cl.CL_size_t(PRODUCER_GROUP_SIZE)
producerGlobalSize := cl.CL_size_t(PRODUCER_GLOBAL_SIZE)
consumerGroupSize := cl.CL_size_t(CONSUMER_GROUP_SIZE)
consumerGlobalSize := cl.CL_size_t(CONSUMER_GLOBAL_SIZE)
var samplePipePkt [2]cl.CL_float
szPipe := cl.CL_uint(PIPE_SIZE)
szPipePkt := cl.CL_uint(unsafe.Sizeof(samplePipePkt))
if szPipe%PRNG_CHANNELS != 0 {
szPipe = (szPipe/PRNG_CHANNELS)*PRNG_CHANNELS + PRNG_CHANNELS
}
consumerGlobalSize = cl.CL_size_t(szPipe)
pipePktPerThread := cl.CL_int(szPipe) / PRNG_CHANNELS
seed := cl.CL_int(SEED)
rngType := cl.CL_int(RV_GAUSSIAN)
var histMin cl.CL_float
var histMax cl.CL_float
if rngType == cl.CL_int(RV_UNIFORM) {
histMin = 0.0
histMax = 1.0
} else {
histMin = -10.0
histMax = 10.0
}
localDevHist := make([]cl.CL_int, MAX_HIST_BINS)
cpuHist := make([]cl.CL_int, MAX_HIST_BINS)
//Create and initialize memory objects
rngPipe := cl.CLCreatePipe(context,
cl.CL_MEM_READ_WRITE,
szPipePkt,
szPipe,
nil,
&status)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clCreatePipe")
devHist := cl.CLCreateBuffer(context,
cl.CL_MEM_READ_WRITE|cl.CL_MEM_COPY_HOST_PTR,
MAX_HIST_BINS*cl.CL_size_t(unsafe.Sizeof(localDevHist[0])),
unsafe.Pointer(&localDevHist[0]),
&status)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clCreateBuffer")
//-----------------------------------------------------
// STEP 6: Create and compile the program
//-----------------------------------------------------
programSource, programeSize := utils.Load_programsource("pipe.cl")
// Create a program using clCreateProgramWithSource()
program := cl.CLCreateProgramWithSource(context,
1,
programSource[:],
programeSize[:],
&status)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLCreateProgramWithSource")
defer cl.CLReleaseProgram(program)
// Build (compile) the program for the devices with
// clBuildProgram()
options := "-cl-std=CL2.0"
status = cl.CLBuildProgram(program,
numDevices,
devices,
[]byte(options),
nil,
nil)
if status != cl.CL_SUCCESS {
var program_log interface{}
var log_size cl.CL_size_t
/* Find size of log and print to std output */
cl.CLGetProgramBuildInfo(program, devices[0], cl.CL_PROGRAM_BUILD_LOG,
0, nil, &log_size)
cl.CLGetProgramBuildInfo(program, devices[0], cl.CL_PROGRAM_BUILD_LOG,
log_size, &program_log, nil)
fmt.Printf("%s\n", program_log)
return
}
//utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLBuildProgram")
//-----------------------------------------------------
// STEP 7: Create the kernel
//-----------------------------------------------------
// Use clCreateKernel() to create a kernel
produceKernel := cl.CLCreateKernel(program, []byte("pipe_producer"), &status)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLCreateKernel")
defer cl.CLReleaseKernel(produceKernel)
consumeKernel := cl.CLCreateKernel(program, []byte("pipe_consumer"), &status)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLCreateKernel")
defer cl.CLReleaseKernel(consumeKernel)
//-----------------------------------------------------
// STEP 8: Set the kernel arguments
//-----------------------------------------------------
// Associate the input and output buffers with the
// kernel
// using clSetKernelArg()
// Set appropriate arguments to the kernel
status = cl.CLSetKernelArg(produceKernel,
0,
cl.CL_size_t(unsafe.Sizeof(rngPipe)),
unsafe.Pointer(&rngPipe))
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clSetKernelArg(rngPipe)")
status = cl.CLSetKernelArg(produceKernel,
1,
cl.CL_size_t(unsafe.Sizeof(pipePktPerThread)),
unsafe.Pointer(&pipePktPerThread))
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clSetKernelArg(pipePktPerThread)")
status = cl.CLSetKernelArg(produceKernel,
2,
cl.CL_size_t(unsafe.Sizeof(seed)),
unsafe.Pointer(&seed))
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clSetKernelArg(seed)")
status = cl.CLSetKernelArg(produceKernel,
3,
cl.CL_size_t(unsafe.Sizeof(rngType)),
unsafe.Pointer(&rngType))
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clSetKernelArg(rngType)")
//-----------------------------------------------------
// STEP 9: Configure the work-item structure
//-----------------------------------------------------
// Define an index space (global work size) of work
// items for
// execution. A workgroup size (local work size) is not
// required,
// but can be used.
// Enqueue both the kernels.
var globalThreads = []cl.CL_size_t{producerGlobalSize}
var localThreads = []cl.CL_size_t{producerGroupSize}
//-----------------------------------------------------
// STEP 10: Enqueue the kernel for execution
//-----------------------------------------------------
// Execute the kernel by using
// clEnqueueNDRangeKernel().
// 'globalWorkSize' is the 1D dimension of the
// work-items
var produceEvt [1]cl.CL_event
status = cl.CLEnqueueNDRangeKernel(commandQueue[0],
produceKernel,
1,
nil,
globalThreads,
localThreads,
0,
nil,
&produceEvt[0])
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clEnqueueNDRangeKernel")
/*
launch consumer kernel only after producer has finished.
This is done to avoid concurrent kernels execution as the
memory consistency of pipe is guaranteed only across
synchronization points.
*/
status = cl.CLWaitForEvents(1, produceEvt[:])
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clWaitForEvents(produceEvt)")
//-----------------------------------------------------
// STEP 8: Set the kernel arguments
//-----------------------------------------------------
// Associate the input and output buffers with the
// kernel
// using clSetKernelArg()
// Set appropriate arguments to the kernel
status = cl.CLSetKernelArg(consumeKernel,
0,
cl.CL_size_t(unsafe.Sizeof(rngPipe)),
unsafe.Pointer(&rngPipe))
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clSetKernelArg(rngPipe)")
status = cl.CLSetKernelArg(consumeKernel,
1,
cl.CL_size_t(unsafe.Sizeof(devHist)),
unsafe.Pointer(&devHist))
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clSetKernelArg(devHist)")
status = cl.CLSetKernelArg(consumeKernel,
2,
cl.CL_size_t(unsafe.Sizeof(histMin)),
unsafe.Pointer(&histMin))
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clSetKernelArg(histMin)")
status = cl.CLSetKernelArg(consumeKernel,
3,
cl.CL_size_t(unsafe.Sizeof(histMax)),
unsafe.Pointer(&histMax))
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clSetKernelArg(histMax)")
//-----------------------------------------------------
// STEP 9: Configure the work-item structure
//-----------------------------------------------------
// Define an index space (global work size) of work
// items for
// execution. A workgroup size (local work size) is not
// required,
// but can be used.
globalThreads[0] = consumerGlobalSize
localThreads[0] = consumerGroupSize
//-----------------------------------------------------
// STEP 10: Enqueue the kernel for execution
//-----------------------------------------------------
// Execute the kernel by using
// clEnqueueNDRangeKernel().
// 'globalWorkSize' is the 1D dimension of the
// work-items
var consumeEvt [1]cl.CL_event
status = cl.CLEnqueueNDRangeKernel(
commandQueue[1],
consumeKernel,
1,
nil,
globalThreads,
localThreads,
0,
nil,
&consumeEvt[0])
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clEnqueueNDRangeKernel")
status = cl.CLFlush(commandQueue[0])
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clFlush(0)")
status = cl.CLFlush(commandQueue[1])
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clFlush(1)")
//wait for kernels to finish
status = cl.CLFinish(commandQueue[0])
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clFinish(0)")
status = cl.CLFinish(commandQueue[1])
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clFinish(1)")
//-----------------------------------------------------
// STEP 11: Read the output buffer back to the host
//-----------------------------------------------------
// Use clEnqueueReadBuffer() to read the OpenCL output
// buffer (bufferC)
// to the host output array (C)
//copy the data back to host buffer
var readEvt cl.CL_event
status = cl.CLEnqueueReadBuffer(commandQueue[1],
devHist,
cl.CL_TRUE,
0,
(MAX_HIST_BINS)*cl.CL_size_t(unsafe.Sizeof(localDevHist[0])),
unsafe.Pointer(&localDevHist[0]),
0,
nil,
&readEvt)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clEnqueueReadBuffer")
//-----------------------------------------------------
// STEP 12: Verify the results
//-----------------------------------------------------
//Find the tolerance limit
fTol := (float32)(CONSUMER_GLOBAL_SIZE) * (float32)(COMP_TOL) / (float32)(100.0)
iTol := (int)(fTol)
if iTol == 0 {
iTol = 1
}
//CPU side histogram computation
CPUReference(seed, pipePktPerThread, rngType, cpuHist, histMax, histMin)
//Compare
for bin := 0; bin < MAX_HIST_BINS; bin++ {
diff := int(localDevHist[bin] - cpuHist[bin])
if diff < 0 {
diff = -diff
}
if diff > iTol {
println("Failed!")
return
}
}
println("Passed!")
}
func main() {
/* OpenCL data structures */
var device []cl.CL_device_id
var context cl.CL_context
var queue cl.CL_command_queue
var program *cl.CL_program
var kernel cl.CL_kernel
var err cl.CL_int
/* Data and buffers */
var a float32 = 6.0
var b float32 = 2.0
var result float32
var a_buffer, b_buffer, output_buffer cl.CL_mem
/* Extension data */
var sizeofuint cl.CL_uint
var addr_data interface{}
var ext_data interface{}
fp64_ext := "cl_khr_fp64"
var ext_size cl.CL_size_t
var options []byte
/* Create a device and context */
device = utils.Create_device()
context = cl.CLCreateContext(nil, 1, device[:], nil, nil, &err)
if err < 0 {
println("Couldn't create a context")
return
}
/* Obtain the device data */
if cl.CLGetDeviceInfo(device[0], cl.CL_DEVICE_ADDRESS_BITS,
cl.CL_size_t(unsafe.Sizeof(sizeofuint)), &addr_data, nil) < 0 {
println("Couldn't read extension data")
return
}
fmt.Printf("Address width: %v\n", addr_data.(cl.CL_uint))
/* Define "FP_64" option if doubles are supported */
cl.CLGetDeviceInfo(device[0], cl.CL_DEVICE_EXTENSIONS,
0, nil, &ext_size)
// ext_data = (char*)malloc(ext_size + 1);
// ext_data[ext_size] = '\0';
cl.CLGetDeviceInfo(device[0], cl.CL_DEVICE_EXTENSIONS,
ext_size, &ext_data, nil)
if strings.Contains(ext_data.(string), fp64_ext) {
fmt.Printf("The %s extension is supported.\n", fp64_ext)
options = []byte("-DFP_64 ")
} else {
fmt.Printf("The %s extension is not supported. %s\n", fp64_ext, ext_data.(string))
}
/* Build the program and create the kernel */
program = utils.Build_program(context, device[:], PROGRAM_FILE, options)
kernel = cl.CLCreateKernel(*program, KERNEL_FUNC, &err)
if err < 0 {
println("Couldn't create a kernel")
return
}
/* Create CL buffers to hold input and output data */
a_buffer = cl.CLCreateBuffer(context, cl.CL_MEM_READ_ONLY|
cl.CL_MEM_COPY_HOST_PTR, cl.CL_size_t(unsafe.Sizeof(a)), unsafe.Pointer(&a), &err)
if err < 0 {
println("Couldn't create a memory object")
return
}
b_buffer = cl.CLCreateBuffer(context, cl.CL_MEM_READ_ONLY|
cl.CL_MEM_COPY_HOST_PTR, cl.CL_size_t(unsafe.Sizeof(b)), unsafe.Pointer(&b), nil)
output_buffer = cl.CLCreateBuffer(context, cl.CL_MEM_WRITE_ONLY,
cl.CL_size_t(unsafe.Sizeof(b)), nil, nil)
/* Create kernel arguments */
err = cl.CLSetKernelArg(kernel, 0, cl.CL_size_t(unsafe.Sizeof(a_buffer)), unsafe.Pointer(&a_buffer))
if err < 0 {
println("Couldn't set a kernel argument")
return
}
cl.CLSetKernelArg(kernel, 1, cl.CL_size_t(unsafe.Sizeof(b_buffer)), unsafe.Pointer(&b_buffer))
cl.CLSetKernelArg(kernel, 2, cl.CL_size_t(unsafe.Sizeof(output_buffer)), unsafe.Pointer(&output_buffer))
/* Create a command queue */
queue = cl.CLCreateCommandQueue(context, device[0], 0, &err)
if err < 0 {
println("Couldn't create a command queue")
return
}
/* Enqueue kernel */
err = cl.CLEnqueueTask(queue, kernel, 0, nil, nil)
if err < 0 {
println("Couldn't enqueue the kernel")
return
}
/* Read and print the result */
err = cl.CLEnqueueReadBuffer(queue, output_buffer, cl.CL_TRUE, 0,
cl.CL_size_t(unsafe.Sizeof(result)), unsafe.Pointer(&result), 0, nil, nil)
if err < 0 {
println("Couldn't read the output buffer")
return
}
fmt.Printf("The kernel result is %f\n", result)
/* Deallocate resources */
cl.CLReleaseMemObject(a_buffer)
cl.CLReleaseMemObject(b_buffer)
cl.CLReleaseMemObject(output_buffer)
cl.CLReleaseKernel(kernel)
cl.CLReleaseCommandQueue(queue)
cl.CLReleaseProgram(*program)
cl.CLReleaseContext(context)
}
func main() {
/* OpenCL data structures */
var device []cl.CL_device_id
var context cl.CL_context
var queue cl.CL_command_queue
var program *cl.CL_program
var kernel cl.CL_kernel
var err cl.CL_int
var offset, global_size, local_size [1]cl.CL_size_t
/* Data and events */
var data [2]cl.CL_int
var data_buffer cl.CL_mem
/* Create a device and context */
device = utils.Create_device()
context = cl.CLCreateContext(nil, 1, device[:], nil, nil, &err)
if err < 0 {
println("Couldn't create a context")
return
}
/* Build the program and create a kernel */
program = utils.Build_program(context, device[:], PROGRAM_FILE, nil)
kernel = cl.CLCreateKernel(*program, KERNEL_FUNC, &err)
if err < 0 {
println("Couldn't create a kernel")
return
}
/* Create a buffer to hold data */
data_buffer = cl.CLCreateBuffer(context, cl.CL_MEM_WRITE_ONLY,
cl.CL_size_t(unsafe.Sizeof(data[0]))*2, nil, &err)
if err < 0 {
println("Couldn't create a buffer")
return
}
/* Create kernel argument */
err = cl.CLSetKernelArg(kernel, 0, cl.CL_size_t(unsafe.Sizeof(data_buffer)), unsafe.Pointer(&data_buffer))
if err < 0 {
println("Couldn't set a kernel argument")
return
}
/* Create a command queue */
queue = cl.CLCreateCommandQueue(context, device[0], 0, &err)
if err < 0 {
println("Couldn't create a command queue")
return
}
/* Enqueue kernel */
offset[0] = 0
global_size[0] = 8
local_size[0] = 4
err = cl.CLEnqueueNDRangeKernel(queue, kernel, 1, offset[:], global_size[:], local_size[:], 0, nil, nil)
if err < 0 {
println("Couldn't enqueue the kernel")
return
}
/* Read the buffer */
err = cl.CLEnqueueReadBuffer(queue, data_buffer, cl.CL_TRUE, 0,
cl.CL_size_t(unsafe.Sizeof(data[0]))*2, unsafe.Pointer(&data[0]), 0, nil, nil)
if err < 0 {
println("Couldn't read the buffer")
return
}
fmt.Printf("Increment: %d\n", data[0])
fmt.Printf("Atomic increment: %d\n", data[1])
/* Deallocate resources */
cl.CLReleaseMemObject(data_buffer)
cl.CLReleaseKernel(kernel)
cl.CLReleaseCommandQueue(queue)
cl.CLReleaseProgram(*program)
cl.CLReleaseContext(context)
}
func main() {
/* OpenCL data structures */
var device []cl.CL_device_id
var context cl.CL_context
var queue cl.CL_command_queue
var program *cl.CL_program
var kernel cl.CL_kernel
var err cl.CL_int
/* Data and buffers */
var data_one, data_two, result_array [100]float32
var buffer_one, buffer_two cl.CL_mem
var mapped_memory unsafe.Pointer
/* Initialize arrays */
for i := 0; i < 100; i++ {
data_one[i] = 1.0 * float32(i)
data_two[i] = -1.0 * float32(i)
result_array[i] = 0.0
}
/* Create a device and context */
device = utils.Create_device()
context = cl.CLCreateContext(nil, 1, device[:], nil, nil, &err)
if err < 0 {
println("Couldn't create a context")
return
}
/* Build the program and create the kernel */
program = utils.Build_program(context, device[:], PROGRAM_FILE, nil)
kernel = cl.CLCreateKernel(*program, []byte(KERNEL_FUNC), &err)
if err < 0 {
println("Couldn't create a kernel")
return
}
/* Create buffers */
buffer_one = cl.CLCreateBuffer(context, cl.CL_MEM_READ_WRITE|
cl.CL_MEM_COPY_HOST_PTR, cl.CL_size_t(unsafe.Sizeof(data_one)), unsafe.Pointer(&data_one[0]), &err)
if err < 0 {
println("Couldn't create buffer object 1")
return
}
buffer_two = cl.CLCreateBuffer(context, cl.CL_MEM_READ_WRITE|
cl.CL_MEM_COPY_HOST_PTR, cl.CL_size_t(unsafe.Sizeof(data_two)), unsafe.Pointer(&data_two), &err)
if err < 0 {
println("Couldn't create buffer object 2")
return
}
/* Set buffers as arguments to the kernel */
err = cl.CLSetKernelArg(kernel, 0, cl.CL_size_t(unsafe.Sizeof(buffer_one)), unsafe.Pointer(&buffer_one))
err |= cl.CLSetKernelArg(kernel, 1, cl.CL_size_t(unsafe.Sizeof(buffer_two)), unsafe.Pointer(&buffer_two))
if err < 0 {
println("Couldn't set the buffer as the kernel argument")
return
}
/* Create a command queue */
queue = cl.CLCreateCommandQueue(context, device[0], 0, &err)
if err < 0 {
println("Couldn't create a command queue")
return
}
/* Enqueue kernel */
err = cl.CLEnqueueTask(queue, kernel, 0, nil, nil)
if err < 0 {
println("Couldn't enqueue the kernel")
return
}
/* Enqueue command to copy buffer one to buffer two */
err = cl.CLEnqueueCopyBuffer(queue, buffer_one, buffer_two, 0, 0,
cl.CL_size_t(unsafe.Sizeof(data_one)), 0, nil, nil)
if err < 0 {
println("Couldn't perform the buffer copy")
return
}
/* Enqueue command to map buffer two to host memory */
mapped_memory = cl.CLEnqueueMapBuffer(queue, buffer_two, cl.CL_TRUE,
cl.CL_MAP_READ, 0, cl.CL_size_t(unsafe.Sizeof(data_two)), 0, nil, nil, &err)
if err < 0 {
println("Couldn't map the buffer to host memory")
return
}
/* Transfer memory and unmap the buffer */
C.memcpy(unsafe.Pointer(&result_array[0]), mapped_memory, C.size_t(unsafe.Sizeof(data_two)))
err = cl.CLEnqueueUnmapMemObject(queue, buffer_two, mapped_memory,
0, nil, nil)
if err < 0 {
println("Couldn't unmap the buffer")
return
}
/* Display updated buffer */
for i := 0; i < 10; i++ {
for j := 0; j < 10; j++ {
fmt.Printf("%6.1f", result_array[j+i*10])
}
fmt.Printf("\n")
}
/* Deallocate resources */
cl.CLReleaseMemObject(buffer_one)
cl.CLReleaseMemObject(buffer_two)
cl.CLReleaseKernel(kernel)
cl.CLReleaseCommandQueue(queue)
cl.CLReleaseProgram(*program)
cl.CLReleaseContext(context)
}
func main() {
/* OpenCL data structures */
var device []cl.CL_device_id
var context cl.CL_context
var queue cl.CL_command_queue
var program *cl.CL_program
var kernel cl.CL_kernel
var err cl.CL_int
/* Data and events */
var num_ints cl.CL_int
var num_items [1]cl.CL_size_t
var data [NUM_INTS]cl.CL_int
var data_buffer cl.CL_mem
var prof_event cl.CL_event
var total_time cl.CL_ulong
var time_start, time_end interface{}
/* Initialize data */
for i := 0; i < NUM_INTS; i++ {
data[i] = cl.CL_int(i)
}
/* Set number of data points and work-items */
num_ints = NUM_INTS
num_items[0] = NUM_ITEMS
/* Create a device and context */
device = utils.Create_device()
context = cl.CLCreateContext(nil, 1, device[:], nil, nil, &err)
if err < 0 {
println("Couldn't create a context")
return
}
/* Build the program and create a kernel */
program = utils.Build_program(context, device[:], PROGRAM_FILE, nil)
kernel = cl.CLCreateKernel(*program, KERNEL_FUNC, &err)
if err < 0 {
println("Couldn't create a kernel")
return
}
/* Create a buffer to hold data */
data_buffer = cl.CLCreateBuffer(context,
cl.CL_MEM_READ_WRITE|cl.CL_MEM_COPY_HOST_PTR,
cl.CL_size_t(unsafe.Sizeof(data[0]))*NUM_INTS, unsafe.Pointer(&data[0]), &err)
if err < 0 {
println("Couldn't create a buffer")
return
}
/* Create kernel argument */
err = cl.CLSetKernelArg(kernel, 0, cl.CL_size_t(unsafe.Sizeof(data_buffer)), unsafe.Pointer(&data_buffer))
if err < 0 {
println("Couldn't set a kernel argument")
return
}
cl.CLSetKernelArg(kernel, 1, cl.CL_size_t(unsafe.Sizeof(num_ints)), unsafe.Pointer(&num_ints))
/* Create a command queue */
queue = cl.CLCreateCommandQueue(context, device[0],
cl.CL_QUEUE_PROFILING_ENABLE, &err)
if err < 0 {
println("Couldn't create a command queue")
return
}
total_time = 0.0
for i := 0; i < NUM_ITERATIONS; i++ {
/* Enqueue kernel */
cl.CLEnqueueNDRangeKernel(queue, kernel, 1, nil, num_items[:],
nil, 0, nil, &prof_event)
if err < 0 {
println("Couldn't enqueue the kernel")
return
}
/* Finish processing the queue and get profiling information */
cl.CLFinish(queue)
cl.CLGetEventProfilingInfo(prof_event, cl.CL_PROFILING_COMMAND_START,
cl.CL_size_t(unsafe.Sizeof(total_time)), &time_start, nil)
cl.CLGetEventProfilingInfo(prof_event, cl.CL_PROFILING_COMMAND_END,
cl.CL_size_t(unsafe.Sizeof(total_time)), &time_end, nil)
total_time += time_end.(cl.CL_ulong) - time_start.(cl.CL_ulong)
}
fmt.Printf("Average time = %v\n", total_time/NUM_ITERATIONS)
/* Deallocate resources */
cl.CLReleaseEvent(prof_event)
cl.CLReleaseKernel(kernel)
cl.CLReleaseMemObject(data_buffer)
cl.CLReleaseCommandQueue(queue)
cl.CLReleaseProgram(*program)
cl.CLReleaseContext(context)
}
func main() {
/* OpenCL data structures */
var device []cl.CL_device_id
var context cl.CL_context
var queue cl.CL_command_queue
var program *cl.CL_program
var kernel cl.CL_kernel
var err cl.CL_int
/* Data and buffers */
var full_matrix, zero_matrix [80]float32
var sizeoffloat32 = cl.CL_size_t(unsafe.Sizeof(full_matrix[0]))
var buffer_origin = [3]cl.CL_size_t{5 * sizeoffloat32, 3, 0}
var host_origin = [3]cl.CL_size_t{1 * sizeoffloat32, 1, 0}
var region = [3]cl.CL_size_t{4 * sizeoffloat32, 4, 1}
var matrix_buffer cl.CL_mem
/* Initialize data */
for i := 0; i < 80; i++ {
full_matrix[i] = float32(i) * 1.0
zero_matrix[i] = 0.0
}
/* Create a device and context */
device = utils.Create_device()
context = cl.CLCreateContext(nil, 1, device[:], nil, nil, &err)
if err < 0 {
println("Couldn't create a context")
return
}
/* Build the program and create the kernel */
program = utils.Build_program(context, device[:], PROGRAM_FILE, nil)
if program == nil {
println("Couldn't build program")
return
}
kernel = cl.CLCreateKernel(*program, []byte(KERNEL_FUNC), &err)
if err < 0 {
println("Couldn't create a kernel")
return
}
/* Create a buffer to hold 80 floats */
matrix_buffer = cl.CLCreateBuffer(context, cl.CL_MEM_READ_WRITE|
cl.CL_MEM_COPY_HOST_PTR, cl.CL_size_t(unsafe.Sizeof(full_matrix)), unsafe.Pointer(&full_matrix[0]), &err)
if err < 0 {
println("Couldn't create a buffer object")
return
}
/* Set buffer as argument to the kernel */
err = cl.CLSetKernelArg(kernel, 0, cl.CL_size_t(unsafe.Sizeof(matrix_buffer)), unsafe.Pointer(&matrix_buffer))
if err < 0 {
println("Couldn't set the buffer as the kernel argument")
return
}
/* Create a command queue */
queue = cl.CLCreateCommandQueue(context, device[0], 0, &err)
if err < 0 {
println("Couldn't create a command queue")
return
}
/* Enqueue kernel */
err = cl.CLEnqueueTask(queue, kernel, 0, nil, nil)
if err < 0 {
println("Couldn't enqueue the kernel")
return
}
/* Enqueue command to write to buffer */
err = cl.CLEnqueueWriteBuffer(queue, matrix_buffer, cl.CL_TRUE, 0,
cl.CL_size_t(unsafe.Sizeof(full_matrix)), unsafe.Pointer(&full_matrix[0]), 0, nil, nil)
if err < 0 {
println("Couldn't write to the buffer object")
return
}
/* Enqueue command to read rectangle of data */
err = cl.CLEnqueueReadBufferRect(queue, matrix_buffer, cl.CL_TRUE,
buffer_origin, host_origin, region, 10*sizeoffloat32, 0,
10*sizeoffloat32, 0, unsafe.Pointer(&zero_matrix[0]), 0, nil, nil)
if err < 0 {
println("Couldn't read the rectangle from the buffer object")
return
}
/* Display updated buffer */
for i := 0; i < 8; i++ {
for j := 0; j < 10; j++ {
fmt.Printf("%6.1f", zero_matrix[j+i*10])
}
fmt.Printf("\n")
}
/* Deallocate resources */
cl.CLReleaseMemObject(matrix_buffer)
cl.CLReleaseKernel(kernel)
cl.CLReleaseCommandQueue(queue)
cl.CLReleaseProgram(*program)
cl.CLReleaseContext(context)
}
func main() {
// Use this to check the output of each API call
var status cl.CL_int
//-----------------------------------------------------
// STEP 1: Discover and initialize the platforms
//-----------------------------------------------------
var numPlatforms cl.CL_uint
// Use clGetPlatformIDs() to retrieve the number of
// platforms
status = cl.CLGetPlatformIDs(0, nil, &numPlatforms)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLGetPlatformIDs")
// Allocate enough space for each platform
platforms := make([]cl.CL_platform_id, numPlatforms)
// Fill in platforms with clGetPlatformIDs()
status = cl.CLGetPlatformIDs(numPlatforms, platforms, nil)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLGetPlatformIDs")
//-----------------------------------------------------
// STEP 2: Discover and initialize the GPU devices
//-----------------------------------------------------
var numDevices cl.CL_uint
// Use clGetDeviceIDs() to retrieve the number of
// devices present
status = cl.CLGetDeviceIDs(platforms[0],
cl.CL_DEVICE_TYPE_GPU,
0,
nil,
&numDevices)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLGetDeviceIDs")
// Allocate enough space for each device
devices := make([]cl.CL_device_id, numDevices)
// Fill in devices with clGetDeviceIDs()
status = cl.CLGetDeviceIDs(platforms[0],
cl.CL_DEVICE_TYPE_GPU,
numDevices,
devices,
nil)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLGetDeviceIDs")
var caps cl.CL_device_svm_capabilities
var caps_value interface{}
status = cl.CLGetDeviceInfo(
devices[0],
cl.CL_DEVICE_SVM_CAPABILITIES,
cl.CL_size_t(unsafe.Sizeof(caps)),
&caps_value,
nil)
caps = caps_value.(cl.CL_device_svm_capabilities)
// Coarse-grained buffer SVM should be available on any OpenCL 2.0 device.
// So it is either not an OpenCL 2.0 device or it must support coarse-grained buffer SVM:
if !(status == cl.CL_SUCCESS && (caps&cl.CL_DEVICE_SVM_FINE_GRAIN_BUFFER) != 0) {
fmt.Printf("Cannot detect fine-grained buffer SVM capabilities on the device. The device seemingly doesn't support fine-grained buffer SVM. caps=%x\n", caps)
println("")
return
}
//-----------------------------------------------------
// STEP 3: Create a context
//-----------------------------------------------------
// Create a context using clCreateContext() and
// associate it with the devices
context := cl.CLCreateContext(nil,
numDevices,
devices,
nil,
nil,
&status)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLCreateContext")
defer cl.CLReleaseContext(context)
//-----------------------------------------------------
// STEP 4: Create a command queue
//-----------------------------------------------------
// Create a command queue using clCreateCommandQueueWithProperties(),
// and associate it with the device you want to execute
queue := cl.CLCreateCommandQueueWithProperties(context,
devices[0],
nil,
&status)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLCreateCommandQueueWithProperties")
defer cl.CLReleaseCommandQueue(queue)
//-----------------------------------------------------
// STEP 5: Create and compile the program
//-----------------------------------------------------
programSource, programeSize := utils.Load_programsource("svmfg.cl")
// Create a program using clCreateProgramWithSource()
program := cl.CLCreateProgramWithSource(context,
1,
programSource[:],
programeSize[:],
&status)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLCreateProgramWithSource")
defer cl.CLReleaseProgram(program)
// Build (compile) the program for the devices with
// clBuildProgram()
options := "-cl-std=CL2.0"
status = cl.CLBuildProgram(program,
numDevices,
devices,
[]byte(options),
nil,
nil)
if status != cl.CL_SUCCESS {
var log interface{}
var log_size cl.CL_size_t
/* Find size of log and print to std output */
cl.CLGetProgramBuildInfo(program, devices[0], cl.CL_PROGRAM_BUILD_LOG, 0, nil, &log_size)
cl.CLGetProgramBuildInfo(program, devices[0], cl.CL_PROGRAM_BUILD_LOG, log_size, &log, nil)
fmt.Printf("%s\n", log)
return
}
//utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLBuildProgram")
//-----------------------------------------------------
// STEP 7: Create the kernel
//-----------------------------------------------------
// Use clCreateKernel() to create a kernel
kernel := cl.CLCreateKernel(program,
[]byte("svmbasic"),
&status)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLCreateKernel")
defer cl.CLReleaseKernel(kernel)
// Then call the main sample routine - resource allocations, OpenCL kernel
// execution, and so on.
svmbasic(1024*1024, context, queue, kernel)
// All resource deallocations happen in defer.
}
func main() {
/* Host/device data structures */
var device []cl.CL_device_id
var context cl.CL_context
var queue cl.CL_command_queue
var program *cl.CL_program
var kernel cl.CL_kernel
var err cl.CL_int
/* Data and buffers */
var test [16]byte
var test_buffer cl.CL_mem
/* Create a context */
device = utils.Create_device()
context = cl.CLCreateContext(nil, 1, device[:], nil, nil, &err)
if err < 0 {
println("Couldn't create a context")
return
}
/* Build the program and create a kernel */
program = utils.Build_program(context, device[:], PROGRAM_FILE, nil)
kernel = cl.CLCreateKernel(*program, KERNEL_FUNC, &err)
if err < 0 {
println("Couldn't create a kernel")
return
}
/* Create a write-only buffer to hold the output data */
test_buffer = cl.CLCreateBuffer(context, cl.CL_MEM_WRITE_ONLY,
cl.CL_size_t(unsafe.Sizeof(test)), nil, &err)
if err < 0 {
println("Couldn't create a buffer")
return
}
/* Create kernel argument */
err = cl.CLSetKernelArg(kernel, 0, cl.CL_size_t(unsafe.Sizeof(test_buffer)), unsafe.Pointer(&test_buffer))
if err < 0 {
println("Couldn't set a kernel argument")
return
}
/* Create a command queue */
queue = cl.CLCreateCommandQueue(context, device[0], 0, &err)
if err < 0 {
println("Couldn't create a command queue")
return
}
/* Enqueue kernel */
err = cl.CLEnqueueTask(queue, kernel, 0, nil, nil)
if err < 0 {
println("Couldn't enqueue the kernel")
return
}
/* Read and print the result */
err = cl.CLEnqueueReadBuffer(queue, test_buffer, cl.CL_TRUE, 0,
cl.CL_size_t(unsafe.Sizeof(test)), unsafe.Pointer(&test), 0, nil, nil)
if err < 0 {
println("Couldn't read the buffer")
return
}
for i := 0; i < 15; i++ {
fmt.Printf("0x%X, ", test[i])
}
fmt.Printf("0x%X\n", test[15])
/* Deallocate resources */
cl.CLReleaseMemObject(test_buffer)
cl.CLReleaseKernel(kernel)
cl.CLReleaseCommandQueue(queue)
cl.CLReleaseProgram(*program)
cl.CLReleaseContext(context)
}
func main() {
// Use this to check the output of each API call
var status cl.CL_int
//-----------------------------------------------------
// STEP 1: Discover and initialize the platforms
//-----------------------------------------------------
var numPlatforms cl.CL_uint
var platforms []cl.CL_platform_id
// Use clGetPlatformIDs() to retrieve the number of
// platforms
status = cl.CLGetPlatformIDs(0, nil, &numPlatforms)
// Allocate enough space for each platform
platforms = make([]cl.CL_platform_id, numPlatforms)
// Fill in platforms with clGetPlatformIDs()
status = cl.CLGetPlatformIDs(numPlatforms, platforms, nil)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLGetPlatformIDs")
//-----------------------------------------------------
// STEP 2: Discover and initialize the GPU devices
//-----------------------------------------------------
var numDevices cl.CL_uint
var devices []cl.CL_device_id
// Use clGetDeviceIDs() to retrieve the number of
// devices present
status = cl.CLGetDeviceIDs(platforms[0],
cl.CL_DEVICE_TYPE_GPU,
0,
nil,
&numDevices)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLGetDeviceIDs")
// Allocate enough space for each device
devices = make([]cl.CL_device_id, numDevices)
// Fill in devices with clGetDeviceIDs()
status = cl.CLGetDeviceIDs(platforms[0],
cl.CL_DEVICE_TYPE_GPU,
numDevices,
devices,
nil)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLGetDeviceIDs")
//-----------------------------------------------------
// STEP 3: Create a context
//-----------------------------------------------------
var context cl.CL_context
// Create a context using clCreateContext() and
// associate it with the devices
context = cl.CLCreateContext(nil,
numDevices,
devices,
nil,
nil,
&status)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLCreateContext")
defer cl.CLReleaseContext(context)
//-----------------------------------------------------
// STEP 4: Create a command queue
//-----------------------------------------------------
var cmdQueue cl.CL_command_queue
// Create a command queue using clCreateCommandQueueWithProperties(),
// and associate it with the device you want to execute
cmdQueue = cl.CLCreateCommandQueueWithProperties(context,
devices[0],
nil,
&status)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLCreateCommandQueueWithProperties")
defer cl.CLReleaseCommandQueue(cmdQueue)
//-----------------------------------------------------
// STEP 5: Create device buffers
//-----------------------------------------------------
// initialize any device/SVM memory here.
/* svm buffer for binary tree */
svmTreeBuf := cl.CLSVMAlloc(context,
cl.CL_MEM_READ_WRITE,
cl.CL_size_t(NUMBER_OF_NODES*unsafe.Sizeof(sampleNode)),
0)
if nil == svmTreeBuf {
println("clSVMAlloc(svmTreeBuf) failed.")
return
}
defer cl.CLSVMFree(context, svmTreeBuf)
/* svm buffer for search keys */
svmSearchBuf := cl.CLSVMAlloc(context,
cl.CL_MEM_READ_WRITE,
cl.CL_size_t(NUMBER_OF_SEARCH_KEY*unsafe.Sizeof(sampleKey)),
0)
if nil == svmSearchBuf {
println("clSVMAlloc(svmSearchBuf) failed.")
return
}
defer cl.CLSVMFree(context, svmSearchBuf)
//create the binary tree and set the root
/* root node of the binary tree */
svmRoot := cpuCreateBinaryTree(cmdQueue, svmTreeBuf)
//initialize search keys
cpuInitSearchKeys(cmdQueue, svmSearchBuf)
/* if voice is not deliberately muzzled, shout parameters */
fmt.Printf("-------------------------------------------------------------------------\n")
fmt.Printf("Searching %d keys in a BST having %d Nodes...\n", NUMBER_OF_SEARCH_KEY, NUMBER_OF_NODES)
fmt.Printf("-------------------------------------------------------------------------\n")
//-----------------------------------------------------
// STEP 6: Create and compile the program
//-----------------------------------------------------
programSource, programeSize := utils.Load_programsource("bst.cl")
// Create a program using clCreateProgramWithSource()
program := cl.CLCreateProgramWithSource(context,
1,
programSource[:],
programeSize[:],
&status)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLCreateProgramWithSource")
defer cl.CLReleaseProgram(program)
// Build (compile) the program for the devices with
// clBuildProgram()
options := "-cl-std=CL2.0"
status = cl.CLBuildProgram(program,
numDevices,
devices,
[]byte(options),
nil,
nil)
if status != cl.CL_SUCCESS {
var program_log interface{}
var log_size cl.CL_size_t
/* Find size of log and print to std output */
cl.CLGetProgramBuildInfo(program, devices[0], cl.CL_PROGRAM_BUILD_LOG,
0, nil, &log_size)
cl.CLGetProgramBuildInfo(program, devices[0], cl.CL_PROGRAM_BUILD_LOG,
log_size, &program_log, nil)
fmt.Printf("%s\n", program_log)
return
}
//utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLBuildProgram")
//-----------------------------------------------------
// STEP 7: Create the kernel
//-----------------------------------------------------
var kernel cl.CL_kernel
// Use clCreateKernel() to create a kernel
kernel = cl.CLCreateKernel(program, []byte("bst_kernel"), &status)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLCreateKernel")
defer cl.CLReleaseKernel(kernel)
//-----------------------------------------------------
// STEP 8: Set the kernel arguments
//-----------------------------------------------------
// Associate the input and output buffers with the
// kernel
// using clSetKernelArg()
// Set appropriate arguments to the kernel
status = cl.CLSetKernelArgSVMPointer(kernel,
0,
svmTreeBuf)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clSetKernelArgSVMPointer(svmTreeBuf)")
status = cl.CLSetKernelArgSVMPointer(kernel,
1,
svmSearchBuf)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clSetKernelArgSVMPointer(svmSearchBuf)")
//-----------------------------------------------------
// STEP 9: Configure the work-item structure
//-----------------------------------------------------
// Define an index space (global work size) of work
// items for
// execution. A workgroup size (local work size) is not
// required,
// but can be used.
var localWorkSize [1]cl.CL_size_t
var kernelWorkGroupSize interface{}
status = cl.CLGetKernelWorkGroupInfo(kernel,
devices[0],
cl.CL_KERNEL_WORK_GROUP_SIZE,
cl.CL_size_t(unsafe.Sizeof(localWorkSize[0])),
&kernelWorkGroupSize,
nil)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLGetKernelWorkGroupInfo")
localWorkSize[0] = kernelWorkGroupSize.(cl.CL_size_t)
var globalWorkSize [1]cl.CL_size_t
globalWorkSize[0] = NUMBER_OF_SEARCH_KEY
//-----------------------------------------------------
// STEP 10: Enqueue the kernel for execution
//-----------------------------------------------------
// Execute the kernel by using
// clEnqueueNDRangeKernel().
// 'globalWorkSize' is the 1D dimension of the
// work-items
status = cl.CLEnqueueNDRangeKernel(cmdQueue,
kernel,
1,
nil,
globalWorkSize[:],
localWorkSize[:],
0,
nil,
nil)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "CLEnqueueNDRangeKernel")
status = cl.CLFlush(cmdQueue)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clFlush")
status = cl.CLFinish(cmdQueue)
utils.CHECK_STATUS(status, cl.CL_SUCCESS, "clFinish")
//-----------------------------------------------------
// STEP 11: Read the output buffer back to the host
//-----------------------------------------------------
// Use clEnqueueReadBuffer() to read the OpenCL output
// buffer (bufferC)
// to the host output array (C)
//copy the data back to host buffer
//this demo doesn't need clEnqueueReadBuffer due to SVM
//-----------------------------------------------------
// STEP 12: Verify the results
//-----------------------------------------------------
// reference implementation
svmBinaryTreeCPUReference(cmdQueue,
svmRoot,
svmTreeBuf,
svmSearchBuf)
// compare the results and see if they match
pass := svmCompareResults(cmdQueue, svmSearchBuf)
if pass {
println("Passed!")
} else {
println("Failed!")
}
}
func main() {
/* OpenCL data structures */
var device []cl.CL_device_id
var context cl.CL_context
var queue cl.CL_command_queue
var program *cl.CL_program
var kernel cl.CL_kernel
var err cl.CL_int
/* Data and buffers */
dim := cl.CL_uint(2)
var global_offset = [2]cl.CL_size_t{3, 5}
var global_size = [2]cl.CL_size_t{6, 4}
var local_size = [2]cl.CL_size_t{3, 2}
var test [24]float32
var test_buffer cl.CL_mem
/* Create a device and context */
device = utils.Create_device()
context = cl.CLCreateContext(nil, 1, device[:], nil, nil, &err)
if err < 0 {
println("Couldn't create a context")
return
}
/* Build the program and create a kernel */
program = utils.Build_program(context, device[:], PROGRAM_FILE, nil)
kernel = cl.CLCreateKernel(*program, KERNEL_FUNC, &err)
if err < 0 {
println("Couldn't create a kernel")
return
}
/* Create a write-only buffer to hold the output data */
test_buffer = cl.CLCreateBuffer(context, cl.CL_MEM_WRITE_ONLY,
cl.CL_size_t(unsafe.Sizeof(test)), nil, &err)
if err < 0 {
println("Couldn't create a buffer")
return
}
/* Create kernel argument */
err = cl.CLSetKernelArg(kernel, 0, cl.CL_size_t(unsafe.Sizeof(test_buffer)), unsafe.Pointer(&test_buffer))
if err < 0 {
println("Couldn't set a kernel argument")
return
}
/* Create a command queue */
queue = cl.CLCreateCommandQueue(context, device[0], 0, &err)
if err < 0 {
println("Couldn't create a command queue")
return
}
/* Enqueue kernel */
err = cl.CLEnqueueNDRangeKernel(queue, kernel, dim, global_offset[:],
global_size[:], local_size[:], 0, nil, nil)
if err < 0 {
println("Couldn't enqueue the kernel")
return
}
/* Read and print the result */
err = cl.CLEnqueueReadBuffer(queue, test_buffer, cl.CL_TRUE, 0,
cl.CL_size_t(unsafe.Sizeof(test)), unsafe.Pointer(&test), 0, nil, nil)
if err < 0 {
println("Couldn't read the buffer")
return
}
for i := 0; i < 24; i += 6 {
fmt.Printf("%.2f %.2f %.2f %.2f %.2f %.2f\n",
test[i], test[i+1], test[i+2], test[i+3], test[i+4], test[i+5])
}
/* Deallocate resources */
cl.CLReleaseMemObject(test_buffer)
cl.CLReleaseKernel(kernel)
cl.CLReleaseCommandQueue(queue)
cl.CLReleaseProgram(*program)
cl.CLReleaseContext(context)
}
func TestQueue(t *testing.T) {
/* Host/device data structures */
var platform [1]cl.CL_platform_id
var device [1]cl.CL_device_id
var context cl.CL_context
var queue cl.CL_command_queue
var err cl.CL_int
/* Program/kernel data structures */
var program cl.CL_program
var program_buffer [1][]byte
var program_log interface{}
var program_size [1]cl.CL_size_t
var log_size cl.CL_size_t
var kernel cl.CL_kernel
/* Access the first installed platform */
err = cl.CLGetPlatformIDs(1, platform[:], nil)
if err < 0 {
t.Errorf("Couldn't find any platforms")
}
/* Access the first GPU/CPU */
err = cl.CLGetDeviceIDs(platform[0], cl.CL_DEVICE_TYPE_GPU, 1, device[:], nil)
if err == cl.CL_DEVICE_NOT_FOUND {
err = cl.CLGetDeviceIDs(platform[0], cl.CL_DEVICE_TYPE_CPU, 1, device[:], nil)
}
if err < 0 {
t.Errorf("Couldn't find any devices")
}
/* Create a context */
context = cl.CLCreateContext(nil, 1, device[:], nil, nil, &err)
if err < 0 {
t.Errorf("Couldn't create a context")
}
/* Read each program file and place content into buffer array */
program_handle, err1 := os.Open("blank.cl")
if err1 != nil {
t.Errorf("Couldn't find the program file")
}
defer program_handle.Close()
fi, err2 := program_handle.Stat()
if err2 != nil {
t.Errorf("Couldn't find the program stat")
}
program_size[0] = cl.CL_size_t(fi.Size())
program_buffer[0] = make([]byte, program_size[0])
read_size, err3 := program_handle.Read(program_buffer[0])
if err3 != nil || cl.CL_size_t(read_size) != program_size[0] {
t.Errorf("read file error or file size wrong")
}
/* Create program from file */
program = cl.CLCreateProgramWithSource(context, 1,
program_buffer[:], program_size[:], &err)
if err < 0 {
t.Errorf("Couldn't create the program")
}
/* Build program */
err = cl.CLBuildProgram(program, 1, device[:], nil, nil, nil)
if err < 0 {
/* Find size of log and print to std output */
cl.CLGetProgramBuildInfo(program, device[0], cl.CL_PROGRAM_BUILD_LOG,
0, nil, &log_size)
//program_log = (char*) malloc(log_size+1);
//program_log[log_size] = '\0';
cl.CLGetProgramBuildInfo(program, device[0], cl.CL_PROGRAM_BUILD_LOG,
log_size, &program_log, nil)
t.Errorf("%s\n", program_log)
//free(program_log);
}
/* Create the kernel */
kernel = cl.CLCreateKernel(program, []byte("blank"), &err)
if err < 0 {
t.Errorf("Couldn't create the kernel")
}
/* Create the command queue */
queue = cl.CLCreateCommandQueue(context, device[0], 0, &err)
if err < 0 {
t.Errorf("Couldn't create the command queue")
}
/* Enqueue the kernel execution command */
err = cl.CLEnqueueTask(queue, kernel, 0, nil, nil)
if err < 0 {
t.Errorf("Couldn't enqueue the kernel execution command")
} else {
t.Logf("Successfully queued kernel.\n")
}
/* Deallocate resources */
cl.CLReleaseCommandQueue(queue)
cl.CLReleaseKernel(kernel)
cl.CLReleaseProgram(program)
cl.CLReleaseContext(context)
}
func main() {
/* OpenCL data structures */
var device []cl.CL_device_id
var context cl.CL_context
var queue cl.CL_command_queue
var program *cl.CL_program
var kernel cl.CL_kernel
var err cl.CL_int
/* Data and buffers */
var r_coords = [4]float32{2, 1, 3, 4}
var angles = [4]float32{3 * M_PI / 8, 3 * M_PI / 4, 4 * M_PI / 3, 11 * M_PI / 6}
var x_coords, y_coords [4]float32
var r_coords_buffer, angles_buffer,
x_coords_buffer, y_coords_buffer cl.CL_mem
/* Create a device and context */
device = utils.Create_device()
context = cl.CLCreateContext(nil, 1, device[:], nil, nil, &err)
if err < 0 {
println("Couldn't create a context")
return
}
/* Create a kernel */
program = utils.Build_program(context, device[:], PROGRAM_FILE, nil)
kernel = cl.CLCreateKernel(*program, KERNEL_FUNC, &err)
if err < 0 {
println("Couldn't create a kernel")
return
}
/* Create a write-only buffer to hold the output data */
r_coords_buffer = cl.CLCreateBuffer(context, cl.CL_MEM_READ_ONLY|cl.CL_MEM_COPY_HOST_PTR,
cl.CL_size_t(unsafe.Sizeof(r_coords)), unsafe.Pointer(&r_coords[0]), &err)
if err < 0 {
println("Couldn't create a buffer")
return
}
angles_buffer = cl.CLCreateBuffer(context, cl.CL_MEM_READ_ONLY|cl.CL_MEM_COPY_HOST_PTR,
cl.CL_size_t(unsafe.Sizeof(angles)), unsafe.Pointer(&angles[0]), &err)
x_coords_buffer = cl.CLCreateBuffer(context, cl.CL_MEM_READ_WRITE,
cl.CL_size_t(unsafe.Sizeof(x_coords)), nil, &err)
y_coords_buffer = cl.CLCreateBuffer(context, cl.CL_MEM_READ_WRITE,
cl.CL_size_t(unsafe.Sizeof(y_coords)), nil, &err)
/* Create kernel argument */
err = cl.CLSetKernelArg(kernel, 0, cl.CL_size_t(unsafe.Sizeof(r_coords_buffer)), unsafe.Pointer(&r_coords_buffer))
if err < 0 {
println("Couldn't set a kernel argument")
return
}
cl.CLSetKernelArg(kernel, 1, cl.CL_size_t(unsafe.Sizeof(angles_buffer)), unsafe.Pointer(&angles_buffer))
cl.CLSetKernelArg(kernel, 2, cl.CL_size_t(unsafe.Sizeof(x_coords_buffer)), unsafe.Pointer(&x_coords_buffer))
cl.CLSetKernelArg(kernel, 3, cl.CL_size_t(unsafe.Sizeof(y_coords_buffer)), unsafe.Pointer(&y_coords_buffer))
/* Create a command queue */
queue = cl.CLCreateCommandQueue(context, device[0], 0, &err)
if err < 0 {
println("Couldn't create a command queue")
return
}
/* Enqueue kernel */
err = cl.CLEnqueueTask(queue, kernel, 0, nil, nil)
if err < 0 {
println("Couldn't enqueue the kernel")
return
}
/* Read and print the result */
err = cl.CLEnqueueReadBuffer(queue, x_coords_buffer, cl.CL_TRUE, 0,
cl.CL_size_t(unsafe.Sizeof(x_coords)), unsafe.Pointer(&x_coords), 0, nil, nil)
if err < 0 {
println("Couldn't read the buffer")
return
}
cl.CLEnqueueReadBuffer(queue, y_coords_buffer, cl.CL_TRUE, 0,
cl.CL_size_t(unsafe.Sizeof(y_coords)), unsafe.Pointer(&y_coords), 0, nil, nil)
/* Display the results */
for i := 0; i < 4; i++ {
fmt.Printf("(%6.3f, %6.3f)\n", x_coords[i], y_coords[i])
}
/* Deallocate resources */
cl.CLReleaseMemObject(r_coords_buffer)
cl.CLReleaseMemObject(angles_buffer)
cl.CLReleaseMemObject(x_coords_buffer)
cl.CLReleaseMemObject(y_coords_buffer)
cl.CLReleaseKernel(kernel)
cl.CLReleaseCommandQueue(queue)
cl.CLReleaseProgram(*program)
cl.CLReleaseContext(context)
}