谁能解释一下这个 CUDA 内核是如何执行的？

Question

我是 CUDA 的初学者，我尝试了这个示例代码。

int main()
{
    int i;
    cudaError_t cudastatus;
    float in[9]={1,2,3,4,5,6,7,8,9};
    float* h_in=in;
    float* d_in={0};
    cudaMalloc((void**)&d_in,9*sizeof(float));
    cudaCheckErrors("malloc failed");
    cudastatus=cudaMemcpy(d_in,h_in,9*sizeof(float),cudaMemcpyHostToDevice);
    cudaCheckErrors("memcpyh2d failed");
    float* d_out={0};
    cudaMalloc((void**)&d_out,9*sizeof(float));
    float* out[9]={0};
    kernel<<<3,3>>>(d_in,d_out);
    cudastatus=cudaDeviceSynchronize();
    cudaError_t cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess) {
         fprintf(stderr, "Kernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
         getchar();
    }
    cudaStatus = cudaDeviceSynchronize();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching             Kernel!\n", cudaStatus);
    getchar();
      }
    cudastatus=cudaMemcpy(out,d_out,9*sizeof(float),cudaMemcpyDeviceToHost);
    cudaCheckErrors("memcpyd2h failed");
    for(i=0;i<9;i++)
    {
        printf("%f\n",out[i]);
    }
    getchar();
    return 0;

}

内核代码是这样的

__device__ void func(float temp)
{
     float a=2;
     temp=temp*a;
     return;
}
__global__ void kernel(float* d_in, float* d_out)
{
     int tid=(blockIdx.x*blockDim.x)+threadIdx.x;
     float temp=d_in[tid];
     func(temp);
     d_out[tid]=d_out[tid]+temp;
}

但是当我打印 out 数组的值时，这些值全为零。我的问题是是否为每个线程从内核调用设备函数？ GPU 中如何执行此操作？

Answer 1

任何时候您在使用 CUDA 代码时遇到问题，您都应该这样做 proper cuda error checking并使用cuda-memcheck运行您的代码。由于程序中的其他错误，或者计算机上的 CUDA 设置存在问题，内核可能根本没有执行。

这个函数:

__device__ void func(float temp)
{
     float a=2;
     temp=temp*a;
     return;
}

由于与 CUDA 无关的原因，无法做任何有用的事情。

在 C/C++ 中，当我们将参数传递给这样的函数时:

void func(float temp)

temp 参数按值传递。这意味着副本由temp组成，并提供给要使用的函数。该副本与调用环境中的任何temp独立。对 temp 的任何修改都不会显示在调用环境中。因此该函数没有做任何有用的事情。尽管它看起来是将 temp 乘以 a(本例中为 2)，但当函数返回时，修改后的值就会丢失。

为了避免这种情况，您应该了解有关 C/C++ 编程的更多信息，并考虑通过指针传递值(模拟按引用传递):

__device__ void func(float *temp)

或引用:

__device__ void func(float &temp)

这行代码也不是您想要的:

float* out[9]={0};

我认为应该是:

float out[9]={0};

您需要一个float数量数组，而不是一个float指针数组。

这是代码的“固定”版本:

#include <stdio.h>

__device__ void func(float &temp)
{
     float a=2;
     temp=temp*a;
     return;
}
__global__ void kernel(float* d_in, float* d_out)
{
     int tid=(blockIdx.x*blockDim.x)+threadIdx.x;
     float temp=d_in[tid];
     func(temp);
     d_out[tid]=d_out[tid]+temp;
}

int main()
{
    int i;
    cudaError_t cudastatus;
    float in[9]={1,2,3,4,5,6,7,8,9};
    float* h_in=in;
    float* d_in={0};
    cudastatus=cudaMalloc((void**)&d_in,9*sizeof(float));
    if (cudastatus != cudaSuccess) {
         fprintf(stderr, "cm0 fail %s\n", cudaGetErrorString(cudastatus));
    }
    cudastatus=cudaMemcpy(d_in,h_in,9*sizeof(float),cudaMemcpyHostToDevice);
    if (cudastatus != cudaSuccess) {
         fprintf(stderr, "cm1 fail %s\n", cudaGetErrorString(cudastatus));
    }
    float* d_out={0};
    cudastatus=cudaMalloc((void**)&d_out,9*sizeof(float));
    if (cudastatus != cudaSuccess) {
         fprintf(stderr, "cm2 fail %s\n", cudaGetErrorString(cudastatus));
    }
    cudaMemset(d_out, 0, 9*sizeof(float));
    float out[9]={0};
    kernel<<<3,3>>>(d_in,d_out);
    cudaDeviceSynchronize();
    cudastatus = cudaGetLastError();
    if (cudastatus != cudaSuccess) {
         fprintf(stderr, "Kernel launch failed: %s\n", cudaGetErrorString(cudastatus));
    }
    cudastatus=cudaMemcpy(out,d_out,9*sizeof(float),cudaMemcpyDeviceToHost);
    if (cudastatus != cudaSuccess) {
         fprintf(stderr, "cm3 fail %s\n", cudaGetErrorString(cudastatus));
    }
    for(i=0;i<9;i++)
    {
        printf("%f\n",out[i]);
    }
    return 0;
}