OpenCL πρόβλημα.

[Dr.Fuzzy]

Παίδες ο κώδικας είναι ο παρακάτω (κόψε-ράψε και αρκετά πρόχειρος):

//
//  main.c
//  OpenCL_Sum
//
//  Created by Dr.Fuzzy on 6/25/14.
//  Copyright (c) 2014 CUT. All rights reserved.
//
 
#define KERNEL_FUNC "float_sum"
#define ARRAY_SIZE 1000000
 
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
 
#ifdef __APPLE__
#include "OpenCL/opencl.h"
#else
#include "CL/cl.h"
#endif
 
// A utility function that checks that the kernel execution performs the
// requested work over the entire range of data.
static int validate(cl_float* input, cl_float* output) {
    int i;
    for (i = 0; i < ARRAY_SIZE; i++) {
        
        // The kernel was supposed to sum each value.
        if ( output[i] != (input[i] + input[i]) ) {
            fprintf(stdout,
                    "Error: Element %d did not match expected output.\n", i);
            fprintf(stdout,
                    "       Saw %1.4f, expected %1.4f\n", output[i],
                    input[i] + input[i]);
            fflush(stdout);
            return 0;
        }
    }
    return 1;
}
 
int main()
{
    char name[128];
    // OpenCL Data structures
    cl_platform_id platform;
    cl_device_id device;
    cl_context context;
    cl_program program;
    cl_kernel kernel;
    cl_command_queue queue;
    cl_mem vec_buffer, result_buffer;
    
    cl_event prof_event;;
    
    // C Data Structures / Data types
    float *vec, *non_parallel;
    float result[ARRAY_SIZE];
        
    cl_ulong time_start, time_end, total_time;
    
    int i;
    
    // Errors
    cl_int err;
    
    non_parallel = (float*)malloc(ARRAY_SIZE * sizeof(float));
    vec          = (float*)malloc(ARRAY_SIZE * sizeof(float));
    
    // Initialize the vector of floats
    for(i = 0; i < ARRAY_SIZE; i++)
        vec[i] = i + 1;
    
    //************************* C Function **************************************
    clock_t start, end;
    
    start = clock();
    
    for( i = 0; i < ARRAY_SIZE; i++)
    {
        non_parallel[i] = vec[i] * vec[i];
    }
    end = clock();
    printf( "Number of seconds: %f\n", (clock()-start)/(double)CLOCKS_PER_SEC );
    
    free(non_parallel);
    //***************************************************************************
    
    clGetPlatformIDs(1, &platform, NULL);
    clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
    clGetDeviceInfo(device, CL_DEVICE_NAME, 128, name, NULL);
    fprintf(stdout, "Created a dispatch queue using the %s\n", name);
    context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
    
    // Context error?
    if(err)
    {
        perror("Cannot create context");
        return 1;
    }
       
    queue = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, &err);
    
    if(err)
    {
        perror("Cannot create command queue");
        return 1;
    }
    
    vec_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                sizeof(float) * ARRAY_SIZE, vec, &err);
    result_buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(float)*ARRAY_SIZE, NULL, &err);
    
    if(err)
    {
        perror("Cannot create the vector buffer");
        return 1;
    }
    
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &vec_buffer);
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &result_buffer);
    
    size_t global_size = ARRAY_SIZE;
    
    clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, NULL, 0, NULL, &prof_event);
    clEnqueueReadBuffer(queue, result_buffer, CL_TRUE, 0, sizeof(float)*ARRAY_SIZE, &result, 0, NULL, NULL);
    clFinish(queue);
    
    clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_START,
                            sizeof(time_start), &time_start, NULL);
    clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_END,
                            sizeof(time_end), &time_end, NULL);
    total_time += time_end - time_start;
    
    printf("\nAverage time in nanoseconds = %lu\n", total_time/ARRAY_SIZE);
 
    // Check to see if the kernel did what it was supposed to:
    if (validate(vec, &result_buffer)) {
        fprintf(stdout, "All vectors were properly added.\n");
    }
    
    clReleaseMemObject(vec_buffer);
    clReleaseMemObject(result_buffer);
    clReleaseKernel(kernel);
    clReleaseCommandQueue(queue);
    clReleaseContext(context);
    
    free(vec);
    
    return 0;
}

