Question

Я объявляю классы A и B, оба с членом int* m_pDevicePtr, значение которого устанавливается с помощью cudaMalloc().

Я проверяю результаты двух ядер:

Функция ядра, такая, что входными параметрами являются следующие указатели устройства: kernelAdd(int* a, int* b).

При использовании этого ядра выходные данные соответствуют ожидаемым.
Функция ядра с сигнатурой kernelAdd(int* a, B* pB), которая использует функцию-член B с сигнатурой B::GetNumber(int index) для выполнения той же работы.

... но это ядро делаетне работает.

Вот код:

#include <iostream>
#include <fstream>
#include <string>
#include <stack>
#include <cstdarg>

#include <limits.h>
#include <windows.h>
#include <tchar.h>
#include <stdio.h>
#include <stdarg.h>
#include <math.h>
#include <malloc.h>
#include <stdlib.h>

#include "cuda_runtime.h"
#include "vector_types.h"

#define checkCudaErrors(val) check((val), #val, __FILE__, __LINE__)

#ifdef __DRIVER_TYPES_H__
#ifndef DEVICE_RESET
#define DEVICE_RESET cudaDeviceReset();
#endif
#else
#ifndef DEVICE_RESET
#define DEVICE_RESET
#endif
#endif

#ifdef __DRIVER_TYPES_H__
static const char *_cudaGetErrorEnum(cudaError_t error) {
    return cudaGetErrorName(error);
}
#endif

template <typename T> void check(T result, char const *const func, const char *const file,
    int const line) {
    if (result) {
        fprintf(stderr, "CUDA error at %s:%d code=%d(%s) \"%s\" \n", file, line,
            static_cast<unsigned int>(result), _cudaGetErrorEnum(result), func);
        DEVICE_RESET
            // Make sure we call CUDA Device Reset before exiting
            exit(EXIT_FAILURE);
    }
}

__global__ void _kInitialArray(int* thearray)
{
    int iX = threadIdx.x + blockDim.x * blockIdx.x;
    int iY = threadIdx.y + blockDim.y * blockIdx.y;
    int iZ = threadIdx.z + blockDim.z * blockIdx.z;

    thearray[iX * 16 + iY * 4 + iZ] = iX * 16 + iY * 4 + iZ;
}

extern "C" {
    void _cInitialArray(int* thearray)
    {
        dim3 block(1, 1, 1);
        dim3 th(4, 4, 4);

        _kInitialArray << <block, th >> > (thearray);
        checkCudaErrors(cudaGetLastError());
    }
}

class B
{
public:
    B()
    {
        checkCudaErrors(cudaMalloc((void**)&m_pDevicePtr, sizeof(int) * 64));
        _cInitialArray(m_pDevicePtr);
    }
    ~B()
    {
        cudaFree(m_pDevicePtr);
    }
    __device__ int GetNumber(int index)
    {
        m_pDevicePtr[index] = m_pDevicePtr[index] + 1;
        return m_pDevicePtr[index];
    }
    int* m_pDevicePtr;
};

__global__ void _kAddArray(int* thearray1, int* thearray2)
{
    int iX = threadIdx.x + blockDim.x * blockIdx.x;
    int iY = threadIdx.y + blockDim.y * blockIdx.y;
    int iZ = threadIdx.z + blockDim.z * blockIdx.z;

    thearray2[iX * 16 + iY * 4 + iZ] = thearray2[iX * 16 + iY * 4 + iZ] + 1;
    thearray1[iX * 16 + iY * 4 + iZ] = thearray1[iX * 16 + iY * 4 + iZ] + thearray2[iX * 16 + iY * 4 + iZ];
}

__global__ void _kAddArrayB(int* thearray1, B* pB)
{
    int iX = threadIdx.x + blockDim.x * blockIdx.x;
    int iY = threadIdx.y + blockDim.y * blockIdx.y;
    int iZ = threadIdx.z + blockDim.z * blockIdx.z;

    thearray1[iX * 16 + iY * 4 + iZ] = thearray1[iX * 16 + iY * 4 + iZ] + pB->GetNumber(iX * 16 + iY * 4 + iZ);
}

