pycuda 可以像 C++/CUDA 那样将浮点数解析为无符号字符数组吗？

Question

我正在尝试使用 pycuda 进行 base64 以在网络上传输数据。 我需要将浮点数转换为字节或无符号字符，在我发现 memcpy 在 CPU 上运行良好之后，我只是通过 cudamemcpy 完成了它。 我的意思是，我只是做 cuda mem 复制一些浮点值，并通过“unsigend char*”将这些值放入内核中，将其视为字节数组。

我还看到我的 c++/cuda 代码也可以正常工作，但相同的代码在 pycuda 中不起作用。

部分代码截图如下； C++/CUDA

#include <cuda_runtime.h>
#include <stdio.h>
#include <iostream>
 
using namespace std;

#define CHECK(call)                                                            \
{                                                                              \
    const cudaError_t error = call;                                            \
    if (error != cudaSuccess)                                                  \
    {                                                                          \
        fprintf(stderr, "Error: %s:%d, ", __FILE__, __LINE__);                 \
        fprintf(stderr, "code: %d, reason: %s\n", error,                       \
                cudaGetErrorString(error));                                    \
        exit(1);                                                               \
    }                                                                          \
}  
// grid 2D block 2D
__global__ void base64_encode(int N, unsigned char* in, unsigned char* out) //////////////// not float type, but uchar to treat it as byte array!!

{
    unsigned int idx = threadIdx.x + blockIdx.x * blockDim.x; 
    if (idx < N){
        out[idx] = in[idx];
        printf("cuda thread %d : %02x \n",idx, in[idx]);
    }
}

int main(int argc, char **argv)
{
    printf("%s Starting...\n", argv[0]);
 
    int dev = 0;
    cudaDeviceProp deviceProp;
    CHECK(cudaGetDeviceProperties(&deviceProp, dev));
    printf("Using Device %d: %s\n", dev, deviceProp.name);
    CHECK(cudaSetDevice(dev));
 
    int nx = 1 << 2;
    int ny = 1 << 2;

    int nxy = nx * ny;
    int nBytes = nxy * sizeof(float);
    printf("Matrix size: nx %d ny %d\n", nx, ny);
 
    float *h_A, *hostRef;
    unsigned char * gpuRef;
    h_A = (float *)malloc(nBytes); 
    hostRef = (float *)malloc(nBytes);
    gpuRef = (unsigned char *)malloc(nBytes);
    int size= (int)(nxy/ sizeof(float));

    unsigned char b[nxy];//sizeof(float)
    for (int i = 0; i < size; i++)
    {
        h_A[i] = (float)(i & 0xFF);
        cout << h_A[i]  << ", " << endl;
    }
    memset(hostRef, 0, nBytes);   
    memcpy(b, &h_A, nxy);  
    memset(gpuRef, 0, nBytes);
    
    unsigned char *d_input, *d_output;
    CHECK(cudaMalloc((void **)&d_input, nBytes)); 
    CHECK(cudaMalloc((void **)&d_output, nBytes));
 
    CHECK(cudaMemcpy(d_input, h_A, nBytes, cudaMemcpyHostToDevice)); 
 
    int dimx = 4*4; 
    dim3 block(dimx, 1);
    dim3 grid((nxy + block.x - 1) / block.x );
 
    base64_encode<<<grid, block>>>(nxy, d_input, d_output);
    CHECK(cudaDeviceSynchronize());  
    CHECK(cudaGetLastError()); 
    CHECK(cudaMemcpy(gpuRef, d_output, nBytes, cudaMemcpyDeviceToHost));
  
    for (int i = 0; i < nxy; i++) 
        printf("%02x, ",gpuRef[i]); 
 
    CHECK(cudaFree(d_input)); 
    CHECK(cudaFree(d_output));
 
    free(h_A); 
    free(hostRef);
    free(gpuRef);
 
    CHECK(cudaDeviceReset());

    return (0);
}

结果看起来不错

0, 
1, 
2, 
3, 
cuda thread 0 : 00 
cuda thread 1 : 00 
cuda thread 2 : 00 
cuda thread 3 : 00 
cuda thread 4 : 00 
cuda thread 5 : 00 
cuda thread 6 : 80 
cuda thread 7 : 3f 
cuda thread 8 : 00 
cuda thread 9 : 00 
cuda thread 10 : 00 
cuda thread 11 : 40 
cuda thread 12 : 00 
cuda thread 13 : 00 
cuda thread 14 : 40 
cuda thread 15 : 40 
00, 00, 00, 00, 00, 00, 80, 3f, 00, 00, 00, 40, 00, 00, 40, 40

