我有一个 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;
}
´´´
您的程序有未定义的行为。由于阵列衰减，这
__global__ void addKernel(float3 *pixeld_d[])
相当于
__global__ void addKernel(float3 **pixeld_d)
所以你已经声明你的内核函数接受一个指向 a 的指针的指针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]实际上是 afloat3*而不是 afloat3因为您在此处使用的指针类型与您实际尝试使用的缓冲区类型不匹配。
这个故事的精神：在编译器关闭之前不要随意更改代码。尝试理解为什么它拒绝编译代码。通常，原因是人们试图做一些没有意义的事情。只有在您了解问题所在以及如何修复它之后才更改代码……并且不要在 C++ 中使用 C 风格的强制转换……