问题[cuda](coding)

Cuda 没有 4D FFT 的直接实现。因此，我想将 4D FFT 分解为 4 个 1D FFT，分别对应 X、Y、Z 和 W 维度。我理解 cufftPlanMany API 最适合此用途，因为它无需使用任何 for 循环，因此速度更快。

我为此编写了一个程序。但是，4D FFT 的最终结果与 4D FFTW 实现不匹配。

以下是分别使用 FFTW 和 Cuda 库的两种实现。我仔细选择了 cufftPlanMany 函数的批次、步长和分布。但是，我不知道自己哪里做错了。如有任何帮助，我将不胜感激。

FFTW 4D实现

#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <time.h>
#include <fftw3.h>

#define PRINT_FLAG 1
#define NPRINTS 5  // print size

void printf_fftw_cmplx_array(fftw_complex *complex_array, unsigned int size) {
    for (unsigned int i = 0; i < NPRINTS; ++i) {
        printf("  (%2.4f, %2.4fi)\n", complex_array[i][0], complex_array[i][1]);
    }
    printf("...\n");
    for (unsigned int i = size - NPRINTS; i < size; ++i) {
        printf("  (%2.4f, %2.4fi)\n", complex_array[i][0], complex_array[i][1]);
    }
}

float run_test_fftw_4d(unsigned int nx, unsigned int ny, unsigned int nz, unsigned int nw) {
    srand(2025);

    // Declaration
    fftw_complex *complex_data;
    fftw_plan plan;

    unsigned int element_size = nx * ny * nz * nw;
    size_t size = sizeof(fftw_complex) * element_size;

    clock_t start, stop;
    float elapsed_time;

    // Allocate memory for input and output arrays
    complex_data = (fftw_complex *)fftw_malloc(size);

    // Initialize input complex signal
    for (unsigned int i = 0; i < element_size; ++i) {
        complex_data[i][0] = rand() / (float)RAND_MAX;
        complex_data[i][1] = 0;
    }

    // Print input stuff
    if (PRINT_FLAG) {
        printf("Complex data...\n");
        printf_fftw_cmplx_array(complex_data, element_size);
    }

    // Setup the FFT plan
    plan = fftw_plan_dft(4, (int[]){nx, ny, nz, nw}, complex_data, complex_data, FFTW_FORWARD, FFTW_ESTIMATE);

    // Start time
    start = clock();
    
    // Execute the FFT
    fftw_execute(plan);

    // End time
    stop = clock();

    // Print output stuff
    if (PRINT_FLAG) {
        printf("Fourier Coefficients...\n");
        printf_fftw_cmplx_array(complex_data, element_size);
    }

    // Compute elapsed time
    elapsed_time = (double)(stop - start) / CLOCKS_PER_SEC;

    // Clean up
    fftw_destroy_plan(plan);
    fftw_free(complex_data);
    fftw_cleanup();

    return elapsed_time;
}


int main(int argc, char **argv) {
    if (argc != 6) {
        printf("Error: This program requires exactly 5 command-line arguments.\n");
        printf("       %s <arg0> <arg1> <arg2> <arg3> <arg4>\n", argv[0]);
        printf("       arg0, arg1, arg2, arg3: FFT lengths in 4D\n");
        printf("       arg4: Number of iterations\n");
        printf("       e.g.: %s 64 64 64 64 5\n", argv[0]);
        return -1;
    }

    unsigned int nx = atoi(argv[1]);
    unsigned int ny = atoi(argv[2]);
    unsigned int nz = atoi(argv[3]);
    unsigned int nw = atoi(argv[4]);
    unsigned int niter = atoi(argv[5]);

    float sum = 0.0;
    float span_s = 0.0;
    for (unsigned int i = 0; i < niter; ++i) {
        span_s = run_test_fftw_4d(nx, ny, nz, nw);
        if (PRINT_FLAG) printf("[%d]: %.6f s\n", i, span_s);
        sum += span_s;
    }
    printf("%.6f\n", sum/(float)niter);

    return 0;
}

错误的 cuFFT4D 实现

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h> 
#include <cufft.h>
#include <math.h>

#define PRINT_FLAG 1
#define NPRINTS 5  // print size

#define CHECK_CUDA(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(EXIT_FAILURE);                                                    \
    }                                                                          \
}

