CUDA の Texture Memory で float3 のような3要素の型は fetch できない

試した環境

本題

CUDA の Texture Memory は int4 や float2 のような 1,2,4 要素の整数型と単精度浮動小数点型が利用できます。 逆に言うと float3 のような3要素の型は使えません。

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#texture-memory

The type of a texel, which is restricted to the basic integer and single-precision floating-point types and any of the 1-, 2-, and 4-component vector types defined in Built-in Vector Types that are derived from the basic integer and single-precision floating-point types.

以下、 Texture Object を利用した例です。 基本型である float も含めて、以下の4つのカーネル関数はビルドが通ります。

__global__ void addKernel(float *dst, cudaTextureObject_t tex_src)
{
    int i = threadIdx.x;
    float src = tex1Dfetch<float>(tex_src, i);
    dst[i] = src + 1.0f;
}
__global__ void addKernel(float *dst, cudaTextureObject_t tex_src)
{
    int i = threadIdx.x;
    float1 src = tex1Dfetch<float1>(tex_src, i);
    dst[i] = src.x + 1.0f;
}
__global__ void addKernel(float *dst, cudaTextureObject_t tex_src)
{
    int i = threadIdx.x;
    float2 src = tex1Dfetch<float2>(tex_src, i);
    dst[i] = src.x + src.y;
}
__global__ void addKernel(float *dst, cudaTextureObject_t tex_src)
{
    int i = threadIdx.x;
    float4 src = tex1Dfetch<float4>(tex_src, i);
    dst[i] = src.x + src.y + src.z + src.w;
}

一方以下のfloat3 を使った例ではビルドエラーになります。

__global__ void addKernel(float *dst, cudaTextureObject_t tex_src)
{
    int i = threadIdx.x;
    float3 src = tex1Dfetch<float3>(tex_src, i);
    dst[i] = src.x + src.y + src.z;
}

no instance of overloaded function "tex1Dfetch" matches the argument list CudaTextureObject

float2 の場合の Texture Object の設定は以下の通りです。 cudaCreateChannelDesc(32, 32, 0, 0, cudaChannelFormatKindFloat)とします。 cudaCreateChannelDesc(64, 0, 0, 0, cudaChannelFormatKindFloat)などとするとcudaCreateTextureObjectで失敗します。

cudaResourceDesc resDesc;
resDesc.resType = cudaResourceTypeLinear;
resDesc.res.linear.devPtr = dev_src;
resDesc.res.linear.desc = cudaCreateChannelDesc(32, 32, 0, 0, cudaChannelFormatKindFloat);
resDesc.res.linear.sizeInBytes = size * sizeof(float2);

cudaTextureDesc texDesc;
std::memset(&texDesc, 0, sizeof(texDesc));
texDesc.readMode = cudaReadModeElementType;

cudaTextureObject_t tex_src = 0;
cudaCreateTextureObject(&tex_src, &resDesc, &texDesc, NULL);

参考

docs.nvidia.com

Chocolatey で Vim をインストールするときインストール先の指定によっては既存の Vim がシンボリックリンクに変わる

確認した環境

本題

Chocolatey で Vim をインストールするとき、/InstallDirパラメータでインストール先を指定することができます。 しかしそのインストール場所に既に Vim が存在するときにはその既存の Vim が Chocolatey でインストールした Vimシンボリックリンク(ショートカット)に変わります。 なので結論としては、Chocolatey で Vim をインストールするときにはインストール先を既存の Vim があるフォルダに指定しない方がいいでしょう。

以下その記録です。(ユーザー名などコンソール出力を一部書き換えています。)
まず Chocolatey で Vim をインストールする前の状態です。 KaoriYa さんで配布されている Vim が置かれています。

C:\app\vim>dir
 ドライブ C のボリューム ラベルは Windows です
 ボリューム シリアル番号は XXXX-XXXX です

 C:\app\vim のディレクトリ