extern "C" {
    void _cAddArray(int* thearray1, int* thearray2)
    {
        dim3 block(1, 1, 1);
        dim3 th(4, 4, 4);

        _kAddArray << <block, th >> > (thearray1, thearray2);
        checkCudaErrors(cudaGetLastError());
    }

    void _cAddArrayB(int* thearray1, B* pB)
    {
        dim3 block(1, 1, 1);
        dim3 th(4, 4, 4);

        _kAddArrayB << <block, th >> > (thearray1, pB);
        checkCudaErrors(cudaGetLastError());
    }
}

class A
{
public:
    A() 
    { 
        checkCudaErrors(cudaMalloc((void**)&m_pDevicePtr, sizeof(int) * 64));
        _cInitialArray(m_pDevicePtr);
    }
    ~A()
    {
        checkCudaErrors(cudaFree(m_pDevicePtr));
    }

    void Add(int* toAdd)
    {
        _cAddArray(m_pDevicePtr, toAdd);
    }

    void Add(B* toAdd)
    {
        _cAddArrayB(m_pDevicePtr, toAdd);
    }

    int* m_pDevicePtr;
};



int main(int argc, char * argv[])
{
    B* pB = new B();
    A* pA = new A();
    pA->Add(pB->m_pDevicePtr);

    int* res = (int*)malloc(sizeof(int) * 64);
    checkCudaErrors(cudaMemcpy(res, pA->m_pDevicePtr, sizeof(int) * 64, cudaMemcpyDeviceToHost));
    printf("----------- A=");
    for (int i = 0; i < 8; ++i)
    {
        printf("\n");
        for (int j = 0; j < 8; ++j)
            printf("res %d=%d  ", i * 8 + j, res[i * 8 + j]);
    }
    printf("\n");
    checkCudaErrors(cudaMemcpy(res, pB->m_pDevicePtr, sizeof(int) * 64, cudaMemcpyDeviceToHost));
    printf("----------- B=");
    for (int i = 0; i < 8; ++i)
    {
        printf("\n");
        for (int j = 0; j < 8; ++j)
            printf("res %d=%d  ", i * 8 + j, res[i * 8 + j]);
    }
    printf("\n");
    B* pB2 = new B();
    A* pA2 = new A();
    pA2->Add(pB2);
    checkCudaErrors(cudaMemcpy(res, pA2->m_pDevicePtr, sizeof(int) * 64, cudaMemcpyDeviceToHost));
    printf("----------- A2=");
    for (int i = 0; i < 8; ++i)
    {
        printf("\n");
        for (int j = 0; j < 8; ++j)
            printf("res %d=%d  ", i * 8 + j, res[i * 8 + j]);
    }
    printf("\n");
    checkCudaErrors(cudaMemcpy(res, pB2->m_pDevicePtr, sizeof(int) * 64, cudaMemcpyDeviceToHost));
    printf("----------- B2=");
    for (int i = 0; i < 8; ++i)
    {
        printf("\n");
        for (int j = 0; j < 8; ++j)
            printf("res %d=%d  ", i * 8 + j, res[i * 8 + j]);
    }
    printf("\n");
    delete pA;
    delete pB;
    delete pA2;
    delete pB2;

    return 0;
}

Результаты запуска этой программы:

----------- A=
res 0=1  res 1=3  res 2=5  res 3=7  res 4=9  res 5=11  res 6=13  res 7=15
res 8=17  res 9=19  res 10=21  res 11=23  res 12=25  res 13=27  res 14=29  res 15=31
res 16=33  res 17=35  res 18=37  res 19=39  res 20=41  res 21=43  res 22=45  res 23=47
res 24=49  res 25=51  res 26=53  res 27=55  res 28=57  res 29=59  res 30=61  res 31=63
res 32=65  res 33=67  res 34=69  res 35=71  res 36=73  res 37=75  res 38=77  res 39=79
res 40=81  res 41=83  res 42=85  res 43=87  res 44=89  res 45=91  res 46=93  res 47=95
res 48=97  res 49=99  res 50=101  res 51=103  res 52=105  res 53=107  res 54=109  res 55=111
res 56=113  res 57=115  res 58=117  res 59=119  res 60=121  res 61=123  res 62=125  res 63=127
----------- B=
res 0=1  res 1=2  res 2=3  res 3=4  res 4=5  res 5=6  res 6=7  res 7=8
res 8=9  res 9=10  res 10=11  res 11=12  res 12=13  res 13=14  res 14=15  res 15=16
res 16=17  res 17=18  res 18=19  res 19=20  res 20=21  res 21=22  res 22=23  res 23=24
res 24=25  res 25=26  res 26=27  res 27=28  res 28=29  res 29=30  res 30=31  res 31=32
res 32=33  res 33=34  res 34=35  res 35=36  res 36=37  res 37=38  res 38=39  res 39=40
res 40=41  res 41=42  res 42=43  res 43=44  res 44=45  res 45=46  res 46=47  res 47=48
res 48=49  res 49=50  res 50=51  res 51=52  res 52=53  res 53=54  res 54=55  res 55=56
res 56=57  res 57=58  res 58=59  res 59=60  res 60=61  res 61=62  res 62=63  res 63=64
CUDA error at F:/CPPProject/CudaLatticeGauge/CudaLatticeGauge/Code/CudaLibTest/CudaHelper.cu:183 code=77(cudaErrorIllegalAddress) "cudaMemcpy(res, pA2->m_pDevicePtr, sizeof(int) * 64, cudaMemcpyDeviceToHost)"

Похоже, япросто не может манипулировать содержимым памяти устройства в функции-члене.Это потому, что я сделал что-то не так?или я просто не могу этого сделать?

Я использую CUDA 10 в Windows 10 и Visual Studio 2017.

Robert Crovella · Answer 1 · 14 декабря 2018

В CUDA запрещено разыменовывать указатель хоста в коде устройства.Вы делаете это здесь:

__global__ void _kAddArrayB(int* thearray1, B* pB)
{
    int iX = threadIdx.x + blockDim.x * blockIdx.x;
    int iY = threadIdx.y + blockDim.y * blockIdx.y;
    int iZ = threadIdx.z + blockDim.z * blockIdx.z;

    thearray1[iX * 16 + iY * 4 + iZ] = thearray1[iX * 16 + iY * 4 + iZ] + pB->GetNumber(iX * 16 + iY * 4 + iZ);
//                                                                        ^^^^^
}

pB в приведенном выше коде является указателем на B объект.Вы ранее создали этот B объект в памяти хоста:

B* pB2 = new B();

Ничего, на что указывает pB2, не доступно в коде устройства.Совсем.Любая попытка разыменования указателя pB2 в коде устройства приведет к ошибке выполнения кода устройства.

Я уверен, что существует ряд возможных способов исправить это.Я считаю, что простой подход без разрыва всего остального кода - просто скопировать объект, на который указывает pB2, в память устройства и использовать указатель на эту копию объекта в вашем ядре.Это может привести к изменениям только вашей подпрограммы _cAddArrayB(), например:

    void _cAddArrayB(int* thearray1, B* pB)
    {
        dim3 block(1, 1, 1);
        dim3 th(4, 4, 4);
        B *dpB;
        cudaMalloc(&dpB, sizeof(B));
        checkCudaErrors(cudaMemcpy(dpB, pB, sizeof(B), cudaMemcpyHostToDevice));
        _kAddArrayB << <block, th >> > (thearray1, dpB);
        cudaDeviceSynchronize();
        checkCudaErrors(cudaGetLastError());
        checkCudaErrors(cudaFree(dpB));
    }