#define CHECK_CUFFT(call)                                                      \
{                                                                              \
    cufftResult error;                                                         \
    if ( (error = (call)) != CUFFT_SUCCESS)                                      \
    {                                                                          \
        fprintf(stderr, "Got CUFFT error %d at %s:%d\n", error, __FILE__,      \
                __LINE__);                                                     \
        exit(EXIT_FAILURE);                                                    \
    }                                                                          \
}

void printf_cufft_cmplx_array(cufftComplex *complex_array, unsigned int size) {
    for (unsigned int i = 0; i < NPRINTS; ++i) {
        printf("  (%2.4f, %2.4fi)\n", complex_array[i].x, complex_array[i].y);
    }
    printf("...\n");
    for (unsigned int i = size - NPRINTS; i < size; ++i) {
        printf("  (%2.4f, %2.4fi)\n", complex_array[i].x, complex_array[i].y);
    }
}

float run_test_cufft_4d_4x1d(unsigned int nx, unsigned int ny, unsigned int nz, unsigned int nw) {
    srand(2025);

    // Declaration
    cufftComplex *complex_data;
    cufftComplex *d_complex_data;
    cufftHandle plan1d_x, plan1d_y, plan1d_z, plan1d_w;

    unsigned int element_size = nx * ny * nz * nw;
    size_t size = sizeof(cufftComplex) * element_size;

    cudaEvent_t start, stop;
    float elapsed_time;

    // Allocate memory for the variables on the host
    complex_data = (cufftComplex *)malloc(size);

    // Initialize input complex signal
    for (unsigned int i = 0; i < element_size; ++i) {
        complex_data[i].x = rand() / (float)RAND_MAX;
        complex_data[i].y = 0;
    }

    // Print input stuff
    if (PRINT_FLAG) {
        printf("Complex data...\n");
        printf_cufft_cmplx_array(complex_data, element_size);
    }

    // Create CUDA events
    CHECK_CUDA(cudaEventCreate(&start));
    CHECK_CUDA(cudaEventCreate(&stop));

    // Allocate device memory for complex signal and output frequency
    CHECK_CUDA(cudaMalloc((void **)&d_complex_data, size));

    int n[1] = { (int)nx };
    int embed[1] = { (int)nx };
    CHECK_CUFFT(cufftPlanMany(&plan1d_x, 1, n,       // 1D FFT of size nx
                            embed, ny * nz * nw, 1, // inembed, istride, idist
                            embed, ny * nz * nw, 1, // onembed, ostride, odist
                            CUFFT_C2C, ny * nz * nw));
    n[0] = (int)ny;
    embed[0] = (int)ny;
    CHECK_CUFFT(cufftPlanMany(&plan1d_y, 1, n,       // 1D FFT of size ny
                            embed, nz * nw, 1, // inembed, istride, idist
                            embed, nz * nw, 1, // onembed, ostride, odist
                            CUFFT_C2C, nx * nz * nw));
    n[0] = (int)nz;
    embed[0] = (int)nz;
    CHECK_CUFFT(cufftPlanMany(&plan1d_z, 1, n,       // 1D FFT of size nz
                            embed, nw, 1, // inembed, istride, idist
                            embed, nw, 1, // onembed, ostride, odist
                            CUFFT_C2C, nx * ny * nw));
    n[0] = (int)nw;
    embed[0] = (int)nw;
    CHECK_CUFFT(cufftPlanMany(&plan1d_w, 1, n,       // 1D FFT of size nw
                            embed, 1, nw, // inembed, istride, idist
                            embed, 1, nw, // onembed, ostride, odist
                            CUFFT_C2C, nx * ny * nz));

    // Record the start event
    CHECK_CUDA(cudaEventRecord(start, 0));

    // Copy host memory to device
    CHECK_CUDA(cudaMemcpy(d_complex_data, complex_data, size, cudaMemcpyHostToDevice));

    // Perform FFT along each dimension sequentially
    CHECK_CUFFT(cufftExecC2C(plan1d_x, d_complex_data, d_complex_data, CUFFT_FORWARD));
    CHECK_CUFFT(cufftDestroy(plan1d_x));
    CHECK_CUFFT(cufftExecC2C(plan1d_y, d_complex_data, d_complex_data, CUFFT_FORWARD));
    CHECK_CUFFT(cufftDestroy(plan1d_y));
    CHECK_CUFFT(cufftExecC2C(plan1d_z, d_complex_data, d_complex_data, CUFFT_FORWARD));
    CHECK_CUFFT(cufftDestroy(plan1d_z));
    CHECK_CUFFT(cufftExecC2C(plan1d_w, d_complex_data, d_complex_data, CUFFT_FORWARD));
    CHECK_CUFFT(cufftDestroy(plan1d_w));

    // Retrieve the results into host memory
    CHECK_CUDA(cudaMemcpy(complex_data, d_complex_data, size, cudaMemcpyDeviceToHost));

    // Record the stop event
    CHECK_CUDA(cudaEventRecord(stop, 0));
    CHECK_CUDA(cudaEventSynchronize(stop));

    // Print output stuff
    if (PRINT_FLAG) {
        printf("Fourier Coefficients...\n");
        printf_cufft_cmplx_array(complex_data, element_size);
    }

    // Compute elapsed time
    CHECK_CUDA(cudaEventElapsedTime(&elapsed_time, start, stop));

    // Clean up
    CHECK_CUDA(cudaFree(d_complex_data));
    CHECK_CUDA(cudaEventDestroy(start));
    CHECK_CUDA(cudaEventDestroy(stop));
    free(complex_data);

    return elapsed_time * 1e-3;
}


int main(int argc, char **argv) {
    if (argc != 6) {
        printf("Error: This program requires exactly 5 command-line arguments.\n");
        printf("       %s <arg0> <arg1> <arg2> <arg3> <arg4>\n", argv[0]);
        printf("       arg0, arg1, arg2, arg3: FFT lengths in 4D\n");
        printf("       arg4: Number of iterations\n");
        printf("       e.g.: %s 64 64 64 64 5\n", argv[0]);
        return -1;
    }

    unsigned int nx = atoi(argv[1]);
    unsigned int ny = atoi(argv[2]);
    unsigned int nz = atoi(argv[3]);
    unsigned int nw = atoi(argv[4]);
    unsigned int niter = atoi(argv[5]);

    float sum = 0.0;
    float span_s = 0.0;
    for (unsigned int i = 0; i < niter; ++i) {
        span_s = run_test_cufft_4d_4x1d(nx, ny, nz, nw);
        if (PRINT_FLAG) printf("[%d]: %.6f s\n", i, span_s);
        sum += span_s;
    }
    printf("%.6f\n", sum/(float)niter);

    CHECK_CUDA(cudaDeviceReset());
    return 0;
}

尝试两种 4x4x4x4 数组的实现，你会发现只有前几个系数匹配。我知道 FFTW 实现能产生正确的结果，因为我可以用不同的方式得到相同的结果，例如先进行 3D FFT，再进行 1D FFT，或者同时使用 FFTW 和 cuFFT 库进行 2 次 2D FFT。

我最近更新了 CUDA Toolkit 版本 12.8。nvcc 编译器现在发出以下警告：

未来版本将删除对“<compute/sm/lto>_75”之前的架构的离线编译支持（使用 -Wno-deprecated-gpu-targets 来抑制警告）。

我对 CUDA 编译有基本的了解。nvcc 编译器会针对各种架构和计算能力生成 PTX，并且 PTX 会在运行时由 NVidia 驱动程序进行 JIT 编译。

我不明白“离线编译”是什么。

我主要关心的是，这个警告对于 Compute 6.1 及更高版本设备的支持意味着什么。我需要继续支持 Compute 6.1。我的 Visual Studio CUDA 代码生成构建设置是compute_61,sm_61。如果我删除此设置，我的应用将无法在 Compute 6.1 设备上运行。运行时错误是由于缺少内核造成的。我想知道正确的代码生成选项，以便在不禁用警告的情况下支持 Compute 6.1 设备，并且不会发出此警告。

这篇文章有助于解释该选项。

以及官方文档。

void CudaRenderer::render() {

    dim3 blockDim(16, 16, 1);
    dim3 gridDim((image->width + blockDim.x - 1) / blockDim.x,
                  (image->height + blockDim.y - 1) / blockDim.y);

    int *filteredCircles,
        *lastIndices,
        sz = gridDim.y * gridDim.x ;
    cudaMalloc((void **)&filteredCircles, sizeof(int) * sz * 2000);
    cudaMalloc((void **)&lastIndices, sizeof(int) * sz);
    cudaMemset(lastIndices, 0, sizeof(int) * sz);
    filterCircles<<<gridDim, blockDim>>>(filteredCircles, lastIndices);
        for (int i = 0; i < 10; ++i)
                printf("lastIndices[%d] = %d\n", i, lastIndices[i]);
    kernelRenderCircles<<<gridDim, blockDim>>>(filteredCircles, lastIndices);
    cudaFree(filteredCircles);
    cudaDeviceSynchronize();

在添加打印之前，代码编译正常，但产生了错误的结果。在添加打印以检查 lastIndices 之后，它一直给我分段错误。我在这里做错了什么？

我对多节点数据中心还不熟悉。以下情况发生在我身上。

首先，我使用此答案中的程序来检查 CUDA 设备。我构建了它（我在那里遇到了一些问题，但那是另一个问题），可执行文件名为device_info8。

所以我登录到我的数据中心，并从登录节点运行文件

me@login01 test]$ ./device_info8 
Number of devices: 1
Device Number: 0
  Device name: Tesla V100-PCIE-16GB
  Memory Clock Rate (MHz): 856
  Memory Bus Width (bits): 4096
  Peak Memory Bandwidth (GB/s): 898.0
  Total global memory (Gbytes) 15.8
  Shared memory per block (Kbytes) 48.0
  minor-major: 0-7
  Warp-size: 32
  Concurrent kernels: yes
  Concurrent computation/communication: yes

我无法直接访问我想要测试的节点，所以我

me@login01 test]$ srun -p partition1 --nodelist Node-11 --gres=gpu:all   --pty -u bash -i  
[me@Node-11 test]$

现在我明白了

[me@Node-11 test]$./device_info8
Number of devices: 0

然而，当我运行时，nvidia-smi我可以清楚地看到我有 8 个可用的 GPU！

[me@Node-11 test]$ nvidia-smi 
Tue Dec  3 18:16:04 2024       
+-----------------------------------------------------------------------------+
| NVIDIA-SMI 418.67       Driver Version: 418.67       CUDA Version: 10.1     |
|-------------------------------+----------------------+----------------------+
| GPU  Name        Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |
| Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
|===============================+======================+======================|
|   0  Tesla V100-PCIE...  On   | 00000000:2D:00.0 Off |                    0 |
| N/A   28C    P0    26W / 250W |      0MiB / 32480MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+
|   1  Tesla V100-PCIE...  On   | 00000000:31:00.0 Off |                    0 |
| N/A   26C    P0    25W / 250W |      0MiB / 32480MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+
|   2  Tesla V100-PCIE...  On   | 00000000:35:00.0 Off |                    0 |
| N/A   26C    P0    25W / 250W |      0MiB / 32480MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+
|   3  Tesla V100-PCIE...  On   | 00000000:39:00.0 Off |                    0 |
| N/A   27C    P0    24W / 250W |      0MiB / 32480MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+
|   4  Tesla V100-PCIE...  On   | 00000000:A9:00.0 Off |                    0 |
| N/A   26C    P0    26W / 250W |      0MiB / 32480MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+
|   5  Tesla V100-PCIE...  On   | 00000000:AD:00.0 Off |                    0 |
| N/A   29C    P0    25W / 250W |      0MiB / 32480MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+
|   6  Tesla V100-PCIE...  On   | 00000000:B1:00.0 Off |                    0 |
| N/A   27C    P0    24W / 250W |      0MiB / 32480MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+
|   7  Tesla V100-PCIE...  On   | 00000000:B5:00.0 Off |                    0 |
| N/A   28C    P0    27W / 250W |      0MiB / 32480MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+
                                                                               
+-----------------------------------------------------------------------------+
| Processes:                                                       GPU Memory |
|  GPU       PID   Type   Process name                             Usage      |
|=============================================================================|
|  No running processes found                                                 |
+-----------------------------------------------------------------------------+

为什么会发生这种情况？我忽略了什么？如何让 GPU 可用于程序？

Thrust 的算法thrust::reduce_by_key对我的一个问题很有效。我想尝试使用 CUB 来更精细地控制内存和流以及与我自己的内核进行交互，但reduce_by_keyCUB 中似乎没有等效的算法。有，cub::DeviceSegmentedReduce但它似乎没有输出密钥顺序，我需要事先找到每个连续密钥段的开头，这是它自己的问题。