και ο kernel αυτός:

//
//  mykernel.cl
//  OpenCL_Sum
//
//  Created by Dr.Fuzzy on 6/16/14.
//  Copyright (c) 2014 CUT. All rights reserved.
//
 
__kernel void float_sum(__global float* vec,__global float* result){
    int gid = get_global_id(0);
    result[gid] = vec[gid] * vec[gid];
}

Λοιπόν η βλακεία (μάλλον) είναι εδώ που κάνω τη σύγκριση μεταξύ CPU και Kernel calculated values, καπου μπερδεύω τους τύπους (cl_mem, cl_float, κλπ):

// Check to see if the kernel did what it was supposed to:
    if (validate(vec, &result_buffer)) {
        fprintf(stdout, "All vectors were properly added.\n");
    }

ιδού και το execution:

Number of seconds: 0.004516
Created a dispatch queue using the HD Graphics 4000
 
Average time in nanoseconds = 0
Error: Element 0 did not match expected output.
       Saw 0.0000, expected 2.0000
Program ended with exit code: 0

Επεξ/σία 26 Ιουνίου 2014 από Dr.Fuzzy

[Dr.Fuzzy]

κανείς..;

[παπι]

Δεν εχεις debbuger tools? Λογικα εχεις stackoverflow. 1εκ floats? στη stack???

[Dr.Fuzzy]

Αναγκαστικά θα κάνω debug...

Και 100 να βάλω το ίδιο γινεται, δεν νομίζω οτι ειναι θεμα stack. Κάτι με τα types...

[ChRis6]

Δεν θα κάνεις compile τον kernel ;

[Dr.Fuzzy]

Προφανώς και κάνω compile τον Kernel!

[ChRis6]

Προφανώς και κάνω compile τον Kernel!

Σε ποιο σημείο τον κάνεις compile; Εγώ βλέπω οτι φτιάχνεις δυο buffers και μετα περνας ορίσματα.

 

Το 'average time in nanoseconds' είναι μηδέν,κάτι που δεν μπορεί να συμβεί αφού ακόμα και το kernel launch έχει overhead

[Dr.Fuzzy]

Το kernel γινεται compile για να φτιαχτεί το executable εκτός και αν εννοείς κάτι άλλο με 'compile'.

 

To launch δε το μετράω όπως και το transfer απο το kernel στο host.

 

Βγάζει 0 όχι γιατι δεν τρέχει αλλα γιατι ειναι λάθος το timescale μάλλον. Ο kernel τρέχει αλλα προφανώς βγάζει λάθος αποτελέσματα και μάλλον φταίει το type (cl_mem και cl_float).

[ChRis6]

Για να κάνεις το device(στην περίπτωσή σου η GPU) να τρέξει τον κώδικα που έχεις στο mykernel.cl πρέπει να διαβάσεις το αρχείο (text) και να δώσεις το string στον driver,να κάνεις compile και να δώσεις το εκτελέσιμο στη GPU.Όταν λέω να κάνεις compile τον kernel,εννοώ μέσα στον κώδικα σου(π.χ στη main) όχι να φτιάξεις το εκτελέσιμο του main.c

O compiler για το main.c δεν φτιάχνει κώδικα για GPU αλλά για CPU.Ο kernel στο mykernel.cl είναι κώδικας που θα τρέξει στο device(GPU).Μέσα από τον κώδικα σου φροντίζεις να φτιάξεις το εκτελέσιμο για τo device.