2021/02/07  22:03    <DIR>          .
2021/02/07  22:03    <DIR>          ..
2021/02/07  22:03    <DIR>          vim82-kaoriya-win64
               0 個のファイル                   0 バイト
               3 個のディレクトリ  XX,XXX,XXX,XXX バイトの空き領域

このフォルダに Chocolatey で Vim をインストールします。 すると vim82-kaoriya-win64 が Chocolatey でインストールした Vimシンボリックリンクに変わります。

C:\app\vim>choco install vim --params "'/NoDesktopShortcuts /InstallDir:C:\app'" -y
Chocolatey v0.10.15
Installing the following packages:
vim
By installing you accept licenses for the packages.
Progress: Downloading vim 8.2.2467... 100%

vim v8.2.2467 [Approved]
vim package files install completed. Performing other installation steps.
Attempt to use original download file name failed for 'C:\ProgramData\chocolatey\lib\vim\tools\gvim_8.2.2467_x64.zip'.
Copying vim
  from 'C:\ProgramData\chocolatey\lib\vim\tools\gvim_8.2.2467_x64.zip'
Extracting C:\Users\____\AppData\Local\Temp\chocolatey\vim\8.2.2467\vimInstall.zip to C:\app...
C:\app
0
C:\app\vim\vim82-kaoriya-win64
C:\app\vim\vim82-kaoriya-win64 <<===>> C:\app\vim\vim82 のシンボリック リンクが 作成されました
  vim may be able to be automatically uninstalled.
 The install of vim was successful.
  Software installed to 'C:\app'

Chocolatey installed 1/1 packages.
 See the log for details (C:\ProgramData\chocolatey\logs\chocolatey.log).

Enjoy using Chocolatey? Explore more amazing features to take your
experience to the next level at
 https://chocolatey.org/compare


C:\app\vim>dir
 ドライブ C のボリューム ラベルは Windows です
 ボリューム シリアル番号は XXXX-XXXX です

 C:\app\vim のディレクトリ

2021/02/07  22:20    <DIR>          .
2021/02/07  22:20    <DIR>          ..
2021/02/06  08:03    <DIR>          vim82
2021/02/07  22:20    <SYMLINKD>     vim82-kaoriya-win64 [C:\app\vim\vim82]
2021/02/07  22:20             1,309 _vimrc
               1 個のファイル               1,309 バイト
               4 個のディレクトリ  XX,XXX,XXX,XXX バイトの空き領域

f:id:kakashibata:20210208011657p:plain

予想ですが、Chocolatey は既存の Vim を Chocolatey 経由でインストールしたものだと判定してシンボリックリンクに書き換えたのだと思います。
そのせいもあってか、この状態では Vim のアンインストールに失敗します。

C:\app\vim>cd ..

C:\app>choco uninstall vim -y
Chocolatey v0.10.15
Uninstalling the following packages:
vim

vim v8.2.2467
ERROR: 要求で指定したタグと再解析ポイントにあるタグが一致しません。
vim uninstall not successful.
Error while running 'C:\ProgramData\chocolatey\lib\vim\tools\chocolateyuninstall.ps1'.
 See log for details.
vim not uninstalled. An error occurred during uninstall:
 vim uninstall not successful.

Chocolatey uninstalled 0/1 packages. 1 packages failed.
 See the log for details (C:\ProgramData\chocolatey\logs\chocolatey.log).

Failures
 - vim (exited -1) - Error while running 'C:\ProgramData\chocolatey\lib\vim\tools\chocolateyuninstall.ps1'.
 See log for details.

If a package uninstall is failing and/or you've already uninstalled the
 software outside of Chocolatey, you can attempt to run the command
 with `-n` to skip running a chocolateyUninstall script, additionally
 adding `--skip-autouninstaller` to skip an attempt to automatically
 remove system-installed software. This will only remove the packaging
 files and not things like software installed to Programs and Features.

If a package is failing because it is a dependency of another package
 or packages, then you may first need to consider if it needs removed
 as it is typically installed as a dependency for a reason. If you
 decide that you still want to remove it, head into
 `$env:ChocolateyInstall\lib` and find the package folder you want
 removed. Then delete the folder for the package. This option should
 only be used as a last resort.

ただし Vim の実行ファイルが含まれるフォルダは削除されています。

C:\app>dir vim
 ドライブ C のボリューム ラベルは Windows です
 ボリューム シリアル番号は XXXX-XXXX です

 C:\app\vim のディレクトリ

2021/02/08  00:22    <DIR>          .
2021/02/08  00:22    <DIR>          ..
2021/02/08  00:22    <SYMLINKD>     vim82-kaoriya-win64 [C:\app\vim\vim82]
2021/02/08  00:22             1,309 _vimrc
               1 個のファイル               1,309 バイト
               3 個のディレクトリ  XX,XXX,XXX,XXX バイトの空き領域

C:\app>choco uninstall vim -y
Chocolatey v0.10.15
Uninstalling the following packages:
vim

vim v8.2.2467
ERROR: This command cannot be run due to the error: 指定されたファイルが見つかりません。.
vim uninstall not successful.
Error while running 'C:\ProgramData\chocolatey\lib\vim\tools\chocolateyuninstall.ps1'.
 See log for details.
vim not uninstalled. An error occurred during uninstall:
 vim uninstall not successful.

Chocolatey uninstalled 0/1 packages. 1 packages failed.
 See the log for details (C:\ProgramData\chocolatey\logs\chocolatey.log).

Failures
 - vim (exited -1) - Error while running 'C:\ProgramData\chocolatey\lib\vim\tools\chocolateyuninstall.ps1'.
 See log for details.

If a package uninstall is failing and/or you've already uninstalled the
 software outside of Chocolatey, you can attempt to run the command
 with `-n` to skip running a chocolateyUninstall script, additionally
 adding `--skip-autouninstaller` to skip an attempt to automatically
 remove system-installed software. This will only remove the packaging
 files and not things like software installed to Programs and Features.

If a package is failing because it is a dependency of another package
 or packages, then you may first need to consider if it needs removed
 as it is typically installed as a dependency for a reason. If you
 decide that you still want to remove it, head into
 `$env:ChocolateyInstall\lib` and find the package folder you want
 removed. Then delete the folder for the package. This option should
 only be used as a last resort.

Chocolatey の出力にある通りに、パッケージングファイルの削除だけを行うようにして Chocolatey 上での Vim のアンインストールを完了させます。

C:\app>choco uninstall vim -n
Chocolatey v0.10.15
Uninstalling the following packages:
vim

vim v8.2.2467
 Running auto uninstaller...
 Skipping auto uninstaller - 'Vim 8.2 (x64)' appears to have been uninstalled already by other means.
 vim has been successfully uninstalled.

Chocolatey uninstalled 1/1 packages.
 See the log for details (C:\ProgramData\chocolatey\logs\chocolatey.log).

参考

Chocolatey で Vim をインストールするときに指定できるパラメータについては以下のページを参考にしました。

chocolatey.org

Conemu に VS 2019 x64 tools prompt が無かったので Refresh default tasks で追加する

試した環境

本題

Conemu で Visual Studio 2019 の MSBuild が使えるコマンドプロンプトを立ち上げたかったのですが、 下図のように Visual Studio 2015 などの設定はあるのものの Visual Studio 2019 用の設定がありませんでした。

f:id:kakashibata:20210201000208p:plain

以下の手順で Visual Studio 2019 用の設定を追加しました。

  • Conemu のメインウィンドウから Setup tasks... をクリックまたは Win+Alt+T キー入力で先の図の Settings のウィンドウを開きます。
  • Add/refresh default tasks... をクリックします。下図のダイアログボックスが表示されます。

f:id:kakashibata:20210201000211p:plain

  • 上図ダイアログボックスのRefresh default tasks をクリックします。

すると下図のように Visual Studio 2019 用の設定である VS 2019 x64 tools prompt が追加されました。

f:id:kakashibata:20210201000214p:plain

最後に Save settings をクリックして、設定を保存します。

参考

conemu.github.io

令和3年度大学入学共通テストを解いた(2日目)