Вот рабочий пример, демонстрирующий, что это изменение приводит к выполнению кода без ошибок:

$ cat t361.cu
#include <iostream>
#include <fstream>
#include <string>
#include <stack>
#include <cstdarg>

#include <limits.h>
#include <stdio.h>
#include <stdarg.h>
#include <math.h>
#include <malloc.h>
#include <stdlib.h>

#define checkCudaErrors(val) check((val), #val, __FILE__, __LINE__)

#ifdef __DRIVER_TYPES_H__
#ifndef DEVICE_RESET
#define DEVICE_RESET cudaDeviceReset();
#endif
#else
#ifndef DEVICE_RESET
#define DEVICE_RESET
#endif
#endif

#ifdef __DRIVER_TYPES_H__
static const char *_cudaGetErrorEnum(cudaError_t error) {
    return cudaGetErrorName(error);
}
#endif

template <typename T> void check(T result, char const *const func, const char *const file,
    int const line) {
    if (result) {
        fprintf(stderr, "CUDA error at %s:%d code=%d(%s) \"%s\" \n", file, line,
            static_cast<unsigned int>(result), _cudaGetErrorEnum(result), func);
        DEVICE_RESET
            // Make sure we call CUDA Device Reset before exiting
            exit(EXIT_FAILURE);
    }
}

__global__ void _kInitialArray(int* thearray)
{
    int iX = threadIdx.x + blockDim.x * blockIdx.x;
    int iY = threadIdx.y + blockDim.y * blockIdx.y;
    int iZ = threadIdx.z + blockDim.z * blockIdx.z;

    thearray[iX * 16 + iY * 4 + iZ] = iX * 16 + iY * 4 + iZ;
}

extern "C" {
    void _cInitialArray(int* thearray)
    {
        dim3 block(1, 1, 1);
        dim3 th(4, 4, 4);

        _kInitialArray << <block, th >> > (thearray);
        checkCudaErrors(cudaGetLastError());
    }
}

class B
{
public:
    B()
    {
        checkCudaErrors(cudaMalloc((void**)&m_pDevicePtr, sizeof(int) * 64));
        _cInitialArray(m_pDevicePtr);
    }
    ~B()
    {
        cudaFree(m_pDevicePtr);
    }
    __device__ int GetNumber(int index)
    {
        m_pDevicePtr[index] = m_pDevicePtr[index] + 1;
        return m_pDevicePtr[index];
    }
    int* m_pDevicePtr;
};

__global__ void _kAddArray(int* thearray1, int* thearray2)
{
    int iX = threadIdx.x + blockDim.x * blockIdx.x;
    int iY = threadIdx.y + blockDim.y * blockIdx.y;
    int iZ = threadIdx.z + blockDim.z * blockIdx.z;

    thearray2[iX * 16 + iY * 4 + iZ] = thearray2[iX * 16 + iY * 4 + iZ] + 1;
    thearray1[iX * 16 + iY * 4 + iZ] = thearray1[iX * 16 + iY * 4 + iZ] + thearray2[iX * 16 + iY * 4 + iZ];
}

__global__ void _kAddArrayB(int* thearray1, B* pB)
{
    int iX = threadIdx.x + blockDim.x * blockIdx.x;
    int iY = threadIdx.y + blockDim.y * blockIdx.y;
    int iZ = threadIdx.z + blockDim.z * blockIdx.z;

    thearray1[iX * 16 + iY * 4 + iZ] = thearray1[iX * 16 + iY * 4 + iZ] + pB->GetNumber(iX * 16 + iY * 4 + iZ);
}

