DriverAPIはよーわからん

とりあえずRuntimeAPIを使えば何とかなるんだけど、DriverAPI叩けるようになったら幅広がるよねーと思って、マニュアル参照しながらVecAddのやつ作ってみたんだけど、なぜかCUDA_ERROR_LAUNCH_FAILEDが出る。
まぁ私が悪いっていう命題が真なわけですが、どうしたらいいんだろ。

とりあえずソースうｐしてみる。誰かなんかわかったら教えてください。
分からないことといえば結構あって、cubinとかptxを生成すると__global__の関数名が変わって、cuModuleGetFunction()使うときに直接"VecAdd"とかで呼べないんだけど、解決法誰か知らないかな。あと、実行ファイルに埋め込まれたcubinなりptxを呼び出すにはどうしたらいいんだろ。まぁ結局ドキュメントを片っ端から読めやｺﾞﾙｧって事なんでしょうけど、めんどい。

#include <iostream>
#include <vector>

using namespace std;

#include <cuda.h>


__global__ void VecAdd( const float *A, const float *B, float *C )
{
    unsigned long tid = threadIdx.x;
    C[ tid ] = A[ tid ] + B[ tid ];
}

int main( void )
{
    const float h_A[ 10 ] = { 0.0f, 1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f, 9.0f };
    const float h_B[ 10 ] = { 11.0f, 12.0f, 13.0f, 14.0f, 15.0f, 16.0f, 17.0f, 18.0f, 19.0f, 20.0f };

    if ( cuInit( 0 ) != CUDA_SUCCESS )
    {
        return -1;
    }

    int deviceCount = 0;
    cuDeviceGetCount( &deviceCount );
    if ( !deviceCount )
    {
        cerr << "There is no device supporting CUDA." << endl;
        return -1;
    }

    for ( int i = 0; i < deviceCount; ++i )
    {
        CUdevice dev;
        cuDeviceGet( &dev, i );

        CUcontext context;
        cuCtxCreate( &context, 0, dev );

        CUdevice curdev;
        cuCtxGetDevice( &curdev );
        char cur[ BUFSIZ ];
        cuDeviceGetName( cur, sizeof( cur ), curdev );
        cout << "Now running under... [" << cur << "]" << endl;

        CUmodule module;
        cuModuleLoad( &module, "VecAdd.cubin" );

        size_t size = sizeof( h_A );
        CUdeviceptr d_A;
        cuMemAlloc( &d_A, size );
        CUdeviceptr d_B;
        cuMemAlloc( &d_B, size );
        CUdeviceptr d_C;
        cuMemAlloc( &d_C, size );

        cuMemcpyHtoD( d_A, h_A, size );
        cuMemcpyHtoD( d_B, h_B, size );

        CUfunction vecAdd;
        cuModuleGetFunction( &vecAdd, module, "_Z6VecAddPKfS0_Pf" );

#define ALIGN_UP( offset, alignment ) ( offset ) = ( ( offset ) + ( alignment ) - 1 ) & ~( ( alignment ) - 1 )
        int offset = 0;

        void *ptr = ( void * )( size_t )d_A;
        ALIGN_UP( offset, __alignof( ptr ) );
        cuParamSetv( vecAdd, offset, &ptr, sizeof( ptr ) );
        offset += sizeof( ptr );

        ptr = ( void * )( size_t )d_B;
        ALIGN_UP( offset, __alignof( ptr ) );
        cuParamSetv( vecAdd, offset, &ptr, sizeof( ptr ) );
        offset += sizeof( ptr );

        ptr = ( void * )( size_t )d_C;
        ALIGN_UP( offset, __alignof( ptr ) );
        cuParamSetv( vecAdd, offset, &ptr, sizeof( ptr ) );
        offset += sizeof( ptr );
        cout << "offset: " << offset << endl;

        cuParamSetSize( vecAdd, offset );
        cuFuncSetBlockShape( vecAdd, 10, 1, 1 );
        cuLaunchGrid( vecAdd, 1, 1 );

        float h_C[ 10 ];
        cuMemcpyDtoH( h_C, d_C, size );
        for ( int i = 0; i < 10; ++i )
        {
            cout << h_C[ i ] << " " << ends;
        }
        cout << endl;

        cuMemFree( d_A );
        cuMemFree( d_B );
        cuMemFree( d_C );
    }

    return 0;
}