先週に続き、大学入学共通テスト2日目の問題を解きました。
解いたのは『数学Ⅰ・数学A』、『数学Ⅱ・数学B』、『物理』、『化学』の4科目です。
問題は先週と同じく、以下のページで公開されている問題と解答を利用させて頂きました。

edu.chunichi.co.jp

結果は以下の通りです。

  • 数学Ⅰ・数学A 72
  • 数学Ⅱ・数学B 77
  • 物理 74
  • 化学 47

2か月前に解いた時より点数は上がっていますが、問題を解くための基本的な公式を思い出せなくて苦戦しました。

ちなみに得点調整が入るようです。

得点調整について | 大学入試センター

これによると私の得点は、物理は78点、化学は55点に換算されるようですが、受験するわけではないので私にはあまり関係ないでしょう。

令和3年度大学入学共通テストを解いた

昨日実施された大学入学共通テスト1日目の問題を解きました。
解いたのは『英語(リーディング・リスニング)』、『国語』、『地理B』、『倫理、政治・経済』の4科目です。
去年に解いたセンター試験の結果および2か月前に解いたセンター追・再試験の結果はそれぞれ以下の記事です。 ちなみに去年は『地理B』以外の3科目を解いています。

問題

以下のページで公開されている問題と解答を利用させて頂きました。

edu.chunichi.co.jp

結果

結果は以下の通りです。 括弧の中は配点です。 なお、大学入学共通テストから英語の配点が変わりました。

  • 英語(リーディング) 60点(100点)
  • 英語(リスニング) 66点(100点)
  • 国語 171点(200点)
  • 地理B 87点(100点)
  • 倫理、政治・経済 88点(100点)

以下、科目ごとの感想です。 (2019/2/24追記)感想を一部加筆しました。

リーディング 60点

時間内に解き終わりませんでした。 第4問から第6問は適当に答えて適当に当たってた感じです。
出題形式がセンター試験から大幅に変更されましたが、試行調査とほぼ同じ形式だと感じました。 民間の英語テストに近い形式で、今後民間試験を導入した際にギャップを減らしたい意図があるのかもしれません。

リスニング 66点

第3問の音声と音声の間のインターバルが足りませんでした。 音声が流れているときに問題文を読んで音声を聞き逃すミスばかりしてしまいました。

国語 171点

現代文、古文、漢文に分けたときの点数は以下の通りです。

  • 現代文 91点
  • 古文 35点
  • 漢文 45点

自信を持って答えられた問題は少なかったのですが、大体正解になりました。
問題形式はセンター試験と変わらない印象です。 そもそも試行調査の時も、センター試験の形式に記述式問題の大問が1つ追加された印象だったので、その記述式問題が見送られてそうなったのだと思います。

倫理、政治・経済 88点

倫理と政治・経済に分けたときの点数は以下の通りです。

  • 倫理 44点
  • 政治・経済 44点

試行調査の問題よりもセンター試験の問題に近い印象でした。

地理B 87点

センター試験を解いた時には試験時間ぎりぎり使って解き終わってたのですが、今回は10分以上時間が余りました。

全体の感想

どの科目も解き終わったときは点数が取れている手ごたえが全く感じられず7割を下回るだろうと思っていました。 なので英語以外で8割以上取れたことは驚きです。
全体的に運がよかったかもしれません。 自信が半分くらいの問題が大体正解になりました。
一方で英語は運が味方した上での6割なので残念です。

そういえば今日は朝から運が良かったです。

__constant__ で宣言した構造体の変数を初期化する場合は初期化子リストか constexpr で宣言した変数を使用する【CUDA】

試した環境

本題

先に試したコードを載せます。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

struct Point {
    int x;
    int y;
};

__constant__ Point dev_P0 = { 100, 15 };

Point host_P1_1 = { 10, 25 };
const Point host_P1_2 = { 20, 45 };
constexpr Point host_P1_3 = { 30, 65 };

__constant__ Point dev_P1 = host_P1_3;


__global__ void addConstant(Point* dst, const Point* src) {
    int i = threadIdx.x;
    dst[i].x = src[i].x + dev_P0.x + dev_P1.x;
    dst[i].y = src[i].y + dev_P0.y + dev_P1.y;
}

int main()
{
    const int array_size = 5;
    const Point src[array_size] = { {1, -1}, {2, -2}, {3, -3}, {4, -4}, {5, -5} };
    Point dst[array_size] = { 0 };

    Point* dev_src = nullptr;
    Point* dev_dst = nullptr;

    cudaMalloc((void**)&dev_src, array_size * sizeof(Point));
    cudaMalloc((void**)&dev_dst, array_size * sizeof(Point));

    cudaMemcpy(dev_src, src, array_size * sizeof(Point), cudaMemcpyHostToDevice);

    addConstant <<<1, array_size >>> (dev_dst, dev_src);

    cudaMemcpy(dst, dev_dst, array_size * sizeof(Point), cudaMemcpyDeviceToHost);

    printf("{ {1,-1}, {2,-2}, {3,-3}, {4,-4}, {5,-5} } + P0 + P1 = { {%d,%d}, {%d,%d}, {%d,%d}, {%d,%d}, {%d,%d} }\n",
        dst[0].x, dst[0].y, dst[1].x, dst[1].y, dst[2].x, dst[2].y, dst[3].x, dst[3].y, dst[4].x, dst[4].y);

    cudaFree(dev_dst);
    cudaFree(dev_src);

    // cudaDeviceReset must be called before exiting in order for profiling and
    // tracing tools such as Nsight and Visual Profiler to show complete traces.
    cudaError_t cudaStatus = cudaDeviceReset();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceReset failed!");
        return 1;
    }

    return 0;
}

__constant__で宣言した構造体(上記コードではstruct Point)の変数を初期化する場合は、上記コードのdev_P0のように初期化子リストを使うかdev_P1のようにconstexprで宣言した変数を使用します。 __constant__ Point dev_P1 = host_P1_3;はビルドが通りますが、__constant__ Point dev_P1 = host_P1_1;__constant__ Point dev_P1 = host_P1_2;にするとビルドエラーになります。
エラーメッセージは以下の通りです。

dynamic initialization is not supported for a __constant__ variable

__constant__ で宣言した変数はホスト側で参照できない【CUDA】

試した環境

本題

__constant__ で宣言した変数はデバイス側の変数ですが、ホスト側で使用するグローバル変数のような初期化の書き方ができます。 だからと言ってホスト側で直接参照できるわけではありません。
下記のコードでは__constant__ で宣言した変数Nをデバイス側ホスト側両方で参照しようとしています。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

__constant__ int N = 100;

__global__ void addConstant(int* dst, const int* src) {
    int i = threadIdx.x;
    dst[i] = src[i] + N;
}

int main()
{
    const int array_size = 5;
    const int src[array_size] = { 1, 2, 3, 4, 5 };
    int dst[array_size] = { 0 };

    int* dev_src = nullptr;
    int* dev_dst = nullptr;

    cudaMalloc((void**)&dev_src, array_size * sizeof(int));
    cudaMalloc((void**)&dev_dst, array_size * sizeof(int));

    cudaMemcpy(dev_src, src, array_size * sizeof(int), cudaMemcpyHostToDevice);

    addConstant <<<1, array_size >>> (dev_dst, dev_src);

    cudaMemcpy(dst, dev_dst, array_size * sizeof(int), cudaMemcpyDeviceToHost);

    printf("N = %d\n", N);
    printf("{1,2,3,4,5} + 100 = {%d,%d,%d,%d,%d}\n",
        dst[0], dst[1], dst[2], dst[3], dst[4]);

    cudaFree(dev_dst);
    cudaFree(dev_src);

    return 0;
}

上記ソースコードを実行した結果は以下の通りです。

N = 0
{1,2,3,4,5} + 100 = {101,102,103,104,105}

__constant__ で宣言した変数はデバイス側の変数なのでホスト側で参照に失敗します。 当たり前といえば当たり前なのですが、私はミスしてしまいました。