extern "C" {
    void _cAddArray(int* thearray1, int* thearray2)
    {
        dim3 block(1, 1, 1);
        dim3 th(4, 4, 4);

        _kAddArray << <block, th >> > (thearray1, thearray2);
        checkCudaErrors(cudaGetLastError());
    }

    void _cAddArrayB(int* thearray1, B* pB)
    {
        dim3 block(1, 1, 1);
        dim3 th(4, 4, 4);
        B *dpB;
        cudaMalloc(&dpB, sizeof(B));
        checkCudaErrors(cudaMemcpy(dpB, pB, sizeof(B), cudaMemcpyHostToDevice));
        _kAddArrayB << <block, th >> > (thearray1, dpB);
        cudaDeviceSynchronize();
        checkCudaErrors(cudaGetLastError());
        checkCudaErrors(cudaFree(dpB));
    }
}

class A
{
public:
    A()
    {
        checkCudaErrors(cudaMalloc((void**)&m_pDevicePtr, sizeof(int) * 64));
        _cInitialArray(m_pDevicePtr);
    }
    ~A()
    {
        checkCudaErrors(cudaFree(m_pDevicePtr));
    }

    void Add(int* toAdd)
    {
        _cAddArray(m_pDevicePtr, toAdd);
    }

    void Add(B* toAdd)
    {
        _cAddArrayB(m_pDevicePtr, toAdd);
    }

    int* m_pDevicePtr;
};



int main(int argc, char * argv[])
{
    B* pB = new B();
    A* pA = new A();
    pA->Add(pB->m_pDevicePtr);

    int* res = (int*)malloc(sizeof(int) * 64);
    checkCudaErrors(cudaMemcpy(res, pA->m_pDevicePtr, sizeof(int) * 64, cudaMemcpyDeviceToHost));
    printf("----------- A=");
    for (int i = 0; i < 8; ++i)
    {
        printf("\n");
        for (int j = 0; j < 8; ++j)
            printf("res %d=%d  ", i * 8 + j, res[i * 8 + j]);
    }
    printf("\n");
    checkCudaErrors(cudaMemcpy(res, pB->m_pDevicePtr, sizeof(int) * 64, cudaMemcpyDeviceToHost));
    printf("----------- B=");
    for (int i = 0; i < 8; ++i)
    {
        printf("\n");
        for (int j = 0; j < 8; ++j)
            printf("res %d=%d  ", i * 8 + j, res[i * 8 + j]);
    }
    printf("\n");
    B* pB2 = new B();
    A* pA2 = new A();
    pA2->Add(pB2);
    checkCudaErrors(cudaMemcpy(res, pA2->m_pDevicePtr, sizeof(int) * 64, cudaMemcpyDeviceToHost));
    printf("----------- A2=");
    for (int i = 0; i < 8; ++i)
    {
        printf("\n");
        for (int j = 0; j < 8; ++j)
            printf("res %d=%d  ", i * 8 + j, res[i * 8 + j]);
    }
    printf("\n");
    checkCudaErrors(cudaMemcpy(res, pB2->m_pDevicePtr, sizeof(int) * 64, cudaMemcpyDeviceToHost));
    printf("----------- B2=");
    for (int i = 0; i < 8; ++i)
    {
        printf("\n");
        for (int j = 0; j < 8; ++j)
            printf("res %d=%d  ", i * 8 + j, res[i * 8 + j]);
    }
    printf("\n");
    delete pA;
    delete pB;
    delete pA2;
    delete pB2;

    return 0;
}
$ nvcc -o t361 t361.cu
$ cuda-memcheck ./t361
========= CUDA-MEMCHECK
----------- A=
res 0=1  res 1=3  res 2=5  res 3=7  res 4=9  res 5=11  res 6=13  res 7=15
res 8=17  res 9=19  res 10=21  res 11=23  res 12=25  res 13=27  res 14=29  res 15=31
res 16=33  res 17=35  res 18=37  res 19=39  res 20=41  res 21=43  res 22=45  res 23=47
res 24=49  res 25=51  res 26=53  res 27=55  res 28=57  res 29=59  res 30=61  res 31=63
res 32=65  res 33=67  res 34=69  res 35=71  res 36=73  res 37=75  res 38=77  res 39=79
res 40=81  res 41=83  res 42=85  res 43=87  res 44=89  res 45=91  res 46=93  res 47=95
res 48=97  res 49=99  res 50=101  res 51=103  res 52=105  res 53=107  res 54=109  res 55=111
res 56=113  res 57=115  res 58=117  res 59=119  res 60=121  res 61=123  res 62=125  res 63=127
----------- B=
res 0=1  res 1=2  res 2=3  res 3=4  res 4=5  res 5=6  res 6=7  res 7=8
res 8=9  res 9=10  res 10=11  res 11=12  res 12=13  res 13=14  res 14=15  res 15=16
res 16=17  res 17=18  res 18=19  res 19=20  res 20=21  res 21=22  res 22=23  res 23=24
res 24=25  res 25=26  res 26=27  res 27=28  res 28=29  res 29=30  res 30=31  res 31=32
res 32=33  res 33=34  res 34=35  res 35=36  res 36=37  res 37=38  res 38=39  res 39=40
res 40=41  res 41=42  res 42=43  res 43=44  res 44=45  res 45=46  res 46=47  res 47=48
res 48=49  res 49=50  res 50=51  res 51=52  res 52=53  res 53=54  res 54=55  res 55=56
res 56=57  res 57=58  res 58=59  res 59=60  res 60=61  res 61=62  res 62=63  res 63=64
----------- A2=
res 0=1  res 1=3  res 2=5  res 3=7  res 4=9  res 5=11  res 6=13  res 7=15
res 8=17  res 9=19  res 10=21  res 11=23  res 12=25  res 13=27  res 14=29  res 15=31
res 16=33  res 17=35  res 18=37  res 19=39  res 20=41  res 21=43  res 22=45  res 23=47
res 24=49  res 25=51  res 26=53  res 27=55  res 28=57  res 29=59  res 30=61  res 31=63
res 32=65  res 33=67  res 34=69  res 35=71  res 36=73  res 37=75  res 38=77  res 39=79
res 40=81  res 41=83  res 42=85  res 43=87  res 44=89  res 45=91  res 46=93  res 47=95
res 48=97  res 49=99  res 50=101  res 51=103  res 52=105  res 53=107  res 54=109  res 55=111
res 56=113  res 57=115  res 58=117  res 59=119  res 60=121  res 61=123  res 62=125  res 63=127
----------- B2=
res 0=1  res 1=2  res 2=3  res 3=4  res 4=5  res 5=6  res 6=7  res 7=8
res 8=9  res 9=10  res 10=11  res 11=12  res 12=13  res 13=14  res 14=15  res 15=16
res 16=17  res 17=18  res 18=19  res 19=20  res 20=21  res 21=22  res 22=23  res 23=24
res 24=25  res 25=26  res 26=27  res 27=28  res 28=29  res 29=30  res 30=31  res 31=32
res 32=33  res 33=34  res 34=35  res 35=36  res 36=37  res 37=38  res 38=39  res 39=40
res 40=41  res 41=42  res 42=43  res 43=44  res 44=45  res 45=46  res 46=47  res 47=48
res 48=49  res 49=50  res 50=51  res 51=52  res 52=53  res 53=54  res 54=55  res 55=56
res 56=57  res 57=58  res 58=59  res 59=60  res 60=61  res 61=62  res 62=63  res 63=64
========= ERROR SUMMARY: 0 errors
$

CUDA, Поле члена с ptr устройства и функцией члена устройства, чтобы посетить его, возможно ли это и как?

Пожалуйста, войдите или зарегистрируйтесь чтобы ответить на этот вопрос.

1 Ответ

Пожалуйста, войдите или зарегистрируйтесь что бы добавить комментарий.

CUDA, Поле члена с ptr устройства и функцией члена устройства, чтобы посетить его, возможно ли это и как?

Пожалуйста, войдите или зарегистрируйтесь чтобы ответить на этот вопрос.

1 Ответ

Пожалуйста, войдите или зарегистрируйтесь что бы добавить комментарий.

Похожие темы