如何在gpu v100上使用fma？

__global__ void fmaKernel(float* a, float* b, float* c, int size) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < size) {
        c[tid] = __fmaf_rn(a[tid], b[tid], c[tid]);
    }
}

int main() {
    int size = 1024;
    float* a, * b, * c;
    // 分配和初始化输入数据
    // ...

    // 在GPU上分配内存
    float* d_a, * d_b, * d_c;
    cudaMalloc((void**)&d_a, size * sizeof(float));
    cudaMalloc((void**)&d_b, size * sizeof(float));
    cudaMalloc((void**)&d_c, size * sizeof(float));

    // 将输入数据从主机内存复制到GPU内存
    cudaMemcpy(d_a, a, size * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, size * sizeof(float), cudaMemcpyHostToDevice);

    // 启动CUDA核函数
    int blockSize = 256;
    int numBlocks = (size + blockSize - 1) / blockSize;
    fmaKernel<<<numBlocks, blockSize>>>(d_a, d_b, d_c, size);

    // 将计算结果从GPU内存复制到主机内存
    cudaMemcpy(c, d_c, size * sizeof(float), cudaMemcpyDeviceToHost);

    // 清理GPU内存
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);

    // 处理计算结果
    // ...

    return 0;
}