Giovanni Di Grezia - VFX Artist Forum - CUDA C / C++ https://www.xgiovio.com/forum/ The parallel computing. The future en Sat, 04 Apr 2026 21:44:36 GMT vBulletin 60 images/misc/rss.png Giovanni Di Grezia - VFX Artist Forum - CUDA C / C++ https://www.xgiovio.com/forum/ Dot Product - Shared cache, reduction and final partial sum on host https://www.xgiovio.com/forum/forum/programming/cuda-c-c/446-dot-product-shared-cache-reduction-and-final-partial-sum-on-host Sat, 16 Nov 2013 13:02:35 GMT
PHP Code:
#include <iostream>

#define minimo(a,b) a<b?a:b
#define MAXBLOCKS 32
#define NTHREADS 256 // must be a power of 2


__global__ void dot (int * a, int * b, int *c, int size){


    int id = (blockIdx.x * blockDim.x) + threadIdx.x;
    int nextid = gridDim.x * blockDim.x;

    //// reducted output
    __shared__ int shared_cache [NTHREADS];

    int sum = 0;
    for (;id < size ;){
        sum += (*(a+id)) * (*(b+id));
        id+= nextid;
    }

    *(shared_cache + threadIdx.x) = sum;

    __syncthreads();

    ///////// sum of internal cache

    int i;    
    
    for (i=(NTHREADS /2); i>0 ; i= i/2){
            if (threadIdx.x < i){
                *(shared_cache + threadIdx.x) += *(shared_cache + threadIdx.x + i);
            }
            __syncthreads();
    }

    if ( threadIdx.x == 0){
        *(c+ blockIdx.x) = shared_cache[0];
    }
}




int main () {

    
    int n = 100; /// size of arrays to multiply
    int n_blocks = minimo( MAXBLOCKS, ((n+NTHREADS-1)/NTHREADS));

    int * a = new int[n];
    int * b = new int[n];
    int * c = new int[n_blocks];

    ////////////////////////sample data allocated on host
    for (int i=0 ; i< n ; ++i){
        *(a+i)=1;
        *(b+i)=2;
    }

    int * d_a;
    int * d_b;
    int * d_c;

    

        cudaMalloc((void**) &d_a, n * sizeof(int));
        cudaMalloc((void**) &d_b, n * sizeof(int));
        cudaMalloc((void**) &d_c, n_blocks * sizeof(int));

        cudaMemcpy(d_a , a, n * sizeof(int),cudaMemcpyHostToDevice);
        cudaMemcpy(d_b , b, n * sizeof(int),cudaMemcpyHostToDevice);

        dot<<< n_blocks ,NTHREADS>>>(d_a,d_b,d_c,n);

        cudaMemcpy(c,d_c,n_blocks*sizeof(int),cudaMemcpyDeviceToHost);


        //// final sum on host
        int final_result = 0;
        for (int i=0 ; i< n_blocks ; ++i){
            final_result += *(c+i);
        }

        std::cout << final_result << std::endl;

        std::cin.get();

return 0;} 
]]> CUDA C / C++ xgiovio https://www.xgiovio.com/forum/forum/programming/cuda-c-c/446-dot-product-shared-cache-reduction-and-final-partial-sum-on-host Mergesort - 2^n int array - blocks + threads https://www.xgiovio.com/forum/forum/programming/cuda-c-c/445-mergesort-2-n-int-array-blocks-threads Fri, 15 Nov 2013 14:07:13 GMT
PHP Code:
#include <iostream>
#include <stdlib.h>
#include <time.h>
#include <math.h>
#include <fstream>

using namespace std;

