新規ファイルをアップロード

CUDA課題/streamdivergencefft.cu

+#include <stdio.h>
+#include <math.h>
+#include <cufft.h>
+#define NX 512
+#define NY 512
+#define BLOCK_SIZE 4
+
+// 複素数の構造体定義
+struct Complex {
+    float real;
+    float imag;
+};
+
+__global__ void InputHost(Complex *input) {
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    int j = blockIdx.y * blockDim.y + threadIdx.y;
+    if (i < NX && j < NY) {
+       input[i * NY + j].real = sinf(i + j);
+       input[i * NY + j].imag = cosf(i + j);
+    }
+}
+
+int main() {
+    const int size = NX * NY;
+    const int complexSize = size * sizeof(Complex);
+
+    cufftHandle plan;
+
+    // ホスト上の出力データ
+    Complex hostOutput[size];
+
+    // デバイス上の入力データと出力データ
+    Complex *deviceInput;
+    Complex *deviceOutput;
+
+    // デバイスメモリを確保
+    cudaMalloc((void**)&deviceInput, complexSize);
+    cudaMalloc((void**)&deviceOutput, complexSize);
+
+    // ストリームの生成
+    cudaStream_t stream;
+    cudaStreamCreate(&stream);
+
+    // グリッドとブロックの次元を設定
+    dim3 blockDim(BLOCK_SIZE, BLOCK_SIZE);
+    dim3 gridDim((NX + blockDim.x - 1) / blockDim.x, (NY + blockDim.y - 1) / blockDim.y);
+
+    // カーネルを実行して入力データ初期化
+    InputHost<<<gridDim, blockDim, 0, stream>>>(deviceInput);
+
+    // 2次元FFTのプランを作成
+    cufftPlan2d(&plan, NX, NY, CUFFT_C2C);
+
+    // FFTを実行
+    cufftExecC2C(plan, (cufftComplex*)deviceInput, (cufftComplex*)deviceOutput, CUFFT_FORWARD);
+
+    // 出力データ転送
+    cudaMemcpyAsync(hostOutput, deviceOutput, complexSize, cudaMemcpyDeviceToHost, stream);
+
+    // 結果の出力
+    for (int i = 0; i < NX; i++) {
+        for (int j = 0; j < NY; j++) {
+            int idx = i * NY + j;
+            printf("Output: %f + %fI\n", hostOutput[idx].real, hostOutput[idx].imag);
+        }
+    }
+
+    cudaStreamDestroy(stream);
+    cufftDestroy(plan);
+    cudaFree(deviceInput);
+    cudaFree(deviceOutput);
+
+    return 0;
+}
\ No newline at end of file