Βλέπω οτι έχεις ορίσει τα cl_program και cl_kernel,αλλά δεν τα χρησιμοποιείς...Αυτό σημαίνει οτι είτε αυτά που γράφω παραπάνω σου φάινονται κινέζικα και από κάπου τσίμπησες τον κώδικα,είτε σκέφτεσαι "Τι λέει τώρα ο μαβλάκας"

Τσέκαρε αυτό :
http://www.fixstars.com/en/opencl/book/OpenCLProgrammingBook/online-offline-compilation/

[Dr.Fuzzy]

O κώδικας όπως έγραψα στο 1ο πόστ ειναι κόψε ραψε (δεν το έκρυψα!), που σημαίνει οτι ναι εχω πάρει κάποια κομμάτια έτοιμα.

Text; Δε θελω να διαβάσω text με τιμές, κάνω hardcode κάποιες τιμές μέσα στη main αν εννοείς αυτό.

Εχω 2-3 μερες που παίζω με opencl οποτε όντως μπορει και να λέω βλακείες!

Επεξ/σία 28 Ιουνίου 2014 από Dr.Fuzzy

[ChRis6]

Διάβασε το link που δίνω

[Dr.Fuzzy]

Οκ

Ωραία πες μου κάτι αν δε βαριέσαι μια και το γνωρίζεις καλα το θεμα. Στο παρακάτω παράδειγμα το text που κολλάει (το παράδειγμα το εχω δοκιμάσει και μου δουλευει μια χαρα). Ο κώδικας που έκανα πόστ ειναι βάση αυτού αλλα σε c++.

https://developer.apple.com/library/mac/documentation/Performance/Conceptual/OpenCL_MacProgGuide/ExampleHelloWorld/Example_HelloWorld.html#//apple_ref/doc/uid/TP40008312-CH112-SW3

 

https://developer.apple.com/library/mac/documentation/Performance/Conceptual/OpenCL_MacProgGuide/ExampleHelloWorld/Example_HelloWorld.html#//apple_ref/doc/uid/TP40008312-CH112-SW5

[παπι]

Το κανει compile μεσου ide. Εσυ καντο οπως σου λεει ο chris. Δηλαδη on the fly compile. Για να σου ειναι πιο ξεκαθαρο το context

[Dr.Fuzzy]

Μέσω Xcode το κάνω για να πω την αμαρτία μου, αλλά οκ θα το δω να καταλάβω τι εννοείς...Το παράδειγμα των λινκ που έδωσα είναι offline ή online; Πάντως 

 

 

	/* Load kernel source code */
	fp = fopen(fileName, "r");
	if (!fp) {

 

δεν υπάρχει κάπου.

[παπι]

Να στο πω πολυ απλα. Το προγραμμα σου (exe ή οπως λεγεται στο xos ) εκτος απο αυτα που θα κανει, θα κανει και compile ενα αλλο προγραμμα που λεγεται shader.

Αυτο μπορεις να το κανεις και απο την αρχη, πχ μαζι με το προγραμμα σου να εχεις και compiled shaders ειτε hardcoded μεσα στο προγραμμα (στο .text seg) ειτε σα ξεχωριστα αρχεια. Καλο ειναι ομως να κανεις compile κατα το runtime του προγραμματος για λογους συμβατοτητας.

 

Δηλαδη η διαδικασια που θα ακολουθησει το προγραμμα σου οταν τρεχει.

 

1) load shader source (αρχεια ή απλα strings μεσα στο προγραμμα που εχουν αυτες τις kernel functions)

2) compile αυτο το source με τον compiler που εχει το openCL

3) τα στελνεις στη gpu

4) τα κανεις invoke.

 

Να το πω και αλλιως. Αυτες οι kernel function ΔΕΝ ΕΙΝΑΙ C. Ειναι shader language που ειναι c like και γινονται compile απο το openCL. Αλλη γλωσσα, αλλος compiler.

 

 

Αντε αλλη μια. Το προγραμμα σου ειναι σε php, οι shader σε javascript. Τετοια φαση.