// standard merge function
__device__ void merge (int * first_array, int size1, int * second_array, int size2){

    int i,j,k;
    int * third_array;
    //cudaMalloc ((void**)&third_array, (size1 + size2) * sizeof(int));
    third_array = (int *) malloc((size1 + size2) * sizeof(int));

    for(i=0, j=0, k=-1 ; (i<size1) && (j<size2) ;){
        if ( *(first_array + i) <= *(second_array + j)  ){
            ++k;
            *(third_array + k) =  *(first_array + i);
            ++i;
        } else {
            ++k;
            *(third_array + k) =  *(second_array + j);
            ++j;
        }
    }
    if( i>= size1 ){
        for(;j<size2;){
            ++k;
            *(third_array + k) =  *(second_array + j);
            ++j;
        }
    } else {
        for(;i<size1;){
            ++k;
            *(third_array + k) =  *(first_array + i);
            ++i;
        }
    }


    for(i=0; i<size1 ;++i){
        *(first_array + i) = *(third_array + i);
    }
    j=i;
    for(i=0; i<size2 ;++i){
        *(second_array + i) = *(third_array + i + j);
    }
    //cudaFree (third_array) ;
    free (third_array) ;
}


__global__ void mergesort(int * pointer,int size, int factor){

    int id = ( blockIdx.x * blockDim.x) + threadIdx.x;
    int * pointer_to_pass = NULL;
    for (;id< (size/factor); id += gridDim.x * blockDim.x){
        pointer_to_pass = pointer + ( id * factor); // pointer to pass to each merge
        merge( pointer_to_pass, (factor/2), pointer_to_pass + (factor/2), (factor/2) );
    }
}




int main( void ) {
    /*
    fstream fileout ("data.txt",fstream::out); // file to write the non-ordered and ordered array
    fstream fileout_bin ("data_bin.txt",fstream::out | fstream::binary ); // ordered array in binary mode
    */

    srand (time(NULL));

    int n = 65536;
    int * numeri = new int[n];

    for (int i =0; i< n; ++i){
    *(numeri + i) = rand() +  rand() + rand() + rand() +  rand() + rand() ;  //fill the array with random data
    }

    /*
    fileout << "pre" << endl; // print to file the non-ordered array
    for(int i = 0;i<n;++i){
    fileout << numeri[i]<< endl;
    }
    */
    


    int* d_numeri;
    cudaMalloc((void **)&d_numeri,n*sizeof(int)); //allocating n*sizeof(int) bytes on gpu ram

    cudaMemcpy(d_numeri,numeri,n*sizeof(int),cudaMemcpyHostToDevice); // copy array from host to gpu

    for(int i = 2 ; i<=n ;i=i*2) {
        mergesort<<<128,128>>>(d_numeri,n,i); // launch instances of mergesort and run in parallel
    }

    cudaMemcpy(numeri,d_numeri,n*sizeof(int),cudaMemcpyDeviceToHost); // copy the ordered array from device to host
    
    /*
    fileout << "after" << endl; // write in file the ordered array
    for(int i = 0;i<n;++i){
    fileout << numeri[i]<< endl;
    }

    fileout_bin.write((char *) numeri, n * sizeof(int)); // write to file the ordered array in binary mode
    */


    
return 0;
} 
]]> CUDA C / C++ xgiovio https://www.xgiovio.com/forum/forum/programming/cuda-c-c/445-mergesort-2-n-int-array-blocks-threads MergeSort - N/2 to 1 blocks - 1 thread per block - Int array with size = 2^n https://www.xgiovio.com/forum/forum/programming/cuda-c-c/444-mergesort-n-2-to-1-blocks-1-thread-per-block-int-array-with-size-2-n Wed, 13 Nov 2013 20:32:32 GMT
PHP Code:
#include <iostream>
#include <stdlib.h>
#include <time.h>
#include <math.h>
#include <fstream>

using namespace std;