PyCUDA 代码

import numpy as np
import matplotlib.pyplot as plt
import pycuda.autoinit
import pycuda.driver as drv
from pycuda.compiler import SourceModule
import pycuda.gpuarray as gpuarray

kernel = SourceModule("""
#include <stdio.h>  
using namespace std;

__global__ void base64_encode(int N, unsigned char* in, unsigned char* out){
    int idx = threadIdx.x + blockIdx.x * blockDim.x; 
    if (idx < N){
        out[idx] = in[idx];
        printf("cuda thread %d : %02x \\n",idx, in[idx]);
    }
}
""")
def gpu_rgb2gray(): 
    floatValue = np.asarray(1.0).astype(np.float32)
    floatValue_gpu = cuda.mem_alloc(floatValue.nbytes) 
    cuda.memcpy_htod(floatValue_gpu, floatValue)
    
    h_output = np.asarray(0.0).astype(np.float32)      
    d_output = cuda.mem_alloc(h_output.nbytes)
    cuda.memcpy_htod(d_output, h_output) 
    base64_encoder = kernel.get_function("base64_encode") 
    blockDim = (4, 1, 1)  
    gridDim = (1, 1, 1)  
    base64_encoder(4, floatValue_gpu, d_output, block=blockDim, grid=gridDim)
    
    h_output2 = np.array(d_output.get(), dtype=np.ubyte) 
    return 0#h_output

此代码显示错误：TypeError：参数 #0 上的类型无效（基于 0） 我可以请任何人帮助我吗？

Answer 1

首先是这个导入：

import pycuda.driver as drv

与您的其余代码不匹配。为了匹配您的其余代码，应该是：

import pycuda.driver as cuda

关于你的问题。 pycuda 抱怨的参数是这一行中的参数 0（即第一个参数）：

base64_encoder(4, floatValue_gpu, d_output, block=blockDim, grid=gridDim)
               ^

这与 float 的使用或您在问题中提出的任何主题无关。在您的内核定义中，您需要一个 32 位整数：

__global__ void base64_encode(int N, ...
                              ^^^

但是像 python 中这样的裸常量显然是另外一回事。您可以通过像这样修改调用来解决此问题：

base64_encoder(np.int32(4), floatValue_gpu, d_output, block=blockDim, grid=gridDim)
               ^^^^^^^^^^^

当我进行这两项更改并运行您的 gpu_rgb2gray() 函数时，我会从内核中得到看起来合理的 printf 输出。

进行这些更改后，尽管您在发布的代码中并没有真正使用它，但您将遇到的下一个问题是：

h_output2 = np.array(d_output.get(), dtype=np.ubyte) 

您的d_output 是DeviceAllocation object，而不是GPUArray object，因此它没有get 属性/方法。为了以最少的更改来解决这个问题，我将颠倒您用于填充该对象的方法：

h_output2 = np.empty(floatValue.nbytes, dtype=np.ubyte)
cuda.memcpy_dtoh(h_output2, d_output)

这是一个完整的例子：

$ cat t29.py
import numpy as np
import pycuda.autoinit
import pycuda.driver as cuda
from   pycuda.compiler import SourceModule
import pycuda.gpuarray as gpuarray

kernel = SourceModule("""
#include <stdio.h>
using namespace std;

__global__ void base64_encode(int N, unsigned char* in, unsigned char* out){
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N){
        out[idx] = in[idx];
        printf("cuda thread %d : %02x \\n",idx, in[idx]);
    }
}
""")
def gpu_rgb2gray():
    floatValue = np.asarray(1.0).astype(np.float32)
    floatValue_gpu = cuda.mem_alloc(floatValue.nbytes)
    cuda.memcpy_htod(floatValue_gpu, floatValue)

    h_output = np.asarray(0.0).astype(np.float32)
    d_output = cuda.mem_alloc(h_output.nbytes)
    cuda.memcpy_htod(d_output, h_output)
    base64_encoder = kernel.get_function("base64_encode")
    blockDim = (4, 1, 1)
    gridDim = (1, 1, 1)
    base64_encoder(np.int32(4), floatValue_gpu, d_output, block=blockDim, grid=gridDim)
    h_output2 = np.empty(floatValue.nbytes, dtype=np.ubyte)
    cuda.memcpy_dtoh(h_output2, d_output)
    return h_output2

print(gpu_rgb2gray())

$ cuda-memcheck  python t29.py
========= CUDA-MEMCHECK
cuda thread 0 : 00
cuda thread 1 : 00
cuda thread 2 : 80
cuda thread 3 : 3f
[  0   0 128  63]
========= ERROR SUMMARY: 0 errors
$