Fixed illegal memory access in [split|dense]_embedding_nobag_codegen_forward_unweighted_small_kernel (#1610)

Summary:
Pull Request resolved: #1610

`[split|dense]_embedding_nobag_codegen_forward_unweighted_small_kernel`
always accesses the LXU cache even when the cache is not available.
This diff moves cache accesses under a runtime conditional to ensure
that the access is valid.

Reviewed By: jspark1105, mjanderson09

Differential Revision: D43581917

fbshipit-source-id: ab302092ffeb27b9a7fd18205c3856abc1dd8713

Author: sryap

fbgemm_gpu/codegen/embedding_forward_split_template.cu

@@ -106,11 +106,7 @@ __global__ void {{ "dense" if dense else "split" }}_embedding_nobag_codegen_forw
             {% endif %}
 
             {% if not dense %}
-            auto weight_row_cache = WeightRow<emb_t, cache_t, cache_t>(
-                const_cast<emb_t*>(&weights[idx_j * D_emb]),
-                const_cast<cache_t*>(&lxu_cache_weights[cache_idx_j][0]),
-                D,
-                nullptr);
+
             // assume cache is fp16/fp32 which doesn't require qparams
             float2 qparams_cache = make_float2(0.0f, 0.0f);
 
@@ -128,6 +124,11 @@ __global__ void {{ "dense" if dense else "split" }}_embedding_nobag_codegen_forw
             if (d < D) {
                 {% if not dense %}
                 if (placement == PlacementType::MANAGED_CACHING && cache_idx_j != kCacheLocationMissing) {
+                    auto weight_row_cache = WeightRow<emb_t, cache_t, cache_t>(
+                        const_cast<emb_t*>(&weights[idx_j * D_emb]),
+                        const_cast<cache_t*>(&lxu_cache_weights[cache_idx_j][0]),
+                        D,
+                        nullptr);
                     Vec4T<cache_t> weight = weight_row_cache.load(d, qparams_cache);
                     weight.store(&output[output_j][d]);
                 } else {