0

假设我们有；

struct collapsed {
    char **seq;
    int num;
};
...
__device__ *collapsed xdev;
...

collapsed *x_dev

cudaGetSymbolAddress((void **)&x_dev, xdev);
cudaMemcpyToSymbol(x_dev, x, sizeof(collapsed)*size); //x already defined collapsed * , this line gives ERROR

为什么你认为我在最后一行出错了: invalid device symbol？

4

1 回答 1

2

这里的第一个问题是它x_dev不是设备符号。它可能包含设备内存中的地址，但该地址不能传递给cudaMemcpyToSymbol. 调用应该只是：

cudaMemcpyToSymbol(xdev, ......);

这就引出了第二个问题。这样做：

cudaMemcpyToSymbol(xdev, x, sizeof(collapsed)*size);

将是非法的。xdev是一个指针，因此您可以复制到的唯一有效值xdev是设备地址。如果x是 a struct collapsed在设备内存中的地址，则此内存传输操作的唯一有效版本是

cudaMemcpyToSymbol(xdev, &x, sizeof(collapsed *));

IE。x必须先前已设置为设备中分配的内存地址，例如

collapsed *x;
cudaMalloc((void **)&x, sizeof(collapsed)*size);
cudaMemcpy(x, host_src, sizeof(collapsed)*size, cudaMemcpyHostToDevice);

正如所承诺的，这是一个完整的工作示例。首先是代码：

#include <cstdlib>
#include <iostream>
#include <cuda_runtime.h>

struct collapsed {
    char **seq;
    int num;
};

__device__ collapsed xdev;

__global__
void kernel(const size_t item_sz)
{
    if (threadIdx.x < xdev.num) {
        char *p = xdev.seq[threadIdx.x];
        char val = 0x30 + threadIdx.x;
        for(size_t i=0; i<item_sz; i++) {
            p[i] = val;
        }
    }
}

#define gpuQ(ans) { gpu_assert((ans), __FILE__, __LINE__); }
void gpu_assert(cudaError_t code, const char *file, const int line)
{
    if (code != cudaSuccess)
    {
        std::cerr << "gpu_assert: " << cudaGetErrorString(code) << " " 
            << file << " " << line << std::endl;
        exit(code);
    }
}

int main(void)
{

    const int nitems = 32;
    const size_t item_sz = 16;
    const size_t buf_sz = size_t(nitems) * item_sz;

    // Gpu memory for sequences
    char *_buf;
    gpuQ( cudaMalloc((void **)&_buf, buf_sz) );
    gpuQ( cudaMemset(_buf, 0x7a, buf_sz) );

    // Host array for holding sequence device pointers
    char **seq = new char*[nitems];
    size_t offset = 0;
    for(int i=0; i<nitems; i++, offset += item_sz) {
        seq[i] = _buf + offset;
    }

    // Device array holding sequence pointers
    char **_seq;
    size_t seq_sz =  sizeof(char*) * size_t(nitems);
    gpuQ( cudaMalloc((void **)&_seq, seq_sz) );
    gpuQ( cudaMemcpy(_seq, seq, seq_sz, cudaMemcpyHostToDevice) );

    // Host copy of the xdev structure to copy to the device
    collapsed xdev_host;
    xdev_host.num = nitems;
    xdev_host.seq = _seq;

    // Copy to device symbol
    gpuQ( cudaMemcpyToSymbol(xdev, &xdev_host, sizeof(collapsed)) );

    // Run Kernel
    kernel<<<1,nitems>>>(item_sz);

    // Copy back buffer
    char *buf = new char[buf_sz];
    gpuQ( cudaMemcpy(buf, _buf, buf_sz, cudaMemcpyDeviceToHost) );

    // Print out seq values
    // Each string should be ASCII starting from ´0´ (0x30)
    char *seq_vals = buf; 
    for(int i=0; i<nitems; i++, seq_vals += item_sz) {
        std::string s;
        s.append(seq_vals, item_sz);
        std::cout << s << std::endl;
    }

    return 0;
}

在这里它被编译并运行：

$ /usr/local/cuda/bin/nvcc -arch=sm_12 -Xptxas=-v -g -G -o erogol erogol.cu 
./erogol.cu(19): Warning: Cannot tell what pointer points to, assuming global memory space
ptxas info    : 8 bytes gmem, 4 bytes cmem[14]
ptxas info    : Compiling entry function '_Z6kernelm' for 'sm_12'
ptxas info    : Used 5 registers, 20 bytes smem, 4 bytes cmem[1]

$ /usr/local/cuda/bin/cuda-memcheck ./erogol 
========= CUDA-MEMCHECK
0000000000000000
1111111111111111
2222222222222222
3333333333333333
4444444444444444
5555555555555555
6666666666666666
7777777777777777
8888888888888888
9999999999999999
::::::::::::::::
;;;;;;;;;;;;;;;;
<<<<<<<<<<<<<<<<
================
>>>>>>>>>>>>>>>>
????????????????
@@@@@@@@@@@@@@@@
AAAAAAAAAAAAAAAA
BBBBBBBBBBBBBBBB
CCCCCCCCCCCCCCCC
DDDDDDDDDDDDDDDD
EEEEEEEEEEEEEEEE
FFFFFFFFFFFFFFFF
GGGGGGGGGGGGGGGG
HHHHHHHHHHHHHHHH
IIIIIIIIIIIIIIII
JJJJJJJJJJJJJJJJ
KKKKKKKKKKKKKKKK
LLLLLLLLLLLLLLLL
MMMMMMMMMMMMMMMM
NNNNNNNNNNNNNNNN
OOOOOOOOOOOOOOOO
========= ERROR SUMMARY: 0 errors

一些注意事项：

  1. 为了简化一点，我只使用了一个内存分配_buf来保存所有的字符串数据。的每个值seq都设置为 内的不同地址_buf。这在功能上等同于为每个指针运行单独的cudaMalloc调用，但要快得多。
  2. 关键概念是在主机内存中组装您希望在设备上访问的结构的副本，然后将其复制到设备。my 中的所有指针xdev_host都是设备指针。CUDA API 没有任何类型的深拷贝或自动指针转换工具，因此程序员有责任确保这是正确的。
  3. 内核中的每个线程只是用不同的 ASCII 字符填充其序列。请注意，我已将 xdev 声明为结构，而不是指向结构和复制值的指针，而不是对__device__符号的引用（再次稍微简化一下）。但除此之外，您需要使您的设计模式工作的操作顺序。
  4. 因为我只能访问计算 1.x 设备，所以编译器会发出警告。一个计算 2.x 和 3.x 不会发生这种情况，因为这些设备中改进了内存模型。该警告是正常的，可以安全地忽略。
  5. 因为每个序列只是写入 的不同部分_buf，所以我可以通过一个 cudaMemcpy 调用将所有序列传输回主机。
于 2013-06-09T02:10:35.690 回答