Добрый день !

Хочу написать простейшую работающую программу на ptx-ассемблере.

Для начала, написал основную программу (cuTest.cpp), которая передает в вызываемое вычислительное ядро массив,
проинициализированный значением -1:

Код:
#include "stdio.h"
#include "malloc.h"
#include "cuda.h"
#include "cuda_runtime_api.h"

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

int main ( int argc , char** argv ) {

 int rt_version;

 int N = 8;

 int nDevice = atoi ( argv [ 1 ] );
 printf ( "Device number = %d\n", nDevice );

 char* file_name = argv [ 2 ];
 char* func_name = argv [ 3 ];

 CUresult result = cuInit(0);

 int deviceCount = 0;
 cuDeviceGetCount ( &deviceCount );
 printf ( "deviceCount = %d\n", deviceCount );

 CUdevice cuDevice = 0;
 cuDeviceGet ( &cuDevice, nDevice );

 CUcontext cuContext;
 cuCtxCreate ( &cuContext, 0, cuDevice );

 cudaRuntimeGetVersion ( &rt_version );
 printf ( "Runtime version = %d.%d\n", rt_version/1000, rt_version%100 );

 int* A = (int*) malloc ( N * sizeof ( int ) );
 for ( int i = 0; i < N; i++ )
  A [ i ] = -1;
 printf ( "Before ->\n" );
 for ( int i = 0; i < N; i++ )
  printf ( "%d\n", A [ i ] );

 CUdeviceptr dA = cuMemAlloc ( &dA, N * sizeof ( int ) );

 cuMemcpyHtoD(dA, A, N * sizeof ( int ) );

 CUmodule cuModule;
 result = cuModuleLoad ( &cuModule, file_name );
 printf ( "ModuleLoad result = %d\n", result );

 CUfunction cuFunc;
 result = cuModuleGetFunction ( &cuFunc, cuModule, func_name );
 printf ( "GetFunction result = %d\n", result );

 int offset = 0;
 void* ptr;

 ptr = (void*)(size_t)dA;
 ALIGN_UP(offset, __alignof(ptr));

 result = cuParamSetv ( cuFunc, offset, &ptr, sizeof (ptr) );
 printf ( "ParamSetv result = %d\n", result );
 offset += sizeof(ptr);

 result = cuParamSetSize ( cuFunc, offset );
 printf ( "ParamSetSize result = %d\n", result );

 int threadsPerBlock = N;
 int blocksPerGrid   = 1;

 result = cuFuncSetBlockShape ( cuFunc, threadsPerBlock, 1, 1);
 printf ( "FuncSetBlockShape result = %d\n", result );

 result = cuLaunchGrid ( cuFunc, blocksPerGrid, 1 );
 printf ( "LaunchGrid result = %d\n", result );

 cuMemcpyDtoH( A, dA, N * sizeof ( int ) );

 printf ( "After ->\n" );
 for ( int i = 0; i < N; i++ )
  printf ( "%d\n", A [ i ] );

}

Транслируется она обычным образом:

Код:
nvcc -L/usr/lib64 -lcuda -o cuTest cuTest.cpp

Затем написал простейший ptx-код (файл test.ptx, который лежит в одной директории с главной программой), который присваивает первому элементу массива (т.е., 0-му по номеру) значение 0:

Код:
 
    .version 1.4
    .target sm_10

        .entry test ( .param .s32 C )
        {
         .reg .s32 %p<1>;

          ld.param.s32 %p0, [C];
          st.global.s32 [%p0], 0;
          ret;
        }

Запускаю и получаю :

Код:
./cuTest 1 test.ptx test

Device number = 1
deviceCount = 2
Runtime version = 3.20
Before ->
-1
-1
-1
-1
-1
-1
-1
-1
ModuleLoad result = 0
GetFunction result = 0
ParamSetv result = 0
ParamSetSize result = 0
FuncSetBlockShape result = 0
LaunchGrid result = 0
After ->
-1
-1
-1
-1
-1
-1
-1
-1

Т.е., программа, вроде бы отрабатывает, но никаких изменений в массиве не происходит.

Более того, если в основной программе закомментировать часть, связанную с передачей параметров,
а именно

Код:
/*
int offset = 0;
 void* ptr;

 ptr = (void*)(size_t)dA;
 ALIGN_UP(offset, __alignof(ptr));

 result = cuParamSetv ( cuFunc, offset, &ptr, sizeof (ptr) );
 printf ( "ParamSetv result = %d\n", result );
 offset += sizeof(ptr);

 result = cuParamSetSize ( cuFunc, offset );
 printf ( "ParamSetSize result = %d\n", result );
*/

результат получается точно такой же -
т.е., похоже, что ptx-функция "не видит" входных параметров.

Есть ли у кого-нибудь успешный опыт написания своих ptx-программ,
и в чём тут может быть дело ?

Спасибо.