reduction example using cuda and CUB

Question

I'm trying to get my head around CUB, and having a bit of trouble following the (rather incomplete) worked examples. CUB looks like it is a fantastic tool, I just can't make sense of the example code.

I've built a simple proto-warp reduce example:

#include <cub/cub.cuh>
#include <cuda.h>
#include <vector>
using std::vector;
#include <iostream>
using std::cout;
using std::endl;

const int N = 128;

__global__ void sum(float *indata, float *outdata) {
    typedef cub::WarpReduce<float,4> WarpReduce;
    __shared__ typename WarpReduce::TempStorage temp_storage;
    int id = blockIdx.x*blockDim.x+threadIdx.x;
    if( id < 128 ) {
        outdata[id] = WarpReduce(temp_storage).Sum(indata[id]);
    }
}

int main() {
    vector<float> y(N), sol(N);
    float *dev_y, *dev_sol;
    cudaMalloc((void**)&dev_y,N*sizeof(float));
    cudaMalloc((void**)&dev_sol,N*sizeof(float));
    for( int i = 0; i < N; i++ ) {
        y[i] = (float)i;
    }
    cout << "input: ";
    for( int i = 0; i < N; i++ ) cout << y[i] << " ";
    cout << endl;
    cudaMemcpy(&y[0],dev_y,N*sizeof(float),cudaMemcpyHostToDevice);
    sum<<<1,32>>>(dev_y,dev_sol);
    cudaMemcpy(dev_sol,&sol[0],N*sizeof(float),cudaMemcpyDeviceToHost);
    cout << "output: ";
    for( int i = 0; i < N; i++ ) cout << sol[i] << " ";
    cout << endl;
    cudaFree(dev_y);
    cudaFree(dev_sol);
    return 0;
}

which returns all zeros.

I'm aware that this code would return a reduction that was banded with every 32nd element being the sum of a warp and the other elements being undefined - I just want to get a feel for how CUB works. Can someone point out what I'm doing wrong?

(also, does CUB deserve its own tag yet?)

[cuda error checking](http://stackoverflow.com/questions/14038589/what-is-the-canonical-way-to-check-for-errors-using-the-cuda-runtime-api) is always a good idea, especially when you are having trouble with CUDA code. It would have thrown errors on your `cudaMemcpy` calls. — Robert Crovella, Aug 29 '13 at 20:17
please don't edit questions to fix code/solve problems. you are invalidating the.question by doing so. — talonmies, Sep 02 '13 at 06:56

score 1 · Accepted Answer · answered Aug 29 '13 at 08:41

1

Your cudaMemcpy arguments are back to front, the destination comes first (to be consistent with memcpy).

cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count, cudaMemcpyKind kind )

See the API reference for more info.

answered Aug 29 '13 at 08:41

Tom

20,852
4
42
54

reduction example using cuda and CUB

1 Answers1