问题是将一个键值对数组缩减为一个数组，其中所有键都是唯一的，并且每个值都有所有相同键值的缩减。

有没有一种使用 CUB 的简单方法可以做到这一点，而我没有看到？

我是 GPU 加速方面的新手。刚刚尝试了在 CUDA 上使用简单内核进行基本的 LWJGL 绑定，没有共享内存，函数签名如下

__global__ void compute(
unsigned int n,
unsigned long long int* timeMs,
double* a, double* b, double* c, 
double *g_odata)

内核函数基本上是从上述数组（timeMs、a、b、c 等）中检索线程 ID 的数据并进行一些数学运算，然后将结果放入相应线程 ID 的 g_odata 数组中。n 是要计算的线程数（当然，它会检查线程 ID 是否超过 n）。没有共享内存或减少。

现在，当我测量内核完成所需的总时间时，出现了关于 n（总线程大小/并行度）和块大小的奇怪情况（我有一个带有 80 个多处理器的 GPU）

我clock64()在内核函数前后添加了时间戳，并收集了每个线程的总时间，很明显，线程越多，执行相同任务的速度就越慢

现在的问题：

1. 为什么总时间在 100 个线程之后就开始增加？考虑到 80 个多处理器和 10K+ 个 cuda 核心，我预计这个数字会更大，所以可能是某些配置问题？
2. 为什么内核函数在更多线程上花费更多时间？执行是否交错（即调度程序可以在其中一个线程完成之前暂停它并执行另一个线程）
3. 为什么线程达到 100 个之后会出现停滞行为？为什么它会再次起飞
4. 根据块数而变化的性能。我读到网格/块只是开发人员的观点，没有影响（特别是对于没有共享/减少的完全隔离的线程）。那么为什么这很重要，以及如何选择最佳块大小？

我正在使用 cuda 来加速我的代码，其中我循环处理每个图像。每个图像都通过 cuda 在 GPU 上进行处理。

我参考cuda-samples编写了以下代码：

• 文件名：my_cuda.cu

#include "cuda_runtime.h"


int process_one_image(args)
{
    // note that declaration of some params is omitted.
    unsigned char *h_data = (unsigned char *)malloc(size);
    unsigned char *h_rgb = (unsigned char *)malloc(size_result);
    // initialize the host memory as an image info.
    ...

    unsigned char *d_data;
    unsigned char *d_rgb;

    cudaMalloc((void **)&d_data, size);
    cudaMalloc((void **)&d_rgb, size_result);
    cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);

    // process the d_data on GPU
    ...

    // copy the result from device to host.
    cudaMemcpy(h_rgb, d_rgb, size_result, cudaMemcpyDeviceToHost);


    free(h_rgb);
    free(h_data)
    cudaFree(d_rgb);
    cudaFree(d_data);
}

在上面的代码中，cudaMalloc和cudaMemcpy位于同一个函数中process_one_image。并且代码运行正确。

但我想重复运行这段代码，例如循环运行超过 10000 次。所以我不想每次处理图像时都cudaMalloc这样做。cudaFree

所以我想将我的代码更改为以下安排。

• cuda_文件：my_cuda.cu

#include "cuda_runtime.h"

int initCuda(unsigned char *h_data, unsigned char *h_rgb, unsigned char *d_data, unsigned char *d_rgb)
{
    // note that declaration of some params is omitted.
    unsigned char *h_data = (unsigned char *)malloc(size);
    unsigned char *h_rgb = (unsigned char *)malloc(size_result);
    cudaMalloc((void **)&d_data, size);
    cudaMalloc((void **)&d_rgb, size);
}

int FinalizeCuda(unsigned char *h_data, unsigned char *h_rgb, unsigned char *d_data, unsigned char *d_rgb)
{
    cudaFree(d_data);
    cudaFree(d_rgb);
    free(h_data);
    free(h_rgb);
}

int process_one_image(unsigned char *h_data, unsigned char *h_rgb, unsigned char *d_data, unsigned char *d_rgb) // note some args are omitted such as size etc.
{

    cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);

    // process the d_data on GPU
    ...

    // copy the result from device to host.
    cudaMemcpy(h_rgb, d_rgb, size, cudaMemcpyDeviceToHost);
}