// standard merge function
__device__ void merge (int * first_array, int size1, int * second_array, int size2){

    int i,j,k;
    int * third_array;
    //cudaMalloc ((void**)&third_array, (size1 + size2) * sizeof(int));
    third_array = (int *) malloc((size1 + size2) * sizeof(int));

    for(i=0, j=0, k=-1 ; (i<size1) && (j<size2) ;){
        if ( *(first_array + i) <= *(second_array + j)  ){
            ++k;
            *(third_array + k) =  *(first_array + i);
            ++i;
        } else {
            ++k;
            *(third_array + k) =  *(second_array + j);
            ++j;
        }
    }
    if( i>= size1 ){
        for(;j<size2;){
            ++k;
            *(third_array + k) =  *(second_array + j);
            ++j;
        }
    } else {
        for(;i<size1;){
            ++k;
            *(third_array + k) =  *(first_array + i);
            ++i;
        }
    }


    for(i=0; i<size1 ;++i){
        *(first_array + i) = *(third_array + i);
    }
    j=i;
    for(i=0; i<size2 ;++i){
        *(second_array + i) = *(third_array + i + j);
    }
    //cudaFree (third_array) ;
    free (third_array) ;
}


__global__ void mergesort(int * pointer, int factor){

    int * pointer_to_pass = pointer + ( blockIdx.x * factor); // pointer to pass to each merge

    merge( pointer_to_pass, (factor/2), pointer_to_pass + (factor/2), (factor/2) );
}




int main( void ) {

    fstream fileout ("data.txt",fstream::out); // file to write the non-ordered and ordered array
    fstream fileout_bin ("data_bin.txt",fstream::out | fstream::binary ); // ordered array in binary mode

    srand (time(NULL));

    int n = 32768;
    int * numeri = new int[n];

    for (int i =0; i< n; ++i){
    *(numeri + i) = rand() % 2000 +  rand() % 2000 + rand() % 2000;  //fill the array with random data
    }

    fileout << "pre" << endl; // print to file the non-ordered array
    for(int i = 0;i<n;++i){
    fileout << numeri[i]<< endl;
    }


    int* d_numeri;
    cudaMalloc((void **)&d_numeri,n*sizeof(int)); //allocating n*sizeof(int) bytes on gpu ram

    cudaMemcpy(d_numeri,numeri,n*sizeof(int),cudaMemcpyHostToDevice); // copy array from host to gpu

    for(int i = 2 ; i<=n ;i=i*2) {
        mergesort<<<(n/i),1>>>(d_numeri,i); // launch instances of mergesort and run in parallel
    }

    cudaMemcpy(numeri,d_numeri,n*sizeof(int),cudaMemcpyDeviceToHost); // copy the ordered array from device to host
    
    
    fileout << "after" << endl; // write in file the ordered array
    for(int i = 0;i<n;++i){
    fileout << numeri[i]<< endl;
    }

    fileout_bin.write((char *) numeri, n * sizeof(int)); // write to file the ordered array in binary mode


    
return 0;
} 
]]> CUDA C / C++ xgiovio https://www.xgiovio.com/forum/forum/programming/cuda-c-c/444-mergesort-n-2-to-1-blocks-1-thread-per-block-int-array-with-size-2-n A simple Sum https://www.xgiovio.com/forum/forum/programming/cuda-c-c/443-a-simple-sum Mon, 11 Nov 2013 18:03:19 GMT
PHP Code:
#include <iostream>

__global__ void sum (int * a, int *b, int* c){

    *c = (*a + *b);
}


int main( void ) {

    int a,b,c;
    a=10;
    b=20;

    int * d_a_pointer;
    int * d_b_pointer;
    int * d_c_pointer;

    cudaMalloc ((void **)&d_a_pointer, sizeof(int));
    cudaMalloc ((void **)&d_b_pointer, sizeof(int));
    cudaMalloc ((void **)&d_c_pointer, sizeof(int));


    cudaMemcpy( d_a_pointer, &a, sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy( d_b_pointer, &b, sizeof(int), cudaMemcpyHostToDevice);

    sum<<<1,1>>>(d_a_pointer, d_b_pointer, d_c_pointer);
    cudaMemcpy( &c, d_c_pointer , sizeof(int), cudaMemcpyDeviceToHost);


    std::cout << c;

    std::cin.get();


return 0;
} 
]]> CUDA C / C++ xgiovio https://www.xgiovio.com/forum/forum/programming/cuda-c-c/443-a-simple-sum