Проблема с CUDA Driver API и ptx-программой

Добрый день !

Натолкнулся на следующую проблему -
имею простейшую программу test.ptx на ptx-ассемблере

.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;
}

(Эта программа записывает 0 в первый (0-ой по номеру) элемент массива,
переданного в качестве входного параметра).

Программа cuTest.cpp , которая ее вызывает, имеет вид:

#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 ];

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 ) );

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 );

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 ] );

}

Эта программа заполняет исходный массив значениями -1,
ожидая получить после своего выполнения 0 в первом элементе массива.

Компилируется данная программа просто:

nvcc -L/usr/lib64 -lcuda -o cuTest cuTest.cpp

а запускается так:

> ./cuTest 1 test.ptx test
(файл test.ptx должен находиться в той же директории, что и основная программа).

Однако, результат получается неправильным:

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

Т.е., в результирующий массив 0 не записался.
Аналогичный результат получается , когда ptx-программа имеет более сложный вид,
например, каждый поток пишет в соответствующий элемент массива свой номер.

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

/*
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-функция вообще "не видит" входной параметр.

Есть ли у кого-нибудь соображения в чем тут может быть дело ?

Forums: