c++ - Cuda 写入设备上的数组不会更改值

Question

我有一个 1D float3 像素数组，为了测试我正在尝试在设备上设置一个数组值。我没有收到错误，但是当我打印数组值时它显示 0。

这是我的设备代码。

__global__ void addKernel(float3 *pixeld_d[])
{
        pixeld_d[threadIdx.x + W *blockIdx.x] = &make_float3(255, 30, 123);
        printf("\n Block %d Thread %d Pixeld_d %d",blockIdx.x,threadIdx.x, pixeld_d[threadIdx.x + W * blockIdx.x]->x);
}

我的主机代码:

        float3* pixeld = new float3[W*H];
        float3** pixeld_d = new float3*[W*H];
        status = cudaMallocManaged((void **)&pixeld_d,(W*H)*sizeof(float3));
        status = cudaMemcpy(pixeld_d,pixeld, (W*H) * sizeof(float3), cudaMemcpyHostToDevice);
        addKernel << <W,H >> > (pixeld_d);

在控制台中我得到这样的结果:

 Block 811 Thread 25 Pixeld_d 0

我希望 Pixeld_d 为 255，但它是 0。

这里是完整的代码(所有注释的代码都被注释了，因为我从函数调用中删除了一些东西并且 vs 会给我构建错误):

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <SFML/Graphics.hpp>
#include <stdio.h>
#include <iostream>
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
   }
}

#define W 960
#define H 540
int mov;
#define Vector3 float3
//,Sphere sphere,Sphere light
#pragma region MyRegion




__device__ inline double dot(const Vector3& a, const Vector3& b) {
    return (a.x*b.x + a.y*b.y + a.z*b.z);
}


__device__ struct Sphere
{
    Vector3 c;
    float r;
    Sphere(Vector3 i, float j) { c = i, r = j; }
    Vector3 getNormal(const Vector3& pi) const { return (make_float3(make_float3(pi.x - c.x, pi.y - c.y, pi.z - c.z).x / r, make_float3(pi.x - c.x, pi.y - c.y, pi.z - c.z).y / r, make_float3(pi.x - c.x, pi.y - c.y, pi.z - c.z).z / r)); }

};


__device__ __host__ struct Color
{
    int r, g, b;
    Color(float a, float e, float t) { r = a, g = e, b = t; }
};
#pragma endregion
__global__ void addKernel(float3 *pixeld_d[])
{
        pixeld_d[threadIdx.x + W *blockIdx.x] = &make_float3(255, 30, 123);
        printf("\n Block %d Thread %d Pixeld_d %d",blockIdx.x,threadIdx.x, pixeld_d[threadIdx.x + W * blockIdx.x]->x);
        return;/*
        float3 black = make_float3(0, 0, 0);
        float3 red = make_float3(255, 0, 0);
        float3 white = make_float3(255, 255, 255);
        pixeld_d[threadIdx.y] = &black;
        float3 o = make_float3(blockIdx.x, threadIdx.x, 0);
        float3 d = make_float3(0, 0, 1);
        double t = 20000;
        const Vector3 oc = make_float3(o.x - sphere.c.x, o.y - sphere.c.y, o.z - sphere.c.z);
        const double b = 2 * dot(oc, d);
        const double c = dot(oc, oc) - sphere.r * sphere.r;
        double disc = b * b - 4 * c;
        if (!disc < 1e-4)
        {
            disc = sqrt(disc);
            const double t0 = -b - disc;
            const double t1 = -b + disc;
            t = (t0 < t1) ? t0 : t1;
            Vector3 pi = make_float3(o.x + make_float3(d.x * t,d.y * t, d.z * t).x, o.y + make_float3(d.x * t, d.y * t, d.z * t).y,o.z + make_float3(d.x * t, d.y * t, d.z * t).z);
            Vector3 L = make_float3(light.c.x - pi.x, light.c.y - pi.y, light.c.z - pi.z);
            Vector3 N = make_float3(make_float3(pi.x - sphere.c.x, pi.y - sphere.c.y, pi.z - sphere.c.z).x / sphere.r, make_float3(pi.x - sphere.c.x, pi.y - sphere.c.y, pi.z - sphere.c.z).y / sphere.r, make_float3(pi.x - sphere.c.x, pi.y - sphere.c.y, pi.z - sphere.c.z).z / sphere.r);
            double mg = sqrt(L.x*L.x + L.y * L.y + L.z * L.z);
            float3 Lf = make_float3(L.x / mg, L.y / mg, L.z / mg);
            mg = sqrt(N.x*N.x + N.y * N.y + N.z * N.z);
            float3 Nf = make_float3(N.x / mg, N.y / mg, N.z / mg);
            float dt = dot(Lf,Nf);
            int r = (red.x + white.x * dt)*0.5;
            int g = (red.y + white.y * dt)*0.5;
            int b = (red.z + white.z * dt)*0.5;
            if (r < 0)
                r = 0;
            if (g < 0)
                g = 0;
            if (b < 0)
                b = 0;
            pixeld_d[threadIdx.y]->x = r;
            pixeld_d[threadIdx.y]->y = g;
            pixeld_d[threadIdx.y]->z = b;

        }
*/
}

int main()
{   


    sf::RenderWindow window(sf::VideoMode(W, H), "SFML works!");
    sf::Image image;
    image.create(W, H, sf::Color::Black);
    sf::Texture tex;
    sf::Sprite sprite;

    while (window.isOpen())
    {   
        Sphere *sphere;
        Sphere *light;
        cudaMalloc((void **)&sphere, sizeof(Sphere));
        cudaMalloc((void **)&light, sizeof(Sphere));

        if (sf::Keyboard::isKeyPressed(sf::Keyboard::A))
        {
            mov -= 3;
        }
        if (sf::Keyboard::isKeyPressed(sf::Keyboard::D))
        {
            mov += 3;
        }
        window.clear();
        cudaError_t status;
        float3* pixeld = new float3[W*H];
        float3** pixeld_d = new float3*[W*H];
        status = cudaMallocManaged((void **)&pixeld_d,(W*H)*sizeof(float3));
        status = cudaMemcpy(pixeld_d,pixeld, (W*H) * sizeof(float3), cudaMemcpyHostToDevice);
        addKernel << <W,H >> > (pixeld_d);
        std::cout << cudaGetErrorString(status);
        gpuErrchk( cudaPeekAtLastError() );
        gpuErrchk( cudaDeviceSynchronize() );
        cudaMemcpy(pixeld,pixeld_d,(W*H)*sizeof(float3), cudaMemcpyDeviceToHost);
        std::cout << pixeld[399359].x;
        cudaFree(pixeld_d);
        for (int x = 0; x < W; x++)
        {
            for (int y = 0; y < H; y++)
            {
                sf::Color pixel;
                pixel.r = pixeld[x*W*y].x;
                pixel.g = pixeld[x*W*y].y;
                pixel.b = pixeld[x*W*y].z;
                image.setPixel(x, y, pixel);
            }
        }
        tex.loadFromImage(image);
        sprite.setTexture(tex, true);
        window.draw(sprite);
        window.display();
    }

//,*sphere,*light

    return 0;
}
´´´

Answer 1

您的程序有未定义的行为。由于array decay , 这个

__global__ void addKernel(float3 *pixeld_d[])

相当于

__global__ void addKernel(float3 **pixeld_d)

因此，您已声明您的内核函数将指向 float3 的指针作为输入参数。我在这里推测，但我猜这很可能是最初导致您在尝试中引入以下所有问题以使编译器关闭并编译代码的原因。你真正想写的是

__global__ void addKernel(float3 *pixeld_d)

即，向您的内核传递一个指向 float3 数组的指针，它应该将结果写入其中。

在主机端，您有pixeld_d，它是指向指向float3 的指针数组的指针，初始化为指向动态分配的指针数组

    float3** pixeld_d = new float3*[W*H];

我又在猜测了，但很可能，您实际上希望它只是一个 float3*，但编译器不允许您将其用作内核调用中的参数。紧随其后，您立即用设备内存分配的结果覆盖该指针，从而泄漏之前在此过程中分配的主机内存:

    status = cudaMallocManaged((void **)&pixeld_d,(W*H)*sizeof(float3));

请注意，此处的类型不匹配。您为 float3 数组分配缓冲区(大概是因为那是您真正想要的)，而不是 float3* 数组，这是您正在使用的类型这一点将授权。 &pixel_d 实际上是一个float3***。所以编译器会在那里发现你的错误，但你强制编译器用 C 风格的转换关闭。这是您调用未定义行为的第一个地方。不幸的是，这种错误通常不会导致崩溃，您的程序只会继续按预期运行。

然后您继续并启动您的内核，它执行以下操作:

    pixeld_d[threadIdx.x + W *blockIdx.x] = &make_float3(255, 30, 123);

在这里，您试图将一个临时对象的地址(make_float3() 的结果)分配给您的float3 数组的每个元素。我不确定你是如何设法编译这段代码的，因为它不是合法的 C++，任何 C++ 编译器(包括 nvcc)都应该拒绝编译它。即使您确实以某种方式设法编译了它:这些临时对象将在该行的末尾自动销毁，并且您在那里获得的指针将不再指向有效对象。我再次猜测，但我认为这也是为了让编译器因类型不匹配而关闭。 pixeld_d[i] 实际上是 float3* 而不是 float3 因为你在这里使用的指针类型与类型不匹配您实际尝试使用的缓冲区。

故事的士气:在编译器关闭之前，不要随意更改代码。试着理解为什么它拒绝编译代码。通常，原因是一个人试图做一些没有意义的事情。只有在您了解问题所在以及如何解决问题后才更改代码……并且不要在 C++ 中使用 C 风格的转换……