• 我的c代码：c_code.c

#include "my_cuda.cu"


int processing_loop(args)  // specific args are omitted
{
    // declaration of host and device memory
    unsigned char *h_data;
    unsigned char *h_rgb;
    unsigned char *d_data;
    unsigned char *d_rgb;

    initCuda(h_data, h_rgb, d_data, d_rgb);
    while (1)
    {
        int ret = process_one_image(h_data, h_rgb, d_data, d_rgb);
    }
    FinalizeCuda(h_data, h_rgb, d_data, d_rgb);

}

在这里，您可以注意到我只想cudaMalloc在 C 文件中运行一次，以加速此代码，但我发现它无法正常工作。它没有报告任何错误，但我没有从h_rgb.

看来（我猜）在处理时 cudaMemcpy，它无法找到正确的地址d_data并复制到d_data。

那么我该如何修复这个错误，或者cudaMalloc，这是只执行一次的正确方法吗？

整个代码位于ffio key_file_path中：

• ffio/ffio/ffio_c/ffio.c- 对应于c_code.c示例文件。
• ffio/ffio/ffio_c/yuv2rgb.cu- 对应于my_cuda.cu示例。

如何运行整个示例：

• ./compiler.sh构建可执行文件main
• 执行main位于ffio/ffio/ffio_c/test
• ffio->cudaFrame使用 Ctrl+F检查变量

CUDA 定义了一些参数，我不太确定 CUDA 的行为方式，如果我必须尊重它们，或者如果我不尊重它们，它确实会运行，但可能不会获得最佳性能。

查看时cudaGetDeviceProperties我可以看到一些参数（取自NVidia网站：https ://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__DEVICE.html#group__CUDART__DEVICE_1g1bf9d625a931d657e08db2b4391170f0 ：

concurrentKernels：设备可能可以同时执行多个内核

这只是一个标志还是有任何影响？我认为并发内核只有在使用流时才有意义，对吧？（另请参阅 asyncEngineCount）。

asyncEngineCount：异步引擎数量

这是上面给出的可以并行运行的内核数量吗？另外，我认为这只在使用流时才有效，对吗？

这让我想到了一个关于流的更普遍的问题。总共可以使用或分配多少个流？例如，如果asyncEngineCount是 3，那么拥有 3 个以上的流是否有意义？如果我所说的不是真的，那么分配比asyncEngineCount.

如果我分配的流多于 asyncEngineCount 并让多个内核同时运行，会发生什么情况？假设我分配 10 个流并同时运行 10 个内核，但 asyncEngineCount 只有 3 个。CUDA 是否会一次只运行 3 个内核，而其他内核会被阻塞，直到有一个流空闲为止？

我总是将流解释为一个简单的同步对象，当调用 CUDA 异步方法时，该对象不会发出信号，以便其他方法等待该方法退出并且流设置为有信号，但这似乎是我这边的过于简单的解释？

我找不到关于异步执行和流的很好的解释，也许有更深入的见解。CUDA文档似乎没有解释细节。也许网络或书籍上有很好的解释？

• 使用全局内存的合并读取来优化跨步访问的处理

上面的链接说：

__global__ void coalescedMultiply(float *a, float *c, int M)
{
  __shared__ float aTile[TILE_DIM][TILE_DIM],
                   transposedTile[TILE_DIM][TILE_DIM];
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  float sum = 0.0f;
  aTile[threadIdx.y][threadIdx.x] = > a[row*TILE_DIM+threadIdx.x];
  transposedTile[threadIdx.x][threadIdx.y] =
      a[(blockIdx.x*blockDim.x + threadIdx.y)*TILE_DIM +
      threadIdx.x];  
  __syncthreads();
  for (int i = 0; i < TILE_DIM; i++) {
      sum += aTile[threadIdx.y][i]* transposedTile[i][threadIdx.x];
  }
  c[row*M+col] = sum;
}

(...)
这些多方银行冲突的代价非常高昂。简单的补救措施是填充共享内存阵列，使其具有额外的列，如以下代码行所示。

__shared__ float transposedTile[TILE_DIM][TILE_DIM+1];

我的问题是，

为什么它们只填充右侧 ( transposedTile) 共享内存，而不是同时填充 (transposedTile